mirror of
https://github.com/dragonpilot/dragonpilot.git
synced 2026-06-23 23:12:04 +08:00
@@ -127,6 +127,8 @@ bool CameraBuf::acquire() {
|
||||
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};
|
||||
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,
|
||||
0, 0, &debayer_event));
|
||||
#else
|
||||
@@ -278,30 +280,39 @@ static void publish_thumbnail(PubMaster *pm, const CameraBuf *b) {
|
||||
free(thumbnail_buffer);
|
||||
}
|
||||
|
||||
float set_exposure_target(const CameraBuf *b, int x_start, int x_end, int x_skip, int y_start, int y_end, int y_skip) {
|
||||
int lum_med;
|
||||
uint32_t lum_binning[256] = {0};
|
||||
float set_exposure_target(const CameraBuf *b, int x_start, int x_end, int x_skip, int y_start, int y_end, int y_skip, int analog_gain, bool hist_ceil, bool hl_weighted) {
|
||||
const uint8_t *pix_ptr = b->cur_yuv_buf->y;
|
||||
|
||||
uint32_t lum_binning[256] = {0};
|
||||
unsigned int lum_total = 0;
|
||||
for (int y = y_start; y < y_end; y += y_skip) {
|
||||
for (int x = x_start; x < x_end; x += x_skip) {
|
||||
uint8_t lum = pix_ptr[(y * b->rgb_width) + x];
|
||||
if (hist_ceil && lum < 80 && lum_binning[lum] > HISTO_CEIL_K * (y_end - y_start) * (x_end - x_start) / x_skip / y_skip / 256) {
|
||||
continue;
|
||||
}
|
||||
lum_binning[lum]++;
|
||||
lum_total += 1;
|
||||
}
|
||||
}
|
||||
|
||||
|
||||
// Find mean lumimance value
|
||||
unsigned int lum_cur = 0;
|
||||
for (lum_med = 255; lum_med >= 0; lum_med--) {
|
||||
int lum_med = 0;
|
||||
int lum_med_alt = 0;
|
||||
for (lum_med=255; lum_med>=0; lum_med--) {
|
||||
lum_cur += lum_binning[lum_med];
|
||||
|
||||
if (hl_weighted) {
|
||||
int lum_med_tmp = 0;
|
||||
int hb = HLC_THRESH + (10 - analog_gain);
|
||||
if (lum_cur > 0 && lum_med > hb) {
|
||||
lum_med_tmp = (lum_med - hb) + 100;
|
||||
}
|
||||
lum_med_alt = lum_med_alt>lum_med_tmp?lum_med_alt:lum_med_tmp;
|
||||
}
|
||||
if (lum_cur >= lum_total / 2) {
|
||||
break;
|
||||
}
|
||||
}
|
||||
lum_med = lum_med_alt>0 ? lum_med + lum_med/32*lum_cur*abs(lum_med_alt - lum_med)/lum_total:lum_med;
|
||||
|
||||
return lum_med / 256.0;
|
||||
}
|
||||
@@ -344,12 +355,18 @@ static void driver_cam_auto_exposure(CameraState *c, SubMaster &sm) {
|
||||
struct ExpRect {int x1, x2, x_skip, y1, y2, y_skip;};
|
||||
const CameraBuf *b = &c->buf;
|
||||
|
||||
bool hist_ceil = false, hl_weighted = false;
|
||||
int x_offset = 0, y_offset = 0;
|
||||
int frame_width = b->rgb_width, frame_height = b->rgb_height;
|
||||
|
||||
#ifndef QCOM2
|
||||
int analog_gain = -1;
|
||||
#else
|
||||
int analog_gain = c->analog_gain;
|
||||
#endif
|
||||
|
||||
ExpRect def_rect;
|
||||
if (Hardware::TICI()) {
|
||||
hist_ceil = hl_weighted = true;
|
||||
x_offset = 630, y_offset = 156;
|
||||
frame_width = 668, frame_height = frame_width / 1.33;
|
||||
def_rect = {96, 1832, 2, 242, 1148, 4};
|
||||
@@ -373,7 +390,7 @@ static void driver_cam_auto_exposure(CameraState *c, SubMaster &sm) {
|
||||
}
|
||||
}
|
||||
|
||||
camera_autoexposure(c, set_exposure_target(b, rect.x1, rect.x2, rect.x_skip, rect.y1, rect.y2, rect.y_skip));
|
||||
camera_autoexposure(c, set_exposure_target(b, rect.x1, rect.x2, rect.x_skip, rect.y1, rect.y2, rect.y_skip, analog_gain, hist_ceil, hl_weighted));
|
||||
}
|
||||
|
||||
void common_process_driver_camera(SubMaster *sm, PubMaster *pm, CameraState *c, int cnt) {
|
||||
|
||||
@@ -27,7 +27,6 @@
|
||||
#define CAMERA_ID_MAX 9
|
||||
|
||||
#define UI_BUF_COUNT 4
|
||||
|
||||
#define LOG_CAMERA_ID_FCAMERA 0
|
||||
#define LOG_CAMERA_ID_DCAMERA 1
|
||||
#define LOG_CAMERA_ID_ECAMERA 2
|
||||
@@ -135,7 +134,7 @@ typedef void (*process_thread_cb)(MultiCameraState *s, CameraState *c, int cnt);
|
||||
|
||||
void fill_frame_data(cereal::FrameData::Builder &framed, const FrameMetadata &frame_data);
|
||||
kj::Array<uint8_t> get_frame_image(const CameraBuf *b);
|
||||
float set_exposure_target(const CameraBuf *b, int x_start, int x_end, int x_skip, int y_start, int y_end, int y_skip);
|
||||
float set_exposure_target(const CameraBuf *b, int x_start, int x_end, int x_skip, int y_start, int y_end, int y_skip, int analog_gain, bool hist_ceil, bool hl_weighted);
|
||||
std::thread start_process_thread(MultiCameraState *cameras, CameraState *cs, process_thread_cb callback);
|
||||
void common_process_driver_camera(SubMaster *sm, PubMaster *pm, CameraState *c, int cnt);
|
||||
|
||||
|
||||
@@ -1113,7 +1113,7 @@ void process_road_camera(MultiCameraState *s, CameraState *c, int cnt) {
|
||||
if (cnt % 3 == 0) {
|
||||
const int x = 290, y = 322, width = 560, height = 314;
|
||||
const int skip = 1;
|
||||
camera_autoexposure(c, set_exposure_target(b, x, x + width, skip, y, y + height, skip));
|
||||
camera_autoexposure(c, set_exposure_target(b, x, x + width, skip, y, y + height, skip, -1, false, false));
|
||||
}
|
||||
}
|
||||
|
||||
|
||||
@@ -22,13 +22,18 @@
|
||||
#include "selfdrive/common/swaglog.h"
|
||||
#include "selfdrive/camerad/cameras/sensor2_i2c.h"
|
||||
|
||||
#define FRAME_WIDTH 1928
|
||||
#define FRAME_HEIGHT 1208
|
||||
//#define FRAME_STRIDE 1936 // for 8 bit output
|
||||
#define FRAME_STRIDE 2416 // for 10 bit output
|
||||
//#define FRAME_STRIDE 1936 // for 8 bit output
|
||||
|
||||
#define MIPI_SETTLE_CNT 33 // Calculated by camera_freqs.py
|
||||
|
||||
extern ExitHandler do_exit;
|
||||
|
||||
const size_t FRAME_WIDTH = 1928;
|
||||
const size_t FRAME_HEIGHT =1208;
|
||||
const size_t FRAME_STRIDE = 2416; // for 10 bit output
|
||||
|
||||
const int MIPI_SETTLE_CNT = 33; // Calculated by camera_freqs.py
|
||||
// global var for AE ops
|
||||
std::atomic<CameraExpInfo> cam_exp[3] = {{{0}}};
|
||||
|
||||
CameraInfo cameras_supported[CAMERA_ID_MAX] = {
|
||||
[CAMERA_ID_AR0231] = {
|
||||
@@ -41,25 +46,12 @@ CameraInfo cameras_supported[CAMERA_ID_MAX] = {
|
||||
},
|
||||
};
|
||||
|
||||
const bool enable_dc_gain = true;
|
||||
const float DC_GAIN = 2.5;
|
||||
const float sensor_analog_gains[] = {
|
||||
1.0/8.0, 2.0/8.0, 2.0/7.0, 3.0/7.0, // 0, 1, 2, 3
|
||||
3.0/6.0, 4.0/6.0, 4.0/5.0, 5.0/5.0, // 4, 5, 6, 7
|
||||
5.0/4.0, 6.0/4.0, 6.0/3.0, 7.0/3.0, // 8, 9, 10, 11
|
||||
7.0/2.0, 8.0/2.0, 8.0/1.0}; // 12, 13, 14, 15 = bypass
|
||||
|
||||
const int ANALOG_GAIN_MIN_IDX = 0x0; // 0.125x
|
||||
const int ANALOG_GAIN_REC_IDX = 0x7; // 1.0x
|
||||
const int ANALOG_GAIN_MAX_IDX = 0xD; // 4.0x
|
||||
|
||||
const int EXPOSURE_TIME_MIN = 2; // with HDR, fastest ss
|
||||
const int EXPOSURE_TIME_MAX = 1904; // with HDR, slowest ss
|
||||
|
||||
// global var for AE ops
|
||||
std::atomic<CameraExpInfo> cam_exp[3] = {{{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,
|
||||
5.0/4.0, 6.0/4.0, 6.0/3.0, 7.0/3.0,
|
||||
7.0/2.0, 8.0/2.0};
|
||||
|
||||
// ************** low level camera helpers ****************
|
||||
|
||||
int cam_control(int fd, int op_code, void *handle, int size) {
|
||||
struct cam_control camcontrol = {0};
|
||||
camcontrol.op_code = op_code;
|
||||
@@ -528,17 +520,15 @@ static void camera_init(MultiCameraState *multi_cam_state, VisionIpcServer * v,
|
||||
|
||||
s->camera_num = camera_num;
|
||||
|
||||
s->dc_gain_enabled = false;
|
||||
s->analog_gain = 0x5;
|
||||
s->analog_gain_frac = sensor_analog_gains[s->analog_gain];
|
||||
s->exposure_time = 256;
|
||||
s->exposure_time_max = 1.2 * EXPOSURE_TIME_MAX / 2;
|
||||
s->exposure_time_min = 0.75 * EXPOSURE_TIME_MIN * 2;
|
||||
s->request_id_last = 0;
|
||||
s->skipped = true;
|
||||
|
||||
s->min_ev = EXPOSURE_TIME_MIN * sensor_analog_gains[ANALOG_GAIN_MIN_IDX] * (enable_dc_gain ? DC_GAIN : 1);
|
||||
s->max_ev = EXPOSURE_TIME_MAX * sensor_analog_gains[ANALOG_GAIN_MAX_IDX] * DC_GAIN;
|
||||
s->target_grey_fraction = 0.3;
|
||||
|
||||
s->dc_gain_enabled = enable_dc_gain;
|
||||
s->gain_idx = ANALOG_GAIN_REC_IDX;
|
||||
s->exposure_time = 5;
|
||||
s->cur_ev = (s->dc_gain_enabled ? DC_GAIN : 1) * sensor_analog_gains[s->gain_idx] * s->exposure_time;
|
||||
s->ef_filtered = 1.0;
|
||||
|
||||
s->buf.init(device_id, ctx, s, v, FRAME_BUF_COUNT, rgb_type, yuv_type);
|
||||
}
|
||||
@@ -915,7 +905,7 @@ void handle_camera_event(CameraState *s, void *evdat) {
|
||||
meta_data.frame_id = main_id - s->idx_offset;
|
||||
meta_data.timestamp_sof = timestamp;
|
||||
s->exp_lock.lock();
|
||||
meta_data.gain = s->dc_gain_enabled ? s->analog_gain_frac * DC_GAIN : s->analog_gain_frac;
|
||||
meta_data.gain = s->dc_gain_enabled ? s->analog_gain_frac * 2.5 : s->analog_gain_frac;
|
||||
meta_data.high_conversion_gain = s->dc_gain_enabled;
|
||||
meta_data.integ_lines = s->exposure_time;
|
||||
meta_data.measured_grey_fraction = s->measured_grey_fraction;
|
||||
@@ -935,76 +925,105 @@ void handle_camera_event(CameraState *s, void *evdat) {
|
||||
}
|
||||
}
|
||||
|
||||
// ******************* exposure control helpers *******************
|
||||
|
||||
void set_exposure_time_bounds(CameraState *s) {
|
||||
switch (s->analog_gain) {
|
||||
case 0: {
|
||||
s->exposure_time_min = EXPOSURE_TIME_MIN;
|
||||
s->exposure_time_max = EXPOSURE_TIME_MAX; // EXPOSURE_TIME_MIN * 4;
|
||||
break;
|
||||
}
|
||||
case ANALOG_GAIN_MAX_IDX - 1: {
|
||||
s->exposure_time_min = EXPOSURE_TIME_MIN; // EXPOSURE_TIME_MAX / 4;
|
||||
s->exposure_time_max = EXPOSURE_TIME_MAX;
|
||||
break;
|
||||
}
|
||||
default: {
|
||||
// finetune margins on both ends
|
||||
float k_up = sensor_analog_gains[s->analog_gain+1] / sensor_analog_gains[s->analog_gain];
|
||||
float k_down = sensor_analog_gains[s->analog_gain-1] / sensor_analog_gains[s->analog_gain];
|
||||
s->exposure_time_min = k_down * EXPOSURE_TIME_MIN * 2;
|
||||
s->exposure_time_max = k_up * EXPOSURE_TIME_MAX / 2;
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
void switch_conversion_gain(CameraState *s) {
|
||||
if (!s->dc_gain_enabled) {
|
||||
s->dc_gain_enabled = true;
|
||||
s->analog_gain -= 4;
|
||||
} else {
|
||||
s->dc_gain_enabled = false;
|
||||
s->analog_gain += 4;
|
||||
}
|
||||
}
|
||||
|
||||
static void set_camera_exposure(CameraState *s, float grey_frac) {
|
||||
const float dt = 0.15;
|
||||
// TODO: get stats from sensor?
|
||||
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);
|
||||
exposure_factor = std::max(exposure_factor, 0.56f);
|
||||
|
||||
const float ts_grey = 10.0;
|
||||
const float ts_ev = 0.05;
|
||||
|
||||
const float k_grey = (dt / ts_grey) / (1.0 + dt / ts_grey);
|
||||
const float k_ev = (dt / ts_ev) / (1.0 + dt / ts_ev);
|
||||
|
||||
// Scale target grey between 0.1 and 0.4 depending on lighting conditions
|
||||
float new_target_grey = std::clamp(0.4 - 0.3 * log2(1.0 + s->cur_ev) / log2(6000.0), 0.1, 0.4);
|
||||
float target_grey = (1.0 - k_grey) * s->target_grey_fraction + k_grey * new_target_grey;
|
||||
|
||||
float desired_ev = std::clamp(s->cur_ev * target_grey / grey_frac, s->min_ev, s->max_ev);
|
||||
desired_ev = (1.0 - k_ev) * s->cur_ev + k_ev * desired_ev;
|
||||
|
||||
float best_ev_score = 1e6;
|
||||
int new_g = 0;
|
||||
int new_t = 0;
|
||||
|
||||
// Simple brute force optimizer to choose sensor parameters
|
||||
// to reach desired EV
|
||||
for (int g = ANALOG_GAIN_MIN_IDX; g <= ANALOG_GAIN_MAX_IDX; g++) {
|
||||
float gain = sensor_analog_gains[g] * (enable_dc_gain ? DC_GAIN : 1);
|
||||
|
||||
// Compute optimal time for given gain
|
||||
int t = std::clamp(int(std::round(desired_ev / gain)), EXPOSURE_TIME_MIN, EXPOSURE_TIME_MAX);
|
||||
|
||||
// Only go below recomended gain when absolutely necessary to not overexpose
|
||||
// behavior is not completely linear causing visible step changes when switching gain
|
||||
if (g < ANALOG_GAIN_REC_IDX && t > 10) {
|
||||
continue;
|
||||
}
|
||||
|
||||
// Compute error to desired ev
|
||||
float score = std::abs(desired_ev - (t * gain)) * 10;
|
||||
|
||||
// Going below recomended gain needs lower penalty to not overexpose
|
||||
float m = g > ANALOG_GAIN_REC_IDX ? 5.0 : 1.0;
|
||||
score += std::abs(g - (int)ANALOG_GAIN_REC_IDX) * m;
|
||||
|
||||
// Small penalty on changing gain
|
||||
score += std::abs(g - s->gain_idx);
|
||||
|
||||
if (score < best_ev_score) {
|
||||
new_t = t;
|
||||
new_g = g;
|
||||
best_ev_score = score;
|
||||
}
|
||||
if (s->camera_num != 1) {
|
||||
s->ef_filtered = (1 - EF_LOWPASS_K) * s->ef_filtered + EF_LOWPASS_K * exposure_factor;
|
||||
exposure_factor = s->ef_filtered;
|
||||
}
|
||||
|
||||
s->exp_lock.lock();
|
||||
|
||||
s->measured_grey_fraction = grey_frac;
|
||||
s->target_grey_fraction = target_grey;
|
||||
|
||||
s->analog_gain_frac = sensor_analog_gains[new_g];
|
||||
s->gain_idx = new_g;
|
||||
s->exposure_time = new_t;
|
||||
s->dc_gain_enabled = enable_dc_gain;
|
||||
// always prioritize exposure time adjust
|
||||
s->exposure_time *= exposure_factor;
|
||||
|
||||
float gain = s->analog_gain_frac * (s->dc_gain_enabled ? DC_GAIN : 1.0);
|
||||
s->cur_ev = s->exposure_time * gain;
|
||||
// switch gain if max/min exposure time is reached
|
||||
// or always switch down to a lower gain when possible
|
||||
bool kd = false;
|
||||
if (s->analog_gain > 0) {
|
||||
kd = 1.1 * s->exposure_time / (sensor_analog_gains[s->analog_gain-1] / sensor_analog_gains[s->analog_gain]) < EXPOSURE_TIME_MAX / 2;
|
||||
}
|
||||
|
||||
if (s->exposure_time > s->exposure_time_max) {
|
||||
if (s->analog_gain < ANALOG_GAIN_MAX_IDX - 1) {
|
||||
s->exposure_time = EXPOSURE_TIME_MAX / 2;
|
||||
s->analog_gain += 1;
|
||||
if (!s->dc_gain_enabled && sensor_analog_gains[s->analog_gain] >= 4.0) { // switch to HCG
|
||||
switch_conversion_gain(s);
|
||||
}
|
||||
set_exposure_time_bounds(s);
|
||||
} else {
|
||||
s->exposure_time = s->exposure_time_max;
|
||||
}
|
||||
} else if (s->exposure_time < s->exposure_time_min || kd) {
|
||||
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->analog_gain -= 1;
|
||||
if (s->dc_gain_enabled && sensor_analog_gains[s->analog_gain] <= 1.25) { // switch back to LCG
|
||||
switch_conversion_gain(s);
|
||||
}
|
||||
set_exposure_time_bounds(s);
|
||||
} else {
|
||||
s->exposure_time = s->exposure_time_min;
|
||||
}
|
||||
}
|
||||
|
||||
// set up config
|
||||
uint16_t AG = s->analog_gain + 4;
|
||||
AG = 0xFF00 + AG * 16 + AG;
|
||||
s->analog_gain_frac = sensor_analog_gains[s->analog_gain];
|
||||
|
||||
s->exp_lock.unlock();
|
||||
// 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);
|
||||
|
||||
uint16_t analog_gain_reg = 0xFF00 | (new_g << 4) | new_g;
|
||||
struct i2c_random_wr_payload exp_reg_array[] = {
|
||||
{0x3366, analog_gain_reg}, // analog gain
|
||||
{0x3362, (uint16_t)(s->dc_gain_enabled ? 0x1 : 0x0)}, // DC_GAIN
|
||||
{0x3366, AG}, // analog gain
|
||||
{0x3362, (uint16_t)(s->dc_gain_enabled?0x1:0x0)}, // DC_GAIN
|
||||
{0x305A, 0x00F8}, // red gain
|
||||
{0x3058, 0x0122}, // blue gain
|
||||
{0x3056, 0x009A}, // g1 gain
|
||||
{0x305C, 0x009A}, // g2 gain
|
||||
{0x3012, (uint16_t)s->exposure_time}, // integ time
|
||||
};
|
||||
//{0x301A, 0x091C}}; // reset
|
||||
@@ -1062,7 +1081,7 @@ void process_road_camera(MultiCameraState *s, CameraState *c, int cnt) {
|
||||
if (cnt % 3 == 0) {
|
||||
const auto [x, y, w, h] = (c == &s->wide_road_cam) ? std::tuple(96, 250, 1734, 524) : std::tuple(96, 160, 1734, 986);
|
||||
const int skip = 2;
|
||||
camera_autoexposure(c, set_exposure_target(b, x, x + w, skip, y, y + h, skip));
|
||||
camera_autoexposure(c, set_exposure_target(b, x, x + w, skip, y, y + h, skip, (int)c->analog_gain, true, true));
|
||||
}
|
||||
}
|
||||
|
||||
|
||||
@@ -10,23 +10,30 @@
|
||||
#include "selfdrive/common/util.h"
|
||||
|
||||
#define FRAME_BUF_COUNT 4
|
||||
|
||||
#define ANALOG_GAIN_MAX_IDX 10 // 0xF is bypass
|
||||
#define EXPOSURE_TIME_MIN 2 // with HDR, fastest ss
|
||||
#define EXPOSURE_TIME_MAX 1904 // with HDR, slowest ss
|
||||
|
||||
#define EF_LOWPASS_K 0.35
|
||||
|
||||
#define DEBAYER_LOCAL_WORKSIZE 16
|
||||
|
||||
typedef struct CameraState {
|
||||
MultiCameraState *multi_cam_state;
|
||||
CameraInfo ci;
|
||||
|
||||
std::mutex exp_lock;
|
||||
|
||||
int exposure_time;
|
||||
bool dc_gain_enabled;
|
||||
float analog_gain_frac;
|
||||
|
||||
float cur_ev;
|
||||
float min_ev, max_ev;
|
||||
uint16_t analog_gain;
|
||||
bool dc_gain_enabled;
|
||||
int exposure_time;
|
||||
int exposure_time_min;
|
||||
int exposure_time_max;
|
||||
float ef_filtered;
|
||||
|
||||
float measured_grey_fraction;
|
||||
float target_grey_fraction;
|
||||
int gain_idx;
|
||||
|
||||
unique_fd sensor_fd;
|
||||
unique_fd csiphy_fd;
|
||||
|
||||
@@ -26,9 +26,9 @@ half mf(half x, half cp) {
|
||||
}
|
||||
}
|
||||
|
||||
half3 color_correct(half3 rgb) {
|
||||
half3 color_correct(half3 rgb, int ggain) {
|
||||
half3 ret = (0,0,0);
|
||||
half cpx = 0.01;
|
||||
half cpx = 0.01; //clamp(0.01h, 0.05h, cpxb + cpxk * min(10, ggain));
|
||||
ret += (half)rgb.x * color_correction[0];
|
||||
ret += (half)rgb.y * color_correction[1];
|
||||
ret += (half)rgb.z * color_correction[2];
|
||||
@@ -89,7 +89,8 @@ half phi(half x) {
|
||||
|
||||
__kernel void debayer10(const __global uchar * in,
|
||||
__global uchar * out,
|
||||
__local half * cached
|
||||
__local half * cached,
|
||||
uint ggain
|
||||
)
|
||||
{
|
||||
const int x_global = get_global_id(0);
|
||||
@@ -199,9 +200,10 @@ __kernel void debayer10(const __global uchar * in,
|
||||
}
|
||||
|
||||
rgb = clamp(0.0h, 1.0h, rgb);
|
||||
rgb = color_correct(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);
|
||||
|
||||
}
|
||||
|
||||
@@ -82,16 +82,16 @@ struct i2c_random_wr_payload init_array_ar0231[] = {
|
||||
{0x328E, 0x0FA0}, // T2 G2
|
||||
|
||||
// Initial Gains
|
||||
{0x3022, 0x0001}, // GROUPED_PARAMETER_HOLD_
|
||||
{0x3366, 0xFF77}, // ANALOG_GAIN (1x)
|
||||
{0x3022, 0x01}, // GROUPED_PARAMETER_HOLD_
|
||||
{0x3366, 0x5555}, // ANALOG_GAIN
|
||||
{0x3060, 0x3333}, // ANALOG_COLOR_GAIN
|
||||
{0x3362, 0x0001}, // DC GAIN (enabled)
|
||||
{0x305A, 0x00F8}, // red gain
|
||||
{0x3058, 0x0122}, // blue gain
|
||||
{0x3056, 0x009A}, // g1 gain
|
||||
{0x305C, 0x009A}, // g2 gain
|
||||
{0x3362, 0x0000}, // DC GAIN
|
||||
{0x305A, 0x0108}, // RED_GAIN
|
||||
{0x3058, 0x00FB}, // BLUE_GAIN
|
||||
{0x3056, 0x009A}, // GREEN1_GAIN
|
||||
{0x305C, 0x009A}, // GREEN2_GAIN
|
||||
{0x3022, 0x00}, // GROUPED_PARAMETER_HOLD_
|
||||
|
||||
// Initial Integration Time
|
||||
{0x3012, 0x0005},
|
||||
{0x3012, 0x256},
|
||||
};
|
||||
|
||||
@@ -34,33 +34,37 @@ int main() {
|
||||
float rtol = 0.05;
|
||||
// generate pattern and calculate EV
|
||||
int cnt = 0;
|
||||
for (int i_0=0; i_0<TONE_SPLITS; i_0++) {
|
||||
for (int i_1=0; i_1<TONE_SPLITS; i_1++) {
|
||||
for (int i_2=0; i_2<TONE_SPLITS; i_2++) {
|
||||
for (int i_3=0; i_3<TONE_SPLITS; i_3++) {
|
||||
int h_0 = i_0 * H / TONE_SPLITS;
|
||||
int h_1 = i_1 * (H - h_0) / TONE_SPLITS;
|
||||
int h_2 = i_2 * (H - h_0 - h_1) / TONE_SPLITS;
|
||||
int h_3 = i_3 * (H - h_0 - h_1 - h_2) / TONE_SPLITS;
|
||||
int h_4 = H - h_0 - h_1 - h_2 - h_3;
|
||||
memset(&fb_y[0], l[0], h_0*W);
|
||||
memset(&fb_y[h_0*W], l[1], h_1*W);
|
||||
memset(&fb_y[h_0*W+h_1*W], l[2], h_2*W);
|
||||
memset(&fb_y[h_0*W+h_1*W+h_2*W], l[3], h_3*W);
|
||||
memset(&fb_y[h_0*W+h_1*W+h_2*W+h_3*W], l[4], h_4*W);
|
||||
float ev = set_exposure_target((const CameraBuf*) &cb, 0, W-1, 1, 0, H-1, 1);
|
||||
// printf("%d/%d/%d/%d/%d ev is %f\n", h_0, h_1, h_2, h_3, h_4, ev);
|
||||
// printf("%f\n", ev);
|
||||
for (int is_qcom2=0; is_qcom2<2; is_qcom2++) {
|
||||
for (int g=0; g<GAIN_SPLITS; g++) {
|
||||
for (int i_0=0; i_0<TONE_SPLITS; i_0++) {
|
||||
for (int i_1=0; i_1<TONE_SPLITS; i_1++) {
|
||||
for (int i_2=0; i_2<TONE_SPLITS; i_2++) {
|
||||
for (int i_3=0; i_3<TONE_SPLITS; i_3++) {
|
||||
int h_0 = i_0 * H / TONE_SPLITS;
|
||||
int h_1 = i_1 * (H - h_0) / TONE_SPLITS;
|
||||
int h_2 = i_2 * (H - h_0 - h_1) / TONE_SPLITS;
|
||||
int h_3 = i_3 * (H - h_0 - h_1 - h_2) / TONE_SPLITS;
|
||||
int h_4 = H - h_0 - h_1 - h_2 - h_3;
|
||||
memset(&fb_y[0], l[0], h_0*W);
|
||||
memset(&fb_y[h_0*W], l[1], h_1*W);
|
||||
memset(&fb_y[h_0*W+h_1*W], l[2], h_2*W);
|
||||
memset(&fb_y[h_0*W+h_1*W+h_2*W], l[3], h_3*W);
|
||||
memset(&fb_y[h_0*W+h_1*W+h_2*W+h_3*W], l[4], h_4*W);
|
||||
float ev = set_exposure_target((const CameraBuf*) &cb, 0, W-1, 1, 0, H-1, 1, g*10, (bool)is_qcom2, (bool)is_qcom2);
|
||||
// printf("%d/%d/%d/%d/%d ev is %f\n", h_0, h_1, h_2, h_3, h_4, ev);
|
||||
// printf("%f\n", ev);
|
||||
|
||||
// compare to gt
|
||||
float evgt = gts[cnt];
|
||||
if (fabs(ev - evgt) > rtol*evgt) {
|
||||
passed = false;
|
||||
// compare to gt
|
||||
float evgt = gts[cnt];
|
||||
if (fabs(ev - evgt) > rtol*evgt) {
|
||||
passed = false;
|
||||
}
|
||||
|
||||
// report
|
||||
printf("%d/%d/%d/%d/%d/%d/%d: ev %f, gt %f, err %f\n", is_qcom2, g*10, h_0, h_1, h_2, h_3, h_4, ev, evgt, fabs(ev - evgt) / (evgt != 0 ? evgt : 0.00001f));
|
||||
cnt++;
|
||||
}
|
||||
}
|
||||
|
||||
// report
|
||||
printf("%d/%d/%d/%d/%d: ev %f, gt %f, err %f\n", h_0, h_1, h_2, h_3, h_4, ev, evgt, fabs(ev - evgt) / (evgt != 0 ? evgt : 0.00001f));
|
||||
cnt++;
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
@@ -4,8 +4,9 @@
|
||||
#define H 160
|
||||
|
||||
#define TONE_SPLITS 3
|
||||
#define GAIN_SPLITS 2
|
||||
|
||||
float gts[TONE_SPLITS*TONE_SPLITS*TONE_SPLITS*TONE_SPLITS] = {
|
||||
float gts[2*TONE_SPLITS*TONE_SPLITS*TONE_SPLITS*TONE_SPLITS*GAIN_SPLITS] = {
|
||||
0.917969,0.917969,0.375000,0.917969,0.375000,0.375000,0.187500,0.187500,0.187500,0.917969,
|
||||
0.375000,0.375000,0.187500,0.187500,0.187500,0.187500,0.187500,0.187500,0.093750,0.093750,
|
||||
0.093750,0.093750,0.093750,0.093750,0.093750,0.093750,0.093750,0.917969,0.375000,0.375000,
|
||||
@@ -14,5 +15,29 @@ float gts[TONE_SPLITS*TONE_SPLITS*TONE_SPLITS*TONE_SPLITS] = {
|
||||
0.093750,0.093750,0.093750,0.093750,0.000000,0.000000,0.000000,0.000000,0.000000,0.000000,
|
||||
0.000000,0.000000,0.000000,0.000000,0.000000,0.000000,0.000000,0.000000,0.000000,0.000000,
|
||||
0.000000,0.000000,0.000000,0.000000,0.000000,0.000000,0.000000,0.000000,0.000000,0.000000,
|
||||
0.000000
|
||||
0.000000,0.917969,0.917969,0.375000,0.917969,0.375000,0.375000,0.187500,0.187500,0.187500,
|
||||
0.917969,0.375000,0.375000,0.187500,0.187500,0.187500,0.187500,0.187500,0.187500,0.093750,
|
||||
0.093750,0.093750,0.093750,0.093750,0.093750,0.093750,0.093750,0.093750,0.917969,0.375000,
|
||||
0.375000,0.187500,0.187500,0.187500,0.187500,0.187500,0.187500,0.093750,0.093750,0.093750,
|
||||
0.093750,0.093750,0.093750,0.093750,0.093750,0.093750,0.093750,0.093750,0.093750,0.093750,
|
||||
0.093750,0.093750,0.093750,0.093750,0.093750,0.000000,0.000000,0.000000,0.000000,0.000000,
|
||||
0.000000,0.000000,0.000000,0.000000,0.000000,0.000000,0.000000,0.000000,0.000000,0.000000,
|
||||
0.000000,0.000000,0.000000,0.000000,0.000000,0.000000,0.000000,0.000000,0.000000,0.000000,
|
||||
0.000000,0.000000,4.527344,3.324219,0.457031,4.421875,3.265625,0.453125,4.324219,3.167969,
|
||||
0.449219,4.421875,3.265625,0.453125,4.234375,3.113281,0.449219,3.980469,2.929688,0.441406,
|
||||
4.324219,3.167969,0.449219,3.980469,2.929688,0.441406,3.558594,0.433594,0.433594,4.421875,
|
||||
3.265625,0.453125,4.234375,3.113281,0.449219,3.980469,2.929688,0.441406,4.234375,3.113281,
|
||||
0.449219,3.929688,2.902344,0.441406,3.484375,0.429688,0.429688,3.980469,2.929688,0.441406,
|
||||
3.484375,0.429688,0.429688,2.871094,0.417969,0.417969,4.324219,3.167969,0.449219,3.980469,
|
||||
2.929688,0.441406,3.558594,0.433594,0.433594,3.980469,2.929688,0.441406,3.484375,0.429688,
|
||||
0.429688,2.871094,0.417969,0.417969,3.558594,0.433594,0.433594,2.871094,0.417969,0.417969,
|
||||
0.308594,0.308594,0.308594,4.253906,3.140625,0.574219,4.156250,3.085938,0.566406,4.066406,
|
||||
2.996094,0.562500,4.156250,3.085938,0.566406,3.984375,2.945312,0.554688,3.750000,2.777344,
|
||||
0.542969,4.066406,2.996094,0.562500,3.750000,2.777344,0.542969,3.359375,0.519531,0.519531,
|
||||
4.156250,3.085938,0.566406,3.984375,2.945312,0.554688,3.750000,2.777344,0.542969,3.984375,
|
||||
2.945312,0.554688,3.699219,2.753906,0.539062,3.289062,0.515625,0.515625,3.750000,2.777344,
|
||||
0.542969,3.289062,0.515625,0.515625,2.722656,0.480469,0.480469,4.066406,2.996094,0.562500,
|
||||
3.750000,2.777344,0.542969,3.359375,0.519531,0.519531,3.750000,2.777344,0.542969,3.289062,
|
||||
0.515625,0.515625,2.722656,0.480469,0.480469,3.359375,0.519531,0.519531,2.722656,0.480469,
|
||||
0.480469,0.328125,0.328125,0.328125,
|
||||
};
|
||||
|
||||
Reference in New Issue
Block a user