mirror of
https://github.com/dragonpilot/dragonpilot.git
synced 2026-07-23 13:52:05 +08:00
C2: remove some dead code (#24256)
* remove some dead code * still needed Co-authored-by: Comma Device <device@comma.ai>
This commit is contained in:
George Hotz
@@ -43,7 +43,7 @@ public:
|
||||
ci->frame_width, ci->frame_height, ci->frame_stride,
|
||||
b->rgb_width, b->rgb_height, b->rgb_stride,
|
||||
ci->bayer_flip, ci->hdr, s->camera_num);
|
||||
const char *cl_file = Hardware::TICI() ? "cameras/real_debayer.cl" : "cameras/debayer.cl";
|
||||
const char *cl_file = "cameras/real_debayer.cl";
|
||||
cl_program prg_debayer = cl_program_from_file(context, device_id, cl_file, args);
|
||||
krnl_ = CL_CHECK_ERR(clCreateKernel(prg_debayer, "debayer10", &err));
|
||||
CL_CHECK(clReleaseProgram(prg_debayer));
|
||||
@@ -53,30 +53,13 @@ public:
|
||||
CL_CHECK(clSetKernelArg(krnl_, 0, sizeof(cl_mem), &cam_buf_cl));
|
||||
CL_CHECK(clSetKernelArg(krnl_, 1, sizeof(cl_mem), &buf_cl));
|
||||
|
||||
if (Hardware::TICI()) {
|
||||
const int debayer_local_worksize = 16;
|
||||
constexpr int localMemSize = (debayer_local_worksize + 2 * (3 / 2)) * (debayer_local_worksize + 2 * (3 / 2)) * sizeof(short int);
|
||||
const size_t globalWorkSize[] = {size_t(width), size_t(height)};
|
||||
const size_t localWorkSize[] = {debayer_local_worksize, debayer_local_worksize};
|
||||
CL_CHECK(clSetKernelArg(krnl_, 2, localMemSize, 0));
|
||||
CL_CHECK(clSetKernelArg(krnl_, 3, sizeof(float), &black_level));
|
||||
CL_CHECK(clEnqueueNDRangeKernel(q, krnl_, 2, NULL, globalWorkSize, localWorkSize, 0, 0, debayer_event));
|
||||
} else {
|
||||
if (hdr_) {
|
||||
// HDR requires a 1-D kernel due to the DPCM compression
|
||||
const size_t debayer_local_worksize = 128;
|
||||
const size_t debayer_work_size = height; // doesn't divide evenly, is this okay?
|
||||
CL_CHECK(clSetKernelArg(krnl_, 2, sizeof(float), &gain));
|
||||
CL_CHECK(clEnqueueNDRangeKernel(q, krnl_, 1, NULL, &debayer_work_size, &debayer_local_worksize, 0, 0, debayer_event));
|
||||
} else {
|
||||
const int debayer_local_worksize = 32;
|
||||
assert(width % 2 == 0);
|
||||
const size_t globalWorkSize[] = {size_t(height), size_t(width / 2)};
|
||||
const size_t localWorkSize[] = {debayer_local_worksize, debayer_local_worksize};
|
||||
CL_CHECK(clSetKernelArg(krnl_, 2, sizeof(float), &gain));
|
||||
CL_CHECK(clEnqueueNDRangeKernel(q, krnl_, 2, NULL, globalWorkSize, localWorkSize, 0, 0, debayer_event));
|
||||
}
|
||||
}
|
||||
const int debayer_local_worksize = 16;
|
||||
constexpr int localMemSize = (debayer_local_worksize + 2 * (3 / 2)) * (debayer_local_worksize + 2 * (3 / 2)) * sizeof(short int);
|
||||
const size_t globalWorkSize[] = {size_t(width), size_t(height)};
|
||||
const size_t localWorkSize[] = {debayer_local_worksize, debayer_local_worksize};
|
||||
CL_CHECK(clSetKernelArg(krnl_, 2, localMemSize, 0));
|
||||
CL_CHECK(clSetKernelArg(krnl_, 3, sizeof(float), &black_level));
|
||||
CL_CHECK(clEnqueueNDRangeKernel(q, krnl_, 2, NULL, globalWorkSize, localWorkSize, 0, 0, debayer_event));
|
||||
}
|
||||
|
||||
~Debayer() {
|
||||
@@ -173,7 +156,6 @@ bool CameraBuf::acquire() {
|
||||
#else
|
||||
if (camera_state->camera_id == CAMERA_ID_IMX390) black_level = 64.0;
|
||||
#endif
|
||||
|
||||
debayer->queue(q, camrabuf_cl, cur_rgb_buf->buf_cl, rgb_width, rgb_height, gain, black_level, &event);
|
||||
} else {
|
||||
assert(rgb_stride == camera_state->ci.frame_stride);
|
||||
|
||||
@@ -1,140 +0,0 @@
|
||||
const __constant float3 color_correction[3] = {
|
||||
// Matrix from WBraw -> sRGBD65 (normalized)
|
||||
(float3)( 1.62393627, -0.2092988, 0.00119886),
|
||||
(float3)(-0.45734315, 1.5534676, -0.59296798),
|
||||
(float3)(-0.16659312, -0.3441688, 1.59176912),
|
||||
};
|
||||
|
||||
float3 color_correct(float3 x) {
|
||||
float3 ret = (0,0,0);
|
||||
|
||||
// white balance of daylight
|
||||
x /= (float3)(0.4609375, 1.0, 0.546875);
|
||||
x = max(0.0, min(1.0, x));
|
||||
|
||||
// fix up the colors
|
||||
ret += x.x * color_correction[0];
|
||||
ret += x.y * color_correction[1];
|
||||
ret += x.z * color_correction[2];
|
||||
return ret;
|
||||
}
|
||||
|
||||
float3 srgb_gamma(float3 p) {
|
||||
// go all out and add an sRGB gamma curve
|
||||
const float3 ph = (1.0f + 0.055f)*pow(p, 1/2.4f) - 0.055f;
|
||||
const float3 pl = p*12.92f;
|
||||
return select(ph, pl, islessequal(p, 0.0031308f));
|
||||
}
|
||||
|
||||
#if HDR
|
||||
|
||||
__constant int dpcm_lookup[512] = {0, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15, 16, 17, 18, 19, 20, 21, 22, 23, 24, 25, 26, 27, 28, 29, 30, 31, 0, -1, -2, -3, -4, -5, -6, -7, -8, -9, -10, -11, -12, -13, -14, -15, -16, -17, -18, -19, -20, -21, -22, -23, -24, -25, -26, -27, -28, -29, -30, -31, 935, 951, 967, 983, 999, 1015, 1031, 1047, 1063, 1079, 1095, 1111, 1127, 1143, 1159, 1175, 1191, 1207, 1223, 1239, 1255, 1271, 1287, 1303, 1319, 1335, 1351, 1367, 1383, 1399, 1415, 1431, -935, -951, -967, -983, -999, -1015, -1031, -1047, -1063, -1079, -1095, -1111, -1127, -1143, -1159, -1175, -1191, -1207, -1223, -1239, -1255, -1271, -1287, -1303, -1319, -1335, -1351, -1367, -1383, -1399, -1415, -1431, 419, 427, 435, 443, 451, 459, 467, 475, 483, 491, 499, 507, 515, 523, 531, 539, 547, 555, 563, 571, 579, 587, 595, 603, 611, 619, 627, 635, 643, 651, 659, 667, 675, 683, 691, 699, 707, 715, 723, 731, 739, 747, 755, 763, 771, 779, 787, 795, 803, 811, 819, 827, 835, 843, 851, 859, 867, 875, 883, 891, 899, 907, 915, 923, -419, -427, -435, -443, -451, -459, -467, -475, -483, -491, -499, -507, -515, -523, -531, -539, -547, -555, -563, -571, -579, -587, -595, -603, -611, -619, -627, -635, -643, -651, -659, -667, -675, -683, -691, -699, -707, -715, -723, -731, -739, -747, -755, -763, -771, -779, -787, -795, -803, -811, -819, -827, -835, -843, -851, -859, -867, -875, -883, -891, -899, -907, -915, -923, 161, 165, 169, 173, 177, 181, 185, 189, 193, 197, 201, 205, 209, 213, 217, 221, 225, 229, 233, 237, 241, 245, 249, 253, 257, 261, 265, 269, 273, 277, 281, 285, 289, 293, 297, 301, 305, 309, 313, 317, 321, 325, 329, 333, 337, 341, 345, 349, 353, 357, 361, 365, 369, 373, 377, 381, 385, 389, 393, 397, 401, 405, 409, 413, -161, -165, -169, -173, -177, -181, -185, -189, -193, -197, -201, -205, -209, -213, -217, -221, -225, -229, -233, -237, -241, -245, -249, -253, -257, -261, -265, -269, -273, -277, -281, -285, -289, -293, -297, -301, -305, -309, -313, -317, -321, -325, -329, -333, -337, -341, -345, -349, -353, -357, -361, -365, -369, -373, -377, -381, -385, -389, -393, -397, -401, -405, -409, -413, 32, 34, 36, 38, 40, 42, 44, 46, 48, 50, 52, 54, 56, 58, 60, 62, 64, 66, 68, 70, 72, 74, 76, 78, 80, 82, 84, 86, 88, 90, 92, 94, 96, 98, 100, 102, 104, 106, 108, 110, 112, 114, 116, 118, 120, 122, 124, 126, 128, 130, 132, 134, 136, 138, 140, 142, 144, 146, 148, 150, 152, 154, 156, 158, -32, -34, -36, -38, -40, -42, -44, -46, -48, -50, -52, -54, -56, -58, -60, -62, -64, -66, -68, -70, -72, -74, -76, -78, -80, -82, -84, -86, -88, -90, -92, -94, -96, -98, -100, -102, -104, -106, -108, -110, -112, -114, -116, -118, -120, -122, -124, -126, -128, -130, -132, -134, -136, -138, -140, -142, -144, -146, -148, -150, -152, -154, -156, -158};
|
||||
|
||||
inline uint4 decompress(uint4 p, uint4 pl) {
|
||||
uint4 r1 = (pl + (uint4)(dpcm_lookup[p.s0], dpcm_lookup[p.s1], dpcm_lookup[p.s2], dpcm_lookup[p.s3]));
|
||||
uint4 r2 = ((p-0x200)<<5) | 0xF;
|
||||
r2 += select((uint4)(0,0,0,0), (uint4)(1,1,1,1), r2 <= pl);
|
||||
return select(r2, r1, p < 0x200);
|
||||
}
|
||||
|
||||
#endif
|
||||
|
||||
__kernel void debayer10(__global uchar const * const in,
|
||||
__global uchar * out, float digital_gain)
|
||||
{
|
||||
const int oy = get_global_id(0);
|
||||
if (oy >= RGB_HEIGHT) return;
|
||||
const int iy = oy * 2;
|
||||
|
||||
#if HDR
|
||||
uint4 pint_last;
|
||||
for (int ox = 0; ox < RGB_WIDTH; ox += 2) {
|
||||
#else
|
||||
int ox = get_global_id(1) * 2;
|
||||
{
|
||||
#endif
|
||||
const int ix = (ox/2) * 5;
|
||||
|
||||
// TODO: why doesn't this work for the frontview
|
||||
/*const uchar8 v1 = vload8(0, &in[iy * FRAME_STRIDE + ix]);
|
||||
const uchar ex1 = v1.s4;
|
||||
const uchar8 v2 = vload8(0, &in[(iy+1) * FRAME_STRIDE + ix]);
|
||||
const uchar ex2 = v2.s4;*/
|
||||
|
||||
const uchar4 v1 = vload4(0, &in[iy * FRAME_STRIDE + ix]);
|
||||
const uchar ex1 = in[iy * FRAME_STRIDE + ix + 4];
|
||||
const uchar4 v2 = vload4(0, &in[(iy+1) * FRAME_STRIDE + ix]);
|
||||
const uchar ex2 = in[(iy+1) * FRAME_STRIDE + ix + 4];
|
||||
|
||||
uint4 pinta[2];
|
||||
pinta[0] = (uint4)(
|
||||
(((uint)v1.s0 << 2) + ( (ex1 >> 0) & 3)),
|
||||
(((uint)v1.s1 << 2) + ( (ex1 >> 2) & 3)),
|
||||
(((uint)v2.s0 << 2) + ( (ex2 >> 0) & 3)),
|
||||
(((uint)v2.s1 << 2) + ( (ex2 >> 2) & 3)));
|
||||
pinta[1] = (uint4)(
|
||||
(((uint)v1.s2 << 2) + ( (ex1 >> 4) & 3)),
|
||||
(((uint)v1.s3 << 2) + ( (ex1 >> 6) & 3)),
|
||||
(((uint)v2.s2 << 2) + ( (ex2 >> 4) & 3)),
|
||||
(((uint)v2.s3 << 2) + ( (ex2 >> 6) & 3)));
|
||||
|
||||
#pragma unroll
|
||||
for (uint px = 0; px < 2; px++) {
|
||||
uint4 pint = pinta[px];
|
||||
|
||||
#if HDR
|
||||
// decompress HDR
|
||||
pint = (ox == 0 && px == 0) ? ((pint<<4) | 8) : decompress(pint, pint_last);
|
||||
pint_last = pint;
|
||||
#endif
|
||||
|
||||
float4 p = convert_float4(pint);
|
||||
|
||||
// 64 is the black level of the sensor, remove
|
||||
// (changed to 56 for HDR)
|
||||
const float black_level = 56.0f;
|
||||
// TODO: switch to max here?
|
||||
p = (p - black_level);
|
||||
|
||||
// correct vignetting (no pow function?)
|
||||
// see https://www.eecis.udel.edu/~jye/lab_research/09/JiUp.pdf the A (4th order)
|
||||
const float r = ((oy - RGB_HEIGHT/2)*(oy - RGB_HEIGHT/2) + (ox - RGB_WIDTH/2)*(ox - RGB_WIDTH/2));
|
||||
const float fake_f = 700.0f; // should be 910, but this fits...
|
||||
const float lil_a = (1.0f + r/(fake_f*fake_f));
|
||||
p = p * lil_a * lil_a;
|
||||
|
||||
// rescale to 1.0
|
||||
#if HDR
|
||||
p /= (16384.0f-black_level);
|
||||
#else
|
||||
p /= (1024.0f-black_level);
|
||||
#endif
|
||||
|
||||
// digital gain
|
||||
p *= digital_gain;
|
||||
|
||||
// use both green channels
|
||||
#if BAYER_FLIP == 3
|
||||
float3 c1 = (float3)(p.s3, (p.s1+p.s2)/2.0f, p.s0);
|
||||
#elif BAYER_FLIP == 2
|
||||
float3 c1 = (float3)(p.s2, (p.s0+p.s3)/2.0f, p.s1);
|
||||
#elif BAYER_FLIP == 1
|
||||
float3 c1 = (float3)(p.s1, (p.s0+p.s3)/2.0f, p.s2);
|
||||
#elif BAYER_FLIP == 0
|
||||
float3 c1 = (float3)(p.s0, (p.s1+p.s2)/2.0f, p.s3);
|
||||
#endif
|
||||
|
||||
// color correction
|
||||
c1 = color_correct(c1);
|
||||
|
||||
#if HDR
|
||||
// srgb gamma isn't right for YUV, so it's disabled for now
|
||||
c1 = srgb_gamma(c1);
|
||||
#endif
|
||||
|
||||
// output BGR
|
||||
const int ooff = oy * RGB_STRIDE/3 + ox;
|
||||
vstore3(convert_uchar3_sat(c1.zyx * 255.0f), ooff+px, out);
|
||||
}
|
||||
}
|
||||
}
|
||||
@@ -232,13 +232,8 @@ OmxEncoder::OmxEncoder(const char* filename, CameraType type, int in_width, int
|
||||
|
||||
if (h265) {
|
||||
// setup HEVC
|
||||
#ifndef QCOM2
|
||||
OMX_VIDEO_PARAM_HEVCTYPE hevc_type = {0};
|
||||
OMX_INDEXTYPE index_type = (OMX_INDEXTYPE) OMX_IndexParamVideoHevc;
|
||||
#else
|
||||
OMX_VIDEO_PARAM_PROFILELEVELTYPE hevc_type = {0};
|
||||
OMX_INDEXTYPE index_type = OMX_IndexParamVideoProfileLevelCurrent;
|
||||
#endif
|
||||
hevc_type.nSize = sizeof(hevc_type);
|
||||
hevc_type.nPortIndex = (OMX_U32) PORT_INDEX_OUT;
|
||||
OMX_CHECK(OMX_GetParameter(this->handle, index_type, (OMX_PTR) &hevc_type));
|
||||
@@ -566,11 +561,6 @@ void OmxEncoder::encoder_open(const char* path) {
|
||||
if (this->write) {
|
||||
this->of = util::safe_fopen(this->vid_path, "wb");
|
||||
assert(this->of);
|
||||
#ifndef QCOM2
|
||||
if (this->codec_config_len > 0) {
|
||||
util::safe_fwrite(this->codec_config, 1, this->codec_config_len, this->of);
|
||||
}
|
||||
#endif
|
||||
}
|
||||
}
|
||||
|
||||
|
||||
Reference in New Issue
Block a user