IQ.Lvbs/IQ.Pilot
8
7
Fork 3
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

IQ.Pilot Prebuilt Release @ 27f668a

Browse Source
This commit is contained in:
IQ.Lvbs CI [bot]
2026-09-03 18:23:24 -05:00
commit b073c5182b
2554 changed files with 679696 additions and 0 deletions
Download Patch File Download Diff File Expand all files Collapse all files

57
iqpilot/selfdrive/iqmodeld/transforms/warp_geometry.cl Normal file
View File

@@ -0,0 +1,57 @@
/*
Copyright © IQ.Lvbs, apart of Project Teal Lvbs, All Rights Reserved, licensed under https://konn3kt.com/tos/
*/
#define INTER_BITS 5
#define INTER_TAB_SIZE (1 << INTER_BITS)
#define INTER_REMAP_COEF_BITS 15
#define INTER_REMAP_COEF_SCALE (1 << INTER_REMAP_COEF_BITS)
__kernel void projectPlaneBilinear(__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);
int y = rint(y0 * w);
short sx = convert_short_sat(x >> INTER_BITS);
short sy = convert_short_sat(y >> INTER_BITS);
short min_col = (short)0;
short max_col = convert_short_sat(src_cols - 1);
short min_row = (short)0;
short max_row = convert_short_sat(src_rows - 1);
short sx_clamp = clamp(sx, min_col, max_col);
short sx_p1_clamp = clamp((short)(sx + 1), min_col, max_col);
short sy_clamp = clamp(sy, min_row, max_row);
short sy_p1_clamp = clamp((short)(sy + 1), min_row, max_row);
int top_left = convert_int(src[mad24(sy_clamp, src_row_stride, src_offset + sx_clamp * src_px_stride)]);
int top_right = convert_int(src[mad24(sy_clamp, src_row_stride, src_offset + sx_p1_clamp * src_px_stride)]);
int bottom_left = convert_int(src[mad24(sy_p1_clamp, src_row_stride, src_offset + sx_clamp * src_px_stride)]);
int bottom_right = 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 coeff0 = convert_short_sat_rte((1.0f - taby) * (1.0f - tabx) * INTER_REMAP_COEF_SCALE);
int coeff1 = convert_short_sat_rte((1.0f - taby) * tabx * INTER_REMAP_COEF_SCALE);
int coeff2 = convert_short_sat_rte(taby * (1.0f - tabx) * INTER_REMAP_COEF_SCALE);
int coeff3 = convert_short_sat_rte(taby * tabx * INTER_REMAP_COEF_SCALE);
int blended = top_left * coeff0 + top_right * coeff1 + bottom_left * coeff2 + bottom_right * coeff3;
int dst_index = mad24(dy, dst_row_stride, dst_offset + dx);
dst[dst_index] = convert_uchar_sat((blended + (1 << (INTER_REMAP_COEF_BITS - 1))) >> INTER_REMAP_COEF_BITS);
}
}

46
iqpilot/selfdrive/iqmodeld/transforms/yuv.cl Normal file
View File

@@ -0,0 +1,46 @@
/*
Copyright © IQ.Lvbs, apart of Project Teal Lvbs, All Rights Reserved, licensed under https://konn3kt.com/tos/
*/
#define UV_SIZE ((TRANSFORMED_WIDTH/2)*(TRANSFORMED_HEIGHT/2))
__kernel void packLumaHalves(__global uchar8 const * const in_luma,
__global uchar * out_frame,
int out_offset)
{
const int gid = get_global_id(0);
const int output_index_start = gid * 8;
const int row = output_index_start / TRANSFORMED_WIDTH;
const int col = output_index_start % TRANSFORMED_WIDTH;
const uchar8 luma_block = in_luma[gid];
__global uchar *top_or_left;
__global uchar *bottom_or_right;
if ((row & 1) == 0) {
top_or_left = out_frame + out_offset;
bottom_or_right = out_frame + out_offset + UV_SIZE * 2;
} else {
top_or_left = out_frame + out_offset + UV_SIZE;
bottom_or_right = out_frame + out_offset + UV_SIZE * 3;
}
const int row_stride = (row / 2) * (TRANSFORMED_WIDTH / 2) + col / 2;
vstore4(luma_block.s0246, 0, top_or_left + row_stride);
vstore4(luma_block.s1357, 0, bottom_or_right + row_stride);
}
__kernel void packChromaPlane(__global uchar8 const * const in_plane,
__global uchar8 * out_frame,
int out_offset)
{
const int gid = get_global_id(0);
out_frame[gid + out_offset / 8] = in_plane[gid];
}
__kernel void copyPlaneBytes(__global uchar8 * in_plane,
__global uchar8 * out_plane,
int in_offset,
int out_offset)
{
const int gid = get_global_id(0);
out_plane[gid + out_offset / 8] = in_plane[gid + in_offset / 8];
}