mawei/dragonpilot
1
0
Fork 0
mirror of https://github.com/dragonpilot/dragonpilot.git synced 2026-06-22 22:42:05 +08:00
Code Issues Packages Projects Releases Wiki Activity

tici fcam vignette compensation (#19971)

Browse Source 
* simple model

* fix api

* this too
This commit is contained in:
ZwX1616
2021-02-01 15:06:15 -08:00
committed by GitHub
parent bacf8727d7
commit 8824966aad
6 changed files with 24 additions and 14 deletions
Download Patch File Download Diff File Expand all files Collapse all files
selfdrive/camerad/cameras/camera_common.cc
+4 -4
View File
@@ -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));
}
selfdrive/camerad/cameras/camera_frame_stream.cc
+2 -1
View File
@@ -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
},
selfdrive/camerad/cameras/camera_frame_stream.h
+1
View File
@@ -15,6 +15,7 @@
typedef struct CameraState {
int camera_id;
int camera_num;
CameraInfo ci;
int fps;
selfdrive/camerad/cameras/camera_webcam.cc
+1 -1
View File
@@ -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);
}
selfdrive/camerad/cameras/camera_webcam.h
+1
View File
@@ -14,6 +14,7 @@
typedef struct CameraState {
CameraInfo ci;
int camera_num;
int fps;
float digital_gain;
CameraBuf buf;
selfdrive/camerad/cameras/real_debayer.cl
+15 -8
View File
@@ -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);
}
}