diff --git a/system/camerad/cameras/camera_common.cc b/system/camerad/cameras/camera_common.cc index dee7e13ffe..b6e7f0ae66 100644 --- a/system/camerad/cameras/camera_common.cc +++ b/system/camerad/cameras/camera_common.cc @@ -28,12 +28,12 @@ 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 -DIS_OS=%d -DIS_10BIT=%d -DIS_HDR=%d -DHDR_OFFSET=%d -DIS_BGGR=%d -DVIGNETTING=%d ", - ci->frame_width, ci->frame_height, ci->frame_stride * 2, ci->frame_offset, + 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::OX03C10, ci->image_sensor == cereal::FrameData::ImageSensor::OS04C10, ci->mipi_format == CAM_FORMAT_MIPI_RAW_10, - ci->hdr_offset > 0, ci->hdr_offset / 2, ci->bggr, + ci->hdr_offset > 0, ci->hdr_offset, ci->bggr, s->camera_num == 1); const char *cl_file = "cameras/real_debayer.cl"; cl_program prg_debayer = cl_program_from_file(context, device_id, cl_file, args); diff --git a/system/camerad/cameras/real_debayer.cl b/system/camerad/cameras/real_debayer.cl index e5290d24e8..a18898825d 100644 --- a/system/camerad/cameras/real_debayer.cl +++ b/system/camerad/cameras/real_debayer.cl @@ -215,10 +215,10 @@ __kernel void debayer10(const __global uchar * in, __global uchar * out, int exp dat[3] = vload8(0, in + start_idx + FRAME_STRIDE*row_after_offset); #if IS_HDR uchar8 short_dat[4]; - short_dat[0] = vload8(0, in + start_idx + FRAME_STRIDE*(row_before_offset+HDR_OFFSET) + FRAME_STRIDE/2); - short_dat[1] = vload8(0, in + start_idx + FRAME_STRIDE*(1+HDR_OFFSET) + FRAME_STRIDE/2); - short_dat[2] = vload8(0, in + start_idx + FRAME_STRIDE*(2+HDR_OFFSET) + FRAME_STRIDE/2); - short_dat[3] = vload8(0, in + start_idx + FRAME_STRIDE*(row_after_offset+HDR_OFFSET) + FRAME_STRIDE/2); + 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); + short_dat[2] = vload8(0, in + start_idx + FRAME_STRIDE*(2+HDR_OFFSET/2) + FRAME_STRIDE/2); + short_dat[3] = vload8(0, in + start_idx + FRAME_STRIDE*(row_after_offset+HDR_OFFSET/2) + FRAME_STRIDE/2); #endif // need extra bit for 10-bit @@ -233,10 +233,10 @@ __kernel void debayer10(const __global uchar * in, __global uchar * out, int exp #if IS_HDR uchar short_extra[4]; if (!aligned10) { - short_extra[0] = in[start_idx + FRAME_STRIDE*(row_before_offset+HDR_OFFSET) + FRAME_STRIDE/2 + 8]; - short_extra[1] = in[start_idx + FRAME_STRIDE*(1+HDR_OFFSET) + FRAME_STRIDE/2 + 8]; - short_extra[2] = in[start_idx + FRAME_STRIDE*(2+HDR_OFFSET) + FRAME_STRIDE/2 + 8]; - short_extra[3] = in[start_idx + FRAME_STRIDE*(row_after_offset+HDR_OFFSET) + FRAME_STRIDE/2 + 8]; + short_extra[0] = in[start_idx + FRAME_STRIDE*(row_before_offset+HDR_OFFSET/2) + FRAME_STRIDE/2 + 8]; + short_extra[1] = in[start_idx + FRAME_STRIDE*(1+HDR_OFFSET/2) + FRAME_STRIDE/2 + 8]; + short_extra[2] = in[start_idx + FRAME_STRIDE*(2+HDR_OFFSET/2) + FRAME_STRIDE/2 + 8]; + short_extra[3] = in[start_idx + FRAME_STRIDE*(row_after_offset+HDR_OFFSET/2) + FRAME_STRIDE/2 + 8]; } #endif #endif