mawei/dragonpilot
1
0
Fork 0
mirror of https://github.com/dragonpilot/dragonpilot.git synced 2026-03-03 17:43:52 +08:00
Code Issues Packages Projects Releases Wiki Activity

cleanup qcom2 (#19506)

Browse Source 
* cleanup qcom2

* define DEBAYER_LOCAL_WORKSIZE in camera_qcom2.h
This commit is contained in:
Dean Lee
2020-12-18 03:43:50 +08:00
committed by GitHub
parent 43c05aa2c6
commit bc1cfa6d4f
3 changed files with 8 additions and 17 deletions
Download Patch File Download Diff File Expand all files Collapse all files

12
selfdrive/camerad/cameras/camera_common.cc
View File

@@ -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);

8
selfdrive/camerad/cameras/camera_qcom2.cc
View File

@@ -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");
}

5
selfdrive/camerad/cameras/camera_qcom2.h
View File

@@ -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;