mirror of
https://github.com/infiniteCable2/openpilot.git
synced 2026-07-27 04:12:10 +08:00
Jason Wen
acd46aa94b
* modeld: Retain pre-20hz drive model support * Method not available anymore on OP * some fixes * Revert "Long planner get accel: new function args (#34288)" * Revert "Fix low-speed allow_throttle behavior in long planner (#33894)" * Revert "long planner: allow throttle reflects usage (#33792)" * Revert "Gate acceleration on model gas press predictions (#33643)" * Reapply "Gate acceleration on model gas press predictions (#33643)" This reverts commit 76b08e37cb8eb94266ad9f6fed80db227e7c3428. * Reapply "long planner: allow throttle reflects usage (#33792)" This reverts commit c75244ca4e9c48084b0205b7c871e1a4e0f4e693. * Reapply "Fix low-speed allow_throttle behavior in long planner (#33894)" This reverts commit b2b7d21b7b685a2785d1beede3d223f0bb954807. * Reapply "Long planner get accel: new function args (#34288)" This reverts commit 74dca2fccf4da59cc8ac62ba9c0ad10ba3fc264b. * don't need * retain snpe * wrong * they're symlinks * remove * put back into VCS * add back * don't include built * Refactor model runner retrieval with caching support Added caching for active model runner type via `ModelRunnerTypeCache` to enhance performance and avoid redundant checks. Introduced a `force_check` flag to bypass the cache when necessary. Updated related code to handle cache clearing during onroad transitions. * Update model runner determination logic with caching fix Enhances `get_active_model_runner` to utilize caching more effectively by ensuring type consistency and updating cache only when necessary. Also updates `is_snpe_model` to pass the `started` state to the runner determination function, improving behavior for dynamic checks. * default to none * enable in next PR * more --------- Co-authored-by: DevTekVE <devtekve@gmail.com>
55 lines
2.5 KiB
Common Lisp
55 lines
2.5 KiB
Common Lisp
#define INTER_BITS 5
|
|
#define INTER_TAB_SIZE (1 << INTER_BITS)
|
|
#define INTER_SCALE 1.f / INTER_TAB_SIZE
|
|
|
|
#define INTER_REMAP_COEF_BITS 15
|
|
#define INTER_REMAP_COEF_SCALE (1 << INTER_REMAP_COEF_BITS)
|
|
|
|
__kernel void warpPerspective(__global const uchar * src,
|
|
int src_row_stride, int src_px_stride, int src_offset, int src_rows, int src_cols,
|
|
__global uchar * dst,
|
|
int dst_row_stride, int dst_offset, int dst_rows, int dst_cols,
|
|
__constant float * M)
|
|
{
|
|
int dx = get_global_id(0);
|
|
int dy = get_global_id(1);
|
|
|
|
if (dx < dst_cols && dy < dst_rows)
|
|
{
|
|
float X0 = M[0] * dx + M[1] * dy + M[2];
|
|
float Y0 = M[3] * dx + M[4] * dy + M[5];
|
|
float W = M[6] * dx + M[7] * dy + M[8];
|
|
W = W != 0.0f ? INTER_TAB_SIZE / W : 0.0f;
|
|
int X = rint(X0 * W), Y = rint(Y0 * W);
|
|
|
|
int sx = convert_short_sat(X >> INTER_BITS);
|
|
int sy = convert_short_sat(Y >> INTER_BITS);
|
|
|
|
short sx_clamp = clamp(sx, 0, src_cols - 1);
|
|
short sx_p1_clamp = clamp(sx + 1, 0, src_cols - 1);
|
|
short sy_clamp = clamp(sy, 0, src_rows - 1);
|
|
short sy_p1_clamp = clamp(sy + 1, 0, src_rows - 1);
|
|
int v0 = convert_int(src[mad24(sy_clamp, src_row_stride, src_offset + sx_clamp*src_px_stride)]);
|
|
int v1 = convert_int(src[mad24(sy_clamp, src_row_stride, src_offset + sx_p1_clamp*src_px_stride)]);
|
|
int v2 = convert_int(src[mad24(sy_p1_clamp, src_row_stride, src_offset + sx_clamp*src_px_stride)]);
|
|
int v3 = convert_int(src[mad24(sy_p1_clamp, src_row_stride, src_offset + sx_p1_clamp*src_px_stride)]);
|
|
|
|
short ay = (short)(Y & (INTER_TAB_SIZE - 1));
|
|
short ax = (short)(X & (INTER_TAB_SIZE - 1));
|
|
float taby = 1.f/INTER_TAB_SIZE*ay;
|
|
float tabx = 1.f/INTER_TAB_SIZE*ax;
|
|
|
|
int dst_index = mad24(dy, dst_row_stride, dst_offset + dx);
|
|
|
|
int itab0 = convert_short_sat_rte( (1.0f-taby)*(1.0f-tabx) * INTER_REMAP_COEF_SCALE );
|
|
int itab1 = convert_short_sat_rte( (1.0f-taby)*tabx * INTER_REMAP_COEF_SCALE );
|
|
int itab2 = convert_short_sat_rte( taby*(1.0f-tabx) * INTER_REMAP_COEF_SCALE );
|
|
int itab3 = convert_short_sat_rte( taby*tabx * INTER_REMAP_COEF_SCALE );
|
|
|
|
int val = v0 * itab0 + v1 * itab1 + v2 * itab2 + v3 * itab3;
|
|
|
|
uchar pix = convert_uchar_sat((val + (1 << (INTER_REMAP_COEF_BITS-1))) >> INTER_REMAP_COEF_BITS);
|
|
dst[dst_index] = pix;
|
|
}
|
|
}
|