diff --git a/release/files_tici b/release/files_tici index 1771c45138..18860e20af 100644 --- a/release/files_tici +++ b/release/files_tici @@ -7,7 +7,7 @@ system/camerad/cameras/camera_qcom2.cc system/camerad/cameras/camera_qcom2.h system/camerad/cameras/camera_util.cc system/camerad/cameras/camera_util.h -system/camerad/cameras/real_debayer.cl +system/camerad/cameras/process_raw.cl system/qcomgpsd/* diff --git a/selfdrive/test/test_onroad.py b/selfdrive/test/test_onroad.py index 250534bf86..8124073e5f 100755 --- a/selfdrive/test/test_onroad.py +++ b/selfdrive/test/test_onroad.py @@ -304,7 +304,7 @@ class TestOnroad(unittest.TestCase): def test_camera_processing_time(self): result = "\n" result += "------------------------------------------------\n" - result += "-------------- Debayer Timing ------------------\n" + result += "-------------- ImgProc Timing ------------------\n" result += "------------------------------------------------\n" ts = [getattr(m, m.which()).processingTime for m in self.lr if 'CameraState' in m.which()] diff --git a/system/camerad/cameras/camera_common.cc b/system/camerad/cameras/camera_common.cc index b6e7f0ae66..13ee28ed30 100644 --- a/system/camerad/cameras/camera_common.cc +++ b/system/camerad/cameras/camera_common.cc @@ -18,9 +18,9 @@ ExitHandler do_exit; -class Debayer { +class ImgProc { public: - Debayer(cl_device_id device_id, cl_context context, const CameraBuf *b, const CameraState *s, int buf_width, int uv_offset) { + ImgProc(cl_device_id device_id, cl_context context, const CameraBuf *b, const CameraState *s, int buf_width, int uv_offset) { char args[4096]; const SensorInfo *ci = s->ci.get(); snprintf(args, sizeof(args), @@ -35,25 +35,25 @@ public: ci->mipi_format == CAM_FORMAT_MIPI_RAW_10, ci->hdr_offset > 0, ci->hdr_offset, ci->bggr, s->camera_num == 1); - 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)); + const char *cl_file = "cameras/process_raw.cl"; + cl_program prg_imgproc = cl_program_from_file(context, device_id, cl_file, args); + krnl_ = CL_CHECK_ERR(clCreateKernel(prg_imgproc, "process_raw", &err)); + CL_CHECK(clReleaseProgram(prg_imgproc)); } - void queue(cl_command_queue q, cl_mem cam_buf_cl, cl_mem buf_cl, int width, int height, cl_event *debayer_event, int expo_time) { + void queue(cl_command_queue q, cl_mem cam_buf_cl, cl_mem buf_cl, int width, int height, cl_event *imgproc_event, int expo_time) { CL_CHECK(clSetKernelArg(krnl_, 0, sizeof(cl_mem), &cam_buf_cl)); CL_CHECK(clSetKernelArg(krnl_, 1, sizeof(cl_mem), &buf_cl)); CL_CHECK(clSetKernelArg(krnl_, 2, sizeof(cl_int), &expo_time)); const size_t globalWorkSize[] = {size_t(width / 2), size_t(height / 2)}; - const int debayer_local_worksize = 16; - const size_t localWorkSize[] = {debayer_local_worksize, debayer_local_worksize}; - CL_CHECK(clEnqueueNDRangeKernel(q, krnl_, 2, NULL, globalWorkSize, localWorkSize, 0, 0, debayer_event)); + const int imgproc_local_worksize = 16; + const size_t localWorkSize[] = {imgproc_local_worksize, imgproc_local_worksize}; + CL_CHECK(clEnqueueNDRangeKernel(q, krnl_, 2, NULL, globalWorkSize, localWorkSize, 0, 0, imgproc_event)); } - ~Debayer() { + ~ImgProc() { CL_CHECK(clReleaseKernel(krnl_)); } @@ -94,7 +94,7 @@ void CameraBuf::init(cl_device_id device_id, cl_context context, CameraState *s, vipc_server->create_buffers_with_sizes(stream_type, YUV_BUFFER_COUNT, false, rgb_width, rgb_height, nv12_size, nv12_width, nv12_uv_offset); LOGD("created %d YUV vipc buffers with size %dx%d", YUV_BUFFER_COUNT, nv12_width, nv12_height); - debayer = new Debayer(device_id, context, this, s, nv12_width, nv12_uv_offset); + imgproc = new ImgProc(device_id, context, this, s, nv12_width, nv12_uv_offset); const cl_queue_properties props[] = {0}; //CL_QUEUE_PRIORITY_KHR, CL_QUEUE_PRIORITY_HIGH_KHR, 0}; q = CL_CHECK_ERR(clCreateCommandQueueWithProperties(context, device_id, props, &err)); @@ -104,7 +104,7 @@ CameraBuf::~CameraBuf() { for (int i = 0; i < frame_buf_count; i++) { camera_bufs[i].free(); } - if (debayer) delete debayer; + if (imgproc) delete imgproc; if (q) CL_CHECK(clReleaseCommandQueue(q)); } @@ -122,7 +122,7 @@ bool CameraBuf::acquire() { double start_time = millis_since_boot(); cl_event event; - debayer->queue(q, camera_bufs[cur_buf_idx].buf_cl, cur_yuv_buf->buf_cl, rgb_width, rgb_height, &event, cur_frame_data.integ_lines); + imgproc->queue(q, camera_bufs[cur_buf_idx].buf_cl, cur_yuv_buf->buf_cl, rgb_width, rgb_height, &event, cur_frame_data.integ_lines); clWaitForEvents(1, &event); CL_CHECK(clReleaseEvent(event)); cur_frame_data.processing_time = (millis_since_boot() - start_time) / 1000.0; diff --git a/system/camerad/cameras/camera_common.h b/system/camerad/cameras/camera_common.h index 2bdc4f8b17..774ab0b2e7 100644 --- a/system/camerad/cameras/camera_common.h +++ b/system/camerad/cameras/camera_common.h @@ -44,12 +44,12 @@ typedef struct FrameMetadata { struct MultiCameraState; class CameraState; -class Debayer; +class ImgProc; class CameraBuf { private: VisionIpcServer *vipc_server; - Debayer *debayer = nullptr; + ImgProc *imgproc = nullptr; VisionStreamType stream_type; int cur_buf_idx; SafeQueue safe_queue; diff --git a/system/camerad/cameras/real_debayer.cl b/system/camerad/cameras/process_raw.cl similarity index 99% rename from system/camerad/cameras/real_debayer.cl rename to system/camerad/cameras/process_raw.cl index f7fb6a002b..861be0d9a3 100644 --- a/system/camerad/cameras/real_debayer.cl +++ b/system/camerad/cameras/process_raw.cl @@ -168,7 +168,7 @@ float get_k(float a, float b, float c, float d) { return 2.0 - (fabs(a - b) + fabs(c - d)); } -__kernel void debayer10(const __global uchar * in, __global uchar * out, int expo_time) +__kernel void process_raw(const __global uchar * in, __global uchar * out, int expo_time) { const int gid_x = get_global_id(0); const int gid_y = get_global_id(1);