sadmen
/
openpilot_comma
mirror of https://github.com/commaai/openpilot.git
1
0
Fork 0
Code Issues Packages Projects Releases Wiki Activity

cleanup

Browse Source
pull/31674/head
ZwX1616 1 year ago
parent 1f6316470a
commit a8494b6a52
2 changed files with 55 additions and 36 deletions
Whitespace
Show all changes Ignore whitespace when comparing lines Ignore changes in amount of whitespace Ignore changes in whitespace at EOL
Unified View
Diff Options
Show Stats Download Patch File Download Diff File
  1. 8
      system/camerad/cameras/camera_common.cc
  2. 83
      system/camerad/cameras/real_debayer.cl

8
system/camerad/cameras/camera_common.cc
Unescape View File

@ -27,11 +27,13 @@ public:
"-cl-fast-relaxed-math -cl-denorms-are-zero " "-cl-fast-relaxed-math -cl-denorms-are-zero "
"-DFRAME_WIDTH=%d -DFRAME_HEIGHT=%d -DFRAME_STRIDE=%d -DFRAME_OFFSET=%d " "-DFRAME_WIDTH=%d -DFRAME_HEIGHT=%d -DFRAME_STRIDE=%d -DFRAME_OFFSET=%d "
"-DRGB_WIDTH=%d -DRGB_HEIGHT=%d -DYUV_STRIDE=%d -DUV_OFFSET=%d " "-DRGB_WIDTH=%d -DRGB_HEIGHT=%d -DYUV_STRIDE=%d -DUV_OFFSET=%d "
"-DIS_OX=%d -DIS_OS=%d -DCAM_NUM=%d%s", "-DIS_OX=%d -DIS_OS=%d -DIS_BGGR=%d -DCAM_NUM=%d%s",
ci->frame_width, ci->frame_height, ci->frame_stride, ci->frame_offset, ci->frame_width, ci->frame_height, ci->frame_stride, ci->frame_offset,
b->rgb_width, b->rgb_height, buf_width, uv_offset, b->rgb_width, b->rgb_height, buf_width, uv_offset,
ci->image_sensor == cereal::FrameData::ImageSensor::OX03C10, ci->image_sensor == cereal::FrameData::ImageSensor::OS04C10, ci->image_sensor == cereal::FrameData::ImageSensor::OX03C10,
s->camera_num, s->camera_num==1 ? "" : ""); ci->image_sensor == cereal::FrameData::ImageSensor::OS04C10,
ci->image_sensor == cereal::FrameData::ImageSensor::OS04C10,
s->camera_num, s->camera_num==1 ? " -DVIGNETTING" : "");
const char *cl_file = "cameras/real_debayer.cl"; const char *cl_file = "cameras/real_debayer.cl";
cl_program prg_debayer = cl_program_from_file(context, device_id, cl_file, args); cl_program prg_debayer = cl_program_from_file(context, device_id, cl_file, args);
krnl_ = CL_CHECK_ERR(clCreateKernel(prg_debayer, "debayer10", &err)); krnl_ = CL_CHECK_ERR(clCreateKernel(prg_debayer, "debayer10", &err));

83
system/camerad/cameras/real_debayer.cl
Unescape View File

@ -114,39 +114,41 @@ __kernel void debayer10(const __global uchar * in, __global uchar * out)
const int gid_x = get_global_id(0); const int gid_x = get_global_id(0);
const int gid_y = get_global_id(1); const int gid_y = get_global_id(1);
const int y_top_mod = (gid_y == 0) ? 2: 0; const int row_before_offset = (gid_y == 0) ? 2 : 0;
const int y_bot_mod = (gid_y == (RGB_HEIGHT/2 - 1)) ? 1: 3; const int row_after_offset = (gid_y == (RGB_HEIGHT/2 - 1)) ? 1 : 3;
float3 rgb; float3 rgb;
uchar3 rgb_out[4]; uchar3 rgb_out[4];
int start; int start_idx;
#if IS_OS #if IS_10BIT
bool aligned10; bool aligned10;
if (gid_x % 2 == 0) { if (gid_x % 2 == 0) {
aligned10 = true; aligned10 = true;
start = (2 * gid_y - 1) * FRAME_STRIDE + (5 * gid_x / 2 - 2) + (FRAME_STRIDE * FRAME_OFFSET); start_idx = (2 * gid_y - 1) * FRAME_STRIDE + (5 * gid_x / 2 - 2) + (FRAME_STRIDE * FRAME_OFFSET);
} else { } else {
aligned10 = false; aligned10 = false;
start = (2 * gid_y - 1) * FRAME_STRIDE + (5 * (gid_x - 1) / 2 + 1) + (FRAME_STRIDE * FRAME_OFFSET); start_idx = (2 * gid_y - 1) * FRAME_STRIDE + (5 * (gid_x - 1) / 2 + 1) + (FRAME_STRIDE * FRAME_OFFSET);
} }
#else #else
start = (2 * gid_y - 1) * FRAME_STRIDE + (3 * gid_x - 2) + (FRAME_STRIDE * FRAME_OFFSET); start_idx = (2 * gid_y - 1) * FRAME_STRIDE + (3 * gid_x - 2) + (FRAME_STRIDE * FRAME_OFFSET);
#endif #endif
// read in 8x4 chars // read in 8x4 chars
uchar8 dat[4]; uchar8 dat[4];
uchar extra[4]; dat[0] = vload8(0, in + start_idx + FRAME_STRIDE*row_before_offset);
dat[0] = vload8(0, in + start + FRAME_STRIDE*y_top_mod); dat[1] = vload8(0, in + start_idx + FRAME_STRIDE*1);
dat[1] = vload8(0, in + start + FRAME_STRIDE*1); dat[2] = vload8(0, in + start_idx + FRAME_STRIDE*2);
dat[2] = vload8(0, in + start + FRAME_STRIDE*2); dat[3] = vload8(0, in + start_idx + FRAME_STRIDE*row_after_offset);
dat[3] = vload8(0, in + start + FRAME_STRIDE*y_bot_mod);
#if IS_OS // need extra bit for 10-bit
#if IS_10BIT
uchar extra[4];
if (!aligned10) { if (!aligned10) {
extra[0] = in[start + FRAME_STRIDE*y_top_mod + 8]; extra[0] = in[start_idx + FRAME_STRIDE*row_before_offset + 8];
extra[1] = in[start + FRAME_STRIDE*1 + 8]; extra[1] = in[start_idx + FRAME_STRIDE*1 + 8];
extra[2] = in[start + FRAME_STRIDE*2 + 8]; extra[2] = in[start_idx + FRAME_STRIDE*2 + 8];
extra[3] = in[start + FRAME_STRIDE*y_bot_mod + 8]; extra[3] = in[start_idx + FRAME_STRIDE*row_after_offset + 8];
} }
#endif #endif
@ -154,24 +156,39 @@ __kernel void debayer10(const __global uchar * in, __global uchar * out)
#if VIGNETTING #if VIGNETTING
int gx = (gid_x*2 - RGB_WIDTH/2); int gx = (gid_x*2 - RGB_WIDTH/2);
int gy = (gid_y*2 - RGB_HEIGHT/2); int gy = (gid_y*2 - RGB_HEIGHT/2);
const float gain = get_vignetting_s(gx*gx + gy*gy); const float gain = get_vignetting_s(gx*gx + gy*gy); // TODO: os distance
#else #else
const float gain = 1.0; const float gain = 1.0;
#endif #endif
// process them to floats // parse into floats
#if IS_OS #if IS_10BIT
float4 vd = val4_from_10(dat[0], extra[0], aligned10, 1.0); #if IS_BGGR
float4 vc = val4_from_10(dat[1], extra[1], aligned10, 1.0); float4 vd = val4_from_10(dat[0], extra[0], aligned10, 1.0);
float4 vb = val4_from_10(dat[2], extra[2], aligned10, 1.0); float4 vc = val4_from_10(dat[1], extra[1], aligned10, 1.0);
float4 va = val4_from_10(dat[3], extra[3], aligned10, 1.0); float4 vb = val4_from_10(dat[2], extra[2], aligned10, 1.0);
float4 va = val4_from_10(dat[3], extra[3], aligned10, 1.0);
#else
float4 va = val4_from_10(dat[0], extra[0], aligned10, 1.0);
float4 vb = val4_from_10(dat[1], extra[1], aligned10, 1.0);
float4 vc = val4_from_10(dat[2], extra[2], aligned10, 1.0);
float4 vd = val4_from_10(dat[3], extra[3], aligned10, 1.0);
#endif
#else #else
float4 va = val4_from_12(dat[0], gain); #if IS_BGGR
float4 vb = val4_from_12(dat[1], gain); float4 vd = val4_from_12(dat[0], gain);
float4 vc = val4_from_12(dat[2], gain); float4 vc = val4_from_12(dat[1], gain);
float4 vd = val4_from_12(dat[3], gain); float4 vb = val4_from_12(dat[2], gain);
float4 va = val4_from_12(dat[3], gain);
#else
float4 va = val4_from_12(dat[0], gain);
float4 vb = val4_from_12(dat[1], gain);
float4 vc = val4_from_12(dat[2], gain);
float4 vd = val4_from_12(dat[3], gain);
#endif
#endif #endif
// mirror padding
if (gid_x == 0) { if (gid_x == 0) {
va.s0 = va.s2; va.s0 = va.s2;
vb.s0 = vb.s2; vb.s0 = vb.s2;
@ -192,7 +209,7 @@ __kernel void debayer10(const __global uchar * in, __global uchar * out)
rgb.x = (k02*vb.s2+k04*vb.s0)/(k02+k04); // R_G1 rgb.x = (k02*vb.s2+k04*vb.s0)/(k02+k04); // R_G1
rgb.y = vb.s1; // G1(R) rgb.y = vb.s1; // G1(R)
rgb.z = (k01*va.s1+k03*vc.s1)/(k01+k03); // B_G1 rgb.z = (k01*va.s1+k03*vc.s1)/(k01+k03); // B_G1
#if IS_OS #if IS_BGGR
rgb_out[2] = convert_uchar3_sat(color_correct(clamp(rgb, 0.0, 1.0)) * 255.0); rgb_out[2] = convert_uchar3_sat(color_correct(clamp(rgb, 0.0, 1.0)) * 255.0);
#else #else
rgb_out[0] = convert_uchar3_sat(color_correct(clamp(rgb, 0.0, 1.0)) * 255.0); rgb_out[0] = convert_uchar3_sat(color_correct(clamp(rgb, 0.0, 1.0)) * 255.0);
@ -205,7 +222,7 @@ __kernel void debayer10(const __global uchar * in, __global uchar * out)
rgb.x = vb.s2; // R rgb.x = vb.s2; // R
rgb.y = (k11*(va.s2+vc.s2)*0.5+k13*(vb.s3+vb.s1)*0.5)/(k11+k13); // G_R rgb.y = (k11*(va.s2+vc.s2)*0.5+k13*(vb.s3+vb.s1)*0.5)/(k11+k13); // G_R
rgb.z = (k12*(va.s3+vc.s1)*0.5+k14*(va.s1+vc.s3)*0.5)/(k12+k14); // B_R rgb.z = (k12*(va.s3+vc.s1)*0.5+k14*(va.s1+vc.s3)*0.5)/(k12+k14); // B_R
#if IS_OS #if IS_BGGR
rgb_out[3] = convert_uchar3_sat(color_correct(clamp(rgb, 0.0, 1.0)) * 255.0); rgb_out[3] = convert_uchar3_sat(color_correct(clamp(rgb, 0.0, 1.0)) * 255.0);
#else #else
rgb_out[1] = convert_uchar3_sat(color_correct(clamp(rgb, 0.0, 1.0)) * 255.0); rgb_out[1] = convert_uchar3_sat(color_correct(clamp(rgb, 0.0, 1.0)) * 255.0);
@ -218,7 +235,7 @@ __kernel void debayer10(const __global uchar * in, __global uchar * out)
rgb.x = (k22*(vb.s2+vd.s0)*0.5+k24*(vb.s0+vd.s2)*0.5)/(k22+k24); // R_B rgb.x = (k22*(vb.s2+vd.s0)*0.5+k24*(vb.s0+vd.s2)*0.5)/(k22+k24); // R_B
rgb.y = (k21*(vb.s1+vd.s1)*0.5+k23*(vc.s2+vc.s0)*0.5)/(k21+k23); // G_B rgb.y = (k21*(vb.s1+vd.s1)*0.5+k23*(vc.s2+vc.s0)*0.5)/(k21+k23); // G_B
rgb.z = vc.s1; // B rgb.z = vc.s1; // B
#if IS_OS #if IS_BGGR
rgb_out[0] = convert_uchar3_sat(color_correct(clamp(rgb, 0.0, 1.0)) * 255.0); rgb_out[0] = convert_uchar3_sat(color_correct(clamp(rgb, 0.0, 1.0)) * 255.0);
#else #else
rgb_out[2] = convert_uchar3_sat(color_correct(clamp(rgb, 0.0, 1.0)) * 255.0); rgb_out[2] = convert_uchar3_sat(color_correct(clamp(rgb, 0.0, 1.0)) * 255.0);
@ -231,7 +248,7 @@ __kernel void debayer10(const __global uchar * in, __global uchar * out)
rgb.x = (k31*vb.s2+k33*vd.s2)/(k31+k33); // R_G2 rgb.x = (k31*vb.s2+k33*vd.s2)/(k31+k33); // R_G2
rgb.y = vc.s2; // G2(B) rgb.y = vc.s2; // G2(B)
rgb.z = (k32*vc.s3+k34*vc.s1)/(k32+k34); // B_G2 rgb.z = (k32*vc.s3+k34*vc.s1)/(k32+k34); // B_G2
#if IS_OS #if IS_BGGR
rgb_out[1] = convert_uchar3_sat(color_correct(clamp(rgb, 0.0, 1.0)) * 255.0); rgb_out[1] = convert_uchar3_sat(color_correct(clamp(rgb, 0.0, 1.0)) * 255.0);
#else #else
rgb_out[3] = convert_uchar3_sat(color_correct(clamp(rgb, 0.0, 1.0)) * 255.0); rgb_out[3] = convert_uchar3_sat(color_correct(clamp(rgb, 0.0, 1.0)) * 255.0);

Write Preview
Loading…
Cancel
Save