extended tm

selfdrive/camerad/cameras/camera_common.cc

@@ -138,6 +138,7 @@ bool CameraBuf::acquire() {
     const size_t globalWorkSize[] = {size_t(camera_state->ci.frame_width), size_t(camera_state->ci.frame_height)};
     const size_t localWorkSize[] = {DEBAYER_LOCAL_WORKSIZE, DEBAYER_LOCAL_WORKSIZE};
     CL_CHECK(clSetKernelArg(krnl_debayer, 2, localMemSize, 0));
+    CL_CHECK(clSetKernelArg(krnl_debayer, 3, sizeof(float), &camera_state->max_pv));
     CL_CHECK(clEnqueueNDRangeKernel(q, krnl_debayer, 2, NULL, globalWorkSize, localWorkSize,
                                     0, 0, &debayer_event));
 #else
@@ -292,13 +293,17 @@ void set_exposure_target(CameraState *c, const uint8_t *pix_ptr, int x_start, in
   const CameraBuf *b = &c->buf;
 
   uint32_t lum_binning[256] = {0};
+  uint8_t max_pv = 0;
   for (int y = y_start; y < y_end; y += y_skip) {
     for (int x = x_start; x < x_end; x += x_skip) {
       uint8_t lum = pix_ptr[(y * b->rgb_width) + x];
+      max_pv = max_pv > lum ? max_pv : lum;
       lum_binning[lum]++;
     }
   }
-
+#ifdef QCOM2
+  c->max_pv = (float)max_pv/256.0f;
+#endif
   unsigned int lum_total = (y_end - y_start) * (x_end - x_start) / x_skip / y_skip;
   unsigned int lum_cur = 0;
   int lum_med = 0;

selfdrive/camerad/cameras/camera_qcom2.h

@@ -30,6 +30,7 @@ typedef struct CameraState {
   int exposure_time_min;
   int exposure_time_max;
   float ef_filtered;
+  float max_pv;
 
   int device_iommu;
   int cdm_iommu;

selfdrive/camerad/cameras/real_debayer.cl

@@ -9,9 +9,9 @@ const __constant half3 color_correction[3] = {
   (half3)(-0.06943967, -0.01601912, 1.50009016),
 };
 
-half3 color_correct(half3 rgb) {
+half3 color_correct(half3 rgb, half mv) {
   half3 ret = (0,0,0);
-  rgb = 2*rgb / (1.0f + 2*rgb); // reinhard
+  rgb = 2*rgb*(1+(rgb/(2*mv*mv))) / (1.0f + 2*rgb); // reinhard
   ret += rgb.x * color_correction[0];
   ret += rgb.y * color_correction[1];
   ret += rgb.z * color_correction[2];
@@ -67,7 +67,8 @@ half phi(half x) {
 
 __kernel void debayer10(const __global uchar * in,
                         __global uchar * out,
-                        __local half * cached
+                        __local half * cached,
+                        float mv
                        )
 {
   const int x_global = get_global_id(0);
@@ -177,7 +178,7 @@ __kernel void debayer10(const __global uchar * in,
     }
   }
 
-  rgb = color_correct(rgb);
+  rgb = color_correct(rgb, (half)mv);
 
   out[out_idx + 0] = (uchar)(255.0f * rgb.z);
   out[out_idx + 1] = (uchar)(255.0f * rgb.y);