From 5a3c22d804155da72339b0dafbb567caffbc1721 Mon Sep 17 00:00:00 2001 From: Dean Lee Date: Tue, 13 Apr 2021 13:42:53 +0800 Subject: [PATCH] camerad: refactor RGBToYUVState into a class (#20310) * struct RGBToYUVState to class Rgb2Yuv * clFinish * blank line * rebase master * use event --- selfdrive/camerad/cameras/camera_common.cc | 6 ++-- selfdrive/camerad/cameras/camera_common.h | 2 +- selfdrive/camerad/transforms/rgb_to_yuv.cc | 41 ++++++++-------------- selfdrive/camerad/transforms/rgb_to_yuv.h | 27 ++++++-------- 4 files changed, 28 insertions(+), 48 deletions(-) diff --git a/selfdrive/camerad/cameras/camera_common.cc b/selfdrive/camerad/cameras/camera_common.cc index 7f78b2b9..5dcf9e4d 100644 --- a/selfdrive/camerad/cameras/camera_common.cc +++ b/selfdrive/camerad/cameras/camera_common.cc @@ -84,7 +84,7 @@ void CameraBuf::init(cl_device_id device_id, cl_context context, CameraState *s, CL_CHECK(clReleaseProgram(prg_debayer)); } - rgb_to_yuv_init(&rgb_to_yuv_state, context, device_id, rgb_width, rgb_height, rgb_stride); + rgb2yuv = std::make_unique(context, device_id, rgb_width, rgb_height, rgb_stride); #ifdef __APPLE__ q = CL_CHECK_ERR(clCreateCommandQueue(context, device_id, 0, &err)); @@ -99,7 +99,6 @@ CameraBuf::~CameraBuf() { camera_bufs[i].free(); } - if (rgb_to_yuv_state.rgb_to_yuv_krnl) rgb_to_yuv_destroy(&rgb_to_yuv_state); if (krnl_debayer) CL_CHECK(clReleaseKernel(krnl_debayer)); if (q) CL_CHECK(clReleaseCommandQueue(q)); } @@ -114,7 +113,6 @@ bool CameraBuf::acquire() { } cur_frame_data = camera_bufs_metadata[cur_buf_idx]; - cur_rgb_buf = vipc_server->get_buffer(rgb_type); cl_event debayer_event; @@ -151,7 +149,7 @@ bool CameraBuf::acquire() { CL_CHECK(clReleaseEvent(debayer_event)); cur_yuv_buf = vipc_server->get_buffer(yuv_type); - rgb_to_yuv_queue(&rgb_to_yuv_state, q, cur_rgb_buf->buf_cl, cur_yuv_buf->buf_cl); + rgb2yuv->queue(q, cur_rgb_buf->buf_cl, cur_yuv_buf->buf_cl); VisionIpcBufExtra extra = { cur_frame_data.frame_id, diff --git a/selfdrive/camerad/cameras/camera_common.h b/selfdrive/camerad/cameras/camera_common.h index 55c25214..b8e079d8 100644 --- a/selfdrive/camerad/cameras/camera_common.h +++ b/selfdrive/camerad/cameras/camera_common.h @@ -94,7 +94,7 @@ private: CameraState *camera_state; cl_kernel krnl_debayer; - RGBToYUVState rgb_to_yuv_state; + std::unique_ptr rgb2yuv; VisionStreamType rgb_type, yuv_type; diff --git a/selfdrive/camerad/transforms/rgb_to_yuv.cc b/selfdrive/camerad/transforms/rgb_to_yuv.cc index c0ac1920..ce144d34 100644 --- a/selfdrive/camerad/transforms/rgb_to_yuv.cc +++ b/selfdrive/camerad/transforms/rgb_to_yuv.cc @@ -1,17 +1,8 @@ -#include +#include "rgb_to_yuv.h" #include -#include "clutil.h" - -#include "rgb_to_yuv.h" - -void rgb_to_yuv_init(RGBToYUVState* s, cl_context ctx, cl_device_id device_id, int width, int height, int rgb_stride) { - memset(s, 0, sizeof(*s)); - printf("width %d, height %d, rgb_stride %d\n", width, height, rgb_stride); - assert(width % 2 == 0); - assert(height % 2 == 0); - s->width = width; - s->height = height; +Rgb2Yuv::Rgb2Yuv(cl_context ctx, cl_device_id device_id, int width, int height, int rgb_stride) { + assert(width % 2 == 0 && height % 2 == 0); char args[1024]; snprintf(args, sizeof(args), "-cl-fast-relaxed-math -cl-denorms-are-zero " @@ -19,27 +10,25 @@ void rgb_to_yuv_init(RGBToYUVState* s, cl_context ctx, cl_device_id device_id, i "-DCL_DEBUG " #endif "-DWIDTH=%d -DHEIGHT=%d -DUV_WIDTH=%d -DUV_HEIGHT=%d -DRGB_STRIDE=%d -DRGB_SIZE=%d", - width, height, width/ 2, height / 2, rgb_stride, width * height); + width, height, width / 2, height / 2, rgb_stride, width * height); + cl_program prg = cl_program_from_file(ctx, device_id, "transforms/rgb_to_yuv.cl", args); - - s->rgb_to_yuv_krnl = CL_CHECK_ERR(clCreateKernel(prg, "rgb_to_yuv", &err)); - // done with this + krnl = CL_CHECK_ERR(clCreateKernel(prg, "rgb_to_yuv", &err)); CL_CHECK(clReleaseProgram(prg)); + + work_size[0] = (width + (width % 4 == 0 ? 0 : (4 - width % 4))) / 4; + work_size[1] = (height + (height % 4 == 0 ? 0 : (4 - height % 4))) / 4; } -void rgb_to_yuv_destroy(RGBToYUVState* s) { - CL_CHECK(clReleaseKernel(s->rgb_to_yuv_krnl)); +Rgb2Yuv::~Rgb2Yuv() { + CL_CHECK(clReleaseKernel(krnl)); } -void rgb_to_yuv_queue(RGBToYUVState* s, cl_command_queue q, cl_mem rgb_cl, cl_mem yuv_cl) { - CL_CHECK(clSetKernelArg(s->rgb_to_yuv_krnl, 0, sizeof(cl_mem), &rgb_cl)); - CL_CHECK(clSetKernelArg(s->rgb_to_yuv_krnl, 1, sizeof(cl_mem), &yuv_cl)); - const size_t work_size[2] = { - (size_t)(s->width + (s->width % 4 == 0 ? 0 : (4 - s->width % 4))) / 4, - (size_t)(s->height + (s->height % 4 == 0 ? 0 : (4 - s->height % 4))) / 4 - }; +void Rgb2Yuv::queue(cl_command_queue q, cl_mem rgb_cl, cl_mem yuv_cl) { + CL_CHECK(clSetKernelArg(krnl, 0, sizeof(cl_mem), &rgb_cl)); + CL_CHECK(clSetKernelArg(krnl, 1, sizeof(cl_mem), &yuv_cl)); cl_event event; - CL_CHECK(clEnqueueNDRangeKernel(q, s->rgb_to_yuv_krnl, 2, NULL, &work_size[0], NULL, 0, 0, &event)); + CL_CHECK(clEnqueueNDRangeKernel(q, krnl, 2, NULL, &work_size[0], NULL, 0, 0, &event)); CL_CHECK(clWaitForEvents(1, &event)); CL_CHECK(clReleaseEvent(event)); } diff --git a/selfdrive/camerad/transforms/rgb_to_yuv.h b/selfdrive/camerad/transforms/rgb_to_yuv.h index 05d5c5ec..e1de180d 100644 --- a/selfdrive/camerad/transforms/rgb_to_yuv.h +++ b/selfdrive/camerad/transforms/rgb_to_yuv.h @@ -1,21 +1,14 @@ #pragma once -#include -#include +#include "common/clutil.h" -#ifdef __APPLE__ -#include -#else -#include -#endif +class Rgb2Yuv { +public: + Rgb2Yuv(cl_context ctx, cl_device_id device_id, int width, int height, int rgb_stride); + ~Rgb2Yuv(); + void queue(cl_command_queue q, cl_mem rgb_cl, cl_mem yuv_cl); +private: + size_t work_size[2]; + cl_kernel krnl; +}; -typedef struct { - int width, height; - cl_kernel rgb_to_yuv_krnl; -} RGBToYUVState; - -void rgb_to_yuv_init(RGBToYUVState* s, cl_context ctx, cl_device_id device_id, int width, int height, int rgb_stride); - -void rgb_to_yuv_destroy(RGBToYUVState* s); - -void rgb_to_yuv_queue(RGBToYUVState* s, cl_command_queue q, cl_mem rgb_cl, cl_mem yuv_cl);