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

clean up

Browse Source
pull/32112/head
ZwX1616 1 year ago
parent 677cb43131
commit dcf1a8dc23
4 changed files with 113 additions and 108 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. 173
      system/camerad/cameras/process_raw.cl
  2. 7
      system/camerad/sensors/ar0231_cl.h
  3. 34
      system/camerad/sensors/os04c10_cl.h
  4. 7
      system/camerad/sensors/ox03c10_cl.h

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

@ -19,9 +19,6 @@
#endif #endif
float get_vignetting_s(float r) { float get_vignetting_s(float r) {
#if IS_OS
r = r / 2.2545f;
#endif
if (r < 62500) { if (r < 62500) {
return (1.0f + 0.0000008f*r); return (1.0f + 0.0000008f*r);
} else if (r < 490000) { } else if (r < 490000) {
@ -33,74 +30,26 @@ float get_vignetting_s(float r) {
} }
} }
float4 val4_from_12(uchar8 pvs, float gain) { int4 parse_12bit(uchar8 pvs) {
uint4 parsed = (uint4)(((uint)pvs.s0<<4) + (pvs.s1>>4), // is from the previous 10 bit // lower bits scambled?
((uint)pvs.s2<<4) + (pvs.s4&0xF), return (int4)(((int)pvs.s0<<4) + (pvs.s1>>4),
((uint)pvs.s3<<4) + (pvs.s4>>4), ((int)pvs.s2<<4) + (pvs.s4&0xF),
((uint)pvs.s5<<4) + (pvs.s7&0xF)); ((int)pvs.s3<<4) + (pvs.s4>>4),
#if IS_OX ((int)pvs.s5<<4) + (pvs.s7&0xF));
// PWL
//float4 pv = (convert_float4(parsed) - 64.0) / (4096.0 - 64.0);
float4 pv = {ox03c10_lut[parsed.s0], ox03c10_lut[parsed.s1], ox03c10_lut[parsed.s2], ox03c10_lut[parsed.s3]};
// it's a 24 bit signal, center in the middle 8 bits
return clamp(pv*gain*256.0, 0.0, 1.0);
#else // AR
// normalize and scale
float4 pv = (convert_float4(parsed) - 168.0) / (4096.0 - 168.0);
return clamp(pv*gain, 0.0, 1.0);
#endif
} }
float combine_pvs(float lv, float sv, int expo) { int4 parse_10bit(uchar8 pvs, uchar ext, bool aligned) {
float svc = fmax(sv * expo, 61376.0);
float svd = sv * fmin(expo, 8.0) / 8;
if (expo > 64) {
if (lv < 959) {
return lv / (65536.0 - 64.0);
} else {
return (svc / 64) / (65536.0 - 64.0);
}
} else {
if (lv > 32) {
return (lv * 64 / fmax(expo, 8.0)) / (65536.0 - 64.0);
} else {
return svd / (65536.0 - 64.0);
}
}
}
float4 val4_from_10x2(uchar8 long_pvs, uchar long_ext, uchar8 short_pvs, uchar short_ext, bool aligned, float gain, int expo) {
int8 parsed;
if (aligned) { if (aligned) {
parsed = (int8)(((int)long_pvs.s0 << 2) + (long_pvs.s1 & 0b00000011), return (int4)(((int)pvs.s0 << 2) + (pvs.s1 & 0b00000011),
((int)long_pvs.s2 << 2) + ((long_pvs.s6 & 0b11000000) / 64), ((int)pvs.s2 << 2) + ((pvs.s6 & 0b11000000) / 64),
((int)long_pvs.s3 << 2) + ((long_pvs.s6 & 0b00110000) / 16), ((int)pvs.s3 << 2) + ((pvs.s6 & 0b00110000) / 16),
((int)long_pvs.s4 << 2) + ((long_pvs.s6 & 0b00001100) / 4), ((int)pvs.s4 << 2) + ((pvs.s6 & 0b00001100) / 4));
((int)short_pvs.s0 << 2) + (short_pvs.s1 & 0b00000011),
((int)short_pvs.s2 << 2) + ((short_pvs.s6 & 0b11000000) / 64),
((int)short_pvs.s3 << 2) + ((short_pvs.s6 & 0b00110000) / 16),
((int)short_pvs.s4 << 2) + ((short_pvs.s6 & 0b00001100) / 4));
} else { } else {
parsed = (int8)(((int)long_pvs.s0 << 2) + ((long_pvs.s3 & 0b00110000) / 16), return (int4)(((int)pvs.s0 << 2) + ((pvs.s3 & 0b00110000) / 16),
((int)long_pvs.s1 << 2) + ((long_pvs.s3 & 0b00001100) / 4), ((int)pvs.s1 << 2) + ((pvs.s3 & 0b00001100) / 4),
((int)long_pvs.s2 << 2) + ((long_pvs.s3 & 0b00000011)), ((int)pvs.s2 << 2) + ((pvs.s3 & 0b00000011)),
((int)long_pvs.s4 << 2) + ((long_ext & 0b11000000) / 64), ((int)pvs.s4 << 2) + ((ext & 0b11000000) / 64));
((int)short_pvs.s0 << 2) + ((short_pvs.s3 & 0b00110000) / 16),
((int)short_pvs.s1 << 2) + ((short_pvs.s3 & 0b00001100) / 4),
((int)short_pvs.s2 << 2) + ((short_pvs.s3 & 0b00000011)),
((int)short_pvs.s4 << 2) + ((short_ext & 0b11000000) / 64));
} }
float8 pf = convert_float8(parsed - 64);
float4 pv;
pv.s0 = combine_pvs(pf.s0, pf.s4, expo);
pv.s1 = combine_pvs(pf.s1, pf.s5, expo);
pv.s2 = combine_pvs(pf.s2, pf.s6, expo);
pv.s3 = combine_pvs(pf.s3, pf.s7, expo);
return clamp(pv*gain, 0.0, 1.0);
} }
float get_k(float a, float b, float c, float d) { float get_k(float a, float b, float c, float d) {
@ -112,14 +61,23 @@ __kernel void process_raw(const __global uchar * in, __global uchar * out, int e
const int gid_x = get_global_id(0); const int gid_x = get_global_id(0);
const int gid_y = get_global_id(1); const int gid_y = get_global_id(1);
// estimate vignetting
#if VIGNETTING
int gx = (gid_x*2 - RGB_WIDTH/2);
int gy = (gid_y*2 - RGB_HEIGHT/2);
const float vignette_factor = get_vignetting_s((gx*gx + gy*gy) / VIGNETTE_RSZ);
#else
const float vignette_factor = 1.0;
#endif
const int row_before_offset = (gid_y == 0) ? 2 : 0; const int row_before_offset = (gid_y == 0) ? 2 : 0;
const int row_after_offset = (gid_y == (RGB_HEIGHT/2 - 1)) ? 1 : 3; const int row_after_offset = (gid_y == (RGB_HEIGHT/2 - 1)) ? 1 : 3;
float3 rgb_tmp; float3 rgb_tmp;
uchar3 rgb_out[4]; // output is 2x2 window uchar3 rgb_out[4]; // output is 2x2 window
// read offset
int start_idx; int start_idx;
#if BIT_DEPTH == 10 #if BIT_DEPTH == 10
bool aligned10; bool aligned10;
if (gid_x % 2 == 0) { if (gid_x % 2 == 0) {
@ -133,9 +91,11 @@ __kernel void process_raw(const __global uchar * in, __global uchar * out, int e
start_idx = (2 * gid_y - 1) * FRAME_STRIDE + (3 * gid_x - 2) + (FRAME_STRIDE * FRAME_OFFSET); start_idx = (2 * gid_y - 1) * FRAME_STRIDE + (3 * gid_x - 2) + (FRAME_STRIDE * FRAME_OFFSET);
#endif #endif
// read in 8x4 chars // read in 4 rows, 8 uchars each
uchar8 dat[4]; uchar8 dat[4];
// row_before
dat[0] = vload8(0, in + start_idx + FRAME_STRIDE*row_before_offset); dat[0] = vload8(0, in + start_idx + FRAME_STRIDE*row_before_offset);
// row_0
if (gid_x == 0 && gid_y == 0) { if (gid_x == 0 && gid_y == 0) {
// this wasn't a problem due to extra rows // this wasn't a problem due to extra rows
dat[1] = vload8(0, in + start_idx + FRAME_STRIDE*1 + 2); dat[1] = vload8(0, in + start_idx + FRAME_STRIDE*1 + 2);
@ -143,54 +103,53 @@ __kernel void process_raw(const __global uchar * in, __global uchar * out, int e
} else { } else {
dat[1] = vload8(0, in + start_idx + FRAME_STRIDE*1); dat[1] = vload8(0, in + start_idx + FRAME_STRIDE*1);
} }
// row_1
dat[2] = vload8(0, in + start_idx + FRAME_STRIDE*2); dat[2] = vload8(0, in + start_idx + FRAME_STRIDE*2);
// row_after
dat[3] = vload8(0, in + start_idx + FRAME_STRIDE*row_after_offset); 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 #if HDR_OFFSET > 0
uchar8 short_dat[4]; uchar8 short_dat[4];
short_dat[0] = vload8(0, in + start_idx + FRAME_STRIDE*(row_before_offset+HDR_OFFSET/2) + FRAME_STRIDE/2); short_dat[0] = vload8(0, in + start_idx + FRAME_STRIDE*(row_before_offset+HDR_OFFSET/2) + FRAME_STRIDE/2);
short_dat[1] = vload8(0, in + start_idx + FRAME_STRIDE*(1+HDR_OFFSET/2) + FRAME_STRIDE/2); 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[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); 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 #endif
// need extra bit for 10-bit // parse into floats 0.0-1.0
#if BIT_DEPTH == 10
uchar extra[4];
uchar short_extra[4];
if (!aligned10) {
extra[0] = in[start_idx + FRAME_STRIDE*row_before_offset + 8];
extra[1] = in[start_idx + FRAME_STRIDE*1 + 8];
extra[2] = in[start_idx + FRAME_STRIDE*2 + 8];
extra[3] = in[start_idx + FRAME_STRIDE*row_after_offset + 8];
short_extra[0] = in[start_idx + FRAME_STRIDE*(row_before_offset+HDR_OFFSET/2) + FRAME_STRIDE/2 + 8];
short_extra[1] = in[start_idx + FRAME_STRIDE*(1+HDR_OFFSET/2) + FRAME_STRIDE/2 + 8];
short_extra[2] = in[start_idx + FRAME_STRIDE*(2+HDR_OFFSET/2) + FRAME_STRIDE/2 + 8];
short_extra[3] = in[start_idx + FRAME_STRIDE*(row_after_offset+HDR_OFFSET/2) + FRAME_STRIDE/2 + 8];
}
#endif
// correct vignetting
#if VIGNETTING
int gx = (gid_x*2 - RGB_WIDTH/2);
int gy = (gid_y*2 - RGB_HEIGHT/2);
const float gain = get_vignetting_s(gx*gx + gy*gy);
#else
const float gain = 1.0;
#endif
float4 v_rows[4]; float4 v_rows[4];
// parse into floats for (int i=0; i<4; i++) {
#if BIT_DEPTH == 10 #if BIT_DEPTH == 10
v_rows[ROW_READ_ORDER[0]] = val4_from_10x2(dat[0], extra[0], short_dat[0], short_extra[0], aligned10, gain, expo_time); // for now it's always HDR
v_rows[ROW_READ_ORDER[1]] = val4_from_10x2(dat[1], extra[1], short_dat[1], short_extra[1], aligned10, gain, expo_time); int4 parsed = parse_10bit(dat[i], extra_dat[i], aligned10);
v_rows[ROW_READ_ORDER[2]] = val4_from_10x2(dat[2], extra[2], short_dat[2], short_extra[2], aligned10, gain, expo_time); int4 short_parsed = parse_10bit(short_dat[i], short_extra_dat[i], aligned10);
v_rows[ROW_READ_ORDER[3]] = val4_from_10x2(dat[3], extra[3], short_dat[3], short_extra[3], aligned10, gain, expo_time); v_rows[ROW_READ_ORDER[i]] = normalize_pv_hdr(parsed, short_parsed, vignette_factor, expo_time);
#else #else
v_rows[ROW_READ_ORDER[0]] = val4_from_12(dat[0], gain); // no HDR here
v_rows[ROW_READ_ORDER[1]] = val4_from_12(dat[1], gain); int4 parsed = parse_12bit(dat[i]);
v_rows[ROW_READ_ORDER[2]] = val4_from_12(dat[2], gain); v_rows[ROW_READ_ORDER[i]] = normalize_pv(parsed, vignette_factor);
v_rows[ROW_READ_ORDER[3]] = val4_from_12(dat[3], gain); #endif
#endif }
// mirror padding // mirror padding
if (gid_x == 0) { if (gid_x == 0) {
@ -205,6 +164,7 @@ __kernel void process_raw(const __global uchar * in, __global uchar * out, int e
v_rows[3].s3 = v_rows[3].s1; v_rows[3].s3 = v_rows[3].s1;
} }
// debayering
// a simplified version of https://opensignalprocessingjournal.com/contents/volumes/V6/TOSIGPJ-6-1/TOSIGPJ-6-1.pdf // a simplified version of https://opensignalprocessingjournal.com/contents/volumes/V6/TOSIGPJ-6-1/TOSIGPJ-6-1.pdf
const float k01 = get_k(v_rows[0].s0, v_rows[1].s1, v_rows[0].s2, v_rows[1].s1); const float k01 = get_k(v_rows[0].s0, v_rows[1].s1, v_rows[0].s2, v_rows[1].s1);
const float k02 = get_k(v_rows[0].s2, v_rows[1].s1, v_rows[2].s2, v_rows[1].s1); const float k02 = get_k(v_rows[0].s2, v_rows[1].s1, v_rows[2].s2, v_rows[1].s1);
@ -242,7 +202,7 @@ __kernel void process_raw(const __global uchar * in, __global uchar * out, int e
rgb_tmp.z = (k32*v_rows[2].s3+k34*v_rows[2].s1)/(k32+k34); // B_G2 rgb_tmp.z = (k32*v_rows[2].s3+k34*v_rows[2].s1)/(k32+k34); // B_G2
rgb_out[RGB_WRITE_ORDER[3]] = convert_uchar3_sat(apply_gamma(color_correct(clamp(rgb_tmp, 0.0, 1.0)), expo_time) * 255.0); rgb_out[RGB_WRITE_ORDER[3]] = convert_uchar3_sat(apply_gamma(color_correct(clamp(rgb_tmp, 0.0, 1.0)), expo_time) * 255.0);
// write ys // rgb2yuv(nv12)
uchar2 yy = (uchar2)( uchar2 yy = (uchar2)(
RGB_TO_Y(rgb_out[0].s0, rgb_out[0].s1, rgb_out[0].s2), RGB_TO_Y(rgb_out[0].s0, rgb_out[0].s1, rgb_out[0].s2),
RGB_TO_Y(rgb_out[1].s0, rgb_out[1].s1, rgb_out[1].s2) RGB_TO_Y(rgb_out[1].s0, rgb_out[1].s1, rgb_out[1].s2)
@ -254,7 +214,6 @@ __kernel void process_raw(const __global uchar * in, __global uchar * out, int e
); );
vstore2(yy, 0, out + mad24(gid_y * 2 + 1, YUV_STRIDE, gid_x * 2)); vstore2(yy, 0, out + mad24(gid_y * 2 + 1, YUV_STRIDE, gid_x * 2));
// write uvs
const short ar = AVERAGE(rgb_out[0].s0, rgb_out[1].s0, rgb_out[2].s0, rgb_out[3].s0); const short ar = AVERAGE(rgb_out[0].s0, rgb_out[1].s0, rgb_out[2].s0, rgb_out[3].s0);
const short ag = AVERAGE(rgb_out[0].s1, rgb_out[1].s1, rgb_out[2].s1, rgb_out[3].s1); const short ag = AVERAGE(rgb_out[0].s1, rgb_out[1].s1, rgb_out[2].s1, rgb_out[3].s1);
const short ab = AVERAGE(rgb_out[0].s2, rgb_out[1].s2, rgb_out[2].s2, rgb_out[3].s2); const short ab = AVERAGE(rgb_out[0].s2, rgb_out[1].s2, rgb_out[2].s2, rgb_out[3].s2);

7
system/camerad/sensors/ar0231_cl.h
Unescape View File

@ -1,7 +1,14 @@
#if IS_AR #if IS_AR
#define BIT_DEPTH 12 #define BIT_DEPTH 12
#define PV_MAX 4096
#define BLACK_LVL 168 #define BLACK_LVL 168
#define VIGNETTE_RSZ 1.0f
float normalize_pv(int4 parsed, float vignette_factor) {
float4 pv = (convert_float4(parsed) - BLACK_LVL) / (PV_MAX - BLACK_LVL);
return clamp(pv*vignette_factor, 0.0, 1.0);
}
float3 color_correct(float3 rgb) { float3 color_correct(float3 rgb) {
float3 corrected = rgb.x * (float3)(1.82717181, -0.31231438, 0.07307673); float3 corrected = rgb.x * (float3)(1.82717181, -0.31231438, 0.07307673);

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

@ -3,7 +3,39 @@
#define BGGR #define BGGR
#define BIT_DEPTH 10 #define BIT_DEPTH 10
#define PV_MAX 65536 // gamma curve is calibrated to 16bit
#define BLACK_LVL 64 #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, 61376.0);
float svd = sv * fmin(expo_time, 8.0) / 8;
if (expo_time > 64) {
if (lv < 1023 - BLACK_LVL) {
return lv / (PV_MAX - BLACK_LVL);
} else {
return (svc / 64) / (PV_MAX - BLACK_LVL);
}
} else {
if (lv > 32) {
return (lv * 64 / fmax(expo_time, 8.0)) / (PV_MAX - BLACK_LVL);
} else {
return svd / (PV_MAX - BLACK_LVL);
}
}
}
float 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 color_correct(float3 rgb) {
float3 corrected = rgb.x * (float3)(1.55361989, -0.268894615, -0.000593219); float3 corrected = rgb.x * (float3)(1.55361989, -0.268894615, -0.000593219);
@ -16,7 +48,7 @@ float3 apply_gamma(float3 rgb, int expo_time) {
float s = log2((float)expo_time); float s = log2((float)expo_time);
if (s < 6) {s = fmin(12.0 - s, 9.0);} if (s < 6) {s = fmin(12.0 - s, 9.0);}
// log function adaptive to number of bits // log function adaptive to number of bits
return clamp(log(1 + rgb*65472.0) * (0.48*s*s - 12.92*s + 115.0) - (1.08*s*s - 29.2*s + 260.0), 0.0, 255.0) / 255.0; return clamp(log(1 + rgb*(PV_MAX - 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 #endif

7
system/camerad/sensors/ox03c10_cl.h
Unescape View File

@ -2,6 +2,7 @@
#define BIT_DEPTH 12 #define BIT_DEPTH 12
#define BLACK_LVL 64 #define BLACK_LVL 64
#define VIGNETTE_RSZ 1.0f
constant float ox03c10_lut[] = { constant float ox03c10_lut[] = {
0.0000e+00, 5.9488e-08, 1.1898e-07, 1.7846e-07, 2.3795e-07, 2.9744e-07, 3.5693e-07, 4.1642e-07, 4.7591e-07, 5.3539e-07, 5.9488e-07, 6.5437e-07, 7.1386e-07, 7.7335e-07, 8.3284e-07, 8.9232e-07, 9.5181e-07, 1.0113e-06, 1.0708e-06, 1.1303e-06, 1.1898e-06, 1.2493e-06, 1.3087e-06, 1.3682e-06, 1.4277e-06, 1.4872e-06, 1.5467e-06, 1.6062e-06, 1.6657e-06, 1.7252e-06, 1.7846e-06, 1.8441e-06, 1.9036e-06, 1.9631e-06, 2.0226e-06, 2.0821e-06, 2.1416e-06, 2.2011e-06, 2.2606e-06, 2.3200e-06, 2.3795e-06, 2.4390e-06, 2.4985e-06, 2.5580e-06, 2.6175e-06, 2.6770e-06, 2.7365e-06, 2.7959e-06, 2.8554e-06, 2.9149e-06, 2.9744e-06, 3.0339e-06, 3.0934e-06, 3.1529e-06, 3.2124e-06, 3.2719e-06, 3.3313e-06, 3.3908e-06, 3.4503e-06, 3.5098e-06, 3.5693e-06, 3.6288e-06, 3.6883e-06, 3.7478e-06, 3.8072e-06, 3.8667e-06, 3.9262e-06, 3.9857e-06, 4.0452e-06, 4.1047e-06, 4.1642e-06, 4.2237e-06, 4.2832e-06, 4.3426e-06, 4.4021e-06, 4.4616e-06, 4.5211e-06, 4.5806e-06, 4.6401e-06, 4.6996e-06, 4.7591e-06, 4.8185e-06, 4.8780e-06, 4.9375e-06, 4.9970e-06, 5.0565e-06, 5.1160e-06, 5.1755e-06, 5.2350e-06, 5.2945e-06, 5.3539e-06, 5.4134e-06, 5.4729e-06, 5.5324e-06, 5.5919e-06, 5.6514e-06, 5.7109e-06, 5.7704e-06, 5.8298e-06, 5.8893e-06, 5.9488e-06, 6.0083e-06, 6.0678e-06, 6.1273e-06, 6.1868e-06, 6.2463e-06, 6.3058e-06, 6.3652e-06, 6.4247e-06, 6.4842e-06, 6.5437e-06, 6.6032e-06, 6.6627e-06, 6.7222e-06, 6.7817e-06, 6.8411e-06, 6.9006e-06, 6.9601e-06, 7.0196e-06, 7.0791e-06, 7.1386e-06, 7.1981e-06, 7.2576e-06, 7.3171e-06, 7.3765e-06, 7.4360e-06, 7.4955e-06, 7.5550e-06, 7.6145e-06, 7.6740e-06, 7.7335e-06, 7.7930e-06, 7.8524e-06, 7.9119e-06, 7.9714e-06, 8.0309e-06, 8.0904e-06, 8.1499e-06, 8.2094e-06, 8.2689e-06, 8.3284e-06, 8.3878e-06, 8.4473e-06, 8.5068e-06, 8.5663e-06, 8.6258e-06, 8.6853e-06, 8.7448e-06, 8.8043e-06, 8.8637e-06, 8.9232e-06, 8.9827e-06, 9.0422e-06, 9.1017e-06, 9.1612e-06, 9.2207e-06, 9.2802e-06, 9.3397e-06, 9.3991e-06, 9.4586e-06, 9.5181e-06, 9.5776e-06, 9.6371e-06, 9.6966e-06, 9.7561e-06, 9.8156e-06, 9.8750e-06, 9.9345e-06, 9.9940e-06, 1.0054e-05, 1.0113e-05, 1.0172e-05, 1.0232e-05, 1.0291e-05, 1.0351e-05, 1.0410e-05, 1.0470e-05, 1.0529e-05, 1.0589e-05, 1.0648e-05, 1.0708e-05, 1.0767e-05, 1.0827e-05, 1.0886e-05, 1.0946e-05, 1.1005e-05, 1.1065e-05, 1.1124e-05, 1.1184e-05, 1.1243e-05, 1.1303e-05, 1.1362e-05, 1.1422e-05, 1.1481e-05, 1.1541e-05, 1.1600e-05, 1.1660e-05, 1.1719e-05, 1.1779e-05, 1.1838e-05, 1.1898e-05, 1.1957e-05, 1.2017e-05, 1.2076e-05, 1.2136e-05, 1.2195e-05, 1.2255e-05, 1.2314e-05, 1.2374e-05, 1.2433e-05, 1.2493e-05, 1.2552e-05, 1.2612e-05, 1.2671e-05, 1.2730e-05, 1.2790e-05, 1.2849e-05, 1.2909e-05, 1.2968e-05, 1.3028e-05, 1.3087e-05, 1.3147e-05, 1.3206e-05, 1.3266e-05, 1.3325e-05, 1.3385e-05, 1.3444e-05, 1.3504e-05, 1.3563e-05, 1.3623e-05, 1.3682e-05, 1.3742e-05, 1.3801e-05, 1.3861e-05, 1.3920e-05, 1.3980e-05, 1.4039e-05, 1.4099e-05, 1.4158e-05, 1.4218e-05, 1.4277e-05, 1.4337e-05, 1.4396e-05, 1.4456e-05, 1.4515e-05, 1.4575e-05, 1.4634e-05, 1.4694e-05, 1.4753e-05, 1.4813e-05, 1.4872e-05, 1.4932e-05, 1.4991e-05, 1.5051e-05, 1.5110e-05, 1.5169e-05, // NOLINT 0.0000e+00, 5.9488e-08, 1.1898e-07, 1.7846e-07, 2.3795e-07, 2.9744e-07, 3.5693e-07, 4.1642e-07, 4.7591e-07, 5.3539e-07, 5.9488e-07, 6.5437e-07, 7.1386e-07, 7.7335e-07, 8.3284e-07, 8.9232e-07, 9.5181e-07, 1.0113e-06, 1.0708e-06, 1.1303e-06, 1.1898e-06, 1.2493e-06, 1.3087e-06, 1.3682e-06, 1.4277e-06, 1.4872e-06, 1.5467e-06, 1.6062e-06, 1.6657e-06, 1.7252e-06, 1.7846e-06, 1.8441e-06, 1.9036e-06, 1.9631e-06, 2.0226e-06, 2.0821e-06, 2.1416e-06, 2.2011e-06, 2.2606e-06, 2.3200e-06, 2.3795e-06, 2.4390e-06, 2.4985e-06, 2.5580e-06, 2.6175e-06, 2.6770e-06, 2.7365e-06, 2.7959e-06, 2.8554e-06, 2.9149e-06, 2.9744e-06, 3.0339e-06, 3.0934e-06, 3.1529e-06, 3.2124e-06, 3.2719e-06, 3.3313e-06, 3.3908e-06, 3.4503e-06, 3.5098e-06, 3.5693e-06, 3.6288e-06, 3.6883e-06, 3.7478e-06, 3.8072e-06, 3.8667e-06, 3.9262e-06, 3.9857e-06, 4.0452e-06, 4.1047e-06, 4.1642e-06, 4.2237e-06, 4.2832e-06, 4.3426e-06, 4.4021e-06, 4.4616e-06, 4.5211e-06, 4.5806e-06, 4.6401e-06, 4.6996e-06, 4.7591e-06, 4.8185e-06, 4.8780e-06, 4.9375e-06, 4.9970e-06, 5.0565e-06, 5.1160e-06, 5.1755e-06, 5.2350e-06, 5.2945e-06, 5.3539e-06, 5.4134e-06, 5.4729e-06, 5.5324e-06, 5.5919e-06, 5.6514e-06, 5.7109e-06, 5.7704e-06, 5.8298e-06, 5.8893e-06, 5.9488e-06, 6.0083e-06, 6.0678e-06, 6.1273e-06, 6.1868e-06, 6.2463e-06, 6.3058e-06, 6.3652e-06, 6.4247e-06, 6.4842e-06, 6.5437e-06, 6.6032e-06, 6.6627e-06, 6.7222e-06, 6.7817e-06, 6.8411e-06, 6.9006e-06, 6.9601e-06, 7.0196e-06, 7.0791e-06, 7.1386e-06, 7.1981e-06, 7.2576e-06, 7.3171e-06, 7.3765e-06, 7.4360e-06, 7.4955e-06, 7.5550e-06, 7.6145e-06, 7.6740e-06, 7.7335e-06, 7.7930e-06, 7.8524e-06, 7.9119e-06, 7.9714e-06, 8.0309e-06, 8.0904e-06, 8.1499e-06, 8.2094e-06, 8.2689e-06, 8.3284e-06, 8.3878e-06, 8.4473e-06, 8.5068e-06, 8.5663e-06, 8.6258e-06, 8.6853e-06, 8.7448e-06, 8.8043e-06, 8.8637e-06, 8.9232e-06, 8.9827e-06, 9.0422e-06, 9.1017e-06, 9.1612e-06, 9.2207e-06, 9.2802e-06, 9.3397e-06, 9.3991e-06, 9.4586e-06, 9.5181e-06, 9.5776e-06, 9.6371e-06, 9.6966e-06, 9.7561e-06, 9.8156e-06, 9.8750e-06, 9.9345e-06, 9.9940e-06, 1.0054e-05, 1.0113e-05, 1.0172e-05, 1.0232e-05, 1.0291e-05, 1.0351e-05, 1.0410e-05, 1.0470e-05, 1.0529e-05, 1.0589e-05, 1.0648e-05, 1.0708e-05, 1.0767e-05, 1.0827e-05, 1.0886e-05, 1.0946e-05, 1.1005e-05, 1.1065e-05, 1.1124e-05, 1.1184e-05, 1.1243e-05, 1.1303e-05, 1.1362e-05, 1.1422e-05, 1.1481e-05, 1.1541e-05, 1.1600e-05, 1.1660e-05, 1.1719e-05, 1.1779e-05, 1.1838e-05, 1.1898e-05, 1.1957e-05, 1.2017e-05, 1.2076e-05, 1.2136e-05, 1.2195e-05, 1.2255e-05, 1.2314e-05, 1.2374e-05, 1.2433e-05, 1.2493e-05, 1.2552e-05, 1.2612e-05, 1.2671e-05, 1.2730e-05, 1.2790e-05, 1.2849e-05, 1.2909e-05, 1.2968e-05, 1.3028e-05, 1.3087e-05, 1.3147e-05, 1.3206e-05, 1.3266e-05, 1.3325e-05, 1.3385e-05, 1.3444e-05, 1.3504e-05, 1.3563e-05, 1.3623e-05, 1.3682e-05, 1.3742e-05, 1.3801e-05, 1.3861e-05, 1.3920e-05, 1.3980e-05, 1.4039e-05, 1.4099e-05, 1.4158e-05, 1.4218e-05, 1.4277e-05, 1.4337e-05, 1.4396e-05, 1.4456e-05, 1.4515e-05, 1.4575e-05, 1.4634e-05, 1.4694e-05, 1.4753e-05, 1.4813e-05, 1.4872e-05, 1.4932e-05, 1.4991e-05, 1.5051e-05, 1.5110e-05, 1.5169e-05, // NOLINT
@ -22,6 +23,12 @@ constant float ox03c10_lut[] = {
3.7695e-01, 3.7890e-01, 3.8086e-01, 3.8281e-01, 3.8476e-01, 3.8672e-01, 3.8867e-01, 3.9062e-01, 3.9258e-01, 3.9453e-01, 3.9648e-01, 3.9844e-01, 4.0039e-01, 4.0234e-01, 4.0430e-01, 4.0625e-01, 4.0820e-01, 4.1015e-01, 4.1211e-01, 4.1406e-01, 4.1601e-01, 4.1797e-01, 4.1992e-01, 4.2187e-01, 4.2383e-01, 4.2578e-01, 4.2773e-01, 4.2969e-01, 4.3164e-01, 4.3359e-01, 4.3555e-01, 4.3750e-01, 4.3945e-01, 4.4140e-01, 4.4336e-01, 4.4531e-01, 4.4726e-01, 4.4922e-01, 4.5117e-01, 4.5312e-01, 4.5508e-01, 4.5703e-01, 4.5898e-01, 4.6094e-01, 4.6289e-01, 4.6484e-01, 4.6680e-01, 4.6875e-01, 4.7070e-01, 4.7265e-01, 4.7461e-01, 4.7656e-01, 4.7851e-01, 4.8047e-01, 4.8242e-01, 4.8437e-01, 4.8633e-01, 4.8828e-01, 4.9023e-01, 4.9219e-01, 4.9414e-01, 4.9609e-01, 4.9805e-01, 5.0000e-01, 5.0195e-01, 5.0390e-01, 5.0586e-01, 5.0781e-01, 5.0976e-01, 5.1172e-01, 5.1367e-01, 5.1562e-01, 5.1758e-01, 5.1953e-01, 5.2148e-01, 5.2344e-01, 5.2539e-01, 5.2734e-01, 5.2930e-01, 5.3125e-01, 5.3320e-01, 5.3515e-01, 5.3711e-01, 5.3906e-01, 5.4101e-01, 5.4297e-01, 5.4492e-01, 5.4687e-01, 5.4883e-01, 5.5078e-01, 5.5273e-01, 5.5469e-01, 5.5664e-01, 5.5859e-01, 5.6055e-01, 5.6250e-01, 5.6445e-01, 5.6640e-01, 5.6836e-01, 5.7031e-01, 5.7226e-01, 5.7422e-01, 5.7617e-01, 5.7812e-01, 5.8008e-01, 5.8203e-01, 5.8398e-01, 5.8594e-01, 5.8789e-01, 5.8984e-01, 5.9180e-01, 5.9375e-01, 5.9570e-01, 5.9765e-01, 5.9961e-01, 6.0156e-01, 6.0351e-01, 6.0547e-01, 6.0742e-01, 6.0937e-01, 6.1133e-01, 6.1328e-01, 6.1523e-01, 6.1719e-01, 6.1914e-01, 6.2109e-01, 6.2305e-01, 6.2500e-01, 6.2695e-01, 6.2890e-01, 6.3086e-01, 6.3281e-01, 6.3476e-01, 6.3672e-01, 6.3867e-01, 6.4062e-01, 6.4258e-01, 6.4453e-01, 6.4648e-01, 6.4844e-01, 6.5039e-01, 6.5234e-01, 6.5430e-01, 6.5625e-01, 6.5820e-01, 6.6015e-01, 6.6211e-01, 6.6406e-01, 6.6601e-01, 6.6797e-01, 6.6992e-01, 6.7187e-01, 6.7383e-01, 6.7578e-01, 6.7773e-01, 6.7969e-01, 6.8164e-01, 6.8359e-01, 6.8554e-01, 6.8750e-01, 6.8945e-01, 6.9140e-01, 6.9336e-01, 6.9531e-01, 6.9726e-01, 6.9922e-01, 7.0117e-01, 7.0312e-01, 7.0508e-01, 7.0703e-01, 7.0898e-01, 7.1094e-01, 7.1289e-01, 7.1484e-01, 7.1679e-01, 7.1875e-01, 7.2070e-01, 7.2265e-01, 7.2461e-01, 7.2656e-01, 7.2851e-01, 7.3047e-01, 7.3242e-01, 7.3437e-01, 7.3633e-01, 7.3828e-01, 7.4023e-01, 7.4219e-01, 7.4414e-01, 7.4609e-01, 7.4804e-01, 7.5000e-01, 7.5390e-01, 7.5781e-01, 7.6172e-01, 7.6562e-01, 7.6953e-01, 7.7344e-01, 7.7734e-01, 7.8125e-01, 7.8515e-01, 7.8906e-01, 7.9297e-01, 7.9687e-01, 8.0078e-01, 8.0469e-01, 8.0859e-01, 8.1250e-01, 8.1640e-01, 8.2031e-01, 8.2422e-01, 8.2812e-01, 8.3203e-01, 8.3594e-01, 8.3984e-01, 8.4375e-01, 8.4765e-01, 8.5156e-01, 8.5547e-01, 8.5937e-01, 8.6328e-01, 8.6719e-01, 8.7109e-01, 8.7500e-01, 8.7890e-01, 8.8281e-01, 8.8672e-01, 8.9062e-01, 8.9453e-01, 8.9844e-01, 9.0234e-01, 9.0625e-01, 9.1015e-01, 9.1406e-01, 9.1797e-01, 9.2187e-01, 9.2578e-01, 9.2969e-01, 9.3359e-01, 9.3750e-01, 9.4140e-01, 9.4531e-01, 9.4922e-01, 9.5312e-01, 9.5703e-01, 9.6094e-01, 9.6484e-01, 9.6875e-01, 9.7265e-01, 9.7656e-01, 9.8047e-01, 9.8437e-01, 9.8828e-01, 9.9219e-01, 9.9609e-01, 1.0000e+00 // NOLINT 3.7695e-01, 3.7890e-01, 3.8086e-01, 3.8281e-01, 3.8476e-01, 3.8672e-01, 3.8867e-01, 3.9062e-01, 3.9258e-01, 3.9453e-01, 3.9648e-01, 3.9844e-01, 4.0039e-01, 4.0234e-01, 4.0430e-01, 4.0625e-01, 4.0820e-01, 4.1015e-01, 4.1211e-01, 4.1406e-01, 4.1601e-01, 4.1797e-01, 4.1992e-01, 4.2187e-01, 4.2383e-01, 4.2578e-01, 4.2773e-01, 4.2969e-01, 4.3164e-01, 4.3359e-01, 4.3555e-01, 4.3750e-01, 4.3945e-01, 4.4140e-01, 4.4336e-01, 4.4531e-01, 4.4726e-01, 4.4922e-01, 4.5117e-01, 4.5312e-01, 4.5508e-01, 4.5703e-01, 4.5898e-01, 4.6094e-01, 4.6289e-01, 4.6484e-01, 4.6680e-01, 4.6875e-01, 4.7070e-01, 4.7265e-01, 4.7461e-01, 4.7656e-01, 4.7851e-01, 4.8047e-01, 4.8242e-01, 4.8437e-01, 4.8633e-01, 4.8828e-01, 4.9023e-01, 4.9219e-01, 4.9414e-01, 4.9609e-01, 4.9805e-01, 5.0000e-01, 5.0195e-01, 5.0390e-01, 5.0586e-01, 5.0781e-01, 5.0976e-01, 5.1172e-01, 5.1367e-01, 5.1562e-01, 5.1758e-01, 5.1953e-01, 5.2148e-01, 5.2344e-01, 5.2539e-01, 5.2734e-01, 5.2930e-01, 5.3125e-01, 5.3320e-01, 5.3515e-01, 5.3711e-01, 5.3906e-01, 5.4101e-01, 5.4297e-01, 5.4492e-01, 5.4687e-01, 5.4883e-01, 5.5078e-01, 5.5273e-01, 5.5469e-01, 5.5664e-01, 5.5859e-01, 5.6055e-01, 5.6250e-01, 5.6445e-01, 5.6640e-01, 5.6836e-01, 5.7031e-01, 5.7226e-01, 5.7422e-01, 5.7617e-01, 5.7812e-01, 5.8008e-01, 5.8203e-01, 5.8398e-01, 5.8594e-01, 5.8789e-01, 5.8984e-01, 5.9180e-01, 5.9375e-01, 5.9570e-01, 5.9765e-01, 5.9961e-01, 6.0156e-01, 6.0351e-01, 6.0547e-01, 6.0742e-01, 6.0937e-01, 6.1133e-01, 6.1328e-01, 6.1523e-01, 6.1719e-01, 6.1914e-01, 6.2109e-01, 6.2305e-01, 6.2500e-01, 6.2695e-01, 6.2890e-01, 6.3086e-01, 6.3281e-01, 6.3476e-01, 6.3672e-01, 6.3867e-01, 6.4062e-01, 6.4258e-01, 6.4453e-01, 6.4648e-01, 6.4844e-01, 6.5039e-01, 6.5234e-01, 6.5430e-01, 6.5625e-01, 6.5820e-01, 6.6015e-01, 6.6211e-01, 6.6406e-01, 6.6601e-01, 6.6797e-01, 6.6992e-01, 6.7187e-01, 6.7383e-01, 6.7578e-01, 6.7773e-01, 6.7969e-01, 6.8164e-01, 6.8359e-01, 6.8554e-01, 6.8750e-01, 6.8945e-01, 6.9140e-01, 6.9336e-01, 6.9531e-01, 6.9726e-01, 6.9922e-01, 7.0117e-01, 7.0312e-01, 7.0508e-01, 7.0703e-01, 7.0898e-01, 7.1094e-01, 7.1289e-01, 7.1484e-01, 7.1679e-01, 7.1875e-01, 7.2070e-01, 7.2265e-01, 7.2461e-01, 7.2656e-01, 7.2851e-01, 7.3047e-01, 7.3242e-01, 7.3437e-01, 7.3633e-01, 7.3828e-01, 7.4023e-01, 7.4219e-01, 7.4414e-01, 7.4609e-01, 7.4804e-01, 7.5000e-01, 7.5390e-01, 7.5781e-01, 7.6172e-01, 7.6562e-01, 7.6953e-01, 7.7344e-01, 7.7734e-01, 7.8125e-01, 7.8515e-01, 7.8906e-01, 7.9297e-01, 7.9687e-01, 8.0078e-01, 8.0469e-01, 8.0859e-01, 8.1250e-01, 8.1640e-01, 8.2031e-01, 8.2422e-01, 8.2812e-01, 8.3203e-01, 8.3594e-01, 8.3984e-01, 8.4375e-01, 8.4765e-01, 8.5156e-01, 8.5547e-01, 8.5937e-01, 8.6328e-01, 8.6719e-01, 8.7109e-01, 8.7500e-01, 8.7890e-01, 8.8281e-01, 8.8672e-01, 8.9062e-01, 8.9453e-01, 8.9844e-01, 9.0234e-01, 9.0625e-01, 9.1015e-01, 9.1406e-01, 9.1797e-01, 9.2187e-01, 9.2578e-01, 9.2969e-01, 9.3359e-01, 9.3750e-01, 9.4140e-01, 9.4531e-01, 9.4922e-01, 9.5312e-01, 9.5703e-01, 9.6094e-01, 9.6484e-01, 9.6875e-01, 9.7265e-01, 9.7656e-01, 9.8047e-01, 9.8437e-01, 9.8828e-01, 9.9219e-01, 9.9609e-01, 1.0000e+00 // NOLINT
}; };
float normalize_pv(int4 parsed, float vignette_factor) {
// PWL
float4 pv = {ox03c10_lut[parsed.s0], ox03c10_lut[parsed.s1], ox03c10_lut[parsed.s2], ox03c10_lut[parsed.s3]};
return clamp(pv*vignette_factor*256.0, 0.0, 1.0);
}
float3 color_correct(float3 rgb) { float3 color_correct(float3 rgb) {
float3 corrected = rgb.x * (float3)(1.5664815, -0.29808738, -0.03973474); float3 corrected = rgb.x * (float3)(1.5664815, -0.29808738, -0.03973474);
corrected += rgb.y * (float3)(-0.48672447, 1.41914433, -0.40295248); corrected += rgb.y * (float3)(-0.48672447, 1.41914433, -0.40295248);

Write Preview
Loading…
Cancel
Save