From b0cf2d791eda3b2de5cd8c6709a83b74f6c3c37d Mon Sep 17 00:00:00 2001 From: Comma Device Date: Fri, 23 Feb 2024 17:46:15 -0600 Subject: [PATCH] real 10 --- system/camerad/cameras/real_debayer.cl | 58 +++++++++++++++++++++++--- 1 file changed, 52 insertions(+), 6 deletions(-) diff --git a/system/camerad/cameras/real_debayer.cl b/system/camerad/cameras/real_debayer.cl index 3d1c4ba9d7..ad1c6b1bea 100644 --- a/system/camerad/cameras/real_debayer.cl +++ b/system/camerad/cameras/real_debayer.cl @@ -87,6 +87,24 @@ float4 val4_from_12(uchar8 pvs, float gain) { } +float4 val4_from_10(uchar8 pvs, uchar ext, bool aligned, float gain) { + uint4 parsed; + if (aligned) { + parsed = (uint4)(((uint)pvs.s0 << 2) + (pvs.s1 & 0b00000011), + ((uint)pvs.s2 << 2) + ((pvs.s6 & 0b11000000) / 64), + ((uint)pvs.s3 << 2) + ((pvs.s6 & 0b00110000) / 16), + ((uint)pvs.s4 << 2) + ((pvs.s6 & 0b00001100) / 4)); + } else { + parsed = (uint4)(((uint)pvs.s0 << 2) + ((pvs.s3 & 0b00110000) / 16), + ((uint)pvs.s1 << 2) + ((pvs.s3 & 0b00001100) / 4), + ((uint)pvs.s2 << 2) + ((pvs.s3 & 0b00000011)), + ((uint)pvs.s4 << 2) + ((ext & 0b11000000) / 64)); + } + + float4 pv = convert_float4(parsed) / 1024.0; + return clamp(pv*gain, 0.0, 1.0); +} + float get_k(float a, float b, float c, float d) { return 2.0 - (fabs(a - b) + fabs(c - d)); } @@ -102,14 +120,35 @@ __kernel void debayer10(const __global uchar * in, __global uchar * out) float3 rgb; uchar3 rgb_out[4]; - int start = (2 * gid_y - 1) * FRAME_STRIDE + (3 * gid_x - 2) + (FRAME_STRIDE * FRAME_OFFSET); - + int start; + #if IS_OS + bool aligned10; + + if (gid_x % 2 == 0) { + aligned10 = true; + start = (2 * gid_y - 1) * FRAME_STRIDE + (5 * gid_x / 2 - 2) + (FRAME_STRIDE * FRAME_OFFSET); + } else { + aligned10 = false; + start = (2 * gid_y - 1) * FRAME_STRIDE + (5 * (gid_x - 1) / 2 + 1) + (FRAME_STRIDE * FRAME_OFFSET); + } + #else + start = (2 * gid_y - 1) * FRAME_STRIDE + (3 * gid_x - 2) + (FRAME_STRIDE * FRAME_OFFSET); + #endif // read in 8x4 chars uchar8 dat[4]; + uchar extra[4]; dat[0] = vload8(0, in + start + FRAME_STRIDE*y_top_mod); dat[1] = vload8(0, in + start + FRAME_STRIDE*1); dat[2] = vload8(0, in + start + FRAME_STRIDE*2); dat[3] = vload8(0, in + start + FRAME_STRIDE*y_bot_mod); + #if IS_OS + if (!aligned10) { + extra[0] = in[start + FRAME_STRIDE*y_top_mod + 8]; + extra[1] = in[start + FRAME_STRIDE*1 + 8]; + extra[2] = in[start + FRAME_STRIDE*2 + 8]; + extra[3] = in[start + FRAME_STRIDE*y_bot_mod + 8]; + } + #endif // correct vignetting #if VIGNETTING @@ -121,10 +160,17 @@ __kernel void debayer10(const __global uchar * in, __global uchar * out) #endif // process them to floats - float4 va = val4_from_12(dat[0], gain); - float4 vb = val4_from_12(dat[1], gain); - float4 vc = val4_from_12(dat[2], gain); - float4 vd = val4_from_12(dat[3], gain); + #if IS_OS + float4 va = val4_from_10(dat[0], extra[0], aligned10, 1.0); + float4 vb = val4_from_10(dat[1], extra[1], aligned10, 1.0); + float4 vc = val4_from_10(dat[2], extra[2], aligned10, 1.0); + float4 vd = val4_from_10(dat[3], extra[3], aligned10, 1.0); + #else + float4 va = val4_from_12(dat[0], gain); + float4 vb = val4_from_12(dat[1], gain); + float4 vc = val4_from_12(dat[2], gain); + float4 vd = val4_from_12(dat[3], gain); + #endif if (gid_x == 0) { va.s0 = va.s2;