diff --git a/system/camerad/cameras/camera_common.cc b/system/camerad/cameras/camera_common.cc index 91c6d44f84..cf398a3a2c 100644 --- 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; diff --git a/system/camerad/cameras/camera_qcom2.cc b/system/camerad/cameras/camera_qcom2.cc index e17203a7dd..1c8a2457eb 100644 --- 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"); diff --git a/system/camerad/cameras/real_debayer.cl b/system/camerad/cameras/real_debayer.cl index e15a873d6d..38d77e9d14 100644 --- 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; diff --git a/system/camerad/sensors/os04c10.cc b/system/camerad/sensors/os04c10.cc index 449e06be83..3fa4728074 100644 --- a/system/camerad/sensors/os04c10.cc +++ b/system/camerad/sensors/os04c10.cc @@ -77,15 +77,15 @@ std::vector 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}, }; } diff --git a/system/camerad/sensors/os04c10_registers.h b/system/camerad/sensors/os04c10_registers.h index ad91a02950..5f1d1294f9 100644 --- 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}, };