sadmen
/
jyoung
mirror of https://github.com/jyoung8607/openpilot.git
1
0
Fork 0
Code Issues Packages Projects Releases Wiki Activity

tici camerart III: prehistoric visuals (#20012)

Browse Source 
* reworked

* add ref

* f16

* faster

* extended tm

* fix hdr

* Revert "extended tm"

This reverts commit 55ba3ae3205f7e962bf2f36634b23a595700c9a8.

* reduce static/temporal noise

* split kernels to sync max

* no print

* cleanup qcom2

* optimize

* tune DC

* doesnt work

* minor fix

* adaptive

* smoothen HCG switch

* t

* histogram ceiling

* recalulate params

* group_hold

* less contrast

* recalibrate ccm

* better ae

* better ae revised

* made some changes

* clean up

* midtones

* some improvements

* more

* cleanup

* remove more junk

* lgtm

Co-authored-by: Comma Device <device@comma.ai>
pull/59/head
ZwX1616 5 years ago committed by GitHub
parent c747654f05
commit d8ddc1208e
No known key found for this signature in database
GPG Key ID: 4AEE18F83AFDEB23
5 changed files with 220 additions and 754 deletions
Whitespace
Show all changes Ignore whitespace when comparing lines Ignore changes in amount of whitespace Ignore changes in whitespace at EOL
Unified View
Diff Options
Show Stats Download Patch File Download Diff File
  1. 27
      selfdrive/camerad/cameras/camera_common.cc
  2. 25
      selfdrive/camerad/cameras/camera_qcom2.cc
  3. 9
      selfdrive/camerad/cameras/camera_qcom2.h
  4. 227
      selfdrive/camerad/cameras/real_debayer.cl
  5. 638
      selfdrive/camerad/cameras/sensor2_i2c.h

27
selfdrive/camerad/cameras/camera_common.cc
Unescape View File

@ -134,10 +134,12 @@ bool CameraBuf::acquire() {
CL_CHECK(clSetKernelArg(krnl_debayer, 0, sizeof(cl_mem), &camrabuf_cl)); CL_CHECK(clSetKernelArg(krnl_debayer, 0, sizeof(cl_mem), &camrabuf_cl));
CL_CHECK(clSetKernelArg(krnl_debayer, 1, sizeof(cl_mem), &cur_rgb_buf->buf_cl)); CL_CHECK(clSetKernelArg(krnl_debayer, 1, sizeof(cl_mem), &cur_rgb_buf->buf_cl));
#ifdef QCOM2 #ifdef QCOM2
constexpr int localMemSize = (DEBAYER_LOCAL_WORKSIZE + 2 * (3 / 2)) * (DEBAYER_LOCAL_WORKSIZE + 2 * (3 / 2)) * sizeof(float); constexpr int localMemSize = (DEBAYER_LOCAL_WORKSIZE + 2 * (3 / 2)) * (DEBAYER_LOCAL_WORKSIZE + 2 * (3 / 2)) * sizeof(short int);
const size_t globalWorkSize[] = {size_t(camera_state->ci.frame_width), size_t(camera_state->ci.frame_height)}; const size_t globalWorkSize[] = {size_t(camera_state->ci.frame_width), size_t(camera_state->ci.frame_height)};
const size_t localWorkSize[] = {DEBAYER_LOCAL_WORKSIZE, DEBAYER_LOCAL_WORKSIZE}; const size_t localWorkSize[] = {DEBAYER_LOCAL_WORKSIZE, DEBAYER_LOCAL_WORKSIZE};
CL_CHECK(clSetKernelArg(krnl_debayer, 2, localMemSize, 0)); CL_CHECK(clSetKernelArg(krnl_debayer, 2, localMemSize, 0));
int ggain = camera_state->analog_gain + 4*camera_state->dc_gain_enabled;
CL_CHECK(clSetKernelArg(krnl_debayer, 3, sizeof(int), &ggain));
CL_CHECK(clEnqueueNDRangeKernel(q, krnl_debayer, 2, NULL, globalWorkSize, localWorkSize, CL_CHECK(clEnqueueNDRangeKernel(q, krnl_debayer, 2, NULL, globalWorkSize, localWorkSize,
0, 0, &debayer_event)); 0, 0, &debayer_event));
#else #else
@ -292,35 +294,38 @@ void set_exposure_target(CameraState *c, const uint8_t *pix_ptr, int x_start, in
const CameraBuf *b = &c->buf; const CameraBuf *b = &c->buf;
uint32_t lum_binning[256] = {0}; uint32_t lum_binning[256] = {0};
unsigned int lum_total = 0;
for (int y = y_start; y < y_end; y += y_skip) { for (int y = y_start; y < y_end; y += y_skip) {
for (int x = x_start; x < x_end; x += x_skip) { for (int x = x_start; x < x_end; x += x_skip) {
uint8_t lum = pix_ptr[(y * b->rgb_width) + x]; uint8_t lum = pix_ptr[(y * b->rgb_width) + x];
#ifdef QCOM2
if (lum < 80 && lum_binning[lum] > HISTO_CEIL_K * (y_end - y_start) * (x_end - x_start) / x_skip / y_skip / 256) {
continue;
}
#endif
lum_binning[lum]++; lum_binning[lum]++;
lum_total += 1;
} }
} }
unsigned int lum_total = (y_end - y_start) * (x_end - x_start) / x_skip / y_skip;
unsigned int lum_cur = 0; unsigned int lum_cur = 0;
int lum_med = 0; int lum_med = 0;
int lum_med_alt = 0; int lum_med_alt = 0;
for (lum_med=255; lum_med>=0; lum_med--) { for (lum_med=255; lum_med>=0; lum_med--) {
lum_cur += lum_binning[lum_med]; lum_cur += lum_binning[lum_med];
#ifdef QCOM2 #ifdef QCOM2
bool reach_hlc_perc = false; int lum_med_tmp = 0;
if (c->camera_num == 0) { // wide int hb = HLC_THRESH;
reach_hlc_perc = lum_cur > 2*lum_total / (3*HLC_A); if (lum_cur > 0 && lum_med > hb) {
} else { lum_med_tmp = 4 * (lum_med - hb) + 100;
reach_hlc_perc = lum_cur > lum_total / HLC_A;
}
if (reach_hlc_perc && lum_med > HLC_THRESH) {
lum_med_alt = 86;
} }
lum_med_alt = lum_med_alt>lum_med_tmp?lum_med_alt:lum_med_tmp;
#endif #endif
if (lum_cur >= lum_total / 2) { if (lum_cur >= lum_total / 2) {
break; break;
} }
} }
lum_med = lum_med_alt>lum_med?lum_med_alt:lum_med; lum_med = lum_med_alt>0 ? lum_med + lum_med/32*lum_cur*(lum_med_alt - lum_med)/lum_total/2:lum_med;
camera_autoexposure(c, lum_med / 256.0); camera_autoexposure(c, lum_med / 256.0);
} }

25
selfdrive/camerad/cameras/camera_qcom2.cc
Unescape View File

@ -46,10 +46,9 @@ CameraInfo cameras_supported[CAMERA_ID_MAX] = {
}, },
}; };
float sensor_analog_gains[ANALOG_GAIN_MAX_IDX] = {1.0/8.0, 2.0/8.0, 2.0/7.0, 3.0/7.0, float sensor_analog_gains[ANALOG_GAIN_MAX_IDX] = {3.0/6.0, 4.0/6.0, 4.0/5.0, 5.0/5.0,
3.0/6.0, 4.0/6.0, 4.0/5.0, 5.0/5.0,
5.0/4.0, 6.0/4.0, 6.0/3.0, 7.0/3.0, 5.0/4.0, 6.0/4.0, 6.0/3.0, 7.0/3.0,
7.0/2.0, 8.0/2.0, 8.0/1.0}; 7.0/2.0, 8.0/2.0};
// ************** low level camera helpers **************** // ************** low level camera helpers ****************
@ -992,7 +991,7 @@ void set_exposure_time_bounds(CameraState *s) {
void switch_conversion_gain(CameraState *s) { void switch_conversion_gain(CameraState *s) {
if (!s->dc_gain_enabled) { if (!s->dc_gain_enabled) {
s->dc_gain_enabled = true; s->dc_gain_enabled = true;
s->analog_gain -= 5; s->analog_gain -= 4;
} else { } else {
s->dc_gain_enabled = false; s->dc_gain_enabled = false;
s->analog_gain += 4; s->analog_gain += 4;
@ -1001,9 +1000,9 @@ void switch_conversion_gain(CameraState *s) {
static void set_camera_exposure(CameraState *s, float grey_frac) { static void set_camera_exposure(CameraState *s, float grey_frac) {
// TODO: get stats from sensor? // TODO: get stats from sensor?
float target_grey = 0.3 - (s->analog_gain / 105.0); float target_grey = 0.4 - ((float)(s->analog_gain + 4*s->dc_gain_enabled) / 48.0f);
float exposure_factor = 1 + 30 * pow((target_grey - grey_frac), 3); float exposure_factor = 1 + 30 * pow((target_grey - grey_frac), 3);
exposure_factor = std::max(exposure_factor, 0.55f); exposure_factor = std::max(exposure_factor, 0.4f);
if (s->camera_num != 1) { if (s->camera_num != 1) {
s->ef_filtered = (1 - EF_LOWPASS_K) * s->ef_filtered + EF_LOWPASS_K * exposure_factor; s->ef_filtered = (1 - EF_LOWPASS_K) * s->ef_filtered + EF_LOWPASS_K * exposure_factor;
@ -1024,7 +1023,7 @@ static void set_camera_exposure(CameraState *s, float grey_frac) {
if (s->analog_gain < ANALOG_GAIN_MAX_IDX - 1) { if (s->analog_gain < ANALOG_GAIN_MAX_IDX - 1) {
s->exposure_time = EXPOSURE_TIME_MAX / 2; s->exposure_time = EXPOSURE_TIME_MAX / 2;
s->analog_gain += 1; s->analog_gain += 1;
if (!s->dc_gain_enabled && sensor_analog_gains[s->analog_gain] == 1.0) { // switch to HCG at iso 800 if (!s->dc_gain_enabled && sensor_analog_gains[s->analog_gain] >= 4.0) { // switch to HCG
switch_conversion_gain(s); switch_conversion_gain(s);
} }
set_exposure_time_bounds(s); set_exposure_time_bounds(s);
@ -1035,7 +1034,7 @@ static void set_camera_exposure(CameraState *s, float grey_frac) {
if (s->analog_gain > 0) { if (s->analog_gain > 0) {
s->exposure_time = std::max(EXPOSURE_TIME_MIN * 2, (int)(s->exposure_time / (sensor_analog_gains[s->analog_gain-1] / sensor_analog_gains[s->analog_gain]))); s->exposure_time = std::max(EXPOSURE_TIME_MIN * 2, (int)(s->exposure_time / (sensor_analog_gains[s->analog_gain-1] / sensor_analog_gains[s->analog_gain])));
s->analog_gain -= 1; s->analog_gain -= 1;
if (s->dc_gain_enabled && sensor_analog_gains[s->analog_gain] == 0.25) { // switch back to LCG at iso 200 if (s->dc_gain_enabled && sensor_analog_gains[s->analog_gain] <= 1.25) { // switch back to LCG
switch_conversion_gain(s); switch_conversion_gain(s);
} }
set_exposure_time_bounds(s); set_exposure_time_bounds(s);
@ -1045,20 +1044,22 @@ static void set_camera_exposure(CameraState *s, float grey_frac) {
} }
// set up config // set up config
uint16_t AG = s->analog_gain; uint16_t AG = s->analog_gain + 4;
AG = AG * 4096 + AG * 256 + AG * 16 + AG; AG = 0xFF00 + AG * 16 + AG;
s->analog_gain_frac = sensor_analog_gains[s->analog_gain]; s->analog_gain_frac = sensor_analog_gains[s->analog_gain];
// printf("cam %d, min %d, max %d \n", s->camera_num, s->exposure_time_min, s->exposure_time_max); // printf("cam %d, min %d, max %d \n", s->camera_num, s->exposure_time_min, s->exposure_time_max);
// printf("cam %d, set AG to 0x%X, S to %d, dc %d \n", s->camera_num, AG, s->exposure_time, s->dc_gain_enabled); // printf("cam %d, set AG to 0x%X, S to %d, dc %d \n", s->camera_num, AG, s->exposure_time, s->dc_gain_enabled);
struct i2c_random_wr_payload exp_reg_array[] = {{0x3366, AG}, // analog gain struct i2c_random_wr_payload exp_reg_array[] = {
{0x3366, AG}, // analog gain
{0x3362, (uint16_t)(s->dc_gain_enabled?0x1:0x0)}, // DC_GAIN {0x3362, (uint16_t)(s->dc_gain_enabled?0x1:0x0)}, // DC_GAIN
{0x305A, 0x00D8}, // red gain {0x305A, 0x00D8}, // red gain
{0x3058, 0x011B}, // blue gain {0x3058, 0x011B}, // blue gain
{0x3056, 0x009A}, // g1 gain {0x3056, 0x009A}, // g1 gain
{0x305C, 0x009A}, // g2 gain {0x305C, 0x009A}, // g2 gain
{0x3012, (uint16_t)s->exposure_time}}; // integ time {0x3012, (uint16_t)s->exposure_time}, // integ time
};
//{0x301A, 0x091C}}; // reset //{0x301A, 0x091C}}; // reset
sensors_i2c(s, exp_reg_array, sizeof(exp_reg_array)/sizeof(struct i2c_random_wr_payload), sensors_i2c(s, exp_reg_array, sizeof(exp_reg_array)/sizeof(struct i2c_random_wr_payload),
CAM_SENSOR_PACKET_OPCODE_SENSOR_CONFIG); CAM_SENSOR_PACKET_OPCODE_SENSOR_CONFIG);

9
selfdrive/camerad/cameras/camera_qcom2.h
Unescape View File

@ -9,12 +9,13 @@
#define FRAME_BUF_COUNT 4 #define FRAME_BUF_COUNT 4
#define ANALOG_GAIN_MAX_IDX 15 // 0xF is bypass #define ANALOG_GAIN_MAX_IDX 10 // 0xF is bypass
#define EXPOSURE_TIME_MIN 8 // min time limited by HDR exp factor #define EXPOSURE_TIME_MIN 2 // with HDR, fastest ss
#define EXPOSURE_TIME_MAX 1132 // with HDR, no slower than 1/25 sec (1416 lines) #define EXPOSURE_TIME_MAX 1757 // with HDR, slowest ss
#define HLC_THRESH 240 #define HLC_THRESH 222
#define HLC_A 80 #define HLC_A 80
#define HISTO_CEIL_K 5
#define EF_LOWPASS_K 0.35 #define EF_LOWPASS_K 0.35

227
selfdrive/camerad/cameras/real_debayer.cl
Unescape View File

@ -1,73 +1,124 @@
const __constant float3 color_correction[3] = { #pragma OPENCL EXTENSION cl_khr_fp16 : enable
const half black_level = 42.0;
const __constant half3 color_correction[3] = {
// post wb CCM // post wb CCM
(float3)(1.44602146, -0.24727126, -0.0403062), (half3)(1.25985206, -0.378923, -0.21356857),
(float3)(-0.37658179, 1.26329038, -0.45978396), (half3)(-0.11117607, 1.3962182, -0.46342976),
(float3)(-0.06943967, -0.01601912, 1.50009016), (half3)(-0.21523926, -0.13449348, 1.47665819),
};
const __constant half3 base_ccm[3] = {
(half3)(1,0,0),
(half3)(0,1,0),
(half3)(0,0,1),
}; };
float3 color_correct(float r, float g, float b) {
float3 ret = (0,0,0);
ret += r * color_correction[0]; // tone mapping params
ret += g * color_correction[1]; const half cpk = 0.75;
ret += b * color_correction[2]; const half cpb = 0.125;
ret = max(0.0, min(1.0, ret)); const half cpxk = 0.01;
const half cpxb = 0.01;
half mf(half x, half cp) {
half rk = 8.6 - 66*cp;
if (x > cp) {
return (rk * (x-cp) * (1-(cpk*cp+cpb)) * (1+1/(rk*(1-cp))) / (1+rk*(x-cp))) + cpk*cp + cpb;
} else if (x < cp) {
return (rk * (x-cp) * (cpk*cp+cpb) * (1+1/(rk*cp)) / (1-rk*(x-cp))) + cpk*cp + cpb;
} else {
return x;
}
}
half3 color_correct(half3 rgb, int ggain) {
half3 ret = (0,0,0);
half cpx = clamp(0.03h, 0.1h, cpxb + cpxk * min(10, ggain));
rgb.x = mf(rgb.x, cpx);
rgb.y = mf(rgb.y, cpx);
rgb.z = mf(rgb.z, cpx);
ret += (half)rgb.x * (color_correction[0]);
ret += (half)rgb.y * (color_correction[1]);
ret += (half)rgb.z * (color_correction[2]);
ret = clamp(0.0h, 255.0h, ret*255.0h);
return ret; return ret;
} }
uint int_from_10(const uchar * source, uint start, uint offset) { half val_from_10(const uchar * source, int gx, int gy) {
// source: source // parse 10bit
// start: starting address of 0 int start = gy * FRAME_STRIDE + (5 * (gx / 4));
// offset: 0 - 3 int offset = gx % 4;
uint major = (uint)source[start + offset] << 2; uint major = (uint)source[start + offset] << 2;
uint minor = (source[start + 4] >> (2 * offset)) & 3; uint minor = (source[start + 4] >> (2 * offset)) & 3;
return major + minor; half pv = (half)(major + minor);
}
float to_normal(uint x, int gx, int gy) { // normalize
float pv = (float)(x); pv = max(0.0h, pv - black_level);
const float black_level = 42.0; pv *= 0.00101833h; // /= (1024.0f - black_level);
pv = max(0.0, pv - black_level);
pv /= (1024.0f - black_level); // correct vignetting
if (CAM_NUM == 1) { // fcamera if (CAM_NUM == 1) { // fcamera
gx = (gx - RGB_WIDTH/2); gx = (gx - RGB_WIDTH/2);
gy = (gy - RGB_HEIGHT/2); gy = (gy - RGB_HEIGHT/2);
float r = pow(gx*gx + gy*gy, 0.825); float r = gx*gx + gy*gy;
float s = 1 / (1-0.00000733*r); half s;
if (r < 62500) {
s = (half)(1.0f + 0.0000008f*r);
} else if (r < 490000) {
s = (half)(0.9625f + 0.0000014f*r);
} else if (r < 1102500) {
s = (half)(1.26434f + 0.0000000000016f*r*r);
} else {
s = (half)(0.53503625f + 0.0000000000022f*r*r);
}
pv = s * pv; pv = s * pv;
} }
pv = 20*pv / (1.0f + 20*pv); // reinhard
pv = clamp(0.0h, 1.0h, pv);
return pv; return pv;
} }
half fabs_diff(half x, half y) {
return fabs(x-y);
}
half phi(half x) {
// detection funtion
return 2 - x;
// if (x > 1) {
// return 1 / x;
// } else {
// return 2 - x;
// }
}
__kernel void debayer10(const __global uchar * in, __kernel void debayer10(const __global uchar * in,
__global uchar * out, __global uchar * out,
__local float * cached __local half * cached,
uint ggain
) )
{ {
const int x_global = get_global_id(0); const int x_global = get_global_id(0);
const int y_global = get_global_id(1); const int y_global = get_global_id(1);
// const int globalOffset = ;
const int localRowLen = 2 + get_local_size(0); // 2 padding const int localRowLen = 2 + get_local_size(0); // 2 padding
const int x_local = get_local_id(0); const int x_local = get_local_id(0); // 0-15
const int y_local = get_local_id(1); const int y_local = get_local_id(1); // 0-15
const int localOffset = (y_local + 1) * localRowLen + x_local + 1; // max 18x18-1
const int localOffset = (y_local + 1) * localRowLen + x_local + 1; int out_idx = 3 * x_global + 3 * y_global * RGB_WIDTH;
// cache local pixels first half pv = val_from_10(in, x_global, y_global);
// saves memory access and avoids repeated normalization cached[localOffset] = pv;
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, x_global, y_global);
// edges // don't care
if (x_global < 1 || x_global > RGB_WIDTH - 2 || y_global < 1 || y_global > RGB_HEIGHT - 2) { if (x_global < 1 || x_global >= RGB_WIDTH - 1 || y_global < 1 || y_global >= RGB_HEIGHT - 1) {
barrier(CLK_LOCAL_MEM_FENCE);
return; return;
} else { }
// cache padding
int localColOffset = -1; int localColOffset = -1;
int globalColOffset = -1; int globalColOffset = -1;
@ -75,61 +126,91 @@ __kernel void debayer10(const __global uchar * in,
if (x_local < 1) { if (x_local < 1) {
localColOffset = x_local; localColOffset = x_local;
globalColOffset = -1; 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), x_global, y_global); cached[(y_local + 1) * localRowLen + x_local] = val_from_10(in, x_global-1, y_global);
} else if (x_local >= get_local_size(0) - 1) { } else if (x_local >= get_local_size(0) - 1) {
localColOffset = x_local + 2; localColOffset = x_local + 2;
globalColOffset = 1; globalColOffset = 1;
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); cached[localOffset + 1] = val_from_10(in, x_global+1, y_global);
} }
if (y_local < 1) { if (y_local < 1) {
cached[y_local * localRowLen + x_local + 1] = to_normal(int_from_10(in, globalStart_10 - FRAME_STRIDE, offset_10), x_global, y_global); cached[y_local * localRowLen + x_local + 1] = val_from_10(in, x_global, y_global-1);
if (localColOffset != -1) { 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), x_global, y_global); cached[y_local * localRowLen + localColOffset] = val_from_10(in, x_global+globalColOffset, y_global-1);
} }
} else if (y_local >= get_local_size(1) - 1) { } 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), x_global, y_global); cached[(y_local + 2) * localRowLen + x_local + 1] = val_from_10(in, x_global, y_global+1);
if (localColOffset != -1) { 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), x_global, y_global); cached[(y_local + 2) * localRowLen + localColOffset] = val_from_10(in, x_global+globalColOffset, y_global+1);
} }
} }
// sync // sync
barrier(CLK_LOCAL_MEM_FENCE); barrier(CLK_LOCAL_MEM_FENCE);
// perform debayer half d1 = cached[localOffset - localRowLen - 1];
float r; half d2 = cached[localOffset - localRowLen + 1];
float g; half d3 = cached[localOffset + localRowLen - 1];
float b; half d4 = cached[localOffset + localRowLen + 1];
half n1 = cached[localOffset - localRowLen];
half n2 = cached[localOffset + 1];
half n3 = cached[localOffset + localRowLen];
half n4 = cached[localOffset - 1];
half3 rgb;
// a simplified version of https://opensignalprocessingjournal.com/contents/volumes/V6/TOSIGPJ-6-1/TOSIGPJ-6-1.pdf
if (x_global % 2 == 0) { if (x_global % 2 == 0) {
if (y_global % 2 == 0) { // G1 if (y_global % 2 == 0) {
r = (cached[localOffset - 1] + cached[localOffset + 1]) / 2.0f; rgb.y = pv; // G1(R)
g = (cached[localOffset] + cached[localOffset + localRowLen + 1]) / 2.0f; half k1 = phi(fabs_diff(d1, pv) + fabs_diff(d2, pv));
b = (cached[localOffset - localRowLen] + cached[localOffset + localRowLen]) / 2.0f; half k2 = phi(fabs_diff(d2, pv) + fabs_diff(d4, pv));
} else { // B half k3 = phi(fabs_diff(d3, pv) + fabs_diff(d4, pv));
r = (cached[localOffset - localRowLen - 1] + cached[localOffset - localRowLen + 1] + cached[localOffset + localRowLen - 1] + cached[localOffset + localRowLen + 1]) / 4.0f; half k4 = phi(fabs_diff(d1, pv) + fabs_diff(d3, pv));
g = (cached[localOffset - localRowLen] + cached[localOffset + localRowLen] + cached[localOffset - 1] + cached[localOffset + 1]) / 4.0f; // R_G1
b = cached[localOffset]; rgb.x = (k2*n2+k4*n4)/(k2+k4);
// B_G1
rgb.z = (k1*n1+k3*n3)/(k1+k3);
} else {
rgb.z = pv; // B
half k1 = phi(fabs_diff(d1, d3) + fabs_diff(d2, d4));
half k2 = phi(fabs_diff(n1, n4) + fabs_diff(n2, n3));
half k3 = phi(fabs_diff(d1, d2) + fabs_diff(d3, d4));
half k4 = phi(fabs_diff(n1, n2) + fabs_diff(n3, n4));
// G_B
rgb.y = (k1*(n1+n3)*0.5+k3*(n2+n4)*0.5)/(k1+k3);
// R_B
rgb.x = (k2*(d2+d3)*0.5+k4*(d1+d4)*0.5)/(k2+k4);
} }
} else { } else {
if (y_global % 2 == 0) { // R if (y_global % 2 == 0) {
r = cached[localOffset]; rgb.x = pv; // R
g = (cached[localOffset - localRowLen] + cached[localOffset + localRowLen] + cached[localOffset - 1] + cached[localOffset + 1]) / 4.0f; half k1 = phi(fabs_diff(d1, d3) + fabs_diff(d2, d4));
b = (cached[localOffset - localRowLen - 1] + cached[localOffset - localRowLen + 1] + cached[localOffset + localRowLen - 1] + cached[localOffset + localRowLen + 1]) / 4.0f; half k2 = phi(fabs_diff(n1, n4) + fabs_diff(n2, n3));
} else { // G2 half k3 = phi(fabs_diff(d1, d2) + fabs_diff(d3, d4));
r = (cached[localOffset - localRowLen] + cached[localOffset + localRowLen]) / 2.0f; half k4 = phi(fabs_diff(n1, n2) + fabs_diff(n3, n4));
g = (cached[localOffset] + cached[localOffset - localRowLen - 1]) / 2.0f; // G_R
b = (cached[localOffset - 1] + cached[localOffset + 1]) / 2.0f; rgb.y = (k1*(n1+n3)*0.5+k3*(n2+n4)*0.5)/(k1+k3);
// B_R
rgb.z = (k2*(d2+d3)*0.5+k4*(d1+d4)*0.5)/(k2+k4);
} else {
rgb.y = pv; // G2(B)
half k1 = phi(fabs_diff(d1, pv) + fabs_diff(d2, pv));
half k2 = phi(fabs_diff(d2, pv) + fabs_diff(d4, pv));
half k3 = phi(fabs_diff(d3, pv) + fabs_diff(d4, pv));
half k4 = phi(fabs_diff(d1, pv) + fabs_diff(d3, pv));
// R_G2
rgb.x = (k1*n1+k3*n3)/(k1+k3);
// B_G2
rgb.z = (k2*n2+k4*n4)/(k2+k4);
} }
} }
float3 rgb = color_correct(r, g, b); rgb = clamp(0.0h, 1.0h, rgb);
// rgb = srgb_gamma(rgb); rgb = color_correct(rgb, (int)ggain);
out[out_idx + 0] = (uchar)(rgb.z);
out[out_idx + 1] = (uchar)(rgb.y);
out[out_idx + 2] = (uchar)(rgb.x);
// BGR output
out[3 * x_global + 3 * y_global * RGB_WIDTH + 0] = (uchar)(255.0f * rgb.z);
out[3 * x_global + 3 * y_global * RGB_WIDTH + 1] = (uchar)(255.0f * rgb.y);
out[3 * x_global + 3 * y_global * RGB_WIDTH + 2] = (uchar)(255.0f * rgb.x);
}
} }

638
selfdrive/camerad/cameras/sensor2_i2c.h
Unescape View File

@ -1,10 +1,8 @@
struct i2c_random_wr_payload start_reg_array[] = {{0x301a, 0x91c}}; struct i2c_random_wr_payload start_reg_array[] = {{0x301a, 0x91c}};
//struct i2c_random_wr_payload stop_reg_array[] = {{0x301a, 0x10d8}};
struct i2c_random_wr_payload stop_reg_array[] = {{0x301a, 0x918}};; struct i2c_random_wr_payload stop_reg_array[] = {{0x301a, 0x918}};;
struct i2c_random_wr_payload init_array_ar0231[] = { struct i2c_random_wr_payload init_array_ar0231[] = {
{0x301A, 0x0018}, // RESET_REGISTER {0x301A, 0x0018}, // RESET_REGISTER
//
{0x3092, 0x0C24}, // ROW_NOISE_CONTROL {0x3092, 0x0C24}, // ROW_NOISE_CONTROL
{0x337A, 0x0C80}, // DBLC_SCALE0 {0x337A, 0x0C80}, // DBLC_SCALE0
{0x3520, 0x1288}, // RESERVED_MFR_3520 {0x3520, 0x1288}, // RESERVED_MFR_3520
@ -32,222 +30,6 @@ struct i2c_random_wr_payload init_array_ar0231[] = {
{0x329A, 0x0FA0}, // RESERVED_MFR_329A {0x329A, 0x0FA0}, // RESERVED_MFR_329A
{0x329C, 0x0FA0}, // RESERVED_MFR_329C {0x329C, 0x0FA0}, // RESERVED_MFR_329C
{0x329E, 0x0FA0}, // RESERVED_MFR_329E {0x329E, 0x0FA0}, // RESERVED_MFR_329E
{0x2512, 0x8000}, // SEQ_CTRL_PORT
{0x2510, 0x0905}, // SEQ_DATA_PORT
{0x2510, 0x3350}, // SEQ_DATA_PORT
{0x2510, 0x2004}, // SEQ_DATA_PORT
{0x2510, 0x1460}, // SEQ_DATA_PORT
{0x2510, 0x1578}, // SEQ_DATA_PORT
{0x2510, 0x0901}, // SEQ_DATA_PORT
{0x2510, 0x7B24}, // SEQ_DATA_PORT
{0x2510, 0xFF24}, // SEQ_DATA_PORT
{0x2510, 0xFF24}, // SEQ_DATA_PORT
{0x2510, 0xEA24}, // SEQ_DATA_PORT
{0x2510, 0x1022}, // SEQ_DATA_PORT
{0x2510, 0x2410}, // SEQ_DATA_PORT
{0x2510, 0x155A}, // SEQ_DATA_PORT
{0x2510, 0x0901}, // SEQ_DATA_PORT
{0x2510, 0x1400}, // SEQ_DATA_PORT
{0x2510, 0x24FF}, // SEQ_DATA_PORT
{0x2510, 0x24FF}, // SEQ_DATA_PORT
{0x2510, 0x24EA}, // SEQ_DATA_PORT
{0x2510, 0x2324}, // SEQ_DATA_PORT
{0x2510, 0x647A}, // SEQ_DATA_PORT
{0x2510, 0x2404}, // SEQ_DATA_PORT
{0x2510, 0x052C}, // SEQ_DATA_PORT
{0x2510, 0x400A}, // SEQ_DATA_PORT
{0x2510, 0xFF0A}, // SEQ_DATA_PORT
{0x2510, 0xFF0A}, // SEQ_DATA_PORT
{0x2510, 0x1008}, // SEQ_DATA_PORT
{0x2510, 0x3851}, // SEQ_DATA_PORT
{0x2510, 0x1440}, // SEQ_DATA_PORT
{0x2510, 0x0004}, // SEQ_DATA_PORT
{0x2510, 0x0801}, // SEQ_DATA_PORT
{0x2510, 0x0408}, // SEQ_DATA_PORT
{0x2510, 0x1180}, // SEQ_DATA_PORT
{0x2510, 0x2652}, // SEQ_DATA_PORT
{0x2510, 0x1518}, // SEQ_DATA_PORT
{0x2510, 0x0906}, // SEQ_DATA_PORT
{0x2510, 0x1348}, // SEQ_DATA_PORT
{0x2510, 0x1002}, // SEQ_DATA_PORT
{0x2510, 0x1016}, // SEQ_DATA_PORT
{0x2510, 0x1181}, // SEQ_DATA_PORT
{0x2510, 0x1189}, // SEQ_DATA_PORT
{0x2510, 0x1056}, // SEQ_DATA_PORT
{0x2510, 0x1210}, // SEQ_DATA_PORT
{0x2510, 0x0901}, // SEQ_DATA_PORT
{0x2510, 0x0D09}, // SEQ_DATA_PORT
{0x2510, 0x1413}, // SEQ_DATA_PORT
{0x2510, 0x8809}, // SEQ_DATA_PORT
{0x2510, 0x2B15}, // SEQ_DATA_PORT
{0x2510, 0x8809}, // SEQ_DATA_PORT
{0x2510, 0x0311}, // SEQ_DATA_PORT
{0x2510, 0xD909}, // SEQ_DATA_PORT
{0x2510, 0x1214}, // SEQ_DATA_PORT
{0x2510, 0x4109}, // SEQ_DATA_PORT
{0x2510, 0x0312}, // SEQ_DATA_PORT
{0x2510, 0x1409}, // SEQ_DATA_PORT
{0x2510, 0x0110}, // SEQ_DATA_PORT
{0x2510, 0xD612}, // SEQ_DATA_PORT
{0x2510, 0x1012}, // SEQ_DATA_PORT
{0x2510, 0x1212}, // SEQ_DATA_PORT
{0x2510, 0x1011}, // SEQ_DATA_PORT
{0x2510, 0xDD11}, // SEQ_DATA_PORT
{0x2510, 0xD910}, // SEQ_DATA_PORT
{0x2510, 0x5609}, // SEQ_DATA_PORT
{0x2510, 0x1511}, // SEQ_DATA_PORT
{0x2510, 0xDB09}, // SEQ_DATA_PORT
{0x2510, 0x1511}, // SEQ_DATA_PORT
{0x2510, 0x9B09}, // SEQ_DATA_PORT
{0x2510, 0x0F11}, // SEQ_DATA_PORT
{0x2510, 0xBB12}, // SEQ_DATA_PORT
{0x2510, 0x1A12}, // SEQ_DATA_PORT
{0x2510, 0x1014}, // SEQ_DATA_PORT
{0x2510, 0x6012}, // SEQ_DATA_PORT
{0x2510, 0x5010}, // SEQ_DATA_PORT
{0x2510, 0x7610}, // SEQ_DATA_PORT
{0x2510, 0xE609}, // SEQ_DATA_PORT
{0x2510, 0x0812}, // SEQ_DATA_PORT
{0x2510, 0x4012}, // SEQ_DATA_PORT
{0x2510, 0x6009}, // SEQ_DATA_PORT
{0x2510, 0x290B}, // SEQ_DATA_PORT
{0x2510, 0x0904}, // SEQ_DATA_PORT
{0x2510, 0x1440}, // SEQ_DATA_PORT
{0x2510, 0x0923}, // SEQ_DATA_PORT
{0x2510, 0x15C8}, // SEQ_DATA_PORT
{0x2510, 0x13C8}, // SEQ_DATA_PORT
{0x2510, 0x092C}, // SEQ_DATA_PORT
{0x2510, 0x1588}, // SEQ_DATA_PORT
{0x2510, 0x1388}, // SEQ_DATA_PORT
{0x2510, 0x0C09}, // SEQ_DATA_PORT
{0x2510, 0x0C14}, // SEQ_DATA_PORT
{0x2510, 0x4109}, // SEQ_DATA_PORT
{0x2510, 0x1112}, // SEQ_DATA_PORT
{0x2510, 0x6212}, // SEQ_DATA_PORT
{0x2510, 0x6011}, // SEQ_DATA_PORT
{0x2510, 0xBF11}, // SEQ_DATA_PORT
{0x2510, 0xBB10}, // SEQ_DATA_PORT
{0x2510, 0x6611}, // SEQ_DATA_PORT
{0x2510, 0xFB09}, // SEQ_DATA_PORT
{0x2510, 0x3511}, // SEQ_DATA_PORT
{0x2510, 0xBB12}, // SEQ_DATA_PORT
{0x2510, 0x6312}, // SEQ_DATA_PORT
{0x2510, 0x6014}, // SEQ_DATA_PORT
{0x2510, 0x0015}, // SEQ_DATA_PORT
{0x2510, 0x0011}, // SEQ_DATA_PORT
{0x2510, 0xB812}, // SEQ_DATA_PORT
{0x2510, 0xA012}, // SEQ_DATA_PORT
{0x2510, 0x0010}, // SEQ_DATA_PORT
{0x2510, 0x2610}, // SEQ_DATA_PORT
{0x2510, 0x0013}, // SEQ_DATA_PORT
{0x2510, 0x0011}, // SEQ_DATA_PORT
{0x2510, 0x0008}, // SEQ_DATA_PORT
{0x2510, 0x3053}, // SEQ_DATA_PORT
{0x2510, 0x4215}, // SEQ_DATA_PORT
{0x2510, 0x4013}, // SEQ_DATA_PORT
{0x2510, 0x4010}, // SEQ_DATA_PORT
{0x2510, 0x0210}, // SEQ_DATA_PORT
{0x2510, 0x1611}, // SEQ_DATA_PORT
{0x2510, 0x8111}, // SEQ_DATA_PORT
{0x2510, 0x8910}, // SEQ_DATA_PORT
{0x2510, 0x5612}, // SEQ_DATA_PORT
{0x2510, 0x1009}, // SEQ_DATA_PORT
{0x2510, 0x010D}, // SEQ_DATA_PORT
{0x2510, 0x0815}, // SEQ_DATA_PORT
{0x2510, 0xC015}, // SEQ_DATA_PORT
{0x2510, 0xD013}, // SEQ_DATA_PORT
{0x2510, 0x5009}, // SEQ_DATA_PORT
{0x2510, 0x1313}, // SEQ_DATA_PORT
{0x2510, 0xD009}, // SEQ_DATA_PORT
{0x2510, 0x0215}, // SEQ_DATA_PORT
{0x2510, 0xC015}, // SEQ_DATA_PORT
{0x2510, 0xC813}, // SEQ_DATA_PORT
{0x2510, 0xC009}, // SEQ_DATA_PORT
{0x2510, 0x0515}, // SEQ_DATA_PORT
{0x2510, 0x8813}, // SEQ_DATA_PORT
{0x2510, 0x8009}, // SEQ_DATA_PORT
{0x2510, 0x0213}, // SEQ_DATA_PORT
{0x2510, 0x8809}, // SEQ_DATA_PORT
{0x2510, 0x0411}, // SEQ_DATA_PORT
{0x2510, 0xC909}, // SEQ_DATA_PORT
{0x2510, 0x0814}, // SEQ_DATA_PORT
{0x2510, 0x0109}, // SEQ_DATA_PORT
{0x2510, 0x0B11}, // SEQ_DATA_PORT
{0x2510, 0xD908}, // SEQ_DATA_PORT
{0x2510, 0x1400}, // SEQ_DATA_PORT
{0x2510, 0x091A}, // SEQ_DATA_PORT
{0x2510, 0x1440}, // SEQ_DATA_PORT
{0x2510, 0x0903}, // SEQ_DATA_PORT
{0x2510, 0x1214}, // SEQ_DATA_PORT
{0x2510, 0x0901}, // SEQ_DATA_PORT
{0x2510, 0x10D6}, // SEQ_DATA_PORT
{0x2510, 0x1210}, // SEQ_DATA_PORT
{0x2510, 0x1212}, // SEQ_DATA_PORT
{0x2510, 0x1210}, // SEQ_DATA_PORT
{0x2510, 0x11DD}, // SEQ_DATA_PORT
{0x2510, 0x11D9}, // SEQ_DATA_PORT
{0x2510, 0x1056}, // SEQ_DATA_PORT
{0x2510, 0x0917}, // SEQ_DATA_PORT
{0x2510, 0x11DB}, // SEQ_DATA_PORT
{0x2510, 0x0913}, // SEQ_DATA_PORT
{0x2510, 0x11FB}, // SEQ_DATA_PORT
{0x2510, 0x0905}, // SEQ_DATA_PORT
{0x2510, 0x11BB}, // SEQ_DATA_PORT
{0x2510, 0x121A}, // SEQ_DATA_PORT
{0x2510, 0x1210}, // SEQ_DATA_PORT
{0x2510, 0x1460}, // SEQ_DATA_PORT
{0x2510, 0x1250}, // SEQ_DATA_PORT
{0x2510, 0x1076}, // SEQ_DATA_PORT
{0x2510, 0x10E6}, // SEQ_DATA_PORT
{0x2510, 0x0901}, // SEQ_DATA_PORT
{0x2510, 0x15A8}, // SEQ_DATA_PORT
{0x2510, 0x0901}, // SEQ_DATA_PORT
{0x2510, 0x13A8}, // SEQ_DATA_PORT
{0x2510, 0x1240}, // SEQ_DATA_PORT
{0x2510, 0x1260}, // SEQ_DATA_PORT
{0x2510, 0x0925}, // SEQ_DATA_PORT
{0x2510, 0x13AD}, // SEQ_DATA_PORT
{0x2510, 0x0902}, // SEQ_DATA_PORT
{0x2510, 0x0907}, // SEQ_DATA_PORT
{0x2510, 0x1588}, // SEQ_DATA_PORT
{0x2510, 0x0901}, // SEQ_DATA_PORT
{0x2510, 0x138D}, // SEQ_DATA_PORT
{0x2510, 0x0B09}, // SEQ_DATA_PORT
{0x2510, 0x0914}, // SEQ_DATA_PORT
{0x2510, 0x4009}, // SEQ_DATA_PORT
{0x2510, 0x0B13}, // SEQ_DATA_PORT
{0x2510, 0x8809}, // SEQ_DATA_PORT
{0x2510, 0x1C0C}, // SEQ_DATA_PORT
{0x2510, 0x0920}, // SEQ_DATA_PORT
{0x2510, 0x1262}, // SEQ_DATA_PORT
{0x2510, 0x1260}, // SEQ_DATA_PORT
{0x2510, 0x11BF}, // SEQ_DATA_PORT
{0x2510, 0x11BB}, // SEQ_DATA_PORT
{0x2510, 0x1066}, // SEQ_DATA_PORT
{0x2510, 0x090A}, // SEQ_DATA_PORT
{0x2510, 0x11FB}, // SEQ_DATA_PORT
{0x2510, 0x093B}, // SEQ_DATA_PORT
{0x2510, 0x11BB}, // SEQ_DATA_PORT
{0x2510, 0x1263}, // SEQ_DATA_PORT
{0x2510, 0x1260}, // SEQ_DATA_PORT
{0x2510, 0x1400}, // SEQ_DATA_PORT
{0x2510, 0x1508}, // SEQ_DATA_PORT
{0x2510, 0x11B8}, // SEQ_DATA_PORT
{0x2510, 0x12A0}, // SEQ_DATA_PORT
{0x2510, 0x1200}, // SEQ_DATA_PORT
{0x2510, 0x1026}, // SEQ_DATA_PORT
{0x2510, 0x1000}, // SEQ_DATA_PORT
{0x2510, 0x1300}, // SEQ_DATA_PORT
{0x2510, 0x1100}, // SEQ_DATA_PORT
{0x2510, 0x437A}, // SEQ_DATA_PORT
{0x2510, 0x0609}, // SEQ_DATA_PORT
{0x2510, 0x0B05}, // SEQ_DATA_PORT
{0x2510, 0x0708}, // SEQ_DATA_PORT
{0x2510, 0x4137}, // SEQ_DATA_PORT
{0x2510, 0x502C}, // SEQ_DATA_PORT
{0x2510, 0x2CFE}, // SEQ_DATA_PORT
{0x2510, 0x15FE}, // SEQ_DATA_PORT
{0x2510, 0x0C2C}, // SEQ_DATA_PORT
{0x32E6, 0x00E0}, // RESERVED_MFR_32E6 {0x32E6, 0x00E0}, // RESERVED_MFR_32E6
{0x1008, 0x036F}, // RESERVED_PARAM_1008 {0x1008, 0x036F}, // RESERVED_PARAM_1008
{0x100C, 0x058F}, // RESERVED_PARAM_100C {0x100C, 0x058F}, // RESERVED_PARAM_100C
@ -277,21 +59,6 @@ struct i2c_random_wr_payload init_array_ar0231[] = {
{0x30BA, 0x11F2}, // DIGITAL_CTRL {0x30BA, 0x11F2}, // DIGITAL_CTRL
{0x3044, 0x0400}, // DARK_CONTROL {0x3044, 0x0400}, // DARK_CONTROL
{0x3064, 0x1802}, // SMIA_TEST {0x3064, 0x1802}, // SMIA_TEST
/*{0x3064, 0xCC2}, // STATS_EN
{0x3270, 0x10}, //
{0x3272, 0x30}, //
{0x3274, 0x50}, //
{0x3276, 0x10}, //
{0x3278, 0x30}, //
{0x327A, 0x50}, //
{0x3144, 0x0}, //
{0x3146, 0x0}, //
{0x3244, 0x0}, //
{0x3246, 0x0}, //
{0x3268, 0x0}, //
{0x326A, 0x0}, //
*/
{0x33E0, 0x0C80}, // TEST_ASIL_ROWS {0x33E0, 0x0C80}, // TEST_ASIL_ROWS
{0x3180, 0x0080}, // RESERVED_MFR_3180 {0x3180, 0x0080}, // RESERVED_MFR_3180
{0x33E4, 0x0080}, // RESERVED_MFR_33E4 {0x33E4, 0x0080}, // RESERVED_MFR_33E4
@ -307,7 +74,7 @@ struct i2c_random_wr_payload init_array_ar0231[] = {
{0x3402, 0x0F10}, // X_OUTPUT_CONTROL {0x3402, 0x0F10}, // X_OUTPUT_CONTROL
{0x3404, 0x04B8}, // Y_OUTPUT_CONTROL {0x3404, 0x04B8}, // Y_OUTPUT_CONTROL
{0x3404, 0x0970}, // Y_OUTPUT_CONTROL {0x3404, 0x0970}, // Y_OUTPUT_CONTROL
{0x30BA, 0x11F3}, // DIGITAL_CTRL {0x30BA, 0x11F2}, // DIGITAL_CTRL
// SLAV* MODE // SLAV* MODE
{0x30CE, 0x0120}, {0x30CE, 0x0120},
@ -315,8 +82,8 @@ struct i2c_random_wr_payload init_array_ar0231[] = {
{0x340C, 0x802}, // 2 // 0000 0000 0010 {0x340C, 0x802}, // 2 // 0000 0000 0010
// Readout timing // Readout timing
{0x300C, 0x074B}, // LINE_LENGTH_PCK: min for 3-exposure HDR {0x300C, 0x074B}, // LINE_LENGTH_PCK: min for 2-exposure HDR
{0x300A, 0x06EB}, // FRAME_LENGTH_LINES_ 6EB {0x300A, 0x085E}, // FRAME_LENGTH_LINES_ 6EB
{0x3042, 0x0000}, // EXTRA_DELAY {0x3042, 0x0000}, // EXTRA_DELAY
// Readout Settings // Readout Settings
@ -341,24 +108,23 @@ struct i2c_random_wr_payload init_array_ar0231[] = {
// HDR Settings // HDR Settings
{0x3082, 0x0004}, // OPERATION_MODE_CTRL {0x3082, 0x0004}, // OPERATION_MODE_CTRL
{0x3238, 0x0222}, // EXPOSURE_RATIO {0x3238, 0x0004}, // EXPOSURE_RATIO
{0x3014, 0x098E}, // FINE_INTEGRATION_TIME_ {0x3014, 0x098E}, // FINE_INTEGRATION_TIME_
{0x321E, 0x098E}, // FINE_INTEGRATION_TIME2 {0x321E, 0x098E}, // FINE_INTEGRATION_TIME2
{0x3222, 0x098E}, // FINE_INTEGRATION_TIME3
{0x3226, 0x098E}, // FINE_INTEGRATION_TIME4, 098E?
{0x30B0, 0x0800}, // DIGITAL_TEST {0x30B0, 0x0800}, // DIGITAL_TEST
{0x32EA, 0x3C0E}, // RESERVED_MFR_32EA {0x32EA, 0x3C0E}, // RESERVED_MFR_32EA
{0x32EC, 0x72A1}, // RESERVED_MFR_32EC {0x32EC, 0x72A1}, // RESERVED_MFR_32EC
{0x31D0, 0x0000}, // COMPANDING, no good in 10 bit? {0x31D0, 0x0000}, // COMPANDING, no good in 10 bit?
{0x33DA, 0x0000}, // COMPANDING {0x33DA, 0x0000}, // COMPANDING
{0x3362, 0x0000}, // DC GAIN {0x3370, 0x03B1}, // DBLC
{0x3370, 0x0231}, // DBLC {0x31E0, 0x0001}, // PDC
{0x318E, 0x0200}, // PRE_HDR_GAIN_EN {0x318E, 0x0200}, // PRE_HDR_GAIN_EN
// Initial Gains // Initial Gains
{0x3022, 0x01}, // GROUPED_PARAMETER_HOLD_ {0x3022, 0x01}, // GROUPED_PARAMETER_HOLD_
{0x3366, 0x5555}, // ANALOG_GAIN {0x3366, 0x5555}, // ANALOG_GAIN
{0x3060, 0xBBBB}, // ANALOG_COLOR_GAIN {0x3060, 0x3333}, // ANALOG_COLOR_GAIN
{0x3362, 0x0000}, // DC GAIN
{0x305A, 0x00D8}, // RED_GAIN {0x305A, 0x00D8}, // RED_GAIN
{0x3058, 0x011B}, // BLUE_GAIN {0x3058, 0x011B}, // BLUE_GAIN
{0x3056, 0x009A}, // GREEN1_GAIN {0x3056, 0x009A}, // GREEN1_GAIN
@ -367,393 +133,5 @@ struct i2c_random_wr_payload init_array_ar0231[] = {
// Initial Integration Time // Initial Integration Time
{0x3012, 0x256}, {0x3012, 0x256},
};
struct i2c_random_wr_payload poke_array_ov7750[] = {
{0x3208, 0x0}, {0x380e, 0x1a}, {0x380f, 0xf0}, {0x3500, 0x0}, {0x3501, 0x0}, {0x3502, 0x10}, {0x350a, 0x0}, {0x350b, 0x10}, {0x3208, 0x10}, {0x3208, 0xa0},
//{0x3208, 0x0}, {0x380e, 0x1a}, {0x380f, 0xf0}, {0x3500, 0x0}, {0x3501, 0x0}, {0x3502, 0x10}, {0x350a, 0x0}, {0x350b, 0x10}, {0x3208, 0x10}, {0x3208, 0xa0},
};
struct i2c_random_wr_payload preinit_array_ov7750[] = {
{0x103, 0x1},
{0x303b, 0x2},
{0x302b, 0x80},
};
struct i2c_random_wr_payload init_array_ov7750[] = {
// 2nd batch
{0x3005, 0x0},
{0x3012, 0xc0},
{0x3013, 0xd2},
{0x3014, 0x4},
{0x3016, 0xf0},
{0x3017, 0xf0},
{0x3018, 0xf0},
{0x301a, 0xf0},
{0x301b, 0xf0},
{0x301c, 0xf0},
{0x3023, 0x5},
{0x3037, 0xf0},
{0x3098, 0x4},
{0x3099, 0x28},
{0x309a, 0x5},
{0x309b, 0x4},
{0x30b0, 0xa},
{0x30b1, 0x1},
{0x30b3, 0x64},
{0x30b4, 0x3},
{0x30b5, 0x5},
{0x3106, 0xda},
{0x3500, 0x0},
{0x3501, 0x1f},
{0x3502, 0x80},
{0x3503, 0x7},
{0x3509, 0x10},
{0x350b, 0x10},
{0x3600, 0x1c},
{0x3602, 0x62},
{0x3620, 0xb7},
{0x3622, 0x4},
{0x3626, 0x21},
{0x3627, 0x30},
{0x3630, 0x44},
{0x3631, 0x35},
{0x3634, 0x60},
{0x3636, 0x0},
{0x3662, 0x1},
{0x3663, 0x70},
{0x3664, 0xf0},
{0x3666, 0xa},
{0x3669, 0x1a},
{0x366a, 0x0},
{0x366b, 0x50},
{0x3673, 0x1},
{0x3674, 0xff},
{0x3675, 0x3},
{0x3705, 0xc1},
{0x3709, 0x40},
{0x373c, 0x8},
{0x3742, 0x0},
{0x3757, 0xb3},
{0x3788, 0x0},
{0x37a8, 0x1},
{0x37a9, 0xc0},
{0x3800, 0x0},
{0x3801, 0x4},
{0x3802, 0x0},
{0x3803, 0x4},
{0x3804, 0x2},
{0x3805, 0x8b},
{0x3806, 0x1},
{0x3807, 0xeb},
{0x3808, 0x2},
{0x3809, 0x80},
{0x380a, 0x1},
{0x380b, 0xe0},
{0x380c, 0x3},
{0x380d, 0xa0},
{0x380e, 0x6},
{0x380f, 0xbc},
{0x3810, 0x0},
{0x3811, 0x4},
{0x3812, 0x0},
{0x3813, 0x5},
{0x3814, 0x11},
{0x3815, 0x11},
{0x3820, 0x40},
{0x3821, 0x0},
{0x382f, 0xe},
{0x3832, 0x0},
{0x3833, 0x5},
{0x3834, 0x0},
{0x3835, 0xc},
{0x3837, 0x0},
{0x3b80, 0x0},
{0x3b81, 0xa5},
{0x3b82, 0x10},
{0x3b83, 0x0},
{0x3b84, 0x8},
{0x3b85, 0x0},
{0x3b86, 0x1},
{0x3b87, 0x0},
{0x3b88, 0x0},
{0x3b89, 0x0},
{0x3b8a, 0x0},
{0x3b8b, 0x5},
{0x3b8c, 0x0},
{0x3b8d, 0x0},
{0x3b8e, 0x0},
{0x3b8f, 0x1a},
{0x3b94, 0x5},
{0x3b95, 0xf2},
{0x3b96, 0x40},
{0x3c00, 0x89},
{0x3c01, 0x63},
{0x3c02, 0x1},
{0x3c03, 0x0},
{0x3c04, 0x0},
{0x3c05, 0x3},
{0x3c06, 0x0},
{0x3c07, 0x6},
{0x3c0c, 0x1},
{0x3c0d, 0xd0},
{0x3c0e, 0x2},
{0x3c0f, 0xa},
{0x4001, 0x42},
{0x4004, 0x4},
{0x4005, 0x0},
{0x404e, 0x1},
{0x4300, 0xff},
{0x4301, 0x0},
{0x4315, 0x0},
{0x4501, 0x48},
{0x4600, 0x0},
{0x4601, 0x4e},
{0x4801, 0xf},
{0x4806, 0xf},
{0x4819, 0xaa},
{0x4823, 0x3e},
{0x4837, 0x19},
{0x4a0d, 0x0},
{0x4a47, 0x7f},
{0x4a49, 0xf0},
{0x4a4b, 0x30},
{0x5000, 0x85},
{0x5001, 0x80},
};
struct i2c_random_wr_payload init_array_ov8856[] = {
// part 1 184
{0x103, 0x1},
{0x302, 0x3c},
{0x303, 0x1},
{0x31e, 0xc},
{0x3000, 0x0},
{0x300e, 0x0},
{0x3010, 0x0},
{0x3015, 0x84},
{0x3018, 0x72},
{0x3033, 0x24},
{0x3500, 0x0},
{0x3501, 0x4c},
{0x3502, 0xe0},
{0x3503, 0x8},
{0x3505, 0x83},
{0x3508, 0x1},
{0x3509, 0x80},
{0x350c, 0x0},
{0x350d, 0x80},
{0x350e, 0x4},
{0x350f, 0x0},
{0x3510, 0x0},
{0x3511, 0x2},
{0x3512, 0x0},
{0x3600, 0x72},
{0x3601, 0x40},
{0x3602, 0x30},
{0x3610, 0xc5},
{0x3611, 0x58},
{0x3612, 0x5c},
{0x3613, 0x5a},
{0x3614, 0x60},
{0x3628, 0xff},
{0x3629, 0xff},
{0x362a, 0xff},
{0x3633, 0x10},
{0x3634, 0x10},
{0x3635, 0x10},
{0x3636, 0x10},
{0x3663, 0x8},
{0x3669, 0x34},
{0x366e, 0x8},
{0x3706, 0x86},
{0x370b, 0x7e},
{0x3714, 0x27},
{0x3730, 0x12},
{0x3733, 0x10},
{0x3764, 0x0},
{0x3765, 0x0},
{0x3769, 0x62},
{0x376a, 0x2a},
{0x376b, 0x3b},
{0x3780, 0x0},
{0x3781, 0x24},
{0x3782, 0x0},
{0x3783, 0x23},
{0x3798, 0x2f},
{0x37a1, 0x60},
{0x37a8, 0x6a},
{0x37ab, 0x3f},
{0x37c2, 0x14},
{0x37c3, 0xf1},
{0x37c9, 0x80},
{0x37cb, 0x3},
{0x37cc, 0xa},
{0x37cd, 0x16},
{0x37ce, 0x1f},
{0x3800, 0x0},
{0x3801, 0x0},
{0x3802, 0x0},
{0x3803, 0xc},
{0x3804, 0xc},
{0x3805, 0xdf},
{0x3806, 0x9},
{0x3807, 0xa3},
{0x3808, 0x6},
{0x3809, 0x60},
{0x380a, 0x4},
{0x380b, 0xc8},
{0x380c, 0x7},
{0x380d, 0x8c},
{0x380e, 0x9},
{0x380f, 0xb2},
{0x3810, 0x0},
{0x3811, 0x8},
{0x3812, 0x0},
{0x3813, 0x2},
{0x3814, 0x3},
{0x3815, 0x1},
{0x3816, 0x0},
{0x3817, 0x0},
{0x3818, 0x0},
{0x3819, 0x0},
{0x3820, 0x90},
{0x3821, 0x67},
{0x382a, 0x3},
{0x382b, 0x1},
{0x3830, 0x6},
{0x3836, 0x2},
{0x3862, 0x4},
{0x3863, 0x8},
{0x3cc0, 0x33},
{0x3d85, 0x17},
{0x3d8c, 0x73},
{0x3d8d, 0xde},
{0x4001, 0xe0},
{0x4003, 0x40},
{0x4008, 0x0},
{0x4009, 0x5},
{0x400f, 0x80},
{0x4010, 0xf0},
{0x4011, 0xff},
{0x4012, 0x2},
{0x4013, 0x1},
{0x4014, 0x1},
{0x4015, 0x1},
{0x4042, 0x0},
{0x4043, 0x80},
{0x4044, 0x0},
{0x4045, 0x80},
{0x4046, 0x0},
{0x4047, 0x80},
{0x4048, 0x0},
{0x4049, 0x80},
{0x4041, 0x3},
{0x404c, 0x20},
{0x404d, 0x0},
{0x404e, 0x20},
{0x4203, 0x80},
{0x4307, 0x30},
{0x4317, 0x0},
{0x4503, 0x8},
{0x4601, 0x80},
{0x4816, 0x53},
{0x481b, 0x58},
{0x481f, 0x27},
{0x4837, 0x16},
{0x5000, 0x77},
{0x5001, 0xe},
{0x5004, 0x0},
{0x502e, 0x0},
{0x5030, 0x41},
{0x5795, 0x0},
{0x5796, 0x10},
{0x5797, 0x10},
{0x5798, 0x73},
{0x5799, 0x73},
{0x579a, 0x0},
{0x579b, 0x28},
{0x579c, 0x0},
{0x579d, 0x16},
{0x579e, 0x6},
{0x579f, 0x20},
{0x57a0, 0x4},
{0x57a1, 0xa0},
{0x5780, 0x14},
{0x5781, 0xf},
{0x5782, 0x44},
{0x5783, 0x2},
{0x5784, 0x1},
{0x5785, 0x1},
{0x5786, 0x0},
{0x5787, 0x4},
{0x5788, 0x2},
{0x5789, 0xf},
{0x578a, 0xfd},
{0x578b, 0xf5},
{0x578c, 0xf5},
{0x578d, 0x3},
{0x578e, 0x8},
{0x578f, 0xc},
{0x5790, 0x8},
{0x5791, 0x4},
{0x5792, 0x0},
{0x5793, 0x52},
{0x5794, 0xa3},
{0x5a08, 0x2},
{0x5b00, 0x2},
{0x5b01, 0x10},
{0x5b02, 0x3},
{0x5b03, 0xcf},
{0x5b05, 0x6c},
{0x5e00, 0x0},
// part 2 45
{0x3501, 0x9a},
{0x3502, 0x20},
{0x366d, 0x0},
{0x366e, 0x10},
{0x3714, 0x23},
{0x37c2, 0x4},
{0x3800, 0x0},
{0x3801, 0x0},
{0x3802, 0x0},
{0x3803, 0xc},
{0x3804, 0xc},
{0x3805, 0xdf},
{0x3806, 0x9},
{0x3807, 0xa3},
{0x3808, 0xc},
{0x3809, 0xc0},
{0x380a, 0x9},
{0x380b, 0x90},
{0x380c, 0x7},
{0x380d, 0x8c},
{0x380e, 0x9},
{0x380f, 0xb2},
{0x3811, 0x10},
{0x3813, 0x4},
{0x3814, 0x1},
{0x3820, 0xc6},
{0x3821, 0x40},
{0x382a, 0x1},
{0x4009, 0xb},
{0x4601, 0x80},
{0x5003, 0xc8},
{0x5006, 0x0},
{0x5007, 0x0},
{0x5795, 0x2},
{0x5796, 0x20},
{0x5797, 0x20},
{0x5798, 0xd5},
{0x5799, 0xd5},
{0x579b, 0x50},
{0x579d, 0x2c},
{0x579e, 0xc},
{0x579f, 0x40},
{0x57a0, 0x9},
{0x57a1, 0x40},
{0x5e10, 0xfc},
}; };

Write Preview
Loading…
Cancel
Save