diff --git a/system/camerad/cameras/camera_common.cc b/system/camerad/cameras/camera_common.cc index 8ada8d2689..095e015cb9 100644 --- a/system/camerad/cameras/camera_common.cc +++ b/system/camerad/cameras/camera_common.cc @@ -27,14 +27,14 @@ public: "-cl-fast-relaxed-math -cl-denorms-are-zero -Isensors " "-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_AR=%d -DIS_OX=%d -DIS_OS=%d -DIS_10BIT=%d -DIS_HDR=%d -DHDR_OFFSET=%d -DVIGNETTING=%d ", + "-DIS_AR=%d -DIS_OX=%d -DIS_OS=%d -DIS_10BIT=%d -DHDR_OFFSET=%d -DVIGNETTING=%d ", ci->frame_width, ci->frame_height, ci->hdr_offset > 0 ? ci->frame_stride * 2 : ci->frame_stride, ci->frame_offset, b->rgb_width, b->rgb_height, buf_width, uv_offset, ci->image_sensor == cereal::FrameData::ImageSensor::AR0231, ci->image_sensor == cereal::FrameData::ImageSensor::OX03C10, ci->image_sensor == cereal::FrameData::ImageSensor::OS04C10, ci->mipi_format == CAM_FORMAT_MIPI_RAW_10, - ci->hdr_offset > 0, ci->hdr_offset, + ci->hdr_offset, s->camera_num == 1); const char *cl_file = "cameras/process_raw.cl"; cl_program prg_imgproc = cl_program_from_file(context, device_id, cl_file, args); diff --git a/system/camerad/cameras/process_raw.cl b/system/camerad/cameras/process_raw.cl index 3f8b1ea40f..2ea1d0a30c 100644 --- a/system/camerad/cameras/process_raw.cl +++ b/system/camerad/cameras/process_raw.cl @@ -162,7 +162,7 @@ __kernel void process_raw(const __global uchar * in, __global uchar * out, int e } dat[2] = vload8(0, in + start_idx + FRAME_STRIDE*2); dat[3] = vload8(0, in + start_idx + FRAME_STRIDE*row_after_offset); - #if IS_HDR + #if HDR_COMBINE uchar8 short_dat[4]; short_dat[0] = vload8(0, in + start_idx + FRAME_STRIDE*(row_before_offset+HDR_OFFSET/2) + FRAME_STRIDE/2); short_dat[1] = vload8(0, in + start_idx + FRAME_STRIDE*(1+HDR_OFFSET/2) + FRAME_STRIDE/2); @@ -179,7 +179,7 @@ __kernel void process_raw(const __global uchar * in, __global uchar * out, int e extra[2] = in[start_idx + FRAME_STRIDE*2 + 8]; extra[3] = in[start_idx + FRAME_STRIDE*row_after_offset + 8]; } - #if IS_HDR + #if HDR_COMBINE uchar short_extra[4]; if (!aligned10) { short_extra[0] = in[start_idx + FRAME_STRIDE*(row_before_offset+HDR_OFFSET/2) + FRAME_STRIDE/2 + 8]; @@ -202,7 +202,7 @@ __kernel void process_raw(const __global uchar * in, __global uchar * out, int e float4 v_rows[4]; // parse into floats #if IS_10BIT - #if IS_HDR + #if HDR_COMBINE v_rows[ROW_READ_ORDER[0]] = val4_from_10x2(dat[0], extra[0], short_dat[0], short_extra[0], aligned10, gain, expo_time); v_rows[ROW_READ_ORDER[1]] = val4_from_10x2(dat[1], extra[1], short_dat[1], short_extra[1], aligned10, gain, expo_time); v_rows[ROW_READ_ORDER[2]] = val4_from_10x2(dat[2], extra[2], short_dat[2], short_extra[2], aligned10, gain, expo_time); diff --git a/system/camerad/sensors/os04c10_cl.h b/system/camerad/sensors/os04c10_cl.h index 1e2183ebeb..a963b7bb2f 100644 --- a/system/camerad/sensors/os04c10_cl.h +++ b/system/camerad/sensors/os04c10_cl.h @@ -1,6 +1,7 @@ #if IS_OS #define BGGR +#define HDR_COMBINE float3 color_correct(float3 rgb) { float3 corrected = rgb.x * (float3)(1.55361989, -0.268894615, -0.000593219);