cleanup qcom2 (#19506)

* cleanup qcom2

* define DEBAYER_LOCAL_WORKSIZE in camera_qcom2.h
old-commit-hash: bc1cfa6d4f

selfdrive/camerad/cameras/camera_common.cc

@@ -160,10 +160,12 @@ bool CameraBuf::acquire() {
     CL_CHECK(clSetKernelArg(krnl_debayer, 0, sizeof(cl_mem), &camrabuf_cl));
     CL_CHECK(clSetKernelArg(krnl_debayer, 1, sizeof(cl_mem), &cur_rgb_buf->buf_cl));
 #ifdef QCOM2
-    CL_CHECK(clSetKernelArg(krnl_debayer, 2, camera_state->debayer_cl_localMemSize, 0));
-    CL_CHECK(clEnqueueNDRangeKernel(q, krnl_debayer, 2, NULL,
-                                  camera_state->debayer_cl_globalWorkSize, camera_state->debayer_cl_localWorkSize,
-                                  0, 0, &debayer_event));
+    constexpr int localMemSize = (DEBAYER_LOCAL_WORKSIZE + 2 * (3 / 2)) * (DEBAYER_LOCAL_WORKSIZE + 2 * (3 / 2)) * sizeof(float);
+    const size_t globalWorkSize[] = {size_t(camera_state->ci.frame_width), size_t(camera_state->ci.frame_height)};
+    const size_t localWorkSize[] = {DEBAYER_LOCAL_WORKSIZE, DEBAYER_LOCAL_WORKSIZE};
+    CL_CHECK(clSetKernelArg(krnl_debayer, 2, localMemSize, 0));
+    CL_CHECK(clEnqueueNDRangeKernel(q, krnl_debayer, 2, NULL, globalWorkSize, localWorkSize,
+                                    0, 0, &debayer_event));
 #else
     float digital_gain = camera_state->digital_gain;
     if ((int)digital_gain == 0) {
@@ -172,7 +174,7 @@ bool CameraBuf::acquire() {
     CL_CHECK(clSetKernelArg(krnl_debayer, 2, sizeof(float), &digital_gain));
     const size_t debayer_work_size = rgb_height;  // doesn't divide evenly, is this okay?
     CL_CHECK(clEnqueueNDRangeKernel(q, krnl_debayer, 1, NULL,
-                                  &debayer_work_size, NULL, 0, 0, &debayer_event));
+                                    &debayer_work_size, NULL, 0, 0, &debayer_event));
 #endif
   } else {
     assert(cur_rgb_buf->len >= frame_size);

selfdrive/camerad/cameras/camera_qcom2.cc

@@ -25,8 +25,6 @@
 
 #include "sensor2_i2c.h"
 
-#define DEBAYER_LOCAL_WORKSIZE 16
-
 #define FRAME_WIDTH  1928
 #define FRAME_HEIGHT 1208
 //#define FRAME_STRIDE 1936 // for 8 bit output
@@ -580,12 +578,6 @@ static void camera_init(CameraState *s, int camera_id, int camera_num, unsigned
   s->skipped = true;
   s->ef_filtered = 1.0;
 
-  s->debayer_cl_localMemSize = (DEBAYER_LOCAL_WORKSIZE + 2 * (3 / 2)) * (DEBAYER_LOCAL_WORKSIZE + 2 * (3 / 2)) * sizeof(float);
-  s->debayer_cl_globalWorkSize[0] = s->ci.frame_width;
-  s->debayer_cl_globalWorkSize[1] = s->ci.frame_height;
-  s->debayer_cl_localWorkSize[0] = DEBAYER_LOCAL_WORKSIZE;
-  s->debayer_cl_localWorkSize[1] = DEBAYER_LOCAL_WORKSIZE;
-
   s->buf.init(device_id, ctx, s, FRAME_BUF_COUNT, "frame");
 }
 

selfdrive/camerad/cameras/camera_qcom2.h

@@ -18,6 +18,7 @@
 
 #define EF_LOWPASS_K 0.35
 
+#define DEBAYER_LOCAL_WORKSIZE 16
 
 typedef struct CameraState {
   CameraInfo ci;
@@ -62,10 +63,6 @@ typedef struct CameraState {
   int idx_offset;
   bool skipped;
 
-  int debayer_cl_localMemSize;
-  size_t debayer_cl_globalWorkSize[2];
-  size_t debayer_cl_localWorkSize[2];
-
   struct cam_req_mgr_session_info req_mgr_session_info;
 
   CameraBuf buf;