diff --git a/system/camerad/cameras/process_raw.cl b/system/camerad/cameras/process_raw.cl index 6331e62c23..1c80b003c5 100644 --- a/system/camerad/cameras/process_raw.cl +++ b/system/camerad/cameras/process_raw.cl @@ -19,9 +19,6 @@ #endif float get_vignetting_s(float r) { - #if IS_OS - r = r / 2.2545f; - #endif if (r < 62500) { return (1.0f + 0.0000008f*r); } else if (r < 490000) { @@ -33,74 +30,26 @@ float get_vignetting_s(float r) { } } -float4 val4_from_12(uchar8 pvs, float gain) { - uint4 parsed = (uint4)(((uint)pvs.s0<<4) + (pvs.s1>>4), // is from the previous 10 bit - ((uint)pvs.s2<<4) + (pvs.s4&0xF), - ((uint)pvs.s3<<4) + (pvs.s4>>4), - ((uint)pvs.s5<<4) + (pvs.s7&0xF)); - #if IS_OX - // 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 - +int4 parse_12bit(uchar8 pvs) { + // lower bits scambled? + return (int4)(((int)pvs.s0<<4) + (pvs.s1>>4), + ((int)pvs.s2<<4) + (pvs.s4&0xF), + ((int)pvs.s3<<4) + (pvs.s4>>4), + ((int)pvs.s5<<4) + (pvs.s7&0xF)); } -float combine_pvs(float lv, float sv, int expo) { - 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; +int4 parse_10bit(uchar8 pvs, uchar ext, bool aligned) { if (aligned) { - parsed = (int8)(((int)long_pvs.s0 << 2) + (long_pvs.s1 & 0b00000011), - ((int)long_pvs.s2 << 2) + ((long_pvs.s6 & 0b11000000) / 64), - ((int)long_pvs.s3 << 2) + ((long_pvs.s6 & 0b00110000) / 16), - ((int)long_pvs.s4 << 2) + ((long_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)); + return (int4)(((int)pvs.s0 << 2) + (pvs.s1 & 0b00000011), + ((int)pvs.s2 << 2) + ((pvs.s6 & 0b11000000) / 64), + ((int)pvs.s3 << 2) + ((pvs.s6 & 0b00110000) / 16), + ((int)pvs.s4 << 2) + ((pvs.s6 & 0b00001100) / 4)); } else { - parsed = (int8)(((int)long_pvs.s0 << 2) + ((long_pvs.s3 & 0b00110000) / 16), - ((int)long_pvs.s1 << 2) + ((long_pvs.s3 & 0b00001100) / 4), - ((int)long_pvs.s2 << 2) + ((long_pvs.s3 & 0b00000011)), - ((int)long_pvs.s4 << 2) + ((long_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)); + return (int4)(((int)pvs.s0 << 2) + ((pvs.s3 & 0b00110000) / 16), + ((int)pvs.s1 << 2) + ((pvs.s3 & 0b00001100) / 4), + ((int)pvs.s2 << 2) + ((pvs.s3 & 0b00000011)), + ((int)pvs.s4 << 2) + ((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) { @@ -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_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_after_offset = (gid_y == (RGB_HEIGHT/2 - 1)) ? 1 : 3; float3 rgb_tmp; uchar3 rgb_out[4]; // output is 2x2 window + // read offset int start_idx; - #if BIT_DEPTH == 10 bool aligned10; 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); #endif - // read in 8x4 chars + // read in 4 rows, 8 uchars each uchar8 dat[4]; + // row_before dat[0] = vload8(0, in + start_idx + FRAME_STRIDE*row_before_offset); + // row_0 if (gid_x == 0 && gid_y == 0) { // this wasn't a problem due to extra rows 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 { dat[1] = vload8(0, in + start_idx + FRAME_STRIDE*1); } + // row_1 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 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[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 - // need extra bit for 10-bit - #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 - + // parse into floats 0.0-1.0 float4 v_rows[4]; - // parse into floats - #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); - v_rows[ROW_READ_ORDER[1]] = val4_from_10x2(dat[1], extra[1], short_dat[1], short_extra[1], aligned10, gain, expo_time); - v_rows[ROW_READ_ORDER[2]] = val4_from_10x2(dat[2], extra[2], short_dat[2], short_extra[2], aligned10, gain, expo_time); - v_rows[ROW_READ_ORDER[3]] = val4_from_10x2(dat[3], extra[3], short_dat[3], short_extra[3], aligned10, gain, expo_time); - #else - v_rows[ROW_READ_ORDER[0]] = val4_from_12(dat[0], gain); - v_rows[ROW_READ_ORDER[1]] = val4_from_12(dat[1], gain); - v_rows[ROW_READ_ORDER[2]] = val4_from_12(dat[2], gain); - v_rows[ROW_READ_ORDER[3]] = val4_from_12(dat[3], gain); - #endif + for (int i=0; i<4; i++) { + #if BIT_DEPTH == 10 + // for now it's always HDR + int4 parsed = parse_10bit(dat[i], extra_dat[i], aligned10); + int4 short_parsed = parse_10bit(short_dat[i], short_extra_dat[i], aligned10); + v_rows[ROW_READ_ORDER[i]] = normalize_pv_hdr(parsed, short_parsed, vignette_factor, expo_time); + #else + // no HDR here + int4 parsed = parse_12bit(dat[i]); + v_rows[ROW_READ_ORDER[i]] = normalize_pv(parsed, vignette_factor); + #endif + } // mirror padding 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; } + // debayering // 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 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_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)( 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) @@ -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)); - // write uvs 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 ab = AVERAGE(rgb_out[0].s2, rgb_out[1].s2, rgb_out[2].s2, rgb_out[3].s2); diff --git a/system/camerad/sensors/ar0231_cl.h b/system/camerad/sensors/ar0231_cl.h index 7ec3cf0f26..a70e60e6da 100644 --- a/system/camerad/sensors/ar0231_cl.h +++ b/system/camerad/sensors/ar0231_cl.h @@ -1,7 +1,14 @@ #if IS_AR #define BIT_DEPTH 12 +#define PV_MAX 4096 #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 corrected = rgb.x * (float3)(1.82717181, -0.31231438, 0.07307673); diff --git a/system/camerad/sensors/os04c10_cl.h b/system/camerad/sensors/os04c10_cl.h index 91d5c75a2b..95c12942c1 100644 --- a/system/camerad/sensors/os04c10_cl.h +++ b/system/camerad/sensors/os04c10_cl.h @@ -3,7 +3,39 @@ #define BGGR #define BIT_DEPTH 10 +#define PV_MAX 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, 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 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); if (s < 6) {s = fmin(12.0 - s, 9.0);} // 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 \ No newline at end of file diff --git a/system/camerad/sensors/ox03c10_cl.h b/system/camerad/sensors/ox03c10_cl.h index 7baaf23ce3..d7c5bc4a99 100644 --- a/system/camerad/sensors/ox03c10_cl.h +++ b/system/camerad/sensors/ox03c10_cl.h @@ -2,6 +2,7 @@ #define BIT_DEPTH 12 #define BLACK_LVL 64 +#define VIGNETTE_RSZ 1.0f 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 @@ -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 }; +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 corrected = rgb.x * (float3)(1.5664815, -0.29808738, -0.03973474); corrected += rgb.y * (float3)(-0.48672447, 1.41914433, -0.40295248);