From 677cb4313188754c462ef658be9e4f6fd34ab9e2 Mon Sep 17 00:00:00 2001 From: ZwX1616 Date: Wed, 10 Apr 2024 22:43:02 -0700 Subject: [PATCH] p1 --- system/camerad/cameras/camera_common.cc | 3 +- system/camerad/cameras/process_raw.cl | 50 ++++++------------------- system/camerad/sensors/ar0231_cl.h | 3 ++ system/camerad/sensors/os04c10_cl.h | 4 +- system/camerad/sensors/ox03c10_cl.h | 3 ++ 5 files changed, 21 insertions(+), 42 deletions(-) diff --git a/system/camerad/cameras/camera_common.cc b/system/camerad/cameras/camera_common.cc index 095e015cb9..fa09d49be3 100644 --- a/system/camerad/cameras/camera_common.cc +++ b/system/camerad/cameras/camera_common.cc @@ -27,13 +27,12 @@ public: "-cl-fast-relaxed-math -cl-denorms-are-zero -Isensors " "-DFRAME_WIDTH=%d -DFRAME_HEIGHT=%d -DFRAME_STRIDE=%d -DFRAME_OFFSET=%d " "-DRGB_WIDTH=%d -DRGB_HEIGHT=%d -DYUV_STRIDE=%d -DUV_OFFSET=%d " - "-DIS_AR=%d -DIS_OX=%d -DIS_OS=%d -DIS_10BIT=%d -DHDR_OFFSET=%d -DVIGNETTING=%d ", + "-DIS_AR=%d -DIS_OX=%d -DIS_OS=%d -DHDR_OFFSET=%d -DVIGNETTING=%d ", ci->frame_width, ci->frame_height, ci->hdr_offset > 0 ? ci->frame_stride * 2 : ci->frame_stride, ci->frame_offset, b->rgb_width, b->rgb_height, buf_width, uv_offset, ci->image_sensor == cereal::FrameData::ImageSensor::AR0231, ci->image_sensor == cereal::FrameData::ImageSensor::OX03C10, ci->image_sensor == cereal::FrameData::ImageSensor::OS04C10, - ci->mipi_format == CAM_FORMAT_MIPI_RAW_10, ci->hdr_offset, s->camera_num == 1); const char *cl_file = "cameras/process_raw.cl"; diff --git a/system/camerad/cameras/process_raw.cl b/system/camerad/cameras/process_raw.cl index a1903ba322..6331e62c23 100644 --- a/system/camerad/cameras/process_raw.cl +++ b/system/camerad/cameras/process_raw.cl @@ -53,24 +53,6 @@ float4 val4_from_12(uchar8 pvs, float gain) { } -float4 val4_from_10(uchar8 pvs, uchar ext, bool aligned, float gain) { - uint4 parsed; - if (aligned) { - parsed = (uint4)(((uint)pvs.s0 << 2) + (pvs.s1 & 0b00000011), - ((uint)pvs.s2 << 2) + ((pvs.s6 & 0b11000000) / 64), - ((uint)pvs.s3 << 2) + ((pvs.s6 & 0b00110000) / 16), - ((uint)pvs.s4 << 2) + ((pvs.s6 & 0b00001100) / 4)); - } else { - parsed = (uint4)(((uint)pvs.s0 << 2) + ((pvs.s3 & 0b00110000) / 16), - ((uint)pvs.s1 << 2) + ((pvs.s3 & 0b00001100) / 4), - ((uint)pvs.s2 << 2) + ((pvs.s3 & 0b00000011)), - ((uint)pvs.s4 << 2) + ((ext & 0b11000000) / 64)); - } - - float4 pv = (convert_float4(parsed) - 64.0) / (1024.0 - 64.0); - return clamp(pv*gain, 0.0, 1.0); -} - float combine_pvs(float lv, float sv, int expo) { float svc = fmax(sv * expo, 61376.0); float svd = sv * fmin(expo, 8.0) / 8; @@ -137,7 +119,8 @@ __kernel void process_raw(const __global uchar * in, __global uchar * out, int e uchar3 rgb_out[4]; // output is 2x2 window int start_idx; - #if IS_10BIT + + #if BIT_DEPTH == 10 bool aligned10; if (gid_x % 2 == 0) { aligned10 = true; @@ -162,7 +145,7 @@ __kernel void process_raw(const __global uchar * in, __global uchar * out, int e } dat[2] = vload8(0, in + start_idx + FRAME_STRIDE*2); dat[3] = vload8(0, in + start_idx + FRAME_STRIDE*row_after_offset); - #if HDR_COMBINE + #if HDR_OFFSET > 0 uchar8 short_dat[4]; short_dat[0] = vload8(0, in + start_idx + FRAME_STRIDE*(row_before_offset+HDR_OFFSET/2) + FRAME_STRIDE/2); short_dat[1] = vload8(0, in + start_idx + FRAME_STRIDE*(1+HDR_OFFSET/2) + FRAME_STRIDE/2); @@ -171,23 +154,19 @@ __kernel void process_raw(const __global uchar * in, __global uchar * out, int e #endif // need extra bit for 10-bit - #if IS_10BIT + #if BIT_DEPTH == 10 uchar extra[4]; + uchar short_extra[4]; if (!aligned10) { extra[0] = in[start_idx + FRAME_STRIDE*row_before_offset + 8]; extra[1] = in[start_idx + FRAME_STRIDE*1 + 8]; extra[2] = in[start_idx + FRAME_STRIDE*2 + 8]; extra[3] = in[start_idx + FRAME_STRIDE*row_after_offset + 8]; - } - #if HDR_COMBINE - uchar short_extra[4]; - if (!aligned10) { - short_extra[0] = in[start_idx + FRAME_STRIDE*(row_before_offset+HDR_OFFSET/2) + FRAME_STRIDE/2 + 8]; - short_extra[1] = in[start_idx + FRAME_STRIDE*(1+HDR_OFFSET/2) + FRAME_STRIDE/2 + 8]; - short_extra[2] = in[start_idx + FRAME_STRIDE*(2+HDR_OFFSET/2) + FRAME_STRIDE/2 + 8]; - short_extra[3] = in[start_idx + FRAME_STRIDE*(row_after_offset+HDR_OFFSET/2) + FRAME_STRIDE/2 + 8]; - } - #endif + short_extra[0] = in[start_idx + FRAME_STRIDE*(row_before_offset+HDR_OFFSET/2) + FRAME_STRIDE/2 + 8]; + short_extra[1] = in[start_idx + FRAME_STRIDE*(1+HDR_OFFSET/2) + FRAME_STRIDE/2 + 8]; + short_extra[2] = in[start_idx + FRAME_STRIDE*(2+HDR_OFFSET/2) + FRAME_STRIDE/2 + 8]; + short_extra[3] = in[start_idx + FRAME_STRIDE*(row_after_offset+HDR_OFFSET/2) + FRAME_STRIDE/2 + 8]; + } #endif // correct vignetting @@ -201,18 +180,11 @@ __kernel void process_raw(const __global uchar * in, __global uchar * out, int e float4 v_rows[4]; // parse into floats - #if IS_10BIT - #if HDR_COMBINE + #if BIT_DEPTH == 10 v_rows[ROW_READ_ORDER[0]] = val4_from_10x2(dat[0], extra[0], short_dat[0], short_extra[0], aligned10, gain, expo_time); v_rows[ROW_READ_ORDER[1]] = val4_from_10x2(dat[1], extra[1], short_dat[1], short_extra[1], aligned10, gain, expo_time); v_rows[ROW_READ_ORDER[2]] = val4_from_10x2(dat[2], extra[2], short_dat[2], short_extra[2], aligned10, gain, expo_time); v_rows[ROW_READ_ORDER[3]] = val4_from_10x2(dat[3], extra[3], short_dat[3], short_extra[3], aligned10, gain, expo_time); - #else - v_rows[ROW_READ_ORDER[0]] = val4_from_10(dat[0], extra[0], aligned10, gain); - v_rows[ROW_READ_ORDER[1]] = val4_from_10(dat[1], extra[1], aligned10, gain); - v_rows[ROW_READ_ORDER[2]] = val4_from_10(dat[2], extra[2], aligned10, gain); - v_rows[ROW_READ_ORDER[3]] = val4_from_10(dat[3], extra[3], aligned10, gain); - #endif #else v_rows[ROW_READ_ORDER[0]] = val4_from_12(dat[0], gain); v_rows[ROW_READ_ORDER[1]] = val4_from_12(dat[1], gain); diff --git a/system/camerad/sensors/ar0231_cl.h b/system/camerad/sensors/ar0231_cl.h index 6b33fc2e9c..7ec3cf0f26 100644 --- a/system/camerad/sensors/ar0231_cl.h +++ b/system/camerad/sensors/ar0231_cl.h @@ -1,5 +1,8 @@ #if IS_AR +#define BIT_DEPTH 12 +#define BLACK_LVL 168 + float3 color_correct(float3 rgb) { float3 corrected = rgb.x * (float3)(1.82717181, -0.31231438, 0.07307673); corrected += rgb.y * (float3)(-0.5743977, 1.36858544, -0.53183455); diff --git a/system/camerad/sensors/os04c10_cl.h b/system/camerad/sensors/os04c10_cl.h index a963b7bb2f..91d5c75a2b 100644 --- a/system/camerad/sensors/os04c10_cl.h +++ b/system/camerad/sensors/os04c10_cl.h @@ -1,7 +1,9 @@ #if IS_OS #define BGGR -#define HDR_COMBINE + +#define BIT_DEPTH 10 +#define BLACK_LVL 64 float3 color_correct(float3 rgb) { float3 corrected = rgb.x * (float3)(1.55361989, -0.268894615, -0.000593219); diff --git a/system/camerad/sensors/ox03c10_cl.h b/system/camerad/sensors/ox03c10_cl.h index 825c2e4135..7baaf23ce3 100644 --- a/system/camerad/sensors/ox03c10_cl.h +++ b/system/camerad/sensors/ox03c10_cl.h @@ -1,5 +1,8 @@ #if IS_OX +#define BIT_DEPTH 12 +#define BLACK_LVL 64 + constant float ox03c10_lut[] = { 0.0000e+00, 5.9488e-08, 1.1898e-07, 1.7846e-07, 2.3795e-07, 2.9744e-07, 3.5693e-07, 4.1642e-07, 4.7591e-07, 5.3539e-07, 5.9488e-07, 6.5437e-07, 7.1386e-07, 7.7335e-07, 8.3284e-07, 8.9232e-07, 9.5181e-07, 1.0113e-06, 1.0708e-06, 1.1303e-06, 1.1898e-06, 1.2493e-06, 1.3087e-06, 1.3682e-06, 1.4277e-06, 1.4872e-06, 1.5467e-06, 1.6062e-06, 1.6657e-06, 1.7252e-06, 1.7846e-06, 1.8441e-06, 1.9036e-06, 1.9631e-06, 2.0226e-06, 2.0821e-06, 2.1416e-06, 2.2011e-06, 2.2606e-06, 2.3200e-06, 2.3795e-06, 2.4390e-06, 2.4985e-06, 2.5580e-06, 2.6175e-06, 2.6770e-06, 2.7365e-06, 2.7959e-06, 2.8554e-06, 2.9149e-06, 2.9744e-06, 3.0339e-06, 3.0934e-06, 3.1529e-06, 3.2124e-06, 3.2719e-06, 3.3313e-06, 3.3908e-06, 3.4503e-06, 3.5098e-06, 3.5693e-06, 3.6288e-06, 3.6883e-06, 3.7478e-06, 3.8072e-06, 3.8667e-06, 3.9262e-06, 3.9857e-06, 4.0452e-06, 4.1047e-06, 4.1642e-06, 4.2237e-06, 4.2832e-06, 4.3426e-06, 4.4021e-06, 4.4616e-06, 4.5211e-06, 4.5806e-06, 4.6401e-06, 4.6996e-06, 4.7591e-06, 4.8185e-06, 4.8780e-06, 4.9375e-06, 4.9970e-06, 5.0565e-06, 5.1160e-06, 5.1755e-06, 5.2350e-06, 5.2945e-06, 5.3539e-06, 5.4134e-06, 5.4729e-06, 5.5324e-06, 5.5919e-06, 5.6514e-06, 5.7109e-06, 5.7704e-06, 5.8298e-06, 5.8893e-06, 5.9488e-06, 6.0083e-06, 6.0678e-06, 6.1273e-06, 6.1868e-06, 6.2463e-06, 6.3058e-06, 6.3652e-06, 6.4247e-06, 6.4842e-06, 6.5437e-06, 6.6032e-06, 6.6627e-06, 6.7222e-06, 6.7817e-06, 6.8411e-06, 6.9006e-06, 6.9601e-06, 7.0196e-06, 7.0791e-06, 7.1386e-06, 7.1981e-06, 7.2576e-06, 7.3171e-06, 7.3765e-06, 7.4360e-06, 7.4955e-06, 7.5550e-06, 7.6145e-06, 7.6740e-06, 7.7335e-06, 7.7930e-06, 7.8524e-06, 7.9119e-06, 7.9714e-06, 8.0309e-06, 8.0904e-06, 8.1499e-06, 8.2094e-06, 8.2689e-06, 8.3284e-06, 8.3878e-06, 8.4473e-06, 8.5068e-06, 8.5663e-06, 8.6258e-06, 8.6853e-06, 8.7448e-06, 8.8043e-06, 8.8637e-06, 8.9232e-06, 8.9827e-06, 9.0422e-06, 9.1017e-06, 9.1612e-06, 9.2207e-06, 9.2802e-06, 9.3397e-06, 9.3991e-06, 9.4586e-06, 9.5181e-06, 9.5776e-06, 9.6371e-06, 9.6966e-06, 9.7561e-06, 9.8156e-06, 9.8750e-06, 9.9345e-06, 9.9940e-06, 1.0054e-05, 1.0113e-05, 1.0172e-05, 1.0232e-05, 1.0291e-05, 1.0351e-05, 1.0410e-05, 1.0470e-05, 1.0529e-05, 1.0589e-05, 1.0648e-05, 1.0708e-05, 1.0767e-05, 1.0827e-05, 1.0886e-05, 1.0946e-05, 1.1005e-05, 1.1065e-05, 1.1124e-05, 1.1184e-05, 1.1243e-05, 1.1303e-05, 1.1362e-05, 1.1422e-05, 1.1481e-05, 1.1541e-05, 1.1600e-05, 1.1660e-05, 1.1719e-05, 1.1779e-05, 1.1838e-05, 1.1898e-05, 1.1957e-05, 1.2017e-05, 1.2076e-05, 1.2136e-05, 1.2195e-05, 1.2255e-05, 1.2314e-05, 1.2374e-05, 1.2433e-05, 1.2493e-05, 1.2552e-05, 1.2612e-05, 1.2671e-05, 1.2730e-05, 1.2790e-05, 1.2849e-05, 1.2909e-05, 1.2968e-05, 1.3028e-05, 1.3087e-05, 1.3147e-05, 1.3206e-05, 1.3266e-05, 1.3325e-05, 1.3385e-05, 1.3444e-05, 1.3504e-05, 1.3563e-05, 1.3623e-05, 1.3682e-05, 1.3742e-05, 1.3801e-05, 1.3861e-05, 1.3920e-05, 1.3980e-05, 1.4039e-05, 1.4099e-05, 1.4158e-05, 1.4218e-05, 1.4277e-05, 1.4337e-05, 1.4396e-05, 1.4456e-05, 1.4515e-05, 1.4575e-05, 1.4634e-05, 1.4694e-05, 1.4753e-05, 1.4813e-05, 1.4872e-05, 1.4932e-05, 1.4991e-05, 1.5051e-05, 1.5110e-05, 1.5169e-05, // NOLINT 1.5229e-05, 1.5288e-05, 1.5348e-05, 1.5407e-05, 1.5467e-05, 1.5526e-05, 1.5586e-05, 1.5645e-05, 1.5705e-05, 1.5764e-05, 1.5824e-05, 1.5883e-05, 1.5943e-05, 1.6002e-05, 1.6062e-05, 1.6121e-05, 1.6181e-05, 1.6240e-05, 1.6300e-05, 1.6359e-05, 1.6419e-05, 1.6478e-05, 1.6538e-05, 1.6597e-05, 1.6657e-05, 1.6716e-05, 1.6776e-05, 1.6835e-05, 1.6895e-05, 1.6954e-05, 1.7014e-05, 1.7073e-05, 1.7133e-05, 1.7192e-05, 1.7252e-05, 1.7311e-05, 1.7371e-05, 1.7430e-05, 1.7490e-05, 1.7549e-05, 1.7609e-05, 1.7668e-05, 1.7727e-05, 1.7787e-05, 1.7846e-05, 1.7906e-05, 1.7965e-05, 1.8025e-05, 1.8084e-05, 1.8144e-05, 1.8203e-05, 1.8263e-05, 1.8322e-05, 1.8382e-05, 1.8441e-05, 1.8501e-05, 1.8560e-05, 1.8620e-05, 1.8679e-05, 1.8739e-05, 1.8798e-05, 1.8858e-05, 1.8917e-05, 1.8977e-05, 1.9036e-05, 1.9096e-05, 1.9155e-05, 1.9215e-05, 1.9274e-05, 1.9334e-05, 1.9393e-05, 1.9453e-05, 1.9512e-05, 1.9572e-05, 1.9631e-05, 1.9691e-05, 1.9750e-05, 1.9810e-05, 1.9869e-05, 1.9929e-05, 1.9988e-05, 2.0048e-05, 2.0107e-05, 2.0167e-05, 2.0226e-05, 2.0285e-05, 2.0345e-05, 2.0404e-05, 2.0464e-05, 2.0523e-05, 2.0583e-05, 2.0642e-05, 2.0702e-05, 2.0761e-05, 2.0821e-05, 2.0880e-05, 2.0940e-05, 2.0999e-05, 2.1059e-05, 2.1118e-05, 2.1178e-05, 2.1237e-05, 2.1297e-05, 2.1356e-05, 2.1416e-05, 2.1475e-05, 2.1535e-05, 2.1594e-05, 2.1654e-05, 2.1713e-05, 2.1773e-05, 2.1832e-05, 2.1892e-05, 2.1951e-05, 2.2011e-05, 2.2070e-05, 2.2130e-05, 2.2189e-05, 2.2249e-05, 2.2308e-05, 2.2368e-05, 2.2427e-05, 2.2487e-05, 2.2546e-05, 2.2606e-05, 2.2665e-05, 2.2725e-05, 2.2784e-05, 2.2843e-05, 2.2903e-05, 2.2962e-05, 2.3022e-05, 2.3081e-05, 2.3141e-05, 2.3200e-05, 2.3260e-05, 2.3319e-05, 2.3379e-05, 2.3438e-05, 2.3498e-05, 2.3557e-05, 2.3617e-05, 2.3676e-05, 2.3736e-05, 2.3795e-05, 2.3855e-05, 2.3914e-05, 2.3974e-05, 2.4033e-05, 2.4093e-05, 2.4152e-05, 2.4212e-05, 2.4271e-05, 2.4331e-05, 2.4390e-05, 2.4450e-05, 2.4509e-05, 2.4569e-05, 2.4628e-05, 2.4688e-05, 2.4747e-05, 2.4807e-05, 2.4866e-05, 2.4926e-05, 2.4985e-05, 2.5045e-05, 2.5104e-05, 2.5164e-05, 2.5223e-05, 2.5282e-05, 2.5342e-05, 2.5401e-05, 2.5461e-05, 2.5520e-05, 2.5580e-05, 2.5639e-05, 2.5699e-05, 2.5758e-05, 2.5818e-05, 2.5877e-05, 2.5937e-05, 2.5996e-05, 2.6056e-05, 2.6115e-05, 2.6175e-05, 2.6234e-05, 2.6294e-05, 2.6353e-05, 2.6413e-05, 2.6472e-05, 2.6532e-05, 2.6591e-05, 2.6651e-05, 2.6710e-05, 2.6770e-05, 2.6829e-05, 2.6889e-05, 2.6948e-05, 2.7008e-05, 2.7067e-05, 2.7127e-05, 2.7186e-05, 2.7246e-05, 2.7305e-05, 2.7365e-05, 2.7424e-05, 2.7484e-05, 2.7543e-05, 2.7603e-05, 2.7662e-05, 2.7722e-05, 2.7781e-05, 2.7840e-05, 2.7900e-05, 2.7959e-05, 2.8019e-05, 2.8078e-05, 2.8138e-05, 2.8197e-05, 2.8257e-05, 2.8316e-05, 2.8376e-05, 2.8435e-05, 2.8495e-05, 2.8554e-05, 2.8614e-05, 2.8673e-05, 2.8733e-05, 2.8792e-05, 2.8852e-05, 2.8911e-05, 2.8971e-05, 2.9030e-05, 2.9090e-05, 2.9149e-05, 2.9209e-05, 2.9268e-05, 2.9328e-05, 2.9387e-05, 2.9447e-05, 2.9506e-05, 2.9566e-05, 2.9625e-05, 2.9685e-05, 2.9744e-05, 2.9804e-05, 2.9863e-05, 2.9923e-05, 2.9982e-05, 3.0042e-05, 3.0101e-05, 3.0161e-05, 3.0220e-05, 3.0280e-05, 3.0339e-05, 3.0398e-05, // NOLINT