mirror of https://github.com/1okko/openpilot.git
camerad: OS HDR (#32112)
* it's something * backup * 16:10 * cleanup * this is fine * close * remove some junk * no heck * disos * real 10 * for some reason this is flipped * 20hz * no return * ae * tear * need curve laster * correct real gains * fix time * cleanup * why the scam * disable for now * 0.7 * hdr * that doesnt work * what * hugeoof * clean up * cleanup * fix regs * welp cant * is this corrent * it is sq * remove * back * stg10bit * back2ten * Revert "remove" This reverts commit 18712ab7e103c12621c929cd0f772ecb9b348247. * 20hz and swb * correct height * 10bit * ui hack for now * slight * perfect * blk64 * ccm * fix page faults * template * set 4x * is this fine * try * this seems to work * Revert "this seems to work" This reverts commit d3c9023d3f14bd9394fed2d6276dba777ed0e606. * needs to be static * close * 64 is optimal * 2 * take * not 1 * offset * whats going on * i have no idea * less resistence * box defs * no * reduce blur artifacts * simplify * fix * fake short is too much for bright * can be subzero * should not use lsvc * no wasted bit * cont no slow * no less than 10bit * it is based * wrong * right * quart * shift * raise noise floor * 4.5/4.7 * same ballpark * int is fine * shane owes me m4a4 * Revert "shane owes me m4a4" This reverts commit b4283fee18efebedae628a6cfd926ff1416dcfe5. * back * Revert "4.5/4.7" This reverts commit e38f96e90cb5370bd378f6b66def9e7e3ed0ce5d. * default * oof * clean up * simpilfy * from sensorinfo * no div * better name * not the wrong one * not anymore relevant * too * not call it debayer * cl headers * arg is 2nd * gone is is_bggr * define * no is hdr * rgb_tmp * p1 * clean up * 4 * cant for * fix somewhre else * const * ap * rects * just set staruc * nnew tmp * hmm --------- Co-authored-by: Comma Device <device@comma.ai>
This commit is contained in:
ZwX1616
GitHub
parent
ba2538c29c
commit
03d1c48017
|
|
@ -26,10 +26,10 @@ 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 "
|
||||
"-DSENSOR_ID=%hu -DVIGNETTING=%d ",
|
||||
ci->frame_width, ci->frame_height, ci->frame_stride, ci->frame_offset,
|
||||
"-DSENSOR_ID=%hu -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, s->camera_num == 1);
|
||||
ci->image_sensor, 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);
|
||||
krnl_ = CL_CHECK_ERR(clCreateKernel(prg_imgproc, "process_raw", &err));
|
||||
|
|
@ -74,7 +74,7 @@ void CameraBuf::init(cl_device_id device_id, cl_context context, CameraState *s,
|
|||
LOGD("allocated %d CL buffers", frame_buf_count);
|
||||
|
||||
rgb_width = ci->frame_width;
|
||||
rgb_height = ci->frame_height;
|
||||
rgb_height = ci->hdr_offset > 0 ? (ci->frame_height - ci->hdr_offset) / 2 : ci->frame_height;
|
||||
|
||||
int nv12_width = VENUS_Y_STRIDE(COLOR_FMT_NV12, rgb_width);
|
||||
int nv12_height = VENUS_Y_SCANLINES(COLOR_FMT_NV12, rgb_height);
|
||||
|
|
|
|||
|
|
@ -78,7 +78,18 @@ __kernel void process_raw(const __global uchar * in, __global uchar * out, int e
|
|||
|
||||
// read offset
|
||||
int start_idx;
|
||||
start_idx = (2 * gid_y - 1) * FRAME_STRIDE + (3 * gid_x - 2) + (FRAME_STRIDE * FRAME_OFFSET);
|
||||
#if BIT_DEPTH == 10
|
||||
bool aligned10;
|
||||
if (gid_x % 2 == 0) {
|
||||
aligned10 = true;
|
||||
start_idx = (2 * gid_y - 1) * FRAME_STRIDE + (5 * gid_x / 2 - 2) + (FRAME_STRIDE * FRAME_OFFSET);
|
||||
} else {
|
||||
aligned10 = false;
|
||||
start_idx = (2 * gid_y - 1) * FRAME_STRIDE + (5 * (gid_x - 1) / 2 + 1) + (FRAME_STRIDE * FRAME_OFFSET);
|
||||
}
|
||||
#else
|
||||
start_idx = (2 * gid_y - 1) * FRAME_STRIDE + (3 * gid_x - 2) + (FRAME_STRIDE * FRAME_OFFSET);
|
||||
#endif
|
||||
|
||||
// read in 4 rows, 8 uchars each
|
||||
uchar8 dat[4];
|
||||
|
|
@ -96,6 +107,16 @@ __kernel void process_raw(const __global uchar * in, __global uchar * out, int e
|
|||
dat[2] = vload8(0, in + start_idx + FRAME_STRIDE*2);
|
||||
// row_after
|
||||
dat[3] = vload8(0, in + start_idx + FRAME_STRIDE*row_after_offset);
|
||||
// need extra bit for 10-bit, 4 rows, 1 uchar each
|
||||
#if BIT_DEPTH == 10
|
||||
uchar extra_dat[4];
|
||||
if (!aligned10) {
|
||||
extra_dat[0] = in[start_idx + FRAME_STRIDE*row_before_offset + 8];
|
||||
extra_dat[1] = in[start_idx + FRAME_STRIDE*1 + 8];
|
||||
extra_dat[2] = in[start_idx + FRAME_STRIDE*2 + 8];
|
||||
extra_dat[3] = in[start_idx + FRAME_STRIDE*row_after_offset + 8];
|
||||
}
|
||||
#endif
|
||||
|
||||
// read odd rows for staggered second exposure
|
||||
#if HDR_OFFSET > 0
|
||||
|
|
@ -104,19 +125,44 @@ __kernel void process_raw(const __global uchar * in, __global uchar * out, int e
|
|||
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);
|
||||
#if BIT_DEPTH == 10
|
||||
uchar short_extra_dat[4];
|
||||
if (!aligned10) {
|
||||
short_extra_dat[0] = in[start_idx + FRAME_STRIDE*(row_before_offset+HDR_OFFSET/2) + FRAME_STRIDE/2 + 8];
|
||||
short_extra_dat[1] = in[start_idx + FRAME_STRIDE*(1+HDR_OFFSET/2) + FRAME_STRIDE/2 + 8];
|
||||
short_extra_dat[2] = in[start_idx + FRAME_STRIDE*(2+HDR_OFFSET/2) + FRAME_STRIDE/2 + 8];
|
||||
short_extra_dat[3] = in[start_idx + FRAME_STRIDE*(row_after_offset+HDR_OFFSET/2) + FRAME_STRIDE/2 + 8];
|
||||
}
|
||||
#endif
|
||||
#endif
|
||||
|
||||
// parse into floats 0.0-1.0
|
||||
float4 v_rows[4];
|
||||
// no HDR here
|
||||
int4 parsed = parse_12bit(dat[0]);
|
||||
v_rows[ROW_READ_ORDER[0]] = normalize_pv(parsed, vignette_factor);
|
||||
parsed = parse_12bit(dat[1]);
|
||||
v_rows[ROW_READ_ORDER[1]] = normalize_pv(parsed, vignette_factor);
|
||||
parsed = parse_12bit(dat[2]);
|
||||
v_rows[ROW_READ_ORDER[2]] = normalize_pv(parsed, vignette_factor);
|
||||
parsed = parse_12bit(dat[3]);
|
||||
v_rows[ROW_READ_ORDER[3]] = normalize_pv(parsed, vignette_factor);
|
||||
#if BIT_DEPTH == 10
|
||||
// for now it's always HDR
|
||||
int4 parsed = parse_10bit(dat[0], extra_dat[0], aligned10);
|
||||
int4 short_parsed = parse_10bit(short_dat[0], short_extra_dat[0], aligned10);
|
||||
v_rows[ROW_READ_ORDER[0]] = normalize_pv_hdr(parsed, short_parsed, vignette_factor, expo_time);
|
||||
parsed = parse_10bit(dat[1], extra_dat[1], aligned10);
|
||||
short_parsed = parse_10bit(short_dat[1], short_extra_dat[1], aligned10);
|
||||
v_rows[ROW_READ_ORDER[1]] = normalize_pv_hdr(parsed, short_parsed, vignette_factor, expo_time);
|
||||
parsed = parse_10bit(dat[2], extra_dat[2], aligned10);
|
||||
short_parsed = parse_10bit(short_dat[2], short_extra_dat[2], aligned10);
|
||||
v_rows[ROW_READ_ORDER[2]] = normalize_pv_hdr(parsed, short_parsed, vignette_factor, expo_time);
|
||||
parsed = parse_10bit(dat[3], extra_dat[3], aligned10);
|
||||
short_parsed = parse_10bit(short_dat[3], short_extra_dat[3], aligned10);
|
||||
v_rows[ROW_READ_ORDER[3]] = normalize_pv_hdr(parsed, short_parsed, vignette_factor, expo_time);
|
||||
#else
|
||||
// no HDR here
|
||||
int4 parsed = parse_12bit(dat[0]);
|
||||
v_rows[ROW_READ_ORDER[0]] = normalize_pv(parsed, vignette_factor);
|
||||
parsed = parse_12bit(dat[1]);
|
||||
v_rows[ROW_READ_ORDER[1]] = normalize_pv(parsed, vignette_factor);
|
||||
parsed = parse_12bit(dat[2]);
|
||||
v_rows[ROW_READ_ORDER[2]] = normalize_pv(parsed, vignette_factor);
|
||||
parsed = parse_12bit(dat[3]);
|
||||
v_rows[ROW_READ_ORDER[3]] = normalize_pv(parsed, vignette_factor);
|
||||
#endif
|
||||
|
||||
// mirror padding
|
||||
if (gid_x == 0) {
|
||||
|
|
|
|||
|
|
@ -23,9 +23,10 @@ OS04C10::OS04C10() {
|
|||
pixel_size_mm = 0.002;
|
||||
data_word = false;
|
||||
|
||||
hdr_offset = 64 * 2 + 8; // stagger
|
||||
frame_width = 2688;
|
||||
frame_height = 1520;
|
||||
frame_stride = (frame_width * 12 / 8); // no alignment
|
||||
frame_height = 1520 * 2 + hdr_offset;
|
||||
frame_stride = (frame_width * 10 / 8); // no alignment
|
||||
|
||||
extra_height = 0;
|
||||
frame_offset = 0;
|
||||
|
|
@ -34,8 +35,8 @@ OS04C10::OS04C10() {
|
|||
init_reg_array.assign(std::begin(init_array_os04c10), std::end(init_array_os04c10));
|
||||
probe_reg_addr = 0x300a;
|
||||
probe_expected_data = 0x5304;
|
||||
mipi_format = CAM_FORMAT_MIPI_RAW_12;
|
||||
frame_data_type = 0x2c;
|
||||
mipi_format = CAM_FORMAT_MIPI_RAW_10;
|
||||
frame_data_type = 0x2b;
|
||||
mclk_frequency = 24000000; // Hz
|
||||
|
||||
dc_gain_factor = 1;
|
||||
|
|
@ -66,7 +67,7 @@ std::vector<i2c_random_wr_payload> OS04C10::getExposureRegisters(int exposure_ti
|
|||
return {
|
||||
{0x3501, long_time>>8}, {0x3502, long_time&0xFF},
|
||||
{0x3508, real_gain>>8}, {0x3509, real_gain&0xFF},
|
||||
// {0x350c, real_gain>>8}, {0x350d, real_gain&0xFF},
|
||||
{0x350c, real_gain>>8}, {0x350d, real_gain&0xFF},
|
||||
};
|
||||
}
|
||||
|
||||
|
|
@ -81,6 +82,6 @@ float OS04C10::getExposureScore(float desired_ev, int exp_t, int exp_g_idx, floa
|
|||
score += std::abs(exp_g_idx - (int)analog_gain_rec_idx) * m;
|
||||
score += ((1 - analog_gain_cost_delta) +
|
||||
analog_gain_cost_delta * (exp_g_idx - analog_gain_min_idx) / (analog_gain_max_idx - analog_gain_min_idx)) *
|
||||
std::abs(exp_g_idx - gain_idx) * 5.0;
|
||||
std::abs(exp_g_idx - gain_idx) * 3.0;
|
||||
return score;
|
||||
}
|
||||
|
|
|
|||
|
|
@ -2,25 +2,54 @@
|
|||
|
||||
#define BGGR
|
||||
|
||||
#define BIT_DEPTH 12
|
||||
#define PV_MAX 4096
|
||||
#define BIT_DEPTH 10
|
||||
#define PV_MAX10 1023
|
||||
#define PV_MAX16 65536 // gamma curve is calibrated to 16bit
|
||||
#define BLACK_LVL 64
|
||||
#define VIGNETTE_RSZ 2.2545f
|
||||
|
||||
float4 normalize_pv(int4 parsed, float vignette_factor) {
|
||||
float4 pv = (convert_float4(parsed) - BLACK_LVL) / (PV_MAX - BLACK_LVL);
|
||||
float combine_dual_pvs(float lv, float sv, int expo_time) {
|
||||
float svc = fmax(sv * expo_time, (float)(64 * (PV_MAX10 - BLACK_LVL)));
|
||||
float svd = sv * fmin(expo_time, 8.0) / 8;
|
||||
|
||||
if (expo_time > 64) {
|
||||
if (lv < PV_MAX10 - BLACK_LVL) {
|
||||
return lv / (PV_MAX16 - BLACK_LVL);
|
||||
} else {
|
||||
return (svc / 64) / (PV_MAX16 - BLACK_LVL);
|
||||
}
|
||||
} else {
|
||||
if (lv > 32) {
|
||||
return (lv * 64 / fmax(expo_time, 8.0)) / (PV_MAX16 - BLACK_LVL);
|
||||
} else {
|
||||
return svd / (PV_MAX16 - BLACK_LVL);
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
float4 normalize_pv_hdr(int4 parsed, int4 short_parsed, float vignette_factor, int expo_time) {
|
||||
float4 pl = convert_float4(parsed - BLACK_LVL);
|
||||
float4 ps = convert_float4(short_parsed - BLACK_LVL);
|
||||
float4 pv;
|
||||
pv.s0 = combine_dual_pvs(pl.s0, ps.s0, expo_time);
|
||||
pv.s1 = combine_dual_pvs(pl.s1, ps.s1, expo_time);
|
||||
pv.s2 = combine_dual_pvs(pl.s2, ps.s2, expo_time);
|
||||
pv.s3 = combine_dual_pvs(pl.s3, ps.s3, expo_time);
|
||||
return clamp(pv*vignette_factor, 0.0, 1.0);
|
||||
}
|
||||
|
||||
float3 color_correct(float3 rgb) {
|
||||
float3 corrected = rgb.x * (float3)(1.5664815, -0.29808738, -0.03973474);
|
||||
corrected += rgb.y * (float3)(-0.48672447, 1.41914433, -0.40295248);
|
||||
corrected += rgb.z * (float3)(-0.07975703, -0.12105695, 1.44268722);
|
||||
float3 corrected = rgb.x * (float3)(1.55361989, -0.268894615, -0.000593219);
|
||||
corrected += rgb.y * (float3)(-0.421217301, 1.51883144, -0.69760146);
|
||||
corrected += rgb.z * (float3)(-0.132402589, -0.249936825, 1.69819468);
|
||||
return corrected;
|
||||
}
|
||||
|
||||
float3 apply_gamma(float3 rgb, int expo_time) {
|
||||
return powr(rgb, 0.7);
|
||||
float s = log2((float)expo_time);
|
||||
if (s < 6) {s = fmin(12.0 - s, 9.0);}
|
||||
// log function adaptive to number of bits
|
||||
return clamp(log(1 + rgb*(PV_MAX16 - BLACK_LVL)) * (0.48*s*s - 12.92*s + 115.0) - (1.08*s*s - 29.2*s + 260.0), 0.0, 255.0) / 255.0;
|
||||
}
|
||||
|
||||
#endif
|
||||
|
|
|
|||
|
|
@ -4,18 +4,18 @@ const struct i2c_random_wr_payload start_reg_array_os04c10[] = {{0x100, 1}};
|
|||
const struct i2c_random_wr_payload stop_reg_array_os04c10[] = {{0x100, 0}};
|
||||
|
||||
const struct i2c_random_wr_payload init_array_os04c10[] = {
|
||||
// OS04C10_AA_00_02_17_wAO_2688x1524_MIPI728Mbps_Linear12bit_20FPS_4Lane_MCLK24MHz
|
||||
// DP_2688X1520_NEWSTG_MIPI0776Mbps_30FPS_10BIT_FOURLANE
|
||||
{0x0103, 0x01},
|
||||
|
||||
// PLL
|
||||
{0x0301, 0xe4},
|
||||
{0x0301, 0x84},
|
||||
{0x0303, 0x01},
|
||||
{0x0305, 0xb6},
|
||||
{0x0305, 0x61},
|
||||
{0x0306, 0x01},
|
||||
{0x0307, 0x17},
|
||||
{0x0323, 0x04},
|
||||
{0x0324, 0x01},
|
||||
{0x0325, 0x62},
|
||||
{0x0325, 0x7a},
|
||||
|
||||
{0x3012, 0x06},
|
||||
{0x3013, 0x02},
|
||||
|
|
@ -30,40 +30,40 @@ const struct i2c_random_wr_payload init_array_os04c10[] = {
|
|||
{0x3660, 0x04},
|
||||
{0x3666, 0xa5},
|
||||
{0x3667, 0xa5},
|
||||
{0x366a, 0x50},
|
||||
{0x366a, 0x54},
|
||||
{0x3673, 0x0d},
|
||||
{0x3672, 0x0d},
|
||||
{0x3671, 0x0d},
|
||||
{0x3670, 0x0d},
|
||||
{0x3685, 0x00},
|
||||
{0x3685, 0x0a},
|
||||
{0x3694, 0x0d},
|
||||
{0x3693, 0x0d},
|
||||
{0x3692, 0x0d},
|
||||
{0x3691, 0x0d},
|
||||
{0x3696, 0x4c},
|
||||
{0x3697, 0x4c},
|
||||
{0x3698, 0x40},
|
||||
{0x3698, 0x00},
|
||||
{0x3699, 0x80},
|
||||
{0x369a, 0x18},
|
||||
{0x369a, 0x80},
|
||||
{0x369b, 0x1f},
|
||||
{0x369c, 0x14},
|
||||
{0x369c, 0x1f},
|
||||
{0x369d, 0x80},
|
||||
{0x369e, 0x40},
|
||||
{0x369f, 0x21},
|
||||
{0x36a0, 0x12},
|
||||
{0x36a1, 0x5d},
|
||||
{0x36a1, 0xdd},
|
||||
{0x36a2, 0x66},
|
||||
{0x370a, 0x02},
|
||||
{0x370e, 0x0c},
|
||||
{0x370a, 0x00},
|
||||
{0x370e, 0x00},
|
||||
{0x3710, 0x00},
|
||||
{0x3713, 0x00},
|
||||
{0x3713, 0x04},
|
||||
{0x3725, 0x02},
|
||||
{0x372a, 0x03},
|
||||
{0x3738, 0xce},
|
||||
{0x3748, 0x02},
|
||||
{0x374a, 0x02},
|
||||
{0x374c, 0x02},
|
||||
{0x374e, 0x02},
|
||||
{0x3748, 0x00},
|
||||
{0x374a, 0x00},
|
||||
{0x374c, 0x00},
|
||||
{0x374e, 0x00},
|
||||
{0x3756, 0x00},
|
||||
{0x3757, 0x00},
|
||||
{0x3767, 0x00},
|
||||
|
|
@ -81,20 +81,21 @@ const struct i2c_random_wr_payload init_array_os04c10[] = {
|
|||
{0x37ba, 0x03},
|
||||
{0x37bb, 0x00},
|
||||
{0x37bc, 0x04},
|
||||
{0x37be, 0x08},
|
||||
{0x37be, 0x26},
|
||||
{0x37c4, 0x11},
|
||||
{0x37c5, 0x80},
|
||||
{0x37c6, 0x14},
|
||||
{0x37c7, 0x08},
|
||||
{0x37c7, 0xa8},
|
||||
{0x37da, 0x11},
|
||||
{0x381f, 0x08},
|
||||
{0x3829, 0x03},
|
||||
{0x3832, 0x00},
|
||||
{0x3881, 0x00},
|
||||
{0x3888, 0x04},
|
||||
{0x388b, 0x00},
|
||||
{0x3c80, 0x10},
|
||||
{0x3c86, 0x00},
|
||||
{0x3c8c, 0x20},
|
||||
// {0x3c8c, 0x20},
|
||||
{0x3c9f, 0x01},
|
||||
{0x3d85, 0x1b},
|
||||
{0x3d8c, 0x71},
|
||||
|
|
@ -110,7 +111,7 @@ const struct i2c_random_wr_payload init_array_os04c10[] = {
|
|||
{0x4045, 0x7e},
|
||||
{0x4047, 0x7e},
|
||||
{0x4049, 0x7e},
|
||||
{0x4090, 0x04},
|
||||
{0x4090, 0x14},
|
||||
{0x40b0, 0x00},
|
||||
{0x40b1, 0x00},
|
||||
{0x40b2, 0x00},
|
||||
|
|
@ -128,7 +129,7 @@ const struct i2c_random_wr_payload init_array_os04c10[] = {
|
|||
{0x4503, 0x00},
|
||||
{0x4504, 0x06},
|
||||
{0x4506, 0x00},
|
||||
{0x4507, 0x47},
|
||||
{0x4507, 0x57},
|
||||
{0x4803, 0x00},
|
||||
{0x480c, 0x32},
|
||||
{0x480e, 0x04},
|
||||
|
|
@ -138,7 +139,7 @@ const struct i2c_random_wr_payload init_array_os04c10[] = {
|
|||
{0x4823, 0x3f},
|
||||
{0x4825, 0x30},
|
||||
{0x4833, 0x10},
|
||||
{0x484b, 0x27},
|
||||
{0x484b, 0x07},
|
||||
{0x488b, 0x00},
|
||||
{0x4d00, 0x04},
|
||||
{0x4d01, 0xad},
|
||||
|
|
@ -151,7 +152,7 @@ const struct i2c_random_wr_payload init_array_os04c10[] = {
|
|||
{0x4e0d, 0x00},
|
||||
|
||||
// ISP
|
||||
{0x5001, 0x09},
|
||||
{0x5001, 0x00},
|
||||
{0x5004, 0x00},
|
||||
{0x5080, 0x04},
|
||||
{0x5036, 0x80},
|
||||
|
|
@ -172,32 +173,32 @@ const struct i2c_random_wr_payload init_array_os04c10[] = {
|
|||
{0x301c, 0xf8},
|
||||
{0x301e, 0xb4},
|
||||
{0x301f, 0xf0},
|
||||
{0x3022, 0x61},
|
||||
{0x3022, 0x01},
|
||||
{0x3109, 0xe7},
|
||||
{0x3600, 0x00},
|
||||
{0x3610, 0x65},
|
||||
{0x3610, 0x75},
|
||||
{0x3611, 0x85},
|
||||
{0x3613, 0x3a},
|
||||
{0x3615, 0x60},
|
||||
{0x3621, 0xb0},
|
||||
{0x3621, 0x90},
|
||||
{0x3620, 0x0c},
|
||||
{0x3629, 0x00},
|
||||
{0x3661, 0x04},
|
||||
{0x3664, 0x70},
|
||||
{0x3665, 0x00},
|
||||
{0x3681, 0xa6},
|
||||
{0x3682, 0x53},
|
||||
{0x3683, 0x2a},
|
||||
{0x3684, 0x15},
|
||||
{0x3681, 0x80},
|
||||
{0x3682, 0x40},
|
||||
{0x3683, 0x21},
|
||||
{0x3684, 0x12},
|
||||
{0x3700, 0x2a},
|
||||
{0x3701, 0x12},
|
||||
{0x3703, 0x28},
|
||||
{0x3704, 0x0e},
|
||||
{0x3706, 0x9d},
|
||||
{0x3706, 0x4a},
|
||||
{0x3709, 0x4a},
|
||||
{0x370b, 0x48},
|
||||
{0x370b, 0xa2},
|
||||
{0x370c, 0x01},
|
||||
{0x370f, 0x04},
|
||||
{0x370f, 0x00},
|
||||
{0x3714, 0x24},
|
||||
{0x3716, 0x04},
|
||||
{0x3719, 0x11},
|
||||
|
|
@ -205,19 +206,19 @@ const struct i2c_random_wr_payload init_array_os04c10[] = {
|
|||
{0x3720, 0x00},
|
||||
{0x3724, 0x13},
|
||||
{0x373f, 0xb0},
|
||||
{0x3741, 0x9d},
|
||||
{0x3743, 0x9d},
|
||||
{0x3745, 0x9d},
|
||||
{0x3747, 0x9d},
|
||||
{0x3749, 0x48},
|
||||
{0x374b, 0x48},
|
||||
{0x374d, 0x48},
|
||||
{0x374f, 0x48},
|
||||
{0x3741, 0x4a},
|
||||
{0x3743, 0x4a},
|
||||
{0x3745, 0x4a},
|
||||
{0x3747, 0x4a},
|
||||
{0x3749, 0xa2},
|
||||
{0x374b, 0xa2},
|
||||
{0x374d, 0xa2},
|
||||
{0x374f, 0xa2},
|
||||
{0x3755, 0x10},
|
||||
{0x376c, 0x00},
|
||||
{0x378d, 0x3c},
|
||||
{0x3790, 0x01},
|
||||
{0x3791, 0x01},
|
||||
{0x378d, 0x30},
|
||||
{0x3790, 0x4a},
|
||||
{0x3791, 0xa2},
|
||||
{0x3798, 0x40},
|
||||
{0x379e, 0x00},
|
||||
{0x379f, 0x04},
|
||||
|
|
@ -232,17 +233,17 @@ const struct i2c_random_wr_payload init_array_os04c10[] = {
|
|||
{0x37c0, 0x11},
|
||||
{0x37c2, 0x04},
|
||||
{0x37cd, 0x19},
|
||||
{0x37e0, 0x08},
|
||||
{0x37e6, 0x04},
|
||||
// {0x37e0, 0x08},
|
||||
// {0x37e6, 0x04},
|
||||
{0x37e5, 0x02},
|
||||
{0x37e1, 0x0c},
|
||||
{0x3737, 0x04},
|
||||
// {0x37e1, 0x0c},
|
||||
// {0x3737, 0x04},
|
||||
{0x37d8, 0x02},
|
||||
{0x37e2, 0x10},
|
||||
// {0x37e2, 0x10},
|
||||
{0x3739, 0x10},
|
||||
{0x3662, 0x10},
|
||||
{0x37e4, 0x20},
|
||||
{0x37e3, 0x08},
|
||||
// {0x37e4, 0x20},
|
||||
// {0x37e3, 0x08},
|
||||
{0x37d9, 0x08},
|
||||
{0x4040, 0x00},
|
||||
{0x4041, 0x07},
|
||||
|
|
@ -263,51 +264,58 @@ const struct i2c_random_wr_payload init_array_os04c10[] = {
|
|||
{0x3816, 0x01},
|
||||
{0x3817, 0x01},
|
||||
|
||||
{0x380c, 0x08}, {0x380d, 0x5c}, // HTS
|
||||
{0x380e, 0x09}, {0x380f, 0x38}, // VTS
|
||||
{0x380c, 0x04}, {0x380d, 0x2e}, // HTS
|
||||
{0x380e, 0x09}, {0x380f, 0xdb}, // VTS
|
||||
|
||||
{0x3820, 0xb0},
|
||||
{0x3821, 0x00},
|
||||
{0x3880, 0x25},
|
||||
{0x3821, 0x04},
|
||||
{0x3880, 0x00},
|
||||
{0x3882, 0x20},
|
||||
{0x3c91, 0x0b},
|
||||
{0x3c94, 0x45},
|
||||
{0x3cad, 0x00},
|
||||
{0x3cae, 0x00},
|
||||
// {0x3cad, 0x00},
|
||||
// {0x3cae, 0x00},
|
||||
{0x4000, 0xf3},
|
||||
{0x4001, 0x60},
|
||||
{0x4003, 0x80},
|
||||
{0x4003, 0x40},
|
||||
{0x4300, 0xff},
|
||||
{0x4302, 0x0f},
|
||||
{0x4305, 0x83},
|
||||
{0x4305, 0x93},
|
||||
{0x4505, 0x84},
|
||||
{0x4809, 0x0e},
|
||||
{0x480a, 0x04},
|
||||
{0x4837, 0x15},
|
||||
{0x4837, 0x14},
|
||||
{0x4c00, 0x08},
|
||||
{0x4c01, 0x08},
|
||||
{0x4c04, 0x00},
|
||||
{0x4c05, 0x00},
|
||||
{0x5000, 0xf9},
|
||||
{0x3822, 0x14},
|
||||
// {0x0100, 0x01},
|
||||
// {0x320d, 0x00},
|
||||
// {0x3208, 0xa0},
|
||||
// {0x3822, 0x14},
|
||||
|
||||
// initialize exposure
|
||||
{0x3503, 0x88},
|
||||
|
||||
// long
|
||||
{0x3500, 0x00}, {0x3501, 0x00}, {0x3502, 0x80},
|
||||
{0x3500, 0x00}, {0x3501, 0x00}, {0x3502, 0x10},
|
||||
{0x3508, 0x00}, {0x3509, 0x80},
|
||||
{0x350a, 0x04}, {0x350b, 0x00},
|
||||
|
||||
// short
|
||||
// {0x3510, 0x00}, {0x3511, 0x00}, {0x3512, 0x10},
|
||||
// {0x350c, 0x00}, {0x350d, 0x80},
|
||||
// {0x350e, 0x04}, {0x350f, 0x00},
|
||||
{0x3510, 0x00}, {0x3511, 0x00}, {0x3512, 0x40},
|
||||
{0x350c, 0x00}, {0x350d, 0x80},
|
||||
{0x350e, 0x04}, {0x350f, 0x00},
|
||||
|
||||
// wb
|
||||
{0x5100, 0x06}, {0x5101, 0xcb},
|
||||
// b
|
||||
{0x5100, 0x06}, {0x5101, 0x7e},
|
||||
{0x5140, 0x06}, {0x5141, 0x7e},
|
||||
// g
|
||||
{0x5102, 0x04}, {0x5103, 0x00},
|
||||
{0x5104, 0x08}, {0x5105, 0xde},
|
||||
|
||||
{0x5106, 0x02}, {0x5107, 0x00},
|
||||
};
|
||||
{0x5142, 0x04}, {0x5143, 0x00},
|
||||
// r
|
||||
{0x5104, 0x08}, {0x5105, 0xd6},
|
||||
{0x5144, 0x08}, {0x5145, 0xd6},
|
||||
};
|
||||
|
|
|
|||
|
|
@ -29,6 +29,7 @@ public:
|
|||
uint32_t extra_height = 0;
|
||||
int registers_offset = -1;
|
||||
int stats_offset = -1;
|
||||
int hdr_offset = -1;
|
||||
|
||||
int exposure_time_min;
|
||||
int exposure_time_max;
|
||||
|
|
|
|||
Loading…
Reference in New Issue