From d8ddc1208ebe60c4f51c3c2fb669ad21c3aa6d12 Mon Sep 17 00:00:00 2001
From: ZwX1616 <zwx1616@gmail.com>
Date: Sun, 21 Feb 2021 23:23:40 -0800
Subject: [PATCH] tici camerart III: prehistoric visuals (#20012)

* reworked

* add ref

* f16

* faster

* extended tm

* fix hdr

* Revert "extended tm"

This reverts commit 55ba3ae3205f7e962bf2f36634b23a595700c9a8.

* reduce static/temporal noise

* split kernels to sync max

* no print

* cleanup qcom2

* optimize

* tune DC

* doesnt work

* minor fix

* adaptive

* smoothen HCG switch

* t

* histogram ceiling

* recalulate params

* group_hold

* less contrast

* recalibrate ccm

* better ae

* better ae revised

* made some changes

* clean up

* midtones

* some improvements

* more

* cleanup

* remove more junk

* lgtm

Co-authored-by: Comma Device <device@comma.ai>
---
 selfdrive/camerad/cameras/camera_common.cc |  27 +-
 selfdrive/camerad/cameras/camera_qcom2.cc  |  25 +-
 selfdrive/camerad/cameras/camera_qcom2.h   |   9 +-
 selfdrive/camerad/cameras/real_debayer.cl  | 275 +++++----
 selfdrive/camerad/cameras/sensor2_i2c.h    | 638 +--------------------
 5 files changed, 220 insertions(+), 754 deletions(-)

diff --git a/selfdrive/camerad/cameras/camera_common.cc b/selfdrive/camerad/cameras/camera_common.cc
index 2ebdac67c8..df472d5b4e 100644
--- a/selfdrive/camerad/cameras/camera_common.cc
+++ b/selfdrive/camerad/cameras/camera_common.cc
@@ -134,10 +134,12 @@ bool CameraBuf::acquire() {
     CL_CHECK(clSetKernelArg(krnl_debayer, 0, sizeof(cl_mem), &camrabuf_cl));
     CL_CHECK(clSetKernelArg(krnl_debayer, 1, sizeof(cl_mem), &cur_rgb_buf->buf_cl));
 #ifdef QCOM2
-    constexpr int localMemSize = (DEBAYER_LOCAL_WORKSIZE + 2 * (3 / 2)) * (DEBAYER_LOCAL_WORKSIZE + 2 * (3 / 2)) * sizeof(float);
+    constexpr int localMemSize = (DEBAYER_LOCAL_WORKSIZE + 2 * (3 / 2)) * (DEBAYER_LOCAL_WORKSIZE + 2 * (3 / 2)) * sizeof(short int);
     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));
+    int ggain = camera_state->analog_gain + 4*camera_state->dc_gain_enabled;
+    CL_CHECK(clSetKernelArg(krnl_debayer, 3, sizeof(int), &ggain));
     CL_CHECK(clEnqueueNDRangeKernel(q, krnl_debayer, 2, NULL, globalWorkSize, localWorkSize,
                                     0, 0, &debayer_event));
 #else
@@ -292,35 +294,38 @@ 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};
+  unsigned int lum_total = 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];
+#ifdef QCOM2
+      if (lum < 80 && lum_binning[lum] > HISTO_CEIL_K * (y_end - y_start) * (x_end - x_start) / x_skip / y_skip / 256) {
+        continue;
+      }
+#endif
       lum_binning[lum]++;
+      lum_total += 1;
     }
   }
 
-  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;
   int lum_med_alt = 0;
   for (lum_med=255; lum_med>=0; lum_med--) {
     lum_cur += lum_binning[lum_med];
 #ifdef QCOM2
-    bool reach_hlc_perc = false;
-    if (c->camera_num == 0) { // wide
-      reach_hlc_perc = lum_cur > 2*lum_total / (3*HLC_A);
-    } else {
-      reach_hlc_perc = lum_cur > lum_total / HLC_A;
-    }
-    if (reach_hlc_perc && lum_med > HLC_THRESH) {
-      lum_med_alt = 86;
+    int lum_med_tmp = 0;
+    int hb = HLC_THRESH;
+    if (lum_cur > 0 && lum_med > hb) {
+      lum_med_tmp = 4 * (lum_med - hb) + 100;
     }
+    lum_med_alt = lum_med_alt>lum_med_tmp?lum_med_alt:lum_med_tmp;
 #endif
     if (lum_cur >= lum_total / 2) {
       break;
     }
   }
-  lum_med = lum_med_alt>lum_med?lum_med_alt:lum_med;
+  lum_med = lum_med_alt>0 ? lum_med + lum_med/32*lum_cur*(lum_med_alt - lum_med)/lum_total/2:lum_med;
   camera_autoexposure(c, lum_med / 256.0);
 }
 
diff --git a/selfdrive/camerad/cameras/camera_qcom2.cc b/selfdrive/camerad/cameras/camera_qcom2.cc
index fdef5d949c..3fef926759 100644
--- a/selfdrive/camerad/cameras/camera_qcom2.cc
+++ b/selfdrive/camerad/cameras/camera_qcom2.cc
@@ -46,10 +46,9 @@ CameraInfo cameras_supported[CAMERA_ID_MAX] = {
   },
 };
 
-float sensor_analog_gains[ANALOG_GAIN_MAX_IDX] = {1.0/8.0, 2.0/8.0, 2.0/7.0, 3.0/7.0,
-                                                  3.0/6.0, 4.0/6.0, 4.0/5.0, 5.0/5.0,
+float sensor_analog_gains[ANALOG_GAIN_MAX_IDX] = {3.0/6.0, 4.0/6.0, 4.0/5.0, 5.0/5.0,
                                                   5.0/4.0, 6.0/4.0, 6.0/3.0, 7.0/3.0,
-                                                  7.0/2.0, 8.0/2.0, 8.0/1.0};
+                                                  7.0/2.0, 8.0/2.0};
 
 // ************** low level camera helpers ****************
 
@@ -992,7 +991,7 @@ void set_exposure_time_bounds(CameraState *s) {
 void switch_conversion_gain(CameraState *s) {
   if (!s->dc_gain_enabled) {
     s->dc_gain_enabled = true;
-    s->analog_gain -= 5;
+    s->analog_gain -= 4;
   } else {
     s->dc_gain_enabled = false;
     s->analog_gain += 4;
@@ -1001,9 +1000,9 @@ void switch_conversion_gain(CameraState *s) {
 
 static void set_camera_exposure(CameraState *s, float grey_frac) {
   // TODO: get stats from sensor?
-  float target_grey = 0.3 - (s->analog_gain / 105.0);
+  float target_grey = 0.4 - ((float)(s->analog_gain + 4*s->dc_gain_enabled) / 48.0f);
   float exposure_factor = 1 + 30 * pow((target_grey - grey_frac), 3);
-  exposure_factor = std::max(exposure_factor, 0.55f);
+  exposure_factor = std::max(exposure_factor, 0.4f);
 
   if (s->camera_num != 1) {
     s->ef_filtered = (1 - EF_LOWPASS_K) * s->ef_filtered + EF_LOWPASS_K * exposure_factor;
@@ -1024,7 +1023,7 @@ static void set_camera_exposure(CameraState *s, float grey_frac) {
     if (s->analog_gain < ANALOG_GAIN_MAX_IDX - 1) {
       s->exposure_time = EXPOSURE_TIME_MAX / 2;
       s->analog_gain += 1;
-      if (!s->dc_gain_enabled && sensor_analog_gains[s->analog_gain] == 1.0) { // switch to HCG at iso 800
+      if (!s->dc_gain_enabled && sensor_analog_gains[s->analog_gain] >= 4.0) { // switch to HCG
         switch_conversion_gain(s);
       }
       set_exposure_time_bounds(s);
@@ -1035,7 +1034,7 @@ static void set_camera_exposure(CameraState *s, float grey_frac) {
     if (s->analog_gain > 0) {
       s->exposure_time = std::max(EXPOSURE_TIME_MIN * 2, (int)(s->exposure_time / (sensor_analog_gains[s->analog_gain-1] / sensor_analog_gains[s->analog_gain])));
       s->analog_gain -= 1;
-      if (s->dc_gain_enabled && sensor_analog_gains[s->analog_gain] == 0.25) { // switch back to LCG at iso 200
+      if (s->dc_gain_enabled && sensor_analog_gains[s->analog_gain] <= 1.25) { // switch back to LCG
         switch_conversion_gain(s);
       }
       set_exposure_time_bounds(s);
@@ -1045,20 +1044,22 @@ static void set_camera_exposure(CameraState *s, float grey_frac) {
   }
 
   // set up config
-  uint16_t AG = s->analog_gain;
-  AG = AG * 4096 + AG * 256 + AG * 16 + AG;
+  uint16_t AG = s->analog_gain + 4;
+  AG = 0xFF00 + AG * 16 + AG;
   s->analog_gain_frac = sensor_analog_gains[s->analog_gain];
 
   // printf("cam %d, min %d, max %d \n", s->camera_num, s->exposure_time_min, s->exposure_time_max);
   // printf("cam %d, set AG to 0x%X, S to %d, dc %d \n", s->camera_num, AG, s->exposure_time, s->dc_gain_enabled);
 
-  struct i2c_random_wr_payload exp_reg_array[] = {{0x3366, AG}, // analog gain
+  struct i2c_random_wr_payload exp_reg_array[] = {
+                                                  {0x3366, AG}, // analog gain
                                                   {0x3362, (uint16_t)(s->dc_gain_enabled?0x1:0x0)}, // DC_GAIN
                                                   {0x305A, 0x00D8}, // red gain
                                                   {0x3058, 0x011B}, // blue gain
                                                   {0x3056, 0x009A}, // g1 gain
                                                   {0x305C, 0x009A}, // g2 gain
-                                                  {0x3012, (uint16_t)s->exposure_time}}; // integ time
+                                                  {0x3012, (uint16_t)s->exposure_time}, // integ time
+                                                 };
                                                   //{0x301A, 0x091C}}; // reset
   sensors_i2c(s, exp_reg_array, sizeof(exp_reg_array)/sizeof(struct i2c_random_wr_payload),
                CAM_SENSOR_PACKET_OPCODE_SENSOR_CONFIG);
diff --git a/selfdrive/camerad/cameras/camera_qcom2.h b/selfdrive/camerad/cameras/camera_qcom2.h
index d83f3ceb52..104ba49d8e 100644
--- a/selfdrive/camerad/cameras/camera_qcom2.h
+++ b/selfdrive/camerad/cameras/camera_qcom2.h
@@ -9,12 +9,13 @@
 
 #define FRAME_BUF_COUNT 4
 
-#define ANALOG_GAIN_MAX_IDX 15 // 0xF is bypass
-#define EXPOSURE_TIME_MIN 8 // min time limited by HDR exp factor
-#define EXPOSURE_TIME_MAX 1132 // with HDR, no slower than 1/25 sec (1416 lines)
+#define ANALOG_GAIN_MAX_IDX 10 // 0xF is bypass
+#define EXPOSURE_TIME_MIN 2 // with HDR, fastest ss
+#define EXPOSURE_TIME_MAX 1757 // with HDR, slowest ss
 
-#define HLC_THRESH 240
+#define HLC_THRESH 222
 #define HLC_A 80
+#define HISTO_CEIL_K 5
 
 #define EF_LOWPASS_K 0.35
 
diff --git a/selfdrive/camerad/cameras/real_debayer.cl b/selfdrive/camerad/cameras/real_debayer.cl
index bd6556b915..7a3a94e468 100644
--- a/selfdrive/camerad/cameras/real_debayer.cl
+++ b/selfdrive/camerad/cameras/real_debayer.cl
@@ -1,135 +1,216 @@
-const __constant float3 color_correction[3] = {
+#pragma OPENCL EXTENSION cl_khr_fp16 : enable
+
+const half black_level = 42.0;
+
+const __constant half3 color_correction[3] = {
   // post wb CCM
-  (float3)(1.44602146, -0.24727126, -0.0403062),
-  (float3)(-0.37658179, 1.26329038, -0.45978396),
-  (float3)(-0.06943967, -0.01601912, 1.50009016),
+  (half3)(1.25985206, -0.378923, -0.21356857),
+  (half3)(-0.11117607, 1.3962182, -0.46342976),
+  (half3)(-0.21523926, -0.13449348, 1.47665819),
+};
+
+const __constant half3 base_ccm[3] = {
+  (half3)(1,0,0),
+  (half3)(0,1,0),
+  (half3)(0,0,1),
 };
 
-float3 color_correct(float r, float g, float b) {
-  float3 ret = (0,0,0);
 
-  ret += r * color_correction[0];
-  ret += g * color_correction[1];
-  ret += b * color_correction[2];
-  ret = max(0.0, min(1.0, ret));
+// tone mapping params
+const half cpk = 0.75;
+const half cpb = 0.125;
+const half cpxk = 0.01;
+const half cpxb = 0.01;
+
+half mf(half x, half cp) {
+  half rk = 8.6 - 66*cp;
+  if (x > cp) {
+    return (rk * (x-cp) * (1-(cpk*cp+cpb)) * (1+1/(rk*(1-cp))) / (1+rk*(x-cp))) + cpk*cp + cpb;
+  } else if (x < cp) {
+    return (rk * (x-cp) * (cpk*cp+cpb) * (1+1/(rk*cp)) / (1-rk*(x-cp))) + cpk*cp + cpb;
+  } else {
+    return x;
+  }
+}
+
+half3 color_correct(half3 rgb, int ggain) {
+  half3 ret = (0,0,0);
+  half cpx = clamp(0.03h, 0.1h, cpxb + cpxk * min(10, ggain));
+  rgb.x = mf(rgb.x, cpx);
+  rgb.y = mf(rgb.y, cpx);
+  rgb.z = mf(rgb.z, cpx);
+  ret += (half)rgb.x * (color_correction[0]);
+  ret += (half)rgb.y * (color_correction[1]);
+  ret += (half)rgb.z * (color_correction[2]);
+  ret = clamp(0.0h, 255.0h, ret*255.0h);
   return ret;
 }
 
-uint int_from_10(const uchar * source, uint start, uint offset) {
-  // source: source
-  // start: starting address of 0
-  // offset: 0 - 3
+half val_from_10(const uchar * source, int gx, int gy) {
+  // parse 10bit
+  int start = gy * FRAME_STRIDE + (5 * (gx / 4));
+  int offset = gx % 4;
   uint major = (uint)source[start + offset] << 2;
   uint minor = (source[start + 4] >> (2 * offset)) & 3;
-  return major + minor;
-}
+  half pv = (half)(major + minor);
 
-float to_normal(uint x, int gx, int gy) {
-  float pv = (float)(x);
-  const float black_level = 42.0;
-  pv = max(0.0, pv - black_level);
-  pv /= (1024.0f - black_level);
+  // normalize
+  pv = max(0.0h, pv - black_level);
+  pv *= 0.00101833h; // /= (1024.0f - black_level);
+
+  // correct vignetting
   if (CAM_NUM == 1) { // fcamera
     gx = (gx - RGB_WIDTH/2);
     gy = (gy - RGB_HEIGHT/2);
-    float r = pow(gx*gx + gy*gy, 0.825);
-    float s = 1 / (1-0.00000733*r);
+    float r = gx*gx + gy*gy;
+    half s;
+    if (r < 62500) {
+      s = (half)(1.0f + 0.0000008f*r);
+    } else if (r < 490000) {
+      s = (half)(0.9625f + 0.0000014f*r);
+    } else if (r < 1102500) {
+      s = (half)(1.26434f + 0.0000000000016f*r*r);
+    } else {
+      s = (half)(0.53503625f + 0.0000000000022f*r*r);
+    }
     pv = s * pv;
   }
-  pv = 20*pv / (1.0f + 20*pv); // reinhard
+
+  pv = clamp(0.0h, 1.0h, pv);
   return pv;
 }
 
+half fabs_diff(half x, half y) {
+  return fabs(x-y);
+}
+
+half phi(half x) {
+  // detection funtion
+  return 2 - x;
+  // if (x > 1) {
+  //   return 1 / x;
+  // } else {
+  //   return 2 - x;
+  // }
+}
+
 __kernel void debayer10(const __global uchar * in,
                         __global uchar * out,
-                        __local float * cached
+                        __local half * cached,
+                        uint ggain
                        )
 {
   const int x_global = get_global_id(0);
   const int y_global = get_global_id(1);
 
-  // const int globalOffset = ;
-
   const int localRowLen = 2 + get_local_size(0); // 2 padding
-  const int x_local = get_local_id(0);
-  const int y_local = get_local_id(1);
+  const int x_local = get_local_id(0); // 0-15
+  const int y_local = get_local_id(1); // 0-15
+  const int localOffset = (y_local + 1) * localRowLen + x_local + 1; // max 18x18-1
 
-  const int localOffset = (y_local + 1) * localRowLen + x_local + 1;
+  int out_idx = 3 * x_global + 3 * y_global * RGB_WIDTH;
 
-  // cache local pixels first
-  // saves memory access and avoids repeated normalization
-  uint globalStart_10 = y_global * FRAME_STRIDE + (5 * (x_global / 4));
-  uint offset_10 = x_global % 4;
-  uint raw_val = int_from_10(in, globalStart_10, offset_10);
-  cached[localOffset] = to_normal(raw_val, x_global, y_global);
+  half pv = val_from_10(in, x_global, y_global);
+  cached[localOffset] = pv;
 
-  // edges
-  if (x_global < 1 || x_global > RGB_WIDTH - 2 || y_global < 1 || y_global > RGB_HEIGHT - 2) {
-    barrier(CLK_LOCAL_MEM_FENCE);
+  // don't care
+  if (x_global < 1 || x_global >= RGB_WIDTH - 1 || y_global < 1 || y_global >= RGB_HEIGHT - 1) {
     return;
-  } else {
-    int localColOffset = -1;
-    int globalColOffset = -1;
-
-    // cache padding
-    if (x_local < 1) {
-      localColOffset = x_local;
-      globalColOffset = -1;
-      cached[(y_local + 1) * localRowLen + x_local] = to_normal(int_from_10(in, y_global * FRAME_STRIDE + (5 * ((x_global-1) / 4)), (offset_10 + 3) % 4), x_global, y_global);
-    } else if (x_local >= get_local_size(0) - 1) {
-      localColOffset = x_local + 2;
-      globalColOffset = 1;
-      cached[localOffset + 1] = to_normal(int_from_10(in, y_global * FRAME_STRIDE + (5 * ((x_global+1) / 4)), (offset_10 + 1) % 4), x_global, y_global);
-    }
+  }
+
+  // cache padding
+  int localColOffset = -1;
+  int globalColOffset = -1;
+
+  // cache padding
+  if (x_local < 1) {
+    localColOffset = x_local;
+    globalColOffset = -1;
+    cached[(y_local + 1) * localRowLen + x_local] = val_from_10(in, x_global-1, y_global);
+  } else if (x_local >= get_local_size(0) - 1) {
+    localColOffset = x_local + 2;
+    globalColOffset = 1;
+    cached[localOffset + 1] = val_from_10(in, x_global+1, y_global);
+  }
 
-    if (y_local < 1) {
-      cached[y_local * localRowLen + x_local + 1] = to_normal(int_from_10(in, globalStart_10 - FRAME_STRIDE, offset_10), x_global, y_global);
-      if (localColOffset != -1) {
-        cached[y_local * localRowLen + localColOffset] = to_normal(int_from_10(in, (y_global-1) * FRAME_STRIDE + (5 * ((x_global+globalColOffset) / 4)), (offset_10+4+globalColOffset) % 4), x_global, y_global);
-      }
-    } else if (y_local >= get_local_size(1) - 1) {
-      cached[(y_local + 2) * localRowLen + x_local + 1] = to_normal(int_from_10(in, globalStart_10 + FRAME_STRIDE, offset_10), x_global, y_global);
-      if (localColOffset != -1) {
-        cached[(y_local + 2) * localRowLen + localColOffset] = to_normal(int_from_10(in, (y_global+1) * FRAME_STRIDE + (5 * ((x_global+globalColOffset) / 4)), (offset_10+4+globalColOffset) % 4), x_global, y_global);
-      }
+  if (y_local < 1) {
+    cached[y_local * localRowLen + x_local + 1] = val_from_10(in, x_global, y_global-1);
+    if (localColOffset != -1) {
+      cached[y_local * localRowLen + localColOffset] = val_from_10(in, x_global+globalColOffset, y_global-1);
     }
+  } else if (y_local >= get_local_size(1) - 1) {
+    cached[(y_local + 2) * localRowLen + x_local + 1] = val_from_10(in, x_global, y_global+1);
+    if (localColOffset != -1) {
+      cached[(y_local + 2) * localRowLen + localColOffset] = val_from_10(in, x_global+globalColOffset, y_global+1);
+    }
+  }
 
-    // sync
-    barrier(CLK_LOCAL_MEM_FENCE);
-
-    // perform debayer
-    float r;
-    float g;
-    float b;
-
-    if (x_global % 2 == 0) {
-      if (y_global % 2 == 0) { // G1
-        r = (cached[localOffset - 1] + cached[localOffset + 1]) / 2.0f;
-        g = (cached[localOffset] + cached[localOffset + localRowLen + 1]) / 2.0f;
-        b = (cached[localOffset - localRowLen] + cached[localOffset + localRowLen]) / 2.0f;
-      } else { // B
-        r = (cached[localOffset - localRowLen - 1] + cached[localOffset - localRowLen + 1] + cached[localOffset + localRowLen - 1] + cached[localOffset + localRowLen + 1]) / 4.0f;
-        g = (cached[localOffset - localRowLen] + cached[localOffset + localRowLen] + cached[localOffset - 1] + cached[localOffset + 1]) / 4.0f;
-        b = cached[localOffset];
-      }
+  // sync
+  barrier(CLK_LOCAL_MEM_FENCE);
+
+  half d1 = cached[localOffset - localRowLen - 1];
+  half d2 = cached[localOffset - localRowLen + 1];
+  half d3 = cached[localOffset + localRowLen - 1];
+  half d4 = cached[localOffset + localRowLen + 1];
+  half n1 = cached[localOffset - localRowLen];
+  half n2 = cached[localOffset + 1];
+  half n3 = cached[localOffset + localRowLen];
+  half n4 = cached[localOffset - 1];
+
+  half3 rgb;
+
+  // a simplified version of https://opensignalprocessingjournal.com/contents/volumes/V6/TOSIGPJ-6-1/TOSIGPJ-6-1.pdf
+  if (x_global % 2 == 0) {
+    if (y_global % 2 == 0) {
+      rgb.y = pv; // G1(R)
+      half k1 = phi(fabs_diff(d1, pv) + fabs_diff(d2, pv));
+      half k2 = phi(fabs_diff(d2, pv) + fabs_diff(d4, pv));
+      half k3 = phi(fabs_diff(d3, pv) + fabs_diff(d4, pv));
+      half k4 = phi(fabs_diff(d1, pv) + fabs_diff(d3, pv));
+      // R_G1
+      rgb.x = (k2*n2+k4*n4)/(k2+k4);
+      // B_G1
+      rgb.z = (k1*n1+k3*n3)/(k1+k3);
+    } else {
+      rgb.z = pv; // B
+      half k1 = phi(fabs_diff(d1, d3) + fabs_diff(d2, d4));
+      half k2 = phi(fabs_diff(n1, n4) + fabs_diff(n2, n3));
+      half k3 = phi(fabs_diff(d1, d2) + fabs_diff(d3, d4));
+      half k4 = phi(fabs_diff(n1, n2) + fabs_diff(n3, n4));
+      // G_B
+      rgb.y = (k1*(n1+n3)*0.5+k3*(n2+n4)*0.5)/(k1+k3);
+      // R_B
+      rgb.x = (k2*(d2+d3)*0.5+k4*(d1+d4)*0.5)/(k2+k4);
+    }
+  } else {
+    if (y_global % 2 == 0) {
+      rgb.x = pv; // R
+      half k1 = phi(fabs_diff(d1, d3) + fabs_diff(d2, d4));
+      half k2 = phi(fabs_diff(n1, n4) + fabs_diff(n2, n3));
+      half k3 = phi(fabs_diff(d1, d2) + fabs_diff(d3, d4));
+      half k4 = phi(fabs_diff(n1, n2) + fabs_diff(n3, n4));
+      // G_R
+      rgb.y = (k1*(n1+n3)*0.5+k3*(n2+n4)*0.5)/(k1+k3);
+      // B_R
+      rgb.z = (k2*(d2+d3)*0.5+k4*(d1+d4)*0.5)/(k2+k4);
     } else {
-      if (y_global % 2 == 0) { // R
-        r = cached[localOffset];
-        g = (cached[localOffset - localRowLen] + cached[localOffset + localRowLen] + cached[localOffset - 1] + cached[localOffset + 1]) / 4.0f;
-        b = (cached[localOffset - localRowLen - 1] + cached[localOffset - localRowLen + 1] + cached[localOffset + localRowLen - 1] + cached[localOffset + localRowLen + 1]) / 4.0f;
-      } else { // G2
-        r = (cached[localOffset - localRowLen] + cached[localOffset + localRowLen]) / 2.0f;
-        g = (cached[localOffset] + cached[localOffset - localRowLen - 1]) / 2.0f;
-        b = (cached[localOffset - 1] + cached[localOffset + 1]) / 2.0f;
-      }
+      rgb.y = pv; // G2(B)
+      half k1 = phi(fabs_diff(d1, pv) + fabs_diff(d2, pv));
+      half k2 = phi(fabs_diff(d2, pv) + fabs_diff(d4, pv));
+      half k3 = phi(fabs_diff(d3, pv) + fabs_diff(d4, pv));
+      half k4 = phi(fabs_diff(d1, pv) + fabs_diff(d3, pv));
+      // R_G2
+      rgb.x = (k1*n1+k3*n3)/(k1+k3);
+      // B_G2
+      rgb.z = (k2*n2+k4*n4)/(k2+k4);
     }
+  }
 
-    float3 rgb = color_correct(r, g, b);
-    // rgb = srgb_gamma(rgb); 
+  rgb = clamp(0.0h, 1.0h, rgb);
+  rgb = color_correct(rgb, (int)ggain);
+
+  out[out_idx + 0] = (uchar)(rgb.z);
+  out[out_idx + 1] = (uchar)(rgb.y);
+  out[out_idx + 2] = (uchar)(rgb.x);
 
-    // BGR output
-    out[3 * x_global + 3 * y_global * RGB_WIDTH + 0] = (uchar)(255.0f * rgb.z);
-    out[3 * x_global + 3 * y_global * RGB_WIDTH + 1] = (uchar)(255.0f * rgb.y);
-    out[3 * x_global + 3 * y_global * RGB_WIDTH + 2] = (uchar)(255.0f * rgb.x);
-  }
 }
diff --git a/selfdrive/camerad/cameras/sensor2_i2c.h b/selfdrive/camerad/cameras/sensor2_i2c.h
index 82292347b9..8abec280cd 100644
--- a/selfdrive/camerad/cameras/sensor2_i2c.h
+++ b/selfdrive/camerad/cameras/sensor2_i2c.h
@@ -1,10 +1,8 @@
 struct i2c_random_wr_payload start_reg_array[] = {{0x301a, 0x91c}};
-//struct i2c_random_wr_payload stop_reg_array[] = {{0x301a, 0x10d8}};
 struct i2c_random_wr_payload stop_reg_array[] = {{0x301a, 0x918}};;
 
 struct i2c_random_wr_payload init_array_ar0231[] = {
   {0x301A, 0x0018}, // RESET_REGISTER
-  //
   {0x3092, 0x0C24}, // ROW_NOISE_CONTROL
   {0x337A, 0x0C80}, // DBLC_SCALE0
   {0x3520, 0x1288}, // RESERVED_MFR_3520
@@ -32,222 +30,6 @@ struct i2c_random_wr_payload init_array_ar0231[] = {
   {0x329A, 0x0FA0}, // RESERVED_MFR_329A
   {0x329C, 0x0FA0}, // RESERVED_MFR_329C
   {0x329E, 0x0FA0}, // RESERVED_MFR_329E
-  {0x2512, 0x8000}, // SEQ_CTRL_PORT
-  {0x2510, 0x0905}, // SEQ_DATA_PORT
-  {0x2510, 0x3350}, // SEQ_DATA_PORT
-  {0x2510, 0x2004}, // SEQ_DATA_PORT
-  {0x2510, 0x1460}, // SEQ_DATA_PORT
-  {0x2510, 0x1578}, // SEQ_DATA_PORT
-  {0x2510, 0x0901}, // SEQ_DATA_PORT
-  {0x2510, 0x7B24}, // SEQ_DATA_PORT
-  {0x2510, 0xFF24}, // SEQ_DATA_PORT
-  {0x2510, 0xFF24}, // SEQ_DATA_PORT
-  {0x2510, 0xEA24}, // SEQ_DATA_PORT
-  {0x2510, 0x1022}, // SEQ_DATA_PORT
-  {0x2510, 0x2410}, // SEQ_DATA_PORT
-  {0x2510, 0x155A}, // SEQ_DATA_PORT
-  {0x2510, 0x0901}, // SEQ_DATA_PORT
-  {0x2510, 0x1400}, // SEQ_DATA_PORT
-  {0x2510, 0x24FF}, // SEQ_DATA_PORT
-  {0x2510, 0x24FF}, // SEQ_DATA_PORT
-  {0x2510, 0x24EA}, // SEQ_DATA_PORT
-  {0x2510, 0x2324}, // SEQ_DATA_PORT
-  {0x2510, 0x647A}, // SEQ_DATA_PORT
-  {0x2510, 0x2404}, // SEQ_DATA_PORT
-  {0x2510, 0x052C}, // SEQ_DATA_PORT
-  {0x2510, 0x400A}, // SEQ_DATA_PORT
-  {0x2510, 0xFF0A}, // SEQ_DATA_PORT
-  {0x2510, 0xFF0A}, // SEQ_DATA_PORT
-  {0x2510, 0x1008}, // SEQ_DATA_PORT
-  {0x2510, 0x3851}, // SEQ_DATA_PORT
-  {0x2510, 0x1440}, // SEQ_DATA_PORT
-  {0x2510, 0x0004}, // SEQ_DATA_PORT
-  {0x2510, 0x0801}, // SEQ_DATA_PORT
-  {0x2510, 0x0408}, // SEQ_DATA_PORT
-  {0x2510, 0x1180}, // SEQ_DATA_PORT
-  {0x2510, 0x2652}, // SEQ_DATA_PORT
-  {0x2510, 0x1518}, // SEQ_DATA_PORT
-  {0x2510, 0x0906}, // SEQ_DATA_PORT
-  {0x2510, 0x1348}, // SEQ_DATA_PORT
-  {0x2510, 0x1002}, // SEQ_DATA_PORT
-  {0x2510, 0x1016}, // SEQ_DATA_PORT
-  {0x2510, 0x1181}, // SEQ_DATA_PORT
-  {0x2510, 0x1189}, // SEQ_DATA_PORT
-  {0x2510, 0x1056}, // SEQ_DATA_PORT
-  {0x2510, 0x1210}, // SEQ_DATA_PORT
-  {0x2510, 0x0901}, // SEQ_DATA_PORT
-  {0x2510, 0x0D09}, // SEQ_DATA_PORT
-  {0x2510, 0x1413}, // SEQ_DATA_PORT
-  {0x2510, 0x8809}, // SEQ_DATA_PORT
-  {0x2510, 0x2B15}, // SEQ_DATA_PORT
-  {0x2510, 0x8809}, // SEQ_DATA_PORT
-  {0x2510, 0x0311}, // SEQ_DATA_PORT
-  {0x2510, 0xD909}, // SEQ_DATA_PORT
-  {0x2510, 0x1214}, // SEQ_DATA_PORT
-  {0x2510, 0x4109}, // SEQ_DATA_PORT
-  {0x2510, 0x0312}, // SEQ_DATA_PORT
-  {0x2510, 0x1409}, // SEQ_DATA_PORT
-  {0x2510, 0x0110}, // SEQ_DATA_PORT
-  {0x2510, 0xD612}, // SEQ_DATA_PORT
-  {0x2510, 0x1012}, // SEQ_DATA_PORT
-  {0x2510, 0x1212}, // SEQ_DATA_PORT
-  {0x2510, 0x1011}, // SEQ_DATA_PORT
-  {0x2510, 0xDD11}, // SEQ_DATA_PORT
-  {0x2510, 0xD910}, // SEQ_DATA_PORT
-  {0x2510, 0x5609}, // SEQ_DATA_PORT
-  {0x2510, 0x1511}, // SEQ_DATA_PORT
-  {0x2510, 0xDB09}, // SEQ_DATA_PORT
-  {0x2510, 0x1511}, // SEQ_DATA_PORT
-  {0x2510, 0x9B09}, // SEQ_DATA_PORT
-  {0x2510, 0x0F11}, // SEQ_DATA_PORT
-  {0x2510, 0xBB12}, // SEQ_DATA_PORT
-  {0x2510, 0x1A12}, // SEQ_DATA_PORT
-  {0x2510, 0x1014}, // SEQ_DATA_PORT
-  {0x2510, 0x6012}, // SEQ_DATA_PORT
-  {0x2510, 0x5010}, // SEQ_DATA_PORT
-  {0x2510, 0x7610}, // SEQ_DATA_PORT
-  {0x2510, 0xE609}, // SEQ_DATA_PORT
-  {0x2510, 0x0812}, // SEQ_DATA_PORT
-  {0x2510, 0x4012}, // SEQ_DATA_PORT
-  {0x2510, 0x6009}, // SEQ_DATA_PORT
-  {0x2510, 0x290B}, // SEQ_DATA_PORT
-  {0x2510, 0x0904}, // SEQ_DATA_PORT
-  {0x2510, 0x1440}, // SEQ_DATA_PORT
-  {0x2510, 0x0923}, // SEQ_DATA_PORT
-  {0x2510, 0x15C8}, // SEQ_DATA_PORT
-  {0x2510, 0x13C8}, // SEQ_DATA_PORT
-  {0x2510, 0x092C}, // SEQ_DATA_PORT
-  {0x2510, 0x1588}, // SEQ_DATA_PORT
-  {0x2510, 0x1388}, // SEQ_DATA_PORT
-  {0x2510, 0x0C09}, // SEQ_DATA_PORT
-  {0x2510, 0x0C14}, // SEQ_DATA_PORT
-  {0x2510, 0x4109}, // SEQ_DATA_PORT
-  {0x2510, 0x1112}, // SEQ_DATA_PORT
-  {0x2510, 0x6212}, // SEQ_DATA_PORT
-  {0x2510, 0x6011}, // SEQ_DATA_PORT
-  {0x2510, 0xBF11}, // SEQ_DATA_PORT
-  {0x2510, 0xBB10}, // SEQ_DATA_PORT
-  {0x2510, 0x6611}, // SEQ_DATA_PORT
-  {0x2510, 0xFB09}, // SEQ_DATA_PORT
-  {0x2510, 0x3511}, // SEQ_DATA_PORT
-  {0x2510, 0xBB12}, // SEQ_DATA_PORT
-  {0x2510, 0x6312}, // SEQ_DATA_PORT
-  {0x2510, 0x6014}, // SEQ_DATA_PORT
-  {0x2510, 0x0015}, // SEQ_DATA_PORT
-  {0x2510, 0x0011}, // SEQ_DATA_PORT
-  {0x2510, 0xB812}, // SEQ_DATA_PORT
-  {0x2510, 0xA012}, // SEQ_DATA_PORT
-  {0x2510, 0x0010}, // SEQ_DATA_PORT
-  {0x2510, 0x2610}, // SEQ_DATA_PORT
-  {0x2510, 0x0013}, // SEQ_DATA_PORT
-  {0x2510, 0x0011}, // SEQ_DATA_PORT
-  {0x2510, 0x0008}, // SEQ_DATA_PORT
-  {0x2510, 0x3053}, // SEQ_DATA_PORT
-  {0x2510, 0x4215}, // SEQ_DATA_PORT
-  {0x2510, 0x4013}, // SEQ_DATA_PORT
-  {0x2510, 0x4010}, // SEQ_DATA_PORT
-  {0x2510, 0x0210}, // SEQ_DATA_PORT
-  {0x2510, 0x1611}, // SEQ_DATA_PORT
-  {0x2510, 0x8111}, // SEQ_DATA_PORT
-  {0x2510, 0x8910}, // SEQ_DATA_PORT
-  {0x2510, 0x5612}, // SEQ_DATA_PORT
-  {0x2510, 0x1009}, // SEQ_DATA_PORT
-  {0x2510, 0x010D}, // SEQ_DATA_PORT
-  {0x2510, 0x0815}, // SEQ_DATA_PORT
-  {0x2510, 0xC015}, // SEQ_DATA_PORT
-  {0x2510, 0xD013}, // SEQ_DATA_PORT
-  {0x2510, 0x5009}, // SEQ_DATA_PORT
-  {0x2510, 0x1313}, // SEQ_DATA_PORT
-  {0x2510, 0xD009}, // SEQ_DATA_PORT
-  {0x2510, 0x0215}, // SEQ_DATA_PORT
-  {0x2510, 0xC015}, // SEQ_DATA_PORT
-  {0x2510, 0xC813}, // SEQ_DATA_PORT
-  {0x2510, 0xC009}, // SEQ_DATA_PORT
-  {0x2510, 0x0515}, // SEQ_DATA_PORT
-  {0x2510, 0x8813}, // SEQ_DATA_PORT
-  {0x2510, 0x8009}, // SEQ_DATA_PORT
-  {0x2510, 0x0213}, // SEQ_DATA_PORT
-  {0x2510, 0x8809}, // SEQ_DATA_PORT
-  {0x2510, 0x0411}, // SEQ_DATA_PORT
-  {0x2510, 0xC909}, // SEQ_DATA_PORT
-  {0x2510, 0x0814}, // SEQ_DATA_PORT
-  {0x2510, 0x0109}, // SEQ_DATA_PORT
-  {0x2510, 0x0B11}, // SEQ_DATA_PORT
-  {0x2510, 0xD908}, // SEQ_DATA_PORT
-  {0x2510, 0x1400}, // SEQ_DATA_PORT
-  {0x2510, 0x091A}, // SEQ_DATA_PORT
-  {0x2510, 0x1440}, // SEQ_DATA_PORT
-  {0x2510, 0x0903}, // SEQ_DATA_PORT
-  {0x2510, 0x1214}, // SEQ_DATA_PORT
-  {0x2510, 0x0901}, // SEQ_DATA_PORT
-  {0x2510, 0x10D6}, // SEQ_DATA_PORT
-  {0x2510, 0x1210}, // SEQ_DATA_PORT
-  {0x2510, 0x1212}, // SEQ_DATA_PORT
-  {0x2510, 0x1210}, // SEQ_DATA_PORT
-  {0x2510, 0x11DD}, // SEQ_DATA_PORT
-  {0x2510, 0x11D9}, // SEQ_DATA_PORT
-  {0x2510, 0x1056}, // SEQ_DATA_PORT
-  {0x2510, 0x0917}, // SEQ_DATA_PORT
-  {0x2510, 0x11DB}, // SEQ_DATA_PORT
-  {0x2510, 0x0913}, // SEQ_DATA_PORT
-  {0x2510, 0x11FB}, // SEQ_DATA_PORT
-  {0x2510, 0x0905}, // SEQ_DATA_PORT
-  {0x2510, 0x11BB}, // SEQ_DATA_PORT
-  {0x2510, 0x121A}, // SEQ_DATA_PORT
-  {0x2510, 0x1210}, // SEQ_DATA_PORT
-  {0x2510, 0x1460}, // SEQ_DATA_PORT
-  {0x2510, 0x1250}, // SEQ_DATA_PORT
-  {0x2510, 0x1076}, // SEQ_DATA_PORT
-  {0x2510, 0x10E6}, // SEQ_DATA_PORT
-  {0x2510, 0x0901}, // SEQ_DATA_PORT
-  {0x2510, 0x15A8}, // SEQ_DATA_PORT
-  {0x2510, 0x0901}, // SEQ_DATA_PORT
-  {0x2510, 0x13A8}, // SEQ_DATA_PORT
-  {0x2510, 0x1240}, // SEQ_DATA_PORT
-  {0x2510, 0x1260}, // SEQ_DATA_PORT
-  {0x2510, 0x0925}, // SEQ_DATA_PORT
-  {0x2510, 0x13AD}, // SEQ_DATA_PORT
-  {0x2510, 0x0902}, // SEQ_DATA_PORT
-  {0x2510, 0x0907}, // SEQ_DATA_PORT
-  {0x2510, 0x1588}, // SEQ_DATA_PORT
-  {0x2510, 0x0901}, // SEQ_DATA_PORT
-  {0x2510, 0x138D}, // SEQ_DATA_PORT
-  {0x2510, 0x0B09}, // SEQ_DATA_PORT
-  {0x2510, 0x0914}, // SEQ_DATA_PORT
-  {0x2510, 0x4009}, // SEQ_DATA_PORT
-  {0x2510, 0x0B13}, // SEQ_DATA_PORT
-  {0x2510, 0x8809}, // SEQ_DATA_PORT
-  {0x2510, 0x1C0C}, // SEQ_DATA_PORT
-  {0x2510, 0x0920}, // SEQ_DATA_PORT
-  {0x2510, 0x1262}, // SEQ_DATA_PORT
-  {0x2510, 0x1260}, // SEQ_DATA_PORT
-  {0x2510, 0x11BF}, // SEQ_DATA_PORT
-  {0x2510, 0x11BB}, // SEQ_DATA_PORT
-  {0x2510, 0x1066}, // SEQ_DATA_PORT
-  {0x2510, 0x090A}, // SEQ_DATA_PORT
-  {0x2510, 0x11FB}, // SEQ_DATA_PORT
-  {0x2510, 0x093B}, // SEQ_DATA_PORT
-  {0x2510, 0x11BB}, // SEQ_DATA_PORT
-  {0x2510, 0x1263}, // SEQ_DATA_PORT
-  {0x2510, 0x1260}, // SEQ_DATA_PORT
-  {0x2510, 0x1400}, // SEQ_DATA_PORT
-  {0x2510, 0x1508}, // SEQ_DATA_PORT
-  {0x2510, 0x11B8}, // SEQ_DATA_PORT
-  {0x2510, 0x12A0}, // SEQ_DATA_PORT
-  {0x2510, 0x1200}, // SEQ_DATA_PORT
-  {0x2510, 0x1026}, // SEQ_DATA_PORT
-  {0x2510, 0x1000}, // SEQ_DATA_PORT
-  {0x2510, 0x1300}, // SEQ_DATA_PORT
-  {0x2510, 0x1100}, // SEQ_DATA_PORT
-  {0x2510, 0x437A}, // SEQ_DATA_PORT
-  {0x2510, 0x0609}, // SEQ_DATA_PORT
-  {0x2510, 0x0B05}, // SEQ_DATA_PORT
-  {0x2510, 0x0708}, // SEQ_DATA_PORT
-  {0x2510, 0x4137}, // SEQ_DATA_PORT
-  {0x2510, 0x502C}, // SEQ_DATA_PORT
-  {0x2510, 0x2CFE}, // SEQ_DATA_PORT
-  {0x2510, 0x15FE}, // SEQ_DATA_PORT
-  {0x2510, 0x0C2C}, // SEQ_DATA_PORT
   {0x32E6, 0x00E0}, // RESERVED_MFR_32E6
   {0x1008, 0x036F}, // RESERVED_PARAM_1008
   {0x100C, 0x058F}, // RESERVED_PARAM_100C
@@ -277,21 +59,6 @@ struct i2c_random_wr_payload init_array_ar0231[] = {
   {0x30BA, 0x11F2}, // DIGITAL_CTRL
   {0x3044, 0x0400}, // DARK_CONTROL
   {0x3064, 0x1802}, // SMIA_TEST
-  /*{0x3064, 0xCC2}, // STATS_EN
-  {0x3270, 0x10}, //
-  {0x3272, 0x30}, //
-  {0x3274, 0x50}, //
-  {0x3276, 0x10}, //
-  {0x3278, 0x30}, //
-  {0x327A, 0x50}, //
-
-  {0x3144, 0x0}, //
-  {0x3146, 0x0}, //
-  {0x3244, 0x0}, //
-  {0x3246, 0x0}, //
-  {0x3268, 0x0}, //
-  {0x326A, 0x0}, //
-  */
   {0x33E0, 0x0C80}, // TEST_ASIL_ROWS
   {0x3180, 0x0080}, // RESERVED_MFR_3180
   {0x33E4, 0x0080}, // RESERVED_MFR_33E4
@@ -307,7 +74,7 @@ struct i2c_random_wr_payload init_array_ar0231[] = {
   {0x3402, 0x0F10}, // X_OUTPUT_CONTROL
   {0x3404, 0x04B8}, // Y_OUTPUT_CONTROL
   {0x3404, 0x0970}, // Y_OUTPUT_CONTROL
-  {0x30BA, 0x11F3}, // DIGITAL_CTRL
+  {0x30BA, 0x11F2}, // DIGITAL_CTRL
 
   // SLAV* MODE
   {0x30CE, 0x0120},
@@ -315,8 +82,8 @@ struct i2c_random_wr_payload init_array_ar0231[] = {
   {0x340C, 0x802}, // 2 // 0000 0000 0010
 
   // Readout timing
-  {0x300C, 0x074B}, // LINE_LENGTH_PCK: min for 3-exposure HDR
-  {0x300A, 0x06EB}, // FRAME_LENGTH_LINES_ 6EB
+  {0x300C, 0x074B}, // LINE_LENGTH_PCK: min for 2-exposure HDR
+  {0x300A, 0x085E}, // FRAME_LENGTH_LINES_ 6EB
   {0x3042, 0x0000}, // EXTRA_DELAY
 
   // Readout Settings
@@ -341,24 +108,23 @@ struct i2c_random_wr_payload init_array_ar0231[] = {
 
   // HDR Settings
   {0x3082, 0x0004}, // OPERATION_MODE_CTRL
-  {0x3238, 0x0222}, // EXPOSURE_RATIO
+  {0x3238, 0x0004}, // EXPOSURE_RATIO
   {0x3014, 0x098E}, // FINE_INTEGRATION_TIME_
   {0x321E, 0x098E}, // FINE_INTEGRATION_TIME2
-  {0x3222, 0x098E}, // FINE_INTEGRATION_TIME3
-  {0x3226, 0x098E}, // FINE_INTEGRATION_TIME4, 098E?
   {0x30B0, 0x0800}, // DIGITAL_TEST
   {0x32EA, 0x3C0E}, // RESERVED_MFR_32EA
   {0x32EC, 0x72A1}, // RESERVED_MFR_32EC
   {0x31D0, 0x0000}, // COMPANDING, no good in 10 bit?
   {0x33DA, 0x0000}, // COMPANDING
-  {0x3362, 0x0000}, // DC GAIN
-  {0x3370, 0x0231}, // DBLC
+  {0x3370, 0x03B1}, // DBLC
+  {0x31E0, 0x0001}, // PDC
   {0x318E, 0x0200}, // PRE_HDR_GAIN_EN
 
    // Initial Gains
   {0x3022, 0x01}, // GROUPED_PARAMETER_HOLD_
   {0x3366, 0x5555}, // ANALOG_GAIN
-  {0x3060, 0xBBBB}, // ANALOG_COLOR_GAIN
+  {0x3060, 0x3333}, // ANALOG_COLOR_GAIN
+  {0x3362, 0x0000}, // DC GAIN
   {0x305A, 0x00D8}, // RED_GAIN
   {0x3058, 0x011B}, // BLUE_GAIN
   {0x3056, 0x009A}, // GREEN1_GAIN
@@ -367,393 +133,5 @@ struct i2c_random_wr_payload init_array_ar0231[] = {
 
   // Initial Integration Time
   {0x3012, 0x256},
-
-};
-
-struct i2c_random_wr_payload poke_array_ov7750[] = {
-  {0x3208, 0x0}, {0x380e, 0x1a}, {0x380f, 0xf0}, {0x3500, 0x0}, {0x3501, 0x0}, {0x3502, 0x10}, {0x350a, 0x0}, {0x350b, 0x10}, {0x3208, 0x10}, {0x3208, 0xa0},
-  //{0x3208, 0x0}, {0x380e, 0x1a}, {0x380f, 0xf0}, {0x3500, 0x0}, {0x3501, 0x0}, {0x3502, 0x10}, {0x350a, 0x0}, {0x350b, 0x10}, {0x3208, 0x10}, {0x3208, 0xa0},
-};
-
-struct i2c_random_wr_payload preinit_array_ov7750[] = {
-  {0x103, 0x1},
-  {0x303b, 0x2},
-  {0x302b, 0x80},
-};
-
-struct i2c_random_wr_payload init_array_ov7750[] = {
-  // 2nd batch
-  {0x3005, 0x0},
-  {0x3012, 0xc0},
-  {0x3013, 0xd2},
-  {0x3014, 0x4},
-  {0x3016, 0xf0},
-  {0x3017, 0xf0},
-  {0x3018, 0xf0},
-  {0x301a, 0xf0},
-  {0x301b, 0xf0},
-  {0x301c, 0xf0},
-  {0x3023, 0x5},
-  {0x3037, 0xf0},
-  {0x3098, 0x4},
-  {0x3099, 0x28},
-  {0x309a, 0x5},
-  {0x309b, 0x4},
-  {0x30b0, 0xa},
-  {0x30b1, 0x1},
-  {0x30b3, 0x64},
-  {0x30b4, 0x3},
-  {0x30b5, 0x5},
-  {0x3106, 0xda},
-  {0x3500, 0x0},
-  {0x3501, 0x1f},
-  {0x3502, 0x80},
-  {0x3503, 0x7},
-  {0x3509, 0x10},
-  {0x350b, 0x10},
-  {0x3600, 0x1c},
-  {0x3602, 0x62},
-  {0x3620, 0xb7},
-  {0x3622, 0x4},
-  {0x3626, 0x21},
-  {0x3627, 0x30},
-  {0x3630, 0x44},
-  {0x3631, 0x35},
-  {0x3634, 0x60},
-  {0x3636, 0x0},
-  {0x3662, 0x1},
-  {0x3663, 0x70},
-  {0x3664, 0xf0},
-  {0x3666, 0xa},
-  {0x3669, 0x1a},
-  {0x366a, 0x0},
-  {0x366b, 0x50},
-  {0x3673, 0x1},
-  {0x3674, 0xff},
-  {0x3675, 0x3},
-  {0x3705, 0xc1},
-  {0x3709, 0x40},
-  {0x373c, 0x8},
-  {0x3742, 0x0},
-  {0x3757, 0xb3},
-  {0x3788, 0x0},
-  {0x37a8, 0x1},
-  {0x37a9, 0xc0},
-  {0x3800, 0x0},
-  {0x3801, 0x4},
-  {0x3802, 0x0},
-  {0x3803, 0x4},
-  {0x3804, 0x2},
-  {0x3805, 0x8b},
-  {0x3806, 0x1},
-  {0x3807, 0xeb},
-  {0x3808, 0x2},
-  {0x3809, 0x80},
-  {0x380a, 0x1},
-  {0x380b, 0xe0},
-  {0x380c, 0x3},
-  {0x380d, 0xa0},
-  {0x380e, 0x6},
-  {0x380f, 0xbc},
-  {0x3810, 0x0},
-  {0x3811, 0x4},
-  {0x3812, 0x0},
-  {0x3813, 0x5},
-  {0x3814, 0x11},
-  {0x3815, 0x11},
-  {0x3820, 0x40},
-  {0x3821, 0x0},
-  {0x382f, 0xe},
-  {0x3832, 0x0},
-  {0x3833, 0x5},
-  {0x3834, 0x0},
-  {0x3835, 0xc},
-  {0x3837, 0x0},
-  {0x3b80, 0x0},
-  {0x3b81, 0xa5},
-  {0x3b82, 0x10},
-  {0x3b83, 0x0},
-  {0x3b84, 0x8},
-  {0x3b85, 0x0},
-  {0x3b86, 0x1},
-  {0x3b87, 0x0},
-  {0x3b88, 0x0},
-  {0x3b89, 0x0},
-  {0x3b8a, 0x0},
-  {0x3b8b, 0x5},
-  {0x3b8c, 0x0},
-  {0x3b8d, 0x0},
-  {0x3b8e, 0x0},
-  {0x3b8f, 0x1a},
-  {0x3b94, 0x5},
-  {0x3b95, 0xf2},
-  {0x3b96, 0x40},
-  {0x3c00, 0x89},
-  {0x3c01, 0x63},
-  {0x3c02, 0x1},
-  {0x3c03, 0x0},
-  {0x3c04, 0x0},
-  {0x3c05, 0x3},
-  {0x3c06, 0x0},
-  {0x3c07, 0x6},
-  {0x3c0c, 0x1},
-  {0x3c0d, 0xd0},
-  {0x3c0e, 0x2},
-  {0x3c0f, 0xa},
-  {0x4001, 0x42},
-  {0x4004, 0x4},
-  {0x4005, 0x0},
-  {0x404e, 0x1},
-  {0x4300, 0xff},
-  {0x4301, 0x0},
-  {0x4315, 0x0},
-  {0x4501, 0x48},
-  {0x4600, 0x0},
-  {0x4601, 0x4e},
-  {0x4801, 0xf},
-  {0x4806, 0xf},
-  {0x4819, 0xaa},
-  {0x4823, 0x3e},
-  {0x4837, 0x19},
-  {0x4a0d, 0x0},
-  {0x4a47, 0x7f},
-  {0x4a49, 0xf0},
-  {0x4a4b, 0x30},
-  {0x5000, 0x85},
-  {0x5001, 0x80},
-};
-
-struct i2c_random_wr_payload init_array_ov8856[] = {
-  // part 1 184
-  {0x103, 0x1},
-  {0x302, 0x3c},
-  {0x303, 0x1},
-  {0x31e, 0xc},
-  {0x3000, 0x0},
-  {0x300e, 0x0},
-  {0x3010, 0x0},
-  {0x3015, 0x84},
-  {0x3018, 0x72},
-  {0x3033, 0x24},
-  {0x3500, 0x0},
-  {0x3501, 0x4c},
-  {0x3502, 0xe0},
-  {0x3503, 0x8},
-  {0x3505, 0x83},
-  {0x3508, 0x1},
-  {0x3509, 0x80},
-  {0x350c, 0x0},
-  {0x350d, 0x80},
-  {0x350e, 0x4},
-  {0x350f, 0x0},
-  {0x3510, 0x0},
-  {0x3511, 0x2},
-  {0x3512, 0x0},
-  {0x3600, 0x72},
-  {0x3601, 0x40},
-  {0x3602, 0x30},
-  {0x3610, 0xc5},
-  {0x3611, 0x58},
-  {0x3612, 0x5c},
-  {0x3613, 0x5a},
-  {0x3614, 0x60},
-  {0x3628, 0xff},
-  {0x3629, 0xff},
-  {0x362a, 0xff},
-  {0x3633, 0x10},
-  {0x3634, 0x10},
-  {0x3635, 0x10},
-  {0x3636, 0x10},
-  {0x3663, 0x8},
-  {0x3669, 0x34},
-  {0x366e, 0x8},
-  {0x3706, 0x86},
-  {0x370b, 0x7e},
-  {0x3714, 0x27},
-  {0x3730, 0x12},
-  {0x3733, 0x10},
-  {0x3764, 0x0},
-  {0x3765, 0x0},
-  {0x3769, 0x62},
-  {0x376a, 0x2a},
-  {0x376b, 0x3b},
-  {0x3780, 0x0},
-  {0x3781, 0x24},
-  {0x3782, 0x0},
-  {0x3783, 0x23},
-  {0x3798, 0x2f},
-  {0x37a1, 0x60},
-  {0x37a8, 0x6a},
-  {0x37ab, 0x3f},
-  {0x37c2, 0x14},
-  {0x37c3, 0xf1},
-  {0x37c9, 0x80},
-  {0x37cb, 0x3},
-  {0x37cc, 0xa},
-  {0x37cd, 0x16},
-  {0x37ce, 0x1f},
-  {0x3800, 0x0},
-  {0x3801, 0x0},
-  {0x3802, 0x0},
-  {0x3803, 0xc},
-  {0x3804, 0xc},
-  {0x3805, 0xdf},
-  {0x3806, 0x9},
-  {0x3807, 0xa3},
-  {0x3808, 0x6},
-  {0x3809, 0x60},
-  {0x380a, 0x4},
-  {0x380b, 0xc8},
-  {0x380c, 0x7},
-  {0x380d, 0x8c},
-  {0x380e, 0x9},
-  {0x380f, 0xb2},
-  {0x3810, 0x0},
-  {0x3811, 0x8},
-  {0x3812, 0x0},
-  {0x3813, 0x2},
-  {0x3814, 0x3},
-  {0x3815, 0x1},
-  {0x3816, 0x0},
-  {0x3817, 0x0},
-  {0x3818, 0x0},
-  {0x3819, 0x0},
-  {0x3820, 0x90},
-  {0x3821, 0x67},
-  {0x382a, 0x3},
-  {0x382b, 0x1},
-  {0x3830, 0x6},
-  {0x3836, 0x2},
-  {0x3862, 0x4},
-  {0x3863, 0x8},
-  {0x3cc0, 0x33},
-  {0x3d85, 0x17},
-  {0x3d8c, 0x73},
-  {0x3d8d, 0xde},
-  {0x4001, 0xe0},
-  {0x4003, 0x40},
-  {0x4008, 0x0},
-  {0x4009, 0x5},
-  {0x400f, 0x80},
-  {0x4010, 0xf0},
-  {0x4011, 0xff},
-  {0x4012, 0x2},
-  {0x4013, 0x1},
-  {0x4014, 0x1},
-  {0x4015, 0x1},
-  {0x4042, 0x0},
-  {0x4043, 0x80},
-  {0x4044, 0x0},
-  {0x4045, 0x80},
-  {0x4046, 0x0},
-  {0x4047, 0x80},
-  {0x4048, 0x0},
-  {0x4049, 0x80},
-  {0x4041, 0x3},
-  {0x404c, 0x20},
-  {0x404d, 0x0},
-  {0x404e, 0x20},
-  {0x4203, 0x80},
-  {0x4307, 0x30},
-  {0x4317, 0x0},
-  {0x4503, 0x8},
-  {0x4601, 0x80},
-  {0x4816, 0x53},
-  {0x481b, 0x58},
-  {0x481f, 0x27},
-  {0x4837, 0x16},
-  {0x5000, 0x77},
-  {0x5001, 0xe},
-  {0x5004, 0x0},
-  {0x502e, 0x0},
-  {0x5030, 0x41},
-  {0x5795, 0x0},
-  {0x5796, 0x10},
-  {0x5797, 0x10},
-  {0x5798, 0x73},
-  {0x5799, 0x73},
-  {0x579a, 0x0},
-  {0x579b, 0x28},
-  {0x579c, 0x0},
-  {0x579d, 0x16},
-  {0x579e, 0x6},
-  {0x579f, 0x20},
-  {0x57a0, 0x4},
-  {0x57a1, 0xa0},
-  {0x5780, 0x14},
-  {0x5781, 0xf},
-  {0x5782, 0x44},
-  {0x5783, 0x2},
-  {0x5784, 0x1},
-  {0x5785, 0x1},
-  {0x5786, 0x0},
-  {0x5787, 0x4},
-  {0x5788, 0x2},
-  {0x5789, 0xf},
-  {0x578a, 0xfd},
-  {0x578b, 0xf5},
-  {0x578c, 0xf5},
-  {0x578d, 0x3},
-  {0x578e, 0x8},
-  {0x578f, 0xc},
-  {0x5790, 0x8},
-  {0x5791, 0x4},
-  {0x5792, 0x0},
-  {0x5793, 0x52},
-  {0x5794, 0xa3},
-  {0x5a08, 0x2},
-  {0x5b00, 0x2},
-  {0x5b01, 0x10},
-  {0x5b02, 0x3},
-  {0x5b03, 0xcf},
-  {0x5b05, 0x6c},
-  {0x5e00, 0x0},
-
-  // part 2 45
-  {0x3501, 0x9a},
-  {0x3502, 0x20},
-  {0x366d, 0x0},
-  {0x366e, 0x10},
-  {0x3714, 0x23},
-  {0x37c2, 0x4},
-  {0x3800, 0x0},
-  {0x3801, 0x0},
-  {0x3802, 0x0},
-  {0x3803, 0xc},
-  {0x3804, 0xc},
-  {0x3805, 0xdf},
-  {0x3806, 0x9},
-  {0x3807, 0xa3},
-  {0x3808, 0xc},
-  {0x3809, 0xc0},
-  {0x380a, 0x9},
-  {0x380b, 0x90},
-  {0x380c, 0x7},
-  {0x380d, 0x8c},
-  {0x380e, 0x9},
-  {0x380f, 0xb2},
-  {0x3811, 0x10},
-  {0x3813, 0x4},
-  {0x3814, 0x1},
-  {0x3820, 0xc6},
-  {0x3821, 0x40},
-  {0x382a, 0x1},
-  {0x4009, 0xb},
-  {0x4601, 0x80},
-  {0x5003, 0xc8},
-  {0x5006, 0x0},
-  {0x5007, 0x0},
-  {0x5795, 0x2},
-  {0x5796, 0x20},
-  {0x5797, 0x20},
-  {0x5798, 0xd5},
-  {0x5799, 0xd5},
-  {0x579b, 0x50},
-  {0x579d, 0x2c},
-  {0x579e, 0xc},
-  {0x579f, 0x40},
-  {0x57a0, 0x9},
-  {0x57a1, 0x40},
-  {0x5e10, 0xfc},
 };