From 8824966aad4da328cc393cda26dd80d5bbcd7265 Mon Sep 17 00:00:00 2001 From: ZwX1616 Date: Mon, 1 Feb 2021 15:06:15 -0800 Subject: [PATCH] tici fcam vignette compensation (#19971) * simple model * fix api * this too --- selfdrive/camerad/cameras/camera_common.cc | 8 +++---- .../camerad/cameras/camera_frame_stream.cc | 3 ++- .../camerad/cameras/camera_frame_stream.h | 1 + selfdrive/camerad/cameras/camera_webcam.cc | 2 +- selfdrive/camerad/cameras/camera_webcam.h | 1 + selfdrive/camerad/cameras/real_debayer.cl | 23 ++++++++++++------- 6 files changed, 24 insertions(+), 14 deletions(-) diff --git a/selfdrive/camerad/cameras/camera_common.cc b/selfdrive/camerad/cameras/camera_common.cc index ae9c93cb75..68ef891841 100644 --- a/selfdrive/camerad/cameras/camera_common.cc +++ b/selfdrive/camerad/cameras/camera_common.cc @@ -24,16 +24,16 @@ #include "common/util.h" #include "imgproc/utils.h" -static cl_program build_debayer_program(cl_device_id device_id, cl_context context, const CameraInfo *ci, const CameraBuf *b) { +static cl_program build_debayer_program(cl_device_id device_id, cl_context context, const CameraInfo *ci, const CameraBuf *b, const CameraState *s) { char args[4096]; snprintf(args, sizeof(args), "-cl-fast-relaxed-math -cl-denorms-are-zero " "-DFRAME_WIDTH=%d -DFRAME_HEIGHT=%d -DFRAME_STRIDE=%d " "-DRGB_WIDTH=%d -DRGB_HEIGHT=%d -DRGB_STRIDE=%d " - "-DBAYER_FLIP=%d -DHDR=%d", + "-DBAYER_FLIP=%d -DHDR=%d -DCAM_NUM=%d", ci->frame_width, ci->frame_height, ci->frame_stride, b->rgb_width, b->rgb_height, b->rgb_stride, - ci->bayer_flip, ci->hdr); + ci->bayer_flip, ci->hdr, s->camera_num); #ifdef QCOM2 return cl_program_from_file(context, device_id, "cameras/real_debayer.cl", args); #else @@ -86,7 +86,7 @@ void CameraBuf::init(cl_device_id device_id, cl_context context, CameraState *s, vipc_server->create_buffers(yuv_type, YUV_COUNT, false, rgb_width, rgb_height); if (ci->bayer) { - cl_program prg_debayer = build_debayer_program(device_id, context, ci, this); + cl_program prg_debayer = build_debayer_program(device_id, context, ci, this, s); krnl_debayer = CL_CHECK_ERR(clCreateKernel(prg_debayer, "debayer10", &err)); CL_CHECK(clReleaseProgram(prg_debayer)); } diff --git a/selfdrive/camerad/cameras/camera_frame_stream.cc b/selfdrive/camerad/cameras/camera_frame_stream.cc index 4536f82759..e982416c59 100644 --- a/selfdrive/camerad/cameras/camera_frame_stream.cc +++ b/selfdrive/camerad/cameras/camera_frame_stream.cc @@ -20,6 +20,7 @@ void camera_init(VisionIpcServer * v, CameraState *s, int camera_id, unsigned in s->ci = cameras_supported[camera_id]; assert(s->ci.frame_width != 0); + s->camera_num = camera_id; s->fps = fps; s->buf.init(device_id, ctx, s, v, FRAME_BUF_COUNT, rgb_type, yuv_type); } @@ -66,7 +67,7 @@ CameraInfo cameras_supported[CAMERA_ID_MAX] = { .frame_width = 1632, .frame_height = 1224, .frame_stride = 2040, // seems right - .bayer = true, + .bayer = false, .bayer_flip = 3, .hdr = false }, diff --git a/selfdrive/camerad/cameras/camera_frame_stream.h b/selfdrive/camerad/cameras/camera_frame_stream.h index 65734884f2..98469dfbcd 100644 --- a/selfdrive/camerad/cameras/camera_frame_stream.h +++ b/selfdrive/camerad/cameras/camera_frame_stream.h @@ -15,6 +15,7 @@ typedef struct CameraState { int camera_id; + int camera_num; CameraInfo ci; int fps; diff --git a/selfdrive/camerad/cameras/camera_webcam.cc b/selfdrive/camerad/cameras/camera_webcam.cc index fa54f7964f..1543cc5d15 100644 --- a/selfdrive/camerad/cameras/camera_webcam.cc +++ b/selfdrive/camerad/cameras/camera_webcam.cc @@ -39,8 +39,8 @@ void camera_init(VisionIpcServer * v, CameraState *s, int camera_id, unsigned in s->ci = cameras_supported[camera_id]; assert(s->ci.frame_width != 0); + s->camera_num = camera_id; s->fps = fps; - s->buf.init(device_id, ctx, s, v, FRAME_BUF_COUNT, rgb_type, yuv_type); } diff --git a/selfdrive/camerad/cameras/camera_webcam.h b/selfdrive/camerad/cameras/camera_webcam.h index 50fa514288..80e41fa13f 100644 --- a/selfdrive/camerad/cameras/camera_webcam.h +++ b/selfdrive/camerad/cameras/camera_webcam.h @@ -14,6 +14,7 @@ typedef struct CameraState { CameraInfo ci; + int camera_num; int fps; float digital_gain; CameraBuf buf; diff --git a/selfdrive/camerad/cameras/real_debayer.cl b/selfdrive/camerad/cameras/real_debayer.cl index f2cf806331..bd6556b915 100644 --- a/selfdrive/camerad/cameras/real_debayer.cl +++ b/selfdrive/camerad/cameras/real_debayer.cl @@ -24,11 +24,18 @@ uint int_from_10(const uchar * source, uint start, uint offset) { return major + minor; } -float to_normal(uint x) { +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); + 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); + pv = s * pv; + } pv = 20*pv / (1.0f + 20*pv); // reinhard return pv; } @@ -54,7 +61,7 @@ __kernel void debayer10(const __global uchar * in, 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); + cached[localOffset] = to_normal(raw_val, x_global, y_global); // edges if (x_global < 1 || x_global > RGB_WIDTH - 2 || y_global < 1 || y_global > RGB_HEIGHT - 2) { @@ -68,22 +75,22 @@ __kernel void debayer10(const __global uchar * in, 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)); + 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)); + 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); } if (y_local < 1) { - cached[y_local * localRowLen + x_local + 1] = to_normal(int_from_10(in, globalStart_10 - FRAME_STRIDE, offset_10)); + 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)); + 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)); + 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)); + 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); } }