mirror of
https://github.com/firestar5683/StarPilot.git
synced 2026-09-08 09:03:42 +08:00
Jason Wen
f84d27c4ee
* tinygrad with snpe * force with snpe to validate * fix path * fix more paths * Adjust modeld execution logic based on active model runner Introduced a check to conditionally execute `modeld` based on the active model runner. Added support for distinguishing between SNPE and TinyGrad runners using new helper functions and updated `custom.capnp` definitions. This change optimizes process management by ensuring compatibility with the selected model runner. * Refactor modeld process function checks. Introduce `is_stock_model` to clarify logic and replace direct uses of `is_snpe_model` where the stock model condition is needed. Additionally, rename the duplicate "modeld" process in sunnyPilot to "modeld_snpe" for clarity and consistency. * ignore tg * fix process name * ruff * fix thneed paths * mypy * remove our own * use upstream compile3 * fix thneed * try this * Revert "remove our own" This reverts commit 1cf4f57502565c274628e185c66f1eb04c4d9bbe. * try using compile2.py again * add back symlink * fix path * more fix * wrong path again * Revert "wrong path again" This reverts commit f5301c19d594defb23cda2604168bcb02c830f3e. * update * hardcode path to our submodule * force path * try this * fix file name * try this * again * Revert "again" This reverts commit 17c8cd73768a2de5aaf2a34d202aa6f03d59d72d. * Revert "try this" This reverts commit 767f78bbcf17b08edbc20f283d10a1990bf4264f. * Revert "fix file name" This reverts commit 485eef68da981da571fe9a9ebf3aecdb95e5588c. * Revert "try this" This reverts commit 41fef87680cdda1ba8cbd7e5c59256e6c57010a0. * Revert "force path" This reverts commit 5c3b408937bff0f61b150971b328d127501c5bf0. * Revert "hardcode path to our submodule" This reverts commit 5ee1950b6f4fce73e7e73b6971b996a9989d0a93. * Revert "update" This reverts commit fb313bd7fbeb50111f3c4a11137b4cfd9cabc3fa. * Reapply "wrong path again" This reverts commit 309639aeb3575e20b215bbbcc27e48923ee95c87. * Revert "wrong path again" This reverts commit f5301c19d594defb23cda2604168bcb02c830f3e. * Revert "more fix" This reverts commit 23dd423e78e60bd6a0bd73e921bfba9a5aad88bd. * Revert "fix path" This reverts commit 75d338f2bda2dfa77044e33f4c6553390621b13b. * Revert "add back symlink" This reverts commit 9f71ad0b8ae2068e431c455af9d4f954d024eec7. * Revert "try using compile2.py again" This reverts commit 914117d2e1c6c4270f7bef1a5305847a1d7bab17. * Reapply "remove our own" This reverts commit b1996377b346658f2094cba9032a66230784cb3e. * don't even compile anymore * need it for default snpe model * add to lfs * bring onnx back for sim * must add this back * need this --------- 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;
|
|
}
|
|
}
|