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

pull it for now

Browse Source
pull/32249/head
ZwX1616 1 year ago
parent f923289f3f
commit 97caa34d48
2 changed files with 10 additions and 112 deletions
Whitespace
Show all changes Ignore whitespace when comparing lines Ignore changes in amount of whitespace Ignore changes in whitespace at EOL
Split View
Diff Options
Show Stats Download Patch File Download Diff File
  1. 47
      system/camerad/cameras/process_raw.cl
  2. 55
      system/camerad/sensors/os04c10_cl.h

47
system/camerad/cameras/process_raw.cl
Unescape View File

@ -1,6 +1,5 @@
#include "ar0231_cl.h"
#include "ox03c10_cl.h"
#include "os04c10_cl.h"
#define UV_WIDTH RGB_WIDTH / 2
#define UV_HEIGHT RGB_HEIGHT / 2
@ -78,18 +77,7 @@ __kernel void process_raw(const __global uchar * in, __global uchar * out, int e
// read offset
int start_idx;
#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];
@ -107,16 +95,6 @@ __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
@ -125,34 +103,10 @@ __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];
#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);
@ -162,7 +116,6 @@ __kernel void process_raw(const __global uchar * in, __global uchar * out, int e
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) {

55
system/camerad/sensors/os04c10_cl.h
Unescape View File

@ -1,55 +0,0 @@
#if IS_OS
#define BGGR
#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
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.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) {
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
Write Preview
Loading…
Cancel
Save