From 5acd765c1eb9bbed59c94732b1c8c91c4dc70f7b Mon Sep 17 00:00:00 2001 From: Adeeb Shihadeh Date: Fri, 1 Dec 2023 20:10:07 -0800 Subject: [PATCH] camerad cleanup (#30573) * misc cleanup * rm those * rm utils * fix build * rm pool * little more * goodbye imgproc old-commit-hash: e34ee43eeae000d080cb923b722afaf4c06fbc59 --- docs/c_docs.rst | 2 - release/files_common | 5 -- system/camerad/SConscript | 9 +- system/camerad/cameras/camera_common.cc | 6 -- system/camerad/cameras/camera_common.h | 16 +--- system/camerad/cameras/camera_qcom2.cc | 82 +++++++++--------- system/camerad/cameras/camera_qcom2.h | 3 + system/camerad/cameras/camera_util.cc | 14 +-- system/camerad/cameras/camera_util.h | 2 +- system/camerad/imgproc/conv.cl | 110 ------------------------ system/camerad/imgproc/pool.cl | 34 -------- system/camerad/imgproc/utils.cc | 106 ----------------------- system/camerad/imgproc/utils.h | 37 -------- system/camerad/main.cc | 2 +- system/camerad/snapshot/snapshot.py | 1 - 15 files changed, 59 insertions(+), 370 deletions(-) delete mode 100644 system/camerad/imgproc/conv.cl delete mode 100644 system/camerad/imgproc/pool.cl delete mode 100644 system/camerad/imgproc/utils.cc delete mode 100644 system/camerad/imgproc/utils.h diff --git a/docs/c_docs.rst b/docs/c_docs.rst index 5619cc8a51..0b9d23972e 100644 --- a/docs/c_docs.rst +++ b/docs/c_docs.rst @@ -29,8 +29,6 @@ camerad ^^^^^^^ .. autodoxygenindex:: :project: system_camerad_cameras -.. autodoxygenindex:: - :project: system_camerad_imgproc locationd ^^^^^^^^^ diff --git a/release/files_common b/release/files_common index cfc3830e0d..af5d156d41 100644 --- a/release/files_common +++ b/release/files_common @@ -335,11 +335,6 @@ system/camerad/cameras/camera_common.h system/camerad/cameras/camera_common.cc system/camerad/cameras/sensor2_i2c.h -system/camerad/imgproc/conv.cl -system/camerad/imgproc/pool.cl -system/camerad/imgproc/utils.cc -system/camerad/imgproc/utils.h - selfdrive/manager/__init__.py selfdrive/manager/build.py selfdrive/manager/helpers.py diff --git a/system/camerad/SConscript b/system/camerad/SConscript index ffd7278bbb..83f93344d8 100644 --- a/system/camerad/SConscript +++ b/system/camerad/SConscript @@ -3,12 +3,7 @@ Import('env', 'arch', 'cereal', 'messaging', 'common', 'gpucommon', 'visionipc') libs = ['m', 'pthread', common, 'jpeg', 'OpenCL', 'yuv', cereal, messaging, 'zmq', 'capnp', 'kj', visionipc, gpucommon, 'atomic'] camera_obj = env.Object(['cameras/camera_qcom2.cc', 'cameras/camera_common.cc', 'cameras/camera_util.cc']) -env.Program('camerad', [ - 'main.cc', - camera_obj, - ], LIBS=libs) +env.Program('camerad', ['main.cc', camera_obj], LIBS=libs) if GetOption("extras") and arch == "x86_64": - env.Program('test/ae_gray_test', - ['test/ae_gray_test.cc', camera_obj], - LIBS=libs) + env.Program('test/ae_gray_test', ['test/ae_gray_test.cc', camera_obj], LIBS=libs) diff --git a/system/camerad/cameras/camera_common.cc b/system/camerad/cameras/camera_common.cc index d2c7a35f6e..987ccf23da 100644 --- a/system/camerad/cameras/camera_common.cc +++ b/system/camerad/cameras/camera_common.cc @@ -11,7 +11,6 @@ #include "third_party/libyuv/include/libyuv.h" #include -#include "system/camerad/imgproc/utils.h" #include "common/clutil.h" #include "common/swaglog.h" #include "common/util.h" @@ -93,12 +92,8 @@ void CameraBuf::init(cl_device_id device_id, cl_context context, CameraState *s, debayer = new Debayer(device_id, context, this, s, nv12_width, nv12_uv_offset); -#ifdef __APPLE__ - q = CL_CHECK_ERR(clCreateCommandQueue(context, device_id, 0, &err)); -#else 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)); -#endif } CameraBuf::~CameraBuf() { @@ -281,7 +276,6 @@ float set_exposure_target(const CameraBuf *b, int x_start, int x_end, int x_skip } } - // Find mean lumimance value unsigned int lum_cur = 0; for (lum_med = 255; lum_med >= 0; lum_med--) { diff --git a/system/camerad/cameras/camera_common.h b/system/camerad/cameras/camera_common.h index 0f69aa3774..23bfc8b7d5 100644 --- a/system/camerad/cameras/camera_common.h +++ b/system/camerad/cameras/camera_common.h @@ -14,17 +14,9 @@ #include "common/swaglog.h" #include "system/hardware/hw.h" -#define CAMERA_ID_IMX298 0 -#define CAMERA_ID_IMX179 1 -#define CAMERA_ID_S5K3P8SP 2 -#define CAMERA_ID_OV8865 3 -#define CAMERA_ID_IMX298_FLIPPED 4 -#define CAMERA_ID_OV10640 5 -#define CAMERA_ID_LGC920 6 -#define CAMERA_ID_LGC615 7 -#define CAMERA_ID_AR0231 8 -#define CAMERA_ID_OX03C10 9 -#define CAMERA_ID_MAX 10 +#define CAMERA_ID_AR0231 0 +#define CAMERA_ID_OX03C10 1 +#define CAMERA_ID_MAX 2 const int YUV_BUFFER_COUNT = 20; @@ -55,7 +47,7 @@ typedef struct FrameMetadata { uint32_t frame_id; // Timestamps - uint64_t timestamp_sof; // only set on tici + uint64_t timestamp_sof; uint64_t timestamp_eof; // Exposure diff --git a/system/camerad/cameras/camera_qcom2.cc b/system/camerad/cameras/camera_qcom2.cc index 3a9ecb467b..8d20ed27c8 100644 --- a/system/camerad/cameras/camera_qcom2.cc +++ b/system/camerad/cameras/camera_qcom2.cc @@ -665,51 +665,51 @@ void CameraState::camera_open(MultiCameraState *multi_cam_state_, int camera_num if (!enabled) return; struct cam_isp_in_port_info in_port_info = { - .res_type = (uint32_t[]){CAM_ISP_IFE_IN_RES_PHY_0, CAM_ISP_IFE_IN_RES_PHY_1, CAM_ISP_IFE_IN_RES_PHY_2}[camera_num], + .res_type = (uint32_t[]){CAM_ISP_IFE_IN_RES_PHY_0, CAM_ISP_IFE_IN_RES_PHY_1, CAM_ISP_IFE_IN_RES_PHY_2}[camera_num], - .lane_type = CAM_ISP_LANE_TYPE_DPHY, - .lane_num = 4, - .lane_cfg = 0x3210, + .lane_type = CAM_ISP_LANE_TYPE_DPHY, + .lane_num = 4, + .lane_cfg = 0x3210, - .vc = 0x0, - .dt = dt, + .vc = 0x0, + .dt = dt, + .format = CAM_FORMAT_MIPI_RAW_12, + + .test_pattern = 0x2, // 0x3? + .usage_type = 0x0, + + .left_start = 0, + .left_stop = ci.frame_width - 1, + .left_width = ci.frame_width, + + .right_start = 0, + .right_stop = ci.frame_width - 1, + .right_width = ci.frame_width, + + .line_start = 0, + .line_stop = ci.frame_height + ci.extra_height - 1, + .height = ci.frame_height + ci.extra_height, + + .pixel_clk = 0x0, + .batch_size = 0x0, + .dsp_mode = CAM_ISP_DSP_MODE_NONE, + .hbi_cnt = 0x0, + .custom_csid = 0x0, + + .num_out_res = 0x1, + .data[0] = (struct cam_isp_out_port_info){ + .res_type = CAM_ISP_IFE_OUT_RES_RDI_0, .format = CAM_FORMAT_MIPI_RAW_12, - - .test_pattern = 0x2, // 0x3? - .usage_type = 0x0, - - .left_start = 0, - .left_stop = ci.frame_width - 1, - .left_width = ci.frame_width, - - .right_start = 0, - .right_stop = ci.frame_width - 1, - .right_width = ci.frame_width, - - .line_start = 0, - .line_stop = ci.frame_height + ci.extra_height - 1, + .width = ci.frame_width, .height = ci.frame_height + ci.extra_height, - - .pixel_clk = 0x0, - .batch_size = 0x0, - .dsp_mode = CAM_ISP_DSP_MODE_NONE, - .hbi_cnt = 0x0, - .custom_csid = 0x0, - - .num_out_res = 0x1, - .data[0] = (struct cam_isp_out_port_info){ - .res_type = CAM_ISP_IFE_OUT_RES_RDI_0, - .format = CAM_FORMAT_MIPI_RAW_12, - .width = ci.frame_width, - .height = ci.frame_height + ci.extra_height, - .comp_grp_id = 0x0, .split_point = 0x0, .secure_mode = 0x0, - }, + .comp_grp_id = 0x0, .split_point = 0x0, .secure_mode = 0x0, + }, }; struct cam_isp_resource isp_resource = { - .resource_id = CAM_ISP_RES_ID_PORT, - .handle_type = CAM_HANDLE_USER_POINTER, - .res_hdl = (uint64_t)&in_port_info, - .length = sizeof(in_port_info), + .resource_id = CAM_ISP_RES_ID_PORT, + .handle_type = CAM_HANDLE_USER_POINTER, + .res_hdl = (uint64_t)&in_port_info, + .length = sizeof(in_port_info), }; auto isp_dev_handle_ = device_acquire(multi_cam_state->isp_fd, session_handle, &isp_resource); @@ -1098,8 +1098,8 @@ void CameraState::set_camera_exposure(float grey_frac) { std::string gain_bytes, time_bytes; if (env_ctrl_exp_from_params) { - gain_bytes = Params().get("CameraDebugExpGain"); - time_bytes = Params().get("CameraDebugExpTime"); + gain_bytes = params.get("CameraDebugExpGain"); + time_bytes = params.get("CameraDebugExpTime"); } if (gain_bytes.size() > 0 && time_bytes.size() > 0) { diff --git a/system/camerad/cameras/camera_qcom2.h b/system/camerad/cameras/camera_qcom2.h index 9e0109ab20..1e7cbfc104 100644 --- a/system/camerad/cameras/camera_qcom2.h +++ b/system/camerad/cameras/camera_qcom2.h @@ -107,6 +107,9 @@ private: // Register parsing std::map> ar0231_register_lut; std::map> ar0231_build_register_lut(uint8_t *data); + + // for debugging + Params params; }; typedef struct MultiCameraState { diff --git a/system/camerad/cameras/camera_util.cc b/system/camerad/cameras/camera_util.cc index 74c81a878a..b2cf4115dc 100644 --- a/system/camerad/cameras/camera_util.cc +++ b/system/camerad/cameras/camera_util.cc @@ -30,10 +30,10 @@ int do_cam_control(int fd, int op_code, void *handle, int size) { std::optional device_acquire(int fd, int32_t session_handle, void *data, uint32_t num_resources) { struct cam_acquire_dev_cmd cmd = { - .session_handle = session_handle, - .handle_type = CAM_HANDLE_USER_POINTER, - .num_resources = (uint32_t)(data ? num_resources : 0), - .resource_hdl = (uint64_t)data, + .session_handle = session_handle, + .handle_type = CAM_HANDLE_USER_POINTER, + .num_resources = (uint32_t)(data ? num_resources : 0), + .resource_hdl = (uint64_t)data, }; int err = do_cam_control(fd, CAM_ACQUIRE_DEV, &cmd, sizeof(cmd)); return err == 0 ? std::make_optional(cmd.dev_handle) : std::nullopt; @@ -41,9 +41,9 @@ std::optional device_acquire(int fd, int32_t session_handle, void *data int device_config(int fd, int32_t session_handle, int32_t dev_handle, uint64_t packet_handle) { struct cam_config_dev_cmd cmd = { - .session_handle = session_handle, - .dev_handle = dev_handle, - .packet_handle = packet_handle, + .session_handle = session_handle, + .dev_handle = dev_handle, + .packet_handle = packet_handle, }; return do_cam_control(fd, CAM_CONFIG_DEV, &cmd, sizeof(cmd)); } diff --git a/system/camerad/cameras/camera_util.h b/system/camerad/cameras/camera_util.h index b36f404c0f..891ce3e793 100644 --- a/system/camerad/cameras/camera_util.h +++ b/system/camerad/cameras/camera_util.h @@ -1,6 +1,6 @@ #pragma once - #include +#include #include #include #include diff --git a/system/camerad/imgproc/conv.cl b/system/camerad/imgproc/conv.cl deleted file mode 100644 index a7115ae76c..0000000000 --- a/system/camerad/imgproc/conv.cl +++ /dev/null @@ -1,110 +0,0 @@ -// const __constant float3 rgb_weights = (0.299, 0.587, 0.114); // opencv rgb2gray weights -// const __constant float3 bgr_weights = (0.114, 0.587, 0.299); // bgr2gray weights - -// convert input rgb image to single channel then conv -__kernel void rgb2gray_conv2d( - const __global uchar * input, - __global short * output, - __constant short * filter, - __local uchar3 * cached -) -{ - const int rowOffset = get_global_id(1) * IMAGE_W; - const int my = get_global_id(0) + rowOffset; - - const int localRowLen = TWICE_HALF_FILTER_SIZE + get_local_size(0); - const int localRowOffset = ( get_local_id(1) + HALF_FILTER_SIZE ) * localRowLen; - const int myLocal = localRowOffset + get_local_id(0) + HALF_FILTER_SIZE; - - // cache local pixels - cached[ myLocal ].x = input[ my * 3 ]; // r - cached[ myLocal ].y = input[ my * 3 + 1]; // g - cached[ myLocal ].z = input[ my * 3 + 2]; // b - - // pad - if ( - get_global_id(0) < HALF_FILTER_SIZE || - get_global_id(0) > IMAGE_W - HALF_FILTER_SIZE - 1 || - get_global_id(1) < HALF_FILTER_SIZE || - get_global_id(1) > IMAGE_H - HALF_FILTER_SIZE - 1 - ) - { - barrier(CLK_LOCAL_MEM_FENCE); - return; - } - else - { - int localColOffset = -1; - int globalColOffset = -1; - - // cache extra - if ( get_local_id(0) < HALF_FILTER_SIZE ) - { - localColOffset = get_local_id(0); - globalColOffset = -HALF_FILTER_SIZE; - - cached[ localRowOffset + get_local_id(0) ].x = input[ my * 3 - HALF_FILTER_SIZE * 3 ]; - cached[ localRowOffset + get_local_id(0) ].y = input[ my * 3 - HALF_FILTER_SIZE * 3 + 1]; - cached[ localRowOffset + get_local_id(0) ].z = input[ my * 3 - HALF_FILTER_SIZE * 3 + 2]; - } - else if ( get_local_id(0) >= get_local_size(0) - HALF_FILTER_SIZE ) - { - localColOffset = get_local_id(0) + TWICE_HALF_FILTER_SIZE; - globalColOffset = HALF_FILTER_SIZE; - - cached[ myLocal + HALF_FILTER_SIZE ].x = input[ my * 3 + HALF_FILTER_SIZE * 3 ]; - cached[ myLocal + HALF_FILTER_SIZE ].y = input[ my * 3 + HALF_FILTER_SIZE * 3 + 1]; - cached[ myLocal + HALF_FILTER_SIZE ].z = input[ my * 3 + HALF_FILTER_SIZE * 3 + 2]; - } - - - if ( get_local_id(1) < HALF_FILTER_SIZE ) - { - cached[ get_local_id(1) * localRowLen + get_local_id(0) + HALF_FILTER_SIZE ].x = input[ my * 3 - HALF_FILTER_SIZE_IMAGE_W * 3 ]; - cached[ get_local_id(1) * localRowLen + get_local_id(0) + HALF_FILTER_SIZE ].y = input[ my * 3 - HALF_FILTER_SIZE_IMAGE_W * 3 + 1]; - cached[ get_local_id(1) * localRowLen + get_local_id(0) + HALF_FILTER_SIZE ].z = input[ my * 3 - HALF_FILTER_SIZE_IMAGE_W * 3 + 2]; - if (localColOffset > 0) - { - cached[ get_local_id(1) * localRowLen + localColOffset ].x = input[ my * 3 - HALF_FILTER_SIZE_IMAGE_W * 3 + globalColOffset * 3]; - cached[ get_local_id(1) * localRowLen + localColOffset ].y = input[ my * 3 - HALF_FILTER_SIZE_IMAGE_W * 3 + globalColOffset * 3 + 1]; - cached[ get_local_id(1) * localRowLen + localColOffset ].z = input[ my * 3 - HALF_FILTER_SIZE_IMAGE_W * 3 + globalColOffset * 3 + 2]; - } - } - else if ( get_local_id(1) >= get_local_size(1) -HALF_FILTER_SIZE ) - { - int offset = ( get_local_id(1) + TWICE_HALF_FILTER_SIZE ) * localRowLen; - cached[ offset + get_local_id(0) + HALF_FILTER_SIZE ].x = input[ my * 3 + HALF_FILTER_SIZE_IMAGE_W * 3 ]; - cached[ offset + get_local_id(0) + HALF_FILTER_SIZE ].y = input[ my * 3 + HALF_FILTER_SIZE_IMAGE_W * 3 + 1]; - cached[ offset + get_local_id(0) + HALF_FILTER_SIZE ].z = input[ my * 3 + HALF_FILTER_SIZE_IMAGE_W * 3 + 2]; - if (localColOffset > 0) - { - cached[ offset + localColOffset ].x = input[ my * 3 + HALF_FILTER_SIZE_IMAGE_W * 3 + globalColOffset * 3]; - cached[ offset + localColOffset ].y = input[ my * 3 + HALF_FILTER_SIZE_IMAGE_W * 3 + globalColOffset * 3 + 1]; - cached[ offset + localColOffset ].z = input[ my * 3 + HALF_FILTER_SIZE_IMAGE_W * 3 + globalColOffset * 3 + 2]; - } - } - - // sync - barrier(CLK_LOCAL_MEM_FENCE); - - // perform convolution - int fIndex = 0; - short sum = 0; - - for (int r = -HALF_FILTER_SIZE; r <= HALF_FILTER_SIZE; r++) - { - int curRow = r * localRowLen; - for (int c = -HALF_FILTER_SIZE; c <= HALF_FILTER_SIZE; c++, fIndex++) - { - if (!FLIP_RB){ - // sum += dot(rgb_weights, cached[ myLocal + curRow + c ]) * filter[ fIndex ]; - sum += (cached[ myLocal + curRow + c ].x / 3 + cached[ myLocal + curRow + c ].y / 2 + cached[ myLocal + curRow + c ].z / 9) * filter[ fIndex ]; - } else { - // sum += dot(bgr_weights, cached[ myLocal + curRow + c ]) * filter[ fIndex ]; - sum += (cached[ myLocal + curRow + c ].x / 9 + cached[ myLocal + curRow + c ].y / 2 + cached[ myLocal + curRow + c ].z / 3) * filter[ fIndex ]; - } - } - } - output[my] = sum; - } -} \ No newline at end of file diff --git a/system/camerad/imgproc/pool.cl b/system/camerad/imgproc/pool.cl deleted file mode 100644 index d674b5f363..0000000000 --- a/system/camerad/imgproc/pool.cl +++ /dev/null @@ -1,34 +0,0 @@ -// calculate variance in each subregion -__kernel void var_pool( - const __global char * input, - __global ushort * output // should not be larger than 128*128 so uint16 -) -{ - const int xidx = get_global_id(0) + ROI_X_MIN; - const int yidx = get_global_id(1) + ROI_Y_MIN; - - const int size = X_PITCH * Y_PITCH; - - float fsum = 0; - char mean, max; - - for (int i = 0; i < size; i++) { - int x_offset = i % X_PITCH; - int y_offset = i / X_PITCH; - fsum += input[xidx*X_PITCH + yidx*Y_PITCH*FULL_STRIDE_X + x_offset + y_offset*FULL_STRIDE_X]; - max = input[xidx*X_PITCH + yidx*Y_PITCH*FULL_STRIDE_X + x_offset + y_offset*FULL_STRIDE_X]>max ? input[xidx*X_PITCH + yidx*Y_PITCH*FULL_STRIDE_X + x_offset + y_offset*FULL_STRIDE_X]:max; - } - - mean = convert_char_rte(fsum / size); - - float fvar = 0; - for (int i = 0; i < size; i++) { - int x_offset = i % X_PITCH; - int y_offset = i / X_PITCH; - fvar += (input[xidx*X_PITCH + yidx*Y_PITCH*FULL_STRIDE_X + x_offset + y_offset*FULL_STRIDE_X] - mean) * (input[xidx*X_PITCH + yidx*Y_PITCH*FULL_STRIDE_X + x_offset + y_offset*FULL_STRIDE_X] - mean); - } - - fvar = fvar / size; - - output[(xidx-ROI_X_MIN)+(yidx-ROI_Y_MIN)*(ROI_X_MAX-ROI_X_MIN+1)] = convert_ushort_rte(5 * fvar + convert_float_rte(max)); -} \ No newline at end of file diff --git a/system/camerad/imgproc/utils.cc b/system/camerad/imgproc/utils.cc deleted file mode 100644 index a7bbeb9e86..0000000000 --- a/system/camerad/imgproc/utils.cc +++ /dev/null @@ -1,106 +0,0 @@ -#include "system/camerad/imgproc/utils.h" - -#include -#include -#include -#include -#include - -const int16_t lapl_conv_krnl[9] = {0, 1, 0, - 1, -4, 1, - 0, 1, 0}; - -// calculate score based on laplacians in one area -uint16_t get_lapmap_one(const int16_t *lap, int x_pitch, int y_pitch) { - const int size = x_pitch * y_pitch; - // avg and max of roi - int16_t max = 0; - int sum = 0; - for (int i = 0; i < size; ++i) { - const int16_t v = lap[i]; - sum += v; - if (v > max) max = v; - } - - const int16_t mean = sum / size; - - // var of roi - int var = 0; - for (int i = 0; i < size; ++i) { - var += std::pow(lap[i] - mean, 2); - } - - const float fvar = (float)var / size; - return std::min(5 * fvar + max, (float)65535); -} - -bool is_blur(const uint16_t *lapmap, const size_t size) { - float bad_sum = 0; - for (int i = 0; i < size; i++) { - if (lapmap[i] < LM_THRESH) { - bad_sum += 1 / (float)size; - } - } - return (bad_sum > LM_PREC_THRESH); -} - -static cl_program build_conv_program(cl_device_id device_id, cl_context context, int image_w, int image_h, int filter_size) { - char args[4096]; - snprintf(args, sizeof(args), - "-cl-fast-relaxed-math -cl-denorms-are-zero " - "-DIMAGE_W=%d -DIMAGE_H=%d -DFLIP_RB=%d " - "-DFILTER_SIZE=%d -DHALF_FILTER_SIZE=%d -DTWICE_HALF_FILTER_SIZE=%d -DHALF_FILTER_SIZE_IMAGE_W=%d", - image_w, image_h, 1, - filter_size, filter_size/2, (filter_size/2)*2, (filter_size/2)*image_w); - return cl_program_from_file(context, device_id, "imgproc/conv.cl", args); -} - -LapConv::LapConv(cl_device_id device_id, cl_context ctx, int rgb_width, int rgb_height, int rgb_stride, int filter_size) - : width(rgb_width / NUM_SEGMENTS_X), height(rgb_height / NUM_SEGMENTS_Y), rgb_stride(rgb_stride), - roi_buf(width * height * 3), result_buf(width * height) { - - prg = build_conv_program(device_id, ctx, width, height, filter_size); - krnl = CL_CHECK_ERR(clCreateKernel(prg, "rgb2gray_conv2d", &err)); - // TODO: Removed CL_MEM_SVM_FINE_GRAIN_BUFFER, confirm it doesn't matter - roi_cl = CL_CHECK_ERR(clCreateBuffer(ctx, CL_MEM_READ_WRITE, roi_buf.size() * sizeof(roi_buf[0]), NULL, &err)); - result_cl = CL_CHECK_ERR(clCreateBuffer(ctx, CL_MEM_READ_WRITE, result_buf.size() * sizeof(result_buf[0]), NULL, &err)); - filter_cl = CL_CHECK_ERR(clCreateBuffer(ctx, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, - 9 * sizeof(int16_t), (void *)&lapl_conv_krnl, &err)); -} - -LapConv::~LapConv() { - CL_CHECK(clReleaseMemObject(roi_cl)); - CL_CHECK(clReleaseMemObject(result_cl)); - CL_CHECK(clReleaseMemObject(filter_cl)); - CL_CHECK(clReleaseKernel(krnl)); - CL_CHECK(clReleaseProgram(prg)); -} - -uint16_t LapConv::Update(cl_command_queue q, const uint8_t *rgb_buf, const int roi_id) { - // sharpness scores - const int x_offset = ROI_X_MIN + roi_id % (ROI_X_MAX - ROI_X_MIN + 1); - const int y_offset = ROI_Y_MIN + roi_id / (ROI_X_MAX - ROI_X_MIN + 1); - - const uint8_t *rgb_offset = rgb_buf + y_offset * height * rgb_stride + x_offset * width * 3; - for (int i = 0; i < height; ++i) { - memcpy(&roi_buf[i * width * 3], &rgb_offset[i * rgb_stride], width * 3); - } - - constexpr int local_mem_size = (CONV_LOCAL_WORKSIZE + 2 * (3 / 2)) * (CONV_LOCAL_WORKSIZE + 2 * (3 / 2)) * (3 * sizeof(uint8_t)); - const size_t global_work_size[] = {(size_t)width, (size_t)height}; - const size_t local_work_size[] = {CONV_LOCAL_WORKSIZE, CONV_LOCAL_WORKSIZE}; - - CL_CHECK(clEnqueueWriteBuffer(q, roi_cl, CL_TRUE, 0, roi_buf.size() * sizeof(roi_buf[0]), roi_buf.data(), 0, 0, 0)); - CL_CHECK(clSetKernelArg(krnl, 0, sizeof(cl_mem), (void *)&roi_cl)); - CL_CHECK(clSetKernelArg(krnl, 1, sizeof(cl_mem), (void *)&result_cl)); - CL_CHECK(clSetKernelArg(krnl, 2, sizeof(cl_mem), (void *)&filter_cl)); - CL_CHECK(clSetKernelArg(krnl, 3, local_mem_size, 0)); - cl_event conv_event; - CL_CHECK(clEnqueueNDRangeKernel(q, krnl, 2, NULL, global_work_size, local_work_size, 0, 0, &conv_event)); - CL_CHECK(clWaitForEvents(1, &conv_event)); - CL_CHECK(clReleaseEvent(conv_event)); - CL_CHECK(clEnqueueReadBuffer(q, result_cl, CL_TRUE, 0, - result_buf.size() * sizeof(result_buf[0]), result_buf.data(), 0, 0, 0)); - - return get_lapmap_one(result_buf.data(), width, height); -} diff --git a/system/camerad/imgproc/utils.h b/system/camerad/imgproc/utils.h deleted file mode 100644 index 94323b15c5..0000000000 --- a/system/camerad/imgproc/utils.h +++ /dev/null @@ -1,37 +0,0 @@ -#pragma once - -#include -#include -#include - -#include "common/clutil.h" - -#define NUM_SEGMENTS_X 8 -#define NUM_SEGMENTS_Y 6 - -#define ROI_X_MIN 1 -#define ROI_X_MAX 6 -#define ROI_Y_MIN 2 -#define ROI_Y_MAX 3 - -#define LM_THRESH 120 -#define LM_PREC_THRESH 0.9 // 90 perc is blur -#define CONV_LOCAL_WORKSIZE 16 - -class LapConv { -public: - LapConv(cl_device_id device_id, cl_context ctx, int rgb_width, int rgb_height, int rgb_stride, int filter_size); - ~LapConv(); - uint16_t Update(cl_command_queue q, const uint8_t *rgb_buf, const int roi_id); - -private: - cl_mem roi_cl, result_cl, filter_cl; - cl_program prg; - cl_kernel krnl; - const int width, height; - const int rgb_stride; - std::vector roi_buf; - std::vector result_buf; -}; - -bool is_blur(const uint16_t *lapmap, const size_t size); diff --git a/system/camerad/main.cc b/system/camerad/main.cc index 35a3329f30..19de21c9bb 100644 --- a/system/camerad/main.cc +++ b/system/camerad/main.cc @@ -8,7 +8,7 @@ int main(int argc, char *argv[]) { if (Hardware::PC()) { - printf("camerad is not meant to run on PC\n"); + printf("exiting, camerad is not meant to run on PC\n"); return 0; } diff --git a/system/camerad/snapshot/snapshot.py b/system/camerad/snapshot/snapshot.py index ede97d1a79..8c1b6084c7 100755 --- a/system/camerad/snapshot/snapshot.py +++ b/system/camerad/snapshot/snapshot.py @@ -13,7 +13,6 @@ from openpilot.system.hardware import PC from openpilot.selfdrive.controls.lib.alertmanager import set_offroad_alert from openpilot.selfdrive.manager.process_config import managed_processes -LM_THRESH = 120 # defined in system/camerad/imgproc/utils.h VISION_STREAMS = { "roadCameraState": VisionStreamType.VISION_STREAM_ROAD,