adding back black level support was free

selfdrive/camerad/cameras/camera_common.cc

@@ -51,6 +51,7 @@ public:
   void queue(cl_command_queue q, cl_mem cam_buf_cl, cl_mem buf_cl, int width, int height, float gain, float black_level, cl_event *debayer_event) {
     CL_CHECK(clSetKernelArg(krnl_, 0, sizeof(cl_mem), &cam_buf_cl));
     CL_CHECK(clSetKernelArg(krnl_, 1, sizeof(cl_mem), &buf_cl));
+    CL_CHECK(clSetKernelArg(krnl_, 2, sizeof(float), &black_level));
 
     const size_t globalWorkSize[] = {size_t(width / 2), size_t(height / 2)};
     const int debayer_local_worksize = 16;

selfdrive/camerad/cameras/real_debayer.cl

@@ -38,13 +38,13 @@ float get_vignetting_s(float r) {
   }
 }
 
-float4 precolor_correct_12(uchar8 pvs, float scale) {
+float4 precolor_correct_12(uchar8 pvs, float scale, float bias) {
   uint4 parsed = (uint4)(((uint)pvs.s0<<4) + (pvs.s1>>4),  // is from the previous 10 bit
                          ((uint)pvs.s2<<4) + (pvs.s4&0xF),
                          ((uint)pvs.s3<<4) + (pvs.s4>>4),
                          ((uint)pvs.s5<<4) + (pvs.s7&0xF));
   // normalize and scale
-  float4 pv = (convert_float4(parsed) - 168.0) / (4096.0 - 168.0);
+  float4 pv = (convert_float4(parsed) - bias) / (4096.0 - bias);
   return clamp(pv*scale, 0.0, 1.0);
 }
 
@@ -52,7 +52,7 @@ float get_k(float a, float b, float c, float d) {
   return 2.0 - (fabs(a - b) + fabs(c - d));
 }
 
-__kernel void debayer10(const __global uchar * in, __global uchar * out)
+__kernel void debayer10(const __global uchar * in, __global uchar * out, float black_level)
 {
   const int gid_x = get_global_id(0);
   const int gid_y = get_global_id(1);
@@ -81,10 +81,11 @@ __kernel void debayer10(const __global uchar * in, __global uchar * out)
   #endif
 
   float4 va, vb, vc, vd;
-  va = precolor_correct_12(vac, scale);
-  vb = precolor_correct_12(vbc, scale);
-  vc = precolor_correct_12(vcc, scale);
-  vd = precolor_correct_12(vdc, scale);
+  const float bias = black_level * 4.;
+  va = precolor_correct_12(vac, scale, bias);
+  vb = precolor_correct_12(vbc, scale, bias);
+  vc = precolor_correct_12(vcc, scale, bias);
+  vd = precolor_correct_12(vdc, scale, bias);
 
   // a simplified version of https://opensignalprocessingjournal.com/contents/volumes/V6/TOSIGPJ-6-1/TOSIGPJ-6-1.pdf
   const float k01 = get_k(va.s0, vb.s1, va.s2, vb.s1);