it's something

1 year ago · 77a095fd10
parent b94aba6281
commit 77a095fd10
5 changed files with 76 additions and 28 deletions
--- a/system/camerad/cameras/camera_common.cc
+++ b/system/camerad/cameras/camera_common.cc
@ -21,6 +21,8 @@ ExitHandler do_exit;
 class Debayer {
 public:
  Debayer(cl_device_id device_id, cl_context context, const CameraBuf *b, const CameraState *s, int buf_width, int uv_offset) {
+    compat = true;
+
    char args[4096];
    const SensorInfo *ci = s->ci.get();
    snprintf(args, sizeof(args),
@ -28,17 +30,43 @@ public:
             "-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_OX=%d -DCAM_NUM=%d%s",
-             ci->frame_width, ci->frame_height, ci->frame_stride, ci->frame_offset,
+             ci->frame_width, ci->frame_height, compat ? 2880 : ci->frame_stride, ci->frame_offset,
             b->rgb_width, b->rgb_height, buf_width, uv_offset,
             ci->image_sensor == cereal::FrameData::ImageSensor::OX03C10, s->camera_num, s->camera_num==1 ? " -DVIGNETTING" : "");
    const char *cl_file = "cameras/real_debayer.cl";
    cl_program prg_debayer = cl_program_from_file(context, device_id, cl_file, args);
    krnl_ = CL_CHECK_ERR(clCreateKernel(prg_debayer, "debayer10", &err));
    CL_CHECK(clReleaseProgram(prg_debayer));
+
+    twelve.allocate(2880 * 1080);
+    twelve.init_cl(device_id, context);
  }

-  void queue(cl_command_queue q, cl_mem cam_buf_cl, cl_mem buf_cl, int width, int height, cl_event *debayer_event) {
-    CL_CHECK(clSetKernelArg(krnl_, 0, sizeof(cl_mem), &cam_buf_cl));
+  void queue(cl_command_queue q, VisionBuf cam_buff, cl_mem cam_buf_cl, cl_mem buf_cl, int width, int height, cl_event *debayer_event) {
+    // 10 -> 12 bit
+    uint8_t *cam_buf = (uint8_t *)cam_buff.addr;
+    uint8_t *raw12 = (uint8_t *)twelve.addr;
+    size_t inidx = 0, outidx = 0;
+    for (size_t i = 0; i < (width*height)/ 4; ++i) {
+      // 4 RAW10 pixels (5 bytes)
+      uint8_t in0 = cam_buf[inidx++];
+      uint8_t in1 = cam_buf[inidx++];
+      uint8_t in2 = cam_buf[inidx++];
+      uint8_t in3 = cam_buf[inidx++];
+      uint8_t in4 = cam_buf[inidx++];
+
+      // 4 RAW12 pixels (6 bytes)
+      raw12[outidx++] = in0;
+      raw12[outidx++] = (in1 & 0x0F) | (in4 << 4 & 0x30);
+      raw12[outidx++] = in1 >> 4 | (in2 << 4 & 0xF0);
+      raw12[outidx++] = in2 >> 4 | (in4 << 2 & 0xC0);
+      raw12[outidx++] = in3;
+      raw12[outidx++] = in3 >> 4 | (in4 << 4 & 0xF0);
+    }
+
+    CL_CHECK(clSetKernelArg(krnl_, 0, sizeof(cl_mem), &twelve.buf_cl));
+
+    if (!compat) CL_CHECK(clSetKernelArg(krnl_, 0, sizeof(cl_mem), &cam_buf_cl));
    CL_CHECK(clSetKernelArg(krnl_, 1, sizeof(cl_mem), &buf_cl));

    const size_t globalWorkSize[] = {size_t(width / 2), size_t(height / 2)};
@ -53,6 +81,8 @@ public:

 private:
  cl_kernel krnl_;
+  VisionBuf twelve;
+  bool compat = false;
 };

 void CameraBuf::init(cl_device_id device_id, cl_context context, CameraState *s, VisionIpcServer * v, int frame_cnt, VisionStreamType type) {
@ -112,7 +142,7 @@ bool CameraBuf::acquire() {

  double start_time = millis_since_boot();
  cl_event event;
-  debayer->queue(q, camera_bufs[cur_buf_idx].buf_cl, cur_yuv_buf->buf_cl, rgb_width, rgb_height, &event);
+  debayer->queue(q, camera_bufs[cur_buf_idx], camera_bufs[cur_buf_idx].buf_cl, cur_yuv_buf->buf_cl, rgb_width, rgb_height, &event);
  clWaitForEvents(1, &event);
  CL_CHECK(clReleaseEvent(event));
  cur_frame_data.processing_time = (millis_since_boot() - start_time) / 1000.0;
--- a/system/camerad/cameras/camera_qcom2.cc
+++ b/system/camerad/cameras/camera_qcom2.cc
@ -897,7 +897,7 @@ void CameraState::set_camera_exposure(float grey_frac) {
  // LOGE("ae - camera %d, cur_t %.5f, sof %.5f, dt %.5f", camera_num, 1e-9 * nanos_since_boot(), 1e-9 * buf.cur_frame_data.timestamp_sof, 1e-9 * (nanos_since_boot() - buf.cur_frame_data.timestamp_sof));

  auto exp_reg_array = ci->getExposureRegisters(exposure_time, new_exp_g, dc_gain_enabled);
-  sensors_i2c(exp_reg_array.data(), exp_reg_array.size(), CAM_SENSOR_PACKET_OPCODE_SENSOR_CONFIG, ci->data_word);
+  if (exp_reg_array.size()) sensors_i2c(exp_reg_array.data(), exp_reg_array.size(), CAM_SENSOR_PACKET_OPCODE_SENSOR_CONFIG, ci->data_word);
 }

 static void process_driver_camera(MultiCameraState *s, CameraState *c, int cnt) {
@ -918,7 +918,7 @@ void process_road_camera(MultiCameraState *s, CameraState *c, int cnt) {
  MessageBuilder msg;
  auto framed = c == &s->road_cam ? msg.initEvent().initRoadCameraState() : msg.initEvent().initWideRoadCameraState();
  fill_frame_data(framed, b->cur_frame_data, c);
-  if (env_log_raw_frames && c == &s->road_cam && cnt % 100 == 5) {  // no overlap with qlog decimation
+  if (env_log_raw_frames) {  // no overlap with qlog decimation
    framed.setImage(get_raw_frame_image(b));
  }
  LOGT(c->buf.cur_frame_data.frame_id, "%s: Image set", c == &s->road_cam ? "RoadCamera" : "WideRoadCamera");
--- a/system/camerad/cameras/real_debayer.cl
+++ b/system/camerad/cameras/real_debayer.cl
@ -65,7 +65,7 @@ constant float ox03c10_lut[] = {
  3.7695e-01, 3.7890e-01, 3.8086e-01, 3.8281e-01, 3.8476e-01, 3.8672e-01, 3.8867e-01, 3.9062e-01, 3.9258e-01, 3.9453e-01, 3.9648e-01, 3.9844e-01, 4.0039e-01, 4.0234e-01, 4.0430e-01, 4.0625e-01, 4.0820e-01, 4.1015e-01, 4.1211e-01, 4.1406e-01, 4.1601e-01, 4.1797e-01, 4.1992e-01, 4.2187e-01, 4.2383e-01, 4.2578e-01, 4.2773e-01, 4.2969e-01, 4.3164e-01, 4.3359e-01, 4.3555e-01, 4.3750e-01, 4.3945e-01, 4.4140e-01, 4.4336e-01, 4.4531e-01, 4.4726e-01, 4.4922e-01, 4.5117e-01, 4.5312e-01, 4.5508e-01, 4.5703e-01, 4.5898e-01, 4.6094e-01, 4.6289e-01, 4.6484e-01, 4.6680e-01, 4.6875e-01, 4.7070e-01, 4.7265e-01, 4.7461e-01, 4.7656e-01, 4.7851e-01, 4.8047e-01, 4.8242e-01, 4.8437e-01, 4.8633e-01, 4.8828e-01, 4.9023e-01, 4.9219e-01, 4.9414e-01, 4.9609e-01, 4.9805e-01, 5.0000e-01, 5.0195e-01, 5.0390e-01, 5.0586e-01, 5.0781e-01, 5.0976e-01, 5.1172e-01, 5.1367e-01, 5.1562e-01, 5.1758e-01, 5.1953e-01, 5.2148e-01, 5.2344e-01, 5.2539e-01, 5.2734e-01, 5.2930e-01, 5.3125e-01, 5.3320e-01, 5.3515e-01, 5.3711e-01, 5.3906e-01, 5.4101e-01, 5.4297e-01, 5.4492e-01, 5.4687e-01, 5.4883e-01, 5.5078e-01, 5.5273e-01, 5.5469e-01, 5.5664e-01, 5.5859e-01, 5.6055e-01, 5.6250e-01, 5.6445e-01, 5.6640e-01, 5.6836e-01, 5.7031e-01, 5.7226e-01, 5.7422e-01, 5.7617e-01, 5.7812e-01, 5.8008e-01, 5.8203e-01, 5.8398e-01, 5.8594e-01, 5.8789e-01, 5.8984e-01, 5.9180e-01, 5.9375e-01, 5.9570e-01, 5.9765e-01, 5.9961e-01, 6.0156e-01, 6.0351e-01, 6.0547e-01, 6.0742e-01, 6.0937e-01, 6.1133e-01, 6.1328e-01, 6.1523e-01, 6.1719e-01, 6.1914e-01, 6.2109e-01, 6.2305e-01, 6.2500e-01, 6.2695e-01, 6.2890e-01, 6.3086e-01, 6.3281e-01, 6.3476e-01, 6.3672e-01, 6.3867e-01, 6.4062e-01, 6.4258e-01, 6.4453e-01, 6.4648e-01, 6.4844e-01, 6.5039e-01, 6.5234e-01, 6.5430e-01, 6.5625e-01, 6.5820e-01, 6.6015e-01, 6.6211e-01, 6.6406e-01, 6.6601e-01, 6.6797e-01, 6.6992e-01, 6.7187e-01, 6.7383e-01, 6.7578e-01, 6.7773e-01, 6.7969e-01, 6.8164e-01, 6.8359e-01, 6.8554e-01, 6.8750e-01, 6.8945e-01, 6.9140e-01, 6.9336e-01, 6.9531e-01, 6.9726e-01, 6.9922e-01, 7.0117e-01, 7.0312e-01, 7.0508e-01, 7.0703e-01, 7.0898e-01, 7.1094e-01, 7.1289e-01, 7.1484e-01, 7.1679e-01, 7.1875e-01, 7.2070e-01, 7.2265e-01, 7.2461e-01, 7.2656e-01, 7.2851e-01, 7.3047e-01, 7.3242e-01, 7.3437e-01, 7.3633e-01, 7.3828e-01, 7.4023e-01, 7.4219e-01, 7.4414e-01, 7.4609e-01, 7.4804e-01, 7.5000e-01, 7.5390e-01, 7.5781e-01, 7.6172e-01, 7.6562e-01, 7.6953e-01, 7.7344e-01, 7.7734e-01, 7.8125e-01, 7.8515e-01, 7.8906e-01, 7.9297e-01, 7.9687e-01, 8.0078e-01, 8.0469e-01, 8.0859e-01, 8.1250e-01, 8.1640e-01, 8.2031e-01, 8.2422e-01, 8.2812e-01, 8.3203e-01, 8.3594e-01, 8.3984e-01, 8.4375e-01, 8.4765e-01, 8.5156e-01, 8.5547e-01, 8.5937e-01, 8.6328e-01, 8.6719e-01, 8.7109e-01, 8.7500e-01, 8.7890e-01, 8.8281e-01, 8.8672e-01, 8.9062e-01, 8.9453e-01, 8.9844e-01, 9.0234e-01, 9.0625e-01, 9.1015e-01, 9.1406e-01, 9.1797e-01, 9.2187e-01, 9.2578e-01, 9.2969e-01, 9.3359e-01, 9.3750e-01, 9.4140e-01, 9.4531e-01, 9.4922e-01, 9.5312e-01, 9.5703e-01, 9.6094e-01, 9.6484e-01, 9.6875e-01, 9.7265e-01, 9.7656e-01, 9.8047e-01, 9.8437e-01, 9.8828e-01, 9.9219e-01, 9.9609e-01, 1.0000e+00
 };

-float4 val4_from_12(uchar8 pvs, float gain) {
+float4 val4_from_12(uchar8 pvs, float gain, int loc) {
  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),
@ -100,6 +100,7 @@ __kernel void debayer10(const __global uchar * in, __global uchar * out)
  float3 rgb;
  uchar3 rgb_out[4];

+  int loc = gid_x % 5;
  int start = (2 * gid_y - 1) * FRAME_STRIDE + (3 * gid_x - 2) + (FRAME_STRIDE * FRAME_OFFSET);

  // read in 8x4 chars
@ -119,10 +120,10 @@ __kernel void debayer10(const __global uchar * in, __global uchar * out)
  #endif

  // process them to floats
-  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);
+  float4 va = val4_from_12(dat[0], gain, loc);
+  float4 vb = val4_from_12(dat[1], gain, loc);
+  float4 vc = val4_from_12(dat[2], gain, loc);
+  float4 vd = val4_from_12(dat[3], gain, loc);

  if (gid_x == 0) {
    va.s0 = va.s2;
--- a/system/camerad/sensors/os04c10.cc
+++ b/system/camerad/sensors/os04c10.cc
@ -77,15 +77,15 @@ std::vector<i2c_random_wr_payload> OS04C10::getExposureRegisters(int exposure_ti
  uint32_t real_gain = os04c10_analog_gains_reg[new_exp_g];

  hcg_time = 100;
-  real_gain = 0x320;
+  real_gain = 0x0;

  return {
-    {0x3501, hcg_time>>8}, {0x3502, hcg_time&0xFF},
+    //{0x3501, hcg_time>>8}, {0x3502, hcg_time&0xFF},
    //{0x3581, lcg_time>>8}, {0x3582, lcg_time&0xFF},
    //{0x3541, spd_time>>8}, {0x3542, spd_time&0xFF},
    //{0x35c2, vs_time&0xFF},

-    {0x3508, real_gain>>8}, {0x3509, real_gain&0xFF},
+    //{0x3508, real_gain>>8}, {0x3509, real_gain&0xFF},
  };
 }

--- a/system/camerad/sensors/os04c10_registers.h
+++ b/system/camerad/sensors/os04c10_registers.h
@ -4,6 +4,7 @@ const struct i2c_random_wr_payload start_reg_array_os04c10[] = {{0x100, 1}};
 const struct i2c_random_wr_payload stop_reg_array_os04c10[] = {{0x100, 0}};

 const struct i2c_random_wr_payload init_array_os04c10[] = {
+
  // OS04C10_AA_00_02_17_wAO_1920x1080_MIPI728Mbps_Linear12bit_20FPS_4Lane_MCLK24MHz
  {0x0103, 0x01},
  {0x0301, 0x84},
@ -21,20 +22,7 @@ const struct i2c_random_wr_payload init_array_os04c10[] = {
  {0x3106, 0x21},
  {0x3107, 0xa1},
  {0x3500, 0x00},
-  {0x3501, 0x00},
-  {0x3502, 0x40},
  {0x3503, 0x88},
-  {0x3508, 0x07},
-  {0x3509, 0xc0},
-  {0x350a, 0x04},
-  {0x350b, 0x00},
-  {0x350c, 0x07},
-  {0x350d, 0xc0},
-  {0x350e, 0x04},
-  {0x350f, 0x00},
-  {0x3510, 0x00},
-  {0x3511, 0x00},
-  {0x3512, 0x20},
  {0x3624, 0x00},
  {0x3625, 0x4c},
  {0x3660, 0x00},
@ -158,7 +146,7 @@ const struct i2c_random_wr_payload init_array_os04c10[] = {
  {0x4e0d, 0x00},
  {0x5001, 0x09},
  {0x5004, 0x00},
-  {0x5080, 0x04},
+  //{0x5080, 0x00 | (0b1 << 7)},  // test mode
  {0x5036, 0x00},
  {0x5180, 0x70},
  {0x5181, 0x10},
@ -295,4 +283,33 @@ const struct i2c_random_wr_payload init_array_os04c10[] = {
  {0x4c05, 0x00},
  {0x5000, 0xf9},
  {0x3c8c, 0x10},
+
+  // exposure
+  {0x3501, 0x00},
+  {0x3502, 0x40},
+
+  // gain
+  {0x3508, 0x01},
+  {0x3509, 0xc0},
+  {0x350a, 0x0f},
+  {0x350b, 0x00},
+  {0x350c, 0x07},
+  {0x350d, 0xc0},
+  {0x350e, 0x04},
+  {0x350f, 0x00},
+  {0x3510, 0x00},
+  {0x3511, 0x00},
+  {0x3512, 0x20},
+
+  // WB gain
+  {0x5100, 0x04}, {0x5101, 0x00},
+  {0x5102, 0x04}, {0x5103, 0x00},
+  {0x5104, 0x04}, {0x5105, 0x00},
+  {0x5140, 0x04}, {0x5141, 0x00},
+  {0x5142, 0x04}, {0x5143, 0x00},
+  {0x5144, 0x04}, {0x5145, 0x00},
+
+  //
+  {0x5106, 0x00}, {0x5107, 0x00},
+  {0x5146, 0x00}, {0x5147, 0x00},
 };