|
|
|
@ -33,6 +33,7 @@ |
|
|
|
#include <jpeglib.h> |
|
|
|
#include <jpeglib.h> |
|
|
|
|
|
|
|
|
|
|
|
#define UI_BUF_COUNT 4 |
|
|
|
#define UI_BUF_COUNT 4 |
|
|
|
|
|
|
|
#define DEBAYER_LOCAL_WORKSIZE 16 |
|
|
|
#define YUV_COUNT 40 |
|
|
|
#define YUV_COUNT 40 |
|
|
|
#define MAX_CLIENTS 5 |
|
|
|
#define MAX_CLIENTS 5 |
|
|
|
|
|
|
|
|
|
|
|
@ -44,6 +45,10 @@ void set_do_exit(int sig) { |
|
|
|
do_exit = 1; |
|
|
|
do_exit = 1; |
|
|
|
} |
|
|
|
} |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
/*
|
|
|
|
|
|
|
|
TODO: refactor out camera specific things from here |
|
|
|
|
|
|
|
*/ |
|
|
|
|
|
|
|
|
|
|
|
struct VisionState; |
|
|
|
struct VisionState; |
|
|
|
|
|
|
|
|
|
|
|
struct VisionClientState { |
|
|
|
struct VisionClientState { |
|
|
|
@ -74,8 +79,13 @@ struct VisionState { |
|
|
|
|
|
|
|
|
|
|
|
cl_program prg_debayer_rear; |
|
|
|
cl_program prg_debayer_rear; |
|
|
|
cl_program prg_debayer_front; |
|
|
|
cl_program prg_debayer_front; |
|
|
|
|
|
|
|
cl_program prg_debayer_wide; |
|
|
|
cl_kernel krnl_debayer_rear; |
|
|
|
cl_kernel krnl_debayer_rear; |
|
|
|
cl_kernel krnl_debayer_front; |
|
|
|
cl_kernel krnl_debayer_front; |
|
|
|
|
|
|
|
cl_kernel krnl_debayer_wide; |
|
|
|
|
|
|
|
int debayer_cl_localMemSize; |
|
|
|
|
|
|
|
size_t debayer_cl_globalWorkSize[2]; |
|
|
|
|
|
|
|
size_t debayer_cl_localWorkSize[2]; |
|
|
|
|
|
|
|
|
|
|
|
cl_program prg_rgb_laplacian; |
|
|
|
cl_program prg_rgb_laplacian; |
|
|
|
cl_kernel krnl_rgb_laplacian; |
|
|
|
cl_kernel krnl_rgb_laplacian; |
|
|
|
@ -88,10 +98,12 @@ struct VisionState { |
|
|
|
// processing
|
|
|
|
// processing
|
|
|
|
TBuffer ui_tb; |
|
|
|
TBuffer ui_tb; |
|
|
|
TBuffer ui_front_tb; |
|
|
|
TBuffer ui_front_tb; |
|
|
|
|
|
|
|
TBuffer ui_wide_tb; |
|
|
|
|
|
|
|
|
|
|
|
mat3 yuv_transform; |
|
|
|
mat3 yuv_transform; |
|
|
|
TBuffer *yuv_tb; |
|
|
|
TBuffer *yuv_tb; |
|
|
|
TBuffer *yuv_front_tb; |
|
|
|
TBuffer *yuv_front_tb; |
|
|
|
|
|
|
|
TBuffer *yuv_wide_tb; |
|
|
|
|
|
|
|
|
|
|
|
// TODO: refactor for both cameras?
|
|
|
|
// TODO: refactor for both cameras?
|
|
|
|
Pool yuv_pool; |
|
|
|
Pool yuv_pool; |
|
|
|
@ -113,6 +125,15 @@ struct VisionState { |
|
|
|
int yuv_front_width, yuv_front_height; |
|
|
|
int yuv_front_width, yuv_front_height; |
|
|
|
RGBToYUVState front_rgb_to_yuv_state; |
|
|
|
RGBToYUVState front_rgb_to_yuv_state; |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
Pool yuv_wide_pool; |
|
|
|
|
|
|
|
VisionBuf yuv_wide_ion[YUV_COUNT]; |
|
|
|
|
|
|
|
cl_mem yuv_wide_cl[YUV_COUNT]; |
|
|
|
|
|
|
|
YUVBuf yuv_wide_bufs[YUV_COUNT]; |
|
|
|
|
|
|
|
FrameMetadata yuv_wide_metas[YUV_COUNT]; |
|
|
|
|
|
|
|
size_t yuv_wide_buf_size; |
|
|
|
|
|
|
|
int yuv_wide_width, yuv_wide_height; |
|
|
|
|
|
|
|
RGBToYUVState wide_rgb_to_yuv_state; |
|
|
|
|
|
|
|
|
|
|
|
size_t rgb_buf_size; |
|
|
|
size_t rgb_buf_size; |
|
|
|
int rgb_width, rgb_height, rgb_stride; |
|
|
|
int rgb_width, rgb_height, rgb_stride; |
|
|
|
VisionBuf rgb_bufs[UI_BUF_COUNT]; |
|
|
|
VisionBuf rgb_bufs[UI_BUF_COUNT]; |
|
|
|
@ -129,6 +150,10 @@ struct VisionState { |
|
|
|
int front_meteringbox_xmin, front_meteringbox_xmax; |
|
|
|
int front_meteringbox_xmin, front_meteringbox_xmax; |
|
|
|
int front_meteringbox_ymin, front_meteringbox_ymax; |
|
|
|
int front_meteringbox_ymin, front_meteringbox_ymax; |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
size_t rgb_wide_buf_size; |
|
|
|
|
|
|
|
int rgb_wide_width, rgb_wide_height, rgb_wide_stride; |
|
|
|
|
|
|
|
VisionBuf rgb_wide_bufs[UI_BUF_COUNT]; |
|
|
|
|
|
|
|
cl_mem rgb_wide_bufs_cl[UI_BUF_COUNT]; |
|
|
|
|
|
|
|
|
|
|
|
cl_mem camera_bufs_cl[FRAME_BUF_COUNT]; |
|
|
|
cl_mem camera_bufs_cl[FRAME_BUF_COUNT]; |
|
|
|
VisionBuf camera_bufs[FRAME_BUF_COUNT]; |
|
|
|
VisionBuf camera_bufs[FRAME_BUF_COUNT]; |
|
|
|
@ -138,7 +163,11 @@ struct VisionState { |
|
|
|
cl_mem front_camera_bufs_cl[FRAME_BUF_COUNT]; |
|
|
|
cl_mem front_camera_bufs_cl[FRAME_BUF_COUNT]; |
|
|
|
VisionBuf front_camera_bufs[FRAME_BUF_COUNT]; |
|
|
|
VisionBuf front_camera_bufs[FRAME_BUF_COUNT]; |
|
|
|
|
|
|
|
|
|
|
|
DualCameraState cameras; |
|
|
|
cl_mem wide_camera_bufs_cl[FRAME_BUF_COUNT]; |
|
|
|
|
|
|
|
VisionBuf wide_camera_bufs[FRAME_BUF_COUNT]; |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
MultiCameraState cameras; |
|
|
|
|
|
|
|
|
|
|
|
zsock_t *terminate_pub; |
|
|
|
zsock_t *terminate_pub; |
|
|
|
|
|
|
|
|
|
|
|
@ -179,14 +208,20 @@ void* frontview_thread(void *arg) { |
|
|
|
assert(err == 0); |
|
|
|
assert(err == 0); |
|
|
|
err = clSetKernelArg(s->krnl_debayer_front, 1, sizeof(cl_mem), &s->rgb_front_bufs_cl[rgb_idx]); |
|
|
|
err = clSetKernelArg(s->krnl_debayer_front, 1, sizeof(cl_mem), &s->rgb_front_bufs_cl[rgb_idx]); |
|
|
|
assert(err == 0); |
|
|
|
assert(err == 0); |
|
|
|
|
|
|
|
#ifdef QCOM2 |
|
|
|
|
|
|
|
err = clSetKernelArg(s->krnl_debayer_front, 2, s->debayer_cl_localMemSize, 0); |
|
|
|
|
|
|
|
assert(err == 0); |
|
|
|
|
|
|
|
err = clEnqueueNDRangeKernel(q, s->krnl_debayer_front, 2, NULL, |
|
|
|
|
|
|
|
s->debayer_cl_globalWorkSize, s->debayer_cl_localWorkSize, 0, 0, &debayer_event); |
|
|
|
|
|
|
|
#else |
|
|
|
float digital_gain = 1.0; |
|
|
|
float digital_gain = 1.0; |
|
|
|
err = clSetKernelArg(s->krnl_debayer_front, 2, sizeof(float), &digital_gain); |
|
|
|
err = clSetKernelArg(s->krnl_debayer_front, 2, sizeof(float), &digital_gain); |
|
|
|
assert(err == 0); |
|
|
|
assert(err == 0); |
|
|
|
|
|
|
|
|
|
|
|
const size_t debayer_work_size = s->rgb_front_height; // doesn't divide evenly, is this okay?
|
|
|
|
const size_t debayer_work_size = s->rgb_front_height; // doesn't divide evenly, is this okay?
|
|
|
|
const size_t debayer_local_work_size = 128; |
|
|
|
const size_t debayer_local_work_size = 128; |
|
|
|
err = clEnqueueNDRangeKernel(q, s->krnl_debayer_front, 1, NULL, |
|
|
|
err = clEnqueueNDRangeKernel(q, s->krnl_debayer_front, 1, NULL, |
|
|
|
&debayer_work_size, &debayer_local_work_size, 0, 0, &debayer_event); |
|
|
|
&debayer_work_size, &debayer_local_work_size, 0, 0, &debayer_event); |
|
|
|
|
|
|
|
#endif |
|
|
|
assert(err == 0); |
|
|
|
assert(err == 0); |
|
|
|
} else { |
|
|
|
} else { |
|
|
|
assert(s->rgb_front_buf_size >= s->cameras.front.frame_size); |
|
|
|
assert(s->rgb_front_buf_size >= s->cameras.front.frame_size); |
|
|
|
@ -208,6 +243,12 @@ void* frontview_thread(void *arg) { |
|
|
|
s->rhd_front_checked = state.getRhdChecked(); |
|
|
|
s->rhd_front_checked = state.getRhdChecked(); |
|
|
|
} |
|
|
|
} |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
#ifdef NOSCREEN |
|
|
|
|
|
|
|
if (frame_data.frame_id % 4 == 2) { |
|
|
|
|
|
|
|
sendrgb(&s->cameras, (uint8_t*) s->rgb_front_bufs[rgb_idx].addr, s->rgb_front_bufs[rgb_idx].len, 2); |
|
|
|
|
|
|
|
} |
|
|
|
|
|
|
|
#endif |
|
|
|
|
|
|
|
|
|
|
|
if (sm.updated("driverState")) { |
|
|
|
if (sm.updated("driverState")) { |
|
|
|
auto state = sm["driverState"].getDriverState(); |
|
|
|
auto state = sm["driverState"].getDriverState(); |
|
|
|
float face_prob = state.getFaceProb(); |
|
|
|
float face_prob = state.getFaceProb(); |
|
|
|
@ -256,7 +297,12 @@ void* frontview_thread(void *arg) { |
|
|
|
x_start = s->rhd_front ? 0:s->rgb_front_width * 3 / 5; |
|
|
|
x_start = s->rhd_front ? 0:s->rgb_front_width * 3 / 5; |
|
|
|
x_end = s->rhd_front ? s->rgb_front_width * 2 / 5:s->rgb_front_width; |
|
|
|
x_end = s->rhd_front ? s->rgb_front_width * 2 / 5:s->rgb_front_width; |
|
|
|
} |
|
|
|
} |
|
|
|
|
|
|
|
#ifdef QCOM2 |
|
|
|
|
|
|
|
x_start = 0.15*s->rgb_front_width; |
|
|
|
|
|
|
|
x_end = 0.85*s->rgb_front_width; |
|
|
|
|
|
|
|
y_start = 0.15*s->rgb_front_height; |
|
|
|
|
|
|
|
y_end = 0.85*s->rgb_front_height; |
|
|
|
|
|
|
|
#endif |
|
|
|
uint32_t lum_binning[256] = {0,}; |
|
|
|
uint32_t lum_binning[256] = {0,}; |
|
|
|
for (int y = y_start; y < y_end; ++y) { |
|
|
|
for (int y = y_start; y < y_end; ++y) { |
|
|
|
for (int x = x_start; x < x_end; x += 2) { // every 2nd col
|
|
|
|
for (int x = x_start; x < x_end; x += 2) { // every 2nd col
|
|
|
|
@ -336,6 +382,162 @@ void* frontview_thread(void *arg) { |
|
|
|
|
|
|
|
|
|
|
|
return NULL; |
|
|
|
return NULL; |
|
|
|
} |
|
|
|
} |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
#ifdef QCOM2 |
|
|
|
|
|
|
|
// wide
|
|
|
|
|
|
|
|
void* wideview_thread(void *arg) { |
|
|
|
|
|
|
|
int err; |
|
|
|
|
|
|
|
VisionState *s = (VisionState*)arg; |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
set_thread_name("wideview"); |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
err = set_realtime_priority(1); |
|
|
|
|
|
|
|
LOG("setpriority returns %d", err); |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
// init cl stuff
|
|
|
|
|
|
|
|
const cl_queue_properties props[] = {0}; //CL_QUEUE_PRIORITY_KHR, CL_QUEUE_PRIORITY_HIGH_KHR, 0};
|
|
|
|
|
|
|
|
cl_command_queue q = clCreateCommandQueueWithProperties(s->context, s->device_id, props, &err); |
|
|
|
|
|
|
|
assert(err == 0); |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
// init the net
|
|
|
|
|
|
|
|
LOG("wideview start!"); |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
for (int cnt = 0; !do_exit; cnt++) { |
|
|
|
|
|
|
|
int buf_idx = tbuffer_acquire(&s->cameras.wide.camera_tb); |
|
|
|
|
|
|
|
// int buf_idx = camera_acquire_buffer(s);
|
|
|
|
|
|
|
|
if (buf_idx < 0) { |
|
|
|
|
|
|
|
break; |
|
|
|
|
|
|
|
} |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
double t1 = millis_since_boot(); |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
FrameMetadata frame_data = s->cameras.wide.camera_bufs_metadata[buf_idx]; |
|
|
|
|
|
|
|
uint32_t frame_id = frame_data.frame_id; |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
if (frame_id == -1) { |
|
|
|
|
|
|
|
LOGE("no frame data? wtf"); |
|
|
|
|
|
|
|
tbuffer_release(&s->cameras.wide.camera_tb, buf_idx); |
|
|
|
|
|
|
|
continue; |
|
|
|
|
|
|
|
} |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
int ui_idx = tbuffer_select(&s->ui_wide_tb); |
|
|
|
|
|
|
|
int rgb_idx = ui_idx; |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
cl_event debayer_event; |
|
|
|
|
|
|
|
if (s->cameras.wide.ci.bayer) { |
|
|
|
|
|
|
|
err = clSetKernelArg(s->krnl_debayer_wide, 0, sizeof(cl_mem), &s->wide_camera_bufs_cl[buf_idx]); |
|
|
|
|
|
|
|
assert(err == 0); |
|
|
|
|
|
|
|
err = clSetKernelArg(s->krnl_debayer_wide, 1, sizeof(cl_mem), &s->rgb_wide_bufs_cl[rgb_idx]); |
|
|
|
|
|
|
|
assert(err == 0); |
|
|
|
|
|
|
|
err = clSetKernelArg(s->krnl_debayer_wide, 2, s->debayer_cl_localMemSize, 0); |
|
|
|
|
|
|
|
assert(err == 0); |
|
|
|
|
|
|
|
err = clEnqueueNDRangeKernel(q, s->krnl_debayer_wide, 2, NULL, |
|
|
|
|
|
|
|
s->debayer_cl_globalWorkSize, s->debayer_cl_localWorkSize, 0, 0, &debayer_event); |
|
|
|
|
|
|
|
assert(err == 0); |
|
|
|
|
|
|
|
} else { |
|
|
|
|
|
|
|
assert(s->rgb_wide_buf_size >= s->frame_size); |
|
|
|
|
|
|
|
assert(s->rgb_stride == s->frame_stride); |
|
|
|
|
|
|
|
err = clEnqueueCopyBuffer(q, s->wide_camera_bufs_cl[buf_idx], s->rgb_wide_bufs_cl[rgb_idx], |
|
|
|
|
|
|
|
0, 0, s->rgb_wide_buf_size, 0, 0, &debayer_event); |
|
|
|
|
|
|
|
assert(err == 0); |
|
|
|
|
|
|
|
} |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
clWaitForEvents(1, &debayer_event); |
|
|
|
|
|
|
|
clReleaseEvent(debayer_event); |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
tbuffer_release(&s->cameras.wide.camera_tb, buf_idx); |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
visionbuf_sync(&s->rgb_wide_bufs[rgb_idx], VISIONBUF_SYNC_FROM_DEVICE); |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
#ifdef NOSCREEN |
|
|
|
|
|
|
|
if (frame_data.frame_id % 4 == 0) { |
|
|
|
|
|
|
|
sendrgb(&s->cameras, (uint8_t*) s->rgb_wide_bufs[rgb_idx].addr, s->rgb_wide_bufs[rgb_idx].len, 1); |
|
|
|
|
|
|
|
} |
|
|
|
|
|
|
|
#endif |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
double t2 = millis_since_boot(); |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
double yt1 = millis_since_boot(); |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
int yuv_idx = pool_select(&s->yuv_wide_pool); |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
s->yuv_wide_metas[yuv_idx] = frame_data; |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
uint8_t* yuv_ptr_y = s->yuv_wide_bufs[yuv_idx].y; |
|
|
|
|
|
|
|
cl_mem yuv_wide_cl = s->yuv_wide_cl[yuv_idx]; |
|
|
|
|
|
|
|
rgb_to_yuv_queue(&s->wide_rgb_to_yuv_state, q, s->rgb_wide_bufs_cl[rgb_idx], yuv_wide_cl); |
|
|
|
|
|
|
|
visionbuf_sync(&s->yuv_wide_ion[yuv_idx], VISIONBUF_SYNC_FROM_DEVICE); |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
double yt2 = millis_since_boot(); |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
// keep another reference around till were done processing
|
|
|
|
|
|
|
|
pool_acquire(&s->yuv_wide_pool, yuv_idx); |
|
|
|
|
|
|
|
pool_push(&s->yuv_wide_pool, yuv_idx); |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
// send frame event
|
|
|
|
|
|
|
|
{ |
|
|
|
|
|
|
|
if (s->pm != NULL) { |
|
|
|
|
|
|
|
capnp::MallocMessageBuilder msg; |
|
|
|
|
|
|
|
cereal::Event::Builder event = msg.initRoot<cereal::Event>(); |
|
|
|
|
|
|
|
event.setLogMonoTime(nanos_since_boot()); |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
auto framed = event.initWideFrame(); |
|
|
|
|
|
|
|
framed.setFrameId(frame_data.frame_id); |
|
|
|
|
|
|
|
framed.setEncodeId(cnt); |
|
|
|
|
|
|
|
framed.setTimestampEof(frame_data.timestamp_eof); |
|
|
|
|
|
|
|
framed.setFrameLength(frame_data.frame_length); |
|
|
|
|
|
|
|
framed.setIntegLines(frame_data.integ_lines); |
|
|
|
|
|
|
|
framed.setGlobalGain(frame_data.global_gain); |
|
|
|
|
|
|
|
framed.setLensPos(frame_data.lens_pos); |
|
|
|
|
|
|
|
framed.setLensSag(frame_data.lens_sag); |
|
|
|
|
|
|
|
framed.setLensErr(frame_data.lens_err); |
|
|
|
|
|
|
|
framed.setLensTruePos(frame_data.lens_true_pos); |
|
|
|
|
|
|
|
framed.setGainFrac(frame_data.gain_frac); |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
s->pm->send("wideFrame", msg); |
|
|
|
|
|
|
|
} |
|
|
|
|
|
|
|
} |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
tbuffer_dispatch(&s->ui_wide_tb, ui_idx); |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
// auto exposure over big box
|
|
|
|
|
|
|
|
// TODO: fix this? should not use med imo
|
|
|
|
|
|
|
|
const int exposure_x = 240; |
|
|
|
|
|
|
|
const int exposure_y = 300; |
|
|
|
|
|
|
|
const int exposure_height = 600; |
|
|
|
|
|
|
|
const int exposure_width = 1440; |
|
|
|
|
|
|
|
if (cnt % 3 == 0) { |
|
|
|
|
|
|
|
// find median box luminance for AE
|
|
|
|
|
|
|
|
uint32_t lum_binning[256] = {0,}; |
|
|
|
|
|
|
|
for (int y=0; y<exposure_height; y++) { |
|
|
|
|
|
|
|
for (int x=0; x<exposure_width; x++) { |
|
|
|
|
|
|
|
uint8_t lum = yuv_ptr_y[((exposure_y+y)*s->yuv_wide_width) + exposure_x + x]; |
|
|
|
|
|
|
|
lum_binning[lum]++; |
|
|
|
|
|
|
|
} |
|
|
|
|
|
|
|
} |
|
|
|
|
|
|
|
const unsigned int lum_total = exposure_height * exposure_width; |
|
|
|
|
|
|
|
unsigned int lum_cur = 0; |
|
|
|
|
|
|
|
int lum_med = 0; |
|
|
|
|
|
|
|
for (lum_med=0; lum_med<256; lum_med++) { |
|
|
|
|
|
|
|
// shouldn't be any values less than 16 - yuv footroom
|
|
|
|
|
|
|
|
lum_cur += lum_binning[lum_med]; |
|
|
|
|
|
|
|
if (lum_cur >= lum_total / 2) { |
|
|
|
|
|
|
|
break; |
|
|
|
|
|
|
|
} |
|
|
|
|
|
|
|
} |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
camera_autoexposure(&s->cameras.wide, lum_med / 256.0); |
|
|
|
|
|
|
|
} |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
pool_release(&s->yuv_wide_pool, yuv_idx); |
|
|
|
|
|
|
|
double t5 = millis_since_boot(); |
|
|
|
|
|
|
|
LOGD("queued: %.2fms, yuv: %.2f, | processing: %.3fms", (t2-t1), (yt2-yt1), (t5-t1)); |
|
|
|
|
|
|
|
} |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
return NULL; |
|
|
|
|
|
|
|
} |
|
|
|
|
|
|
|
#endif |
|
|
|
|
|
|
|
|
|
|
|
// processing
|
|
|
|
// processing
|
|
|
|
void* processing_thread(void *arg) { |
|
|
|
void* processing_thread(void *arg) { |
|
|
|
int err; |
|
|
|
int err; |
|
|
|
@ -390,13 +592,19 @@ void* processing_thread(void *arg) { |
|
|
|
assert(err == 0); |
|
|
|
assert(err == 0); |
|
|
|
err = clSetKernelArg(s->krnl_debayer_rear, 1, sizeof(cl_mem), &s->rgb_bufs_cl[rgb_idx]); |
|
|
|
err = clSetKernelArg(s->krnl_debayer_rear, 1, sizeof(cl_mem), &s->rgb_bufs_cl[rgb_idx]); |
|
|
|
assert(err == 0); |
|
|
|
assert(err == 0); |
|
|
|
|
|
|
|
#ifdef QCOM2 |
|
|
|
|
|
|
|
err = clSetKernelArg(s->krnl_debayer_rear, 2, s->debayer_cl_localMemSize, 0); |
|
|
|
|
|
|
|
assert(err == 0); |
|
|
|
|
|
|
|
err = clEnqueueNDRangeKernel(q, s->krnl_debayer_rear, 2, NULL, |
|
|
|
|
|
|
|
s->debayer_cl_globalWorkSize, s->debayer_cl_localWorkSize, 0, 0, &debayer_event); |
|
|
|
|
|
|
|
#else |
|
|
|
err = clSetKernelArg(s->krnl_debayer_rear, 2, sizeof(float), &s->cameras.rear.digital_gain); |
|
|
|
err = clSetKernelArg(s->krnl_debayer_rear, 2, sizeof(float), &s->cameras.rear.digital_gain); |
|
|
|
assert(err == 0); |
|
|
|
assert(err == 0); |
|
|
|
|
|
|
|
|
|
|
|
const size_t debayer_work_size = s->rgb_height; // doesn't divide evenly, is this okay?
|
|
|
|
const size_t debayer_work_size = s->rgb_height; // doesn't divide evenly, is this okay?
|
|
|
|
const size_t debayer_local_work_size = 128; |
|
|
|
const size_t debayer_local_work_size = 128; |
|
|
|
err = clEnqueueNDRangeKernel(q, s->krnl_debayer_rear, 1, NULL, |
|
|
|
err = clEnqueueNDRangeKernel(q, s->krnl_debayer_rear, 1, NULL, |
|
|
|
&debayer_work_size, &debayer_local_work_size, 0, 0, &debayer_event); |
|
|
|
&debayer_work_size, &debayer_local_work_size, 0, 0, &debayer_event); |
|
|
|
|
|
|
|
#endif |
|
|
|
assert(err == 0); |
|
|
|
assert(err == 0); |
|
|
|
} else { |
|
|
|
} else { |
|
|
|
assert(s->rgb_buf_size >= s->frame_size); |
|
|
|
assert(s->rgb_buf_size >= s->frame_size); |
|
|
|
@ -405,7 +613,6 @@ void* processing_thread(void *arg) { |
|
|
|
0, 0, s->rgb_buf_size, 0, 0, &debayer_event); |
|
|
|
0, 0, s->rgb_buf_size, 0, 0, &debayer_event); |
|
|
|
assert(err == 0); |
|
|
|
assert(err == 0); |
|
|
|
} |
|
|
|
} |
|
|
|
|
|
|
|
|
|
|
|
clWaitForEvents(1, &debayer_event); |
|
|
|
clWaitForEvents(1, &debayer_event); |
|
|
|
clReleaseEvent(debayer_event); |
|
|
|
clReleaseEvent(debayer_event); |
|
|
|
|
|
|
|
|
|
|
|
@ -413,6 +620,12 @@ void* processing_thread(void *arg) { |
|
|
|
|
|
|
|
|
|
|
|
visionbuf_sync(&s->rgb_bufs[rgb_idx], VISIONBUF_SYNC_FROM_DEVICE); |
|
|
|
visionbuf_sync(&s->rgb_bufs[rgb_idx], VISIONBUF_SYNC_FROM_DEVICE); |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
#ifdef NOSCREEN |
|
|
|
|
|
|
|
if (frame_data.frame_id % 4 == 1) { |
|
|
|
|
|
|
|
sendrgb(&s->cameras, (uint8_t*) s->rgb_bufs[rgb_idx].addr, s->rgb_bufs[rgb_idx].len, 0); |
|
|
|
|
|
|
|
} |
|
|
|
|
|
|
|
#endif |
|
|
|
|
|
|
|
|
|
|
|
#if defined(QCOM) && !defined(QCOM_REPLAY) |
|
|
|
#if defined(QCOM) && !defined(QCOM_REPLAY) |
|
|
|
/*FILE *dump_rgb_file = fopen("/tmp/process_dump.rgb", "wb");
|
|
|
|
/*FILE *dump_rgb_file = fopen("/tmp/process_dump.rgb", "wb");
|
|
|
|
fwrite(s->rgb_bufs[rgb_idx].addr, s->rgb_bufs[rgb_idx].len, sizeof(uint8_t), dump_rgb_file); |
|
|
|
fwrite(s->rgb_bufs[rgb_idx].addr, s->rgb_bufs[rgb_idx].len, sizeof(uint8_t), dump_rgb_file); |
|
|
|
@ -638,10 +851,17 @@ void* processing_thread(void *arg) { |
|
|
|
tbuffer_dispatch(&s->ui_tb, ui_idx); |
|
|
|
tbuffer_dispatch(&s->ui_tb, ui_idx); |
|
|
|
|
|
|
|
|
|
|
|
// auto exposure over big box
|
|
|
|
// auto exposure over big box
|
|
|
|
|
|
|
|
#ifdef QCOM2 |
|
|
|
|
|
|
|
const int exposure_x = 240; |
|
|
|
|
|
|
|
const int exposure_y = 300; |
|
|
|
|
|
|
|
const int exposure_height = 600; |
|
|
|
|
|
|
|
const int exposure_width = 1440; |
|
|
|
|
|
|
|
#else |
|
|
|
const int exposure_x = 290; |
|
|
|
const int exposure_x = 290; |
|
|
|
const int exposure_y = 282 + 40; |
|
|
|
const int exposure_y = 322; |
|
|
|
const int exposure_height = 314; |
|
|
|
const int exposure_height = 314; |
|
|
|
const int exposure_width = 560; |
|
|
|
const int exposure_width = 560; |
|
|
|
|
|
|
|
#endif |
|
|
|
if (cnt % 3 == 0) { |
|
|
|
if (cnt % 3 == 0) { |
|
|
|
// find median box luminance for AE
|
|
|
|
// find median box luminance for AE
|
|
|
|
uint32_t lum_binning[256] = {0,}; |
|
|
|
uint32_t lum_binning[256] = {0,}; |
|
|
|
@ -661,8 +881,6 @@ void* processing_thread(void *arg) { |
|
|
|
break; |
|
|
|
break; |
|
|
|
} |
|
|
|
} |
|
|
|
} |
|
|
|
} |
|
|
|
// double avg = (double)acc / (big_box_width * big_box_height) - 16;
|
|
|
|
|
|
|
|
// printf("avg %d\n", lum_med);
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
camera_autoexposure(&s->cameras.rear, lum_med / 256.0); |
|
|
|
camera_autoexposure(&s->cameras.rear, lum_med / 256.0); |
|
|
|
} |
|
|
|
} |
|
|
|
@ -769,6 +987,20 @@ void* visionserver_client_thread(void* arg) { |
|
|
|
} else { |
|
|
|
} else { |
|
|
|
assert(false); |
|
|
|
assert(false); |
|
|
|
} |
|
|
|
} |
|
|
|
|
|
|
|
} else if (stream_type == VISION_STREAM_RGB_WIDE) { |
|
|
|
|
|
|
|
stream_bufs->width = s->rgb_wide_width; |
|
|
|
|
|
|
|
stream_bufs->height = s->rgb_wide_height; |
|
|
|
|
|
|
|
stream_bufs->stride = s->rgb_wide_stride; |
|
|
|
|
|
|
|
stream_bufs->buf_len = s->rgb_wide_bufs[0].len; |
|
|
|
|
|
|
|
rep.num_fds = UI_BUF_COUNT; |
|
|
|
|
|
|
|
for (int i=0; i<rep.num_fds; i++) { |
|
|
|
|
|
|
|
rep.fds[i] = s->rgb_wide_bufs[i].fd; |
|
|
|
|
|
|
|
} |
|
|
|
|
|
|
|
if (stream->tb) { |
|
|
|
|
|
|
|
stream->tbuffer = &s->ui_wide_tb; |
|
|
|
|
|
|
|
} else { |
|
|
|
|
|
|
|
assert(false); |
|
|
|
|
|
|
|
} |
|
|
|
} else if (stream_type == VISION_STREAM_YUV) { |
|
|
|
} else if (stream_type == VISION_STREAM_YUV) { |
|
|
|
stream_bufs->width = s->yuv_width; |
|
|
|
stream_bufs->width = s->yuv_width; |
|
|
|
stream_bufs->height = s->yuv_height; |
|
|
|
stream_bufs->height = s->yuv_height; |
|
|
|
@ -797,6 +1029,21 @@ void* visionserver_client_thread(void* arg) { |
|
|
|
} else { |
|
|
|
} else { |
|
|
|
stream->queue = pool_get_queue(&s->yuv_front_pool); |
|
|
|
stream->queue = pool_get_queue(&s->yuv_front_pool); |
|
|
|
} |
|
|
|
} |
|
|
|
|
|
|
|
} else if (stream_type == VISION_STREAM_YUV_WIDE) { |
|
|
|
|
|
|
|
stream_bufs->width = s->yuv_wide_width; |
|
|
|
|
|
|
|
stream_bufs->height = s->yuv_wide_height; |
|
|
|
|
|
|
|
stream_bufs->stride = s->yuv_wide_width; |
|
|
|
|
|
|
|
stream_bufs->buf_len = s->yuv_wide_buf_size; |
|
|
|
|
|
|
|
rep.num_fds = YUV_COUNT; |
|
|
|
|
|
|
|
for (int i=0; i<rep.num_fds; i++) { |
|
|
|
|
|
|
|
rep.fds[i] = s->yuv_wide_ion[i].fd; |
|
|
|
|
|
|
|
} |
|
|
|
|
|
|
|
if (stream->tb) { |
|
|
|
|
|
|
|
stream->tbuffer = s->yuv_wide_tb; |
|
|
|
|
|
|
|
} else { |
|
|
|
|
|
|
|
stream->queue = pool_get_queue(&s->yuv_wide_pool); |
|
|
|
|
|
|
|
} |
|
|
|
|
|
|
|
|
|
|
|
} else { |
|
|
|
} else { |
|
|
|
assert(false); |
|
|
|
assert(false); |
|
|
|
} |
|
|
|
} |
|
|
|
@ -849,6 +1096,9 @@ void* visionserver_client_thread(void* arg) { |
|
|
|
} else if (stream_i == VISION_STREAM_YUV_FRONT) { |
|
|
|
} else if (stream_i == VISION_STREAM_YUV_FRONT) { |
|
|
|
rep.d.stream_acq.extra.frame_id = s->yuv_front_metas[idx].frame_id; |
|
|
|
rep.d.stream_acq.extra.frame_id = s->yuv_front_metas[idx].frame_id; |
|
|
|
rep.d.stream_acq.extra.timestamp_eof = s->yuv_front_metas[idx].timestamp_eof; |
|
|
|
rep.d.stream_acq.extra.timestamp_eof = s->yuv_front_metas[idx].timestamp_eof; |
|
|
|
|
|
|
|
} else if (stream_i == VISION_STREAM_YUV_WIDE) { |
|
|
|
|
|
|
|
rep.d.stream_acq.extra.frame_id = s->yuv_wide_metas[idx].frame_id; |
|
|
|
|
|
|
|
rep.d.stream_acq.extra.timestamp_eof = s->yuv_wide_metas[idx].timestamp_eof; |
|
|
|
} |
|
|
|
} |
|
|
|
vipc_send(fd, &rep); |
|
|
|
vipc_send(fd, &rep); |
|
|
|
} |
|
|
|
} |
|
|
|
@ -959,25 +1209,28 @@ cl_program build_debayer_program(VisionState *s, |
|
|
|
int frame_width, int frame_height, int frame_stride, |
|
|
|
int frame_width, int frame_height, int frame_stride, |
|
|
|
int rgb_width, int rgb_height, int rgb_stride, |
|
|
|
int rgb_width, int rgb_height, int rgb_stride, |
|
|
|
int bayer_flip, int hdr) { |
|
|
|
int bayer_flip, int hdr) { |
|
|
|
|
|
|
|
#ifdef QCOM2 |
|
|
|
|
|
|
|
assert(rgb_width == frame_width); |
|
|
|
|
|
|
|
assert(rgb_height == frame_height); |
|
|
|
|
|
|
|
#else |
|
|
|
assert(rgb_width == frame_width/2); |
|
|
|
assert(rgb_width == frame_width/2); |
|
|
|
assert(rgb_height == frame_height/2); |
|
|
|
assert(rgb_height == frame_height/2); |
|
|
|
|
|
|
|
#endif |
|
|
|
#ifdef QCOM2 |
|
|
|
|
|
|
|
int dnew = 1; |
|
|
|
|
|
|
|
#else |
|
|
|
|
|
|
|
int dnew = 0; |
|
|
|
|
|
|
|
#endif |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
char args[4096]; |
|
|
|
char args[4096]; |
|
|
|
snprintf(args, sizeof(args), |
|
|
|
snprintf(args, sizeof(args), |
|
|
|
"-cl-fast-relaxed-math -cl-denorms-are-zero " |
|
|
|
"-cl-fast-relaxed-math -cl-denorms-are-zero " |
|
|
|
"-DFRAME_WIDTH=%d -DFRAME_HEIGHT=%d -DFRAME_STRIDE=%d " |
|
|
|
"-DFRAME_WIDTH=%d -DFRAME_HEIGHT=%d -DFRAME_STRIDE=%d " |
|
|
|
"-DRGB_WIDTH=%d -DRGB_HEIGHT=%d -DRGB_STRIDE=%d " |
|
|
|
"-DRGB_WIDTH=%d -DRGB_HEIGHT=%d -DRGB_STRIDE=%d " |
|
|
|
"-DBAYER_FLIP=%d -DHDR=%d -DNEW=%d", |
|
|
|
"-DBAYER_FLIP=%d -DHDR=%d", |
|
|
|
frame_width, frame_height, frame_stride, |
|
|
|
frame_width, frame_height, frame_stride, |
|
|
|
rgb_width, rgb_height, rgb_stride, |
|
|
|
rgb_width, rgb_height, rgb_stride, |
|
|
|
bayer_flip, hdr, dnew); |
|
|
|
bayer_flip, hdr); |
|
|
|
|
|
|
|
#ifdef QCOM2 |
|
|
|
|
|
|
|
return CLU_LOAD_FROM_FILE(s->context, s->device_id, "cameras/real_debayer.cl", args); |
|
|
|
|
|
|
|
#else |
|
|
|
return CLU_LOAD_FROM_FILE(s->context, s->device_id, "cameras/debayer.cl", args); |
|
|
|
return CLU_LOAD_FROM_FILE(s->context, s->device_id, "cameras/debayer.cl", args); |
|
|
|
|
|
|
|
#endif |
|
|
|
} |
|
|
|
} |
|
|
|
|
|
|
|
|
|
|
|
cl_program build_conv_program(VisionState *s, |
|
|
|
cl_program build_conv_program(VisionState *s, |
|
|
|
@ -1030,7 +1283,7 @@ void init_buffers(VisionState *s) { |
|
|
|
for (int i=0; i<FRAME_BUF_COUNT; i++) { |
|
|
|
for (int i=0; i<FRAME_BUF_COUNT; i++) { |
|
|
|
s->camera_bufs[i] = visionbuf_allocate_cl(s->frame_size, s->device_id, s->context, |
|
|
|
s->camera_bufs[i] = visionbuf_allocate_cl(s->frame_size, s->device_id, s->context, |
|
|
|
&s->camera_bufs_cl[i]); |
|
|
|
&s->camera_bufs_cl[i]); |
|
|
|
#ifndef QCOM2 |
|
|
|
#ifdef QCOM |
|
|
|
// TODO: make lengths correct
|
|
|
|
// TODO: make lengths correct
|
|
|
|
s->focus_bufs[i] = visionbuf_allocate(0xb80); |
|
|
|
s->focus_bufs[i] = visionbuf_allocate(0xb80); |
|
|
|
s->stats_bufs[i] = visionbuf_allocate(0xb80); |
|
|
|
s->stats_bufs[i] = visionbuf_allocate(0xb80); |
|
|
|
@ -1042,11 +1295,22 @@ void init_buffers(VisionState *s) { |
|
|
|
s->device_id, s->context, |
|
|
|
s->device_id, s->context, |
|
|
|
&s->front_camera_bufs_cl[i]); |
|
|
|
&s->front_camera_bufs_cl[i]); |
|
|
|
} |
|
|
|
} |
|
|
|
|
|
|
|
#ifdef QCOM2 |
|
|
|
|
|
|
|
for (int i=0; i<FRAME_BUF_COUNT; i++) { |
|
|
|
|
|
|
|
s->wide_camera_bufs[i] = visionbuf_allocate_cl(s->cameras.wide.frame_size, |
|
|
|
|
|
|
|
s->device_id, s->context, |
|
|
|
|
|
|
|
&s->wide_camera_bufs_cl[i]); |
|
|
|
|
|
|
|
} |
|
|
|
|
|
|
|
#endif |
|
|
|
// processing buffers
|
|
|
|
// processing buffers
|
|
|
|
if (s->cameras.rear.ci.bayer) { |
|
|
|
if (s->cameras.rear.ci.bayer) { |
|
|
|
s->rgb_width = s->frame_width/2; |
|
|
|
#ifdef QCOM2 |
|
|
|
s->rgb_height = s->frame_height/2; |
|
|
|
s->rgb_width = s->frame_width; |
|
|
|
|
|
|
|
s->rgb_height = s->frame_height; |
|
|
|
|
|
|
|
#else |
|
|
|
|
|
|
|
s->rgb_width = s->frame_width / 2; |
|
|
|
|
|
|
|
s->rgb_height = s->frame_height / 2; |
|
|
|
|
|
|
|
#endif |
|
|
|
} else { |
|
|
|
} else { |
|
|
|
s->rgb_width = s->frame_width; |
|
|
|
s->rgb_width = s->frame_width; |
|
|
|
s->rgb_height = s->frame_height; |
|
|
|
s->rgb_height = s->frame_height; |
|
|
|
@ -1062,15 +1326,22 @@ void init_buffers(VisionState *s) { |
|
|
|
} |
|
|
|
} |
|
|
|
tbuffer_init(&s->ui_tb, UI_BUF_COUNT, "rgb"); |
|
|
|
tbuffer_init(&s->ui_tb, UI_BUF_COUNT, "rgb"); |
|
|
|
|
|
|
|
|
|
|
|
//assert(s->cameras.front.ci.bayer);
|
|
|
|
|
|
|
|
if (s->cameras.front.ci.bayer) { |
|
|
|
if (s->cameras.front.ci.bayer) { |
|
|
|
s->rgb_front_width = s->cameras.front.ci.frame_width/2; |
|
|
|
#ifdef QCOM2 |
|
|
|
s->rgb_front_height = s->cameras.front.ci.frame_height/2; |
|
|
|
s->rgb_front_width = s->cameras.front.ci.frame_width; |
|
|
|
|
|
|
|
s->rgb_front_height = s->cameras.front.ci.frame_height; |
|
|
|
|
|
|
|
#else |
|
|
|
|
|
|
|
s->rgb_front_width = s->cameras.front.ci.frame_width / 2; |
|
|
|
|
|
|
|
s->rgb_front_height = s->cameras.front.ci.frame_height / 2; |
|
|
|
|
|
|
|
#endif |
|
|
|
} else { |
|
|
|
} else { |
|
|
|
s->rgb_front_width = s->cameras.front.ci.frame_width; |
|
|
|
s->rgb_front_width = s->cameras.front.ci.frame_width; |
|
|
|
s->rgb_front_height = s->cameras.front.ci.frame_height; |
|
|
|
s->rgb_front_height = s->cameras.front.ci.frame_height; |
|
|
|
} |
|
|
|
} |
|
|
|
|
|
|
|
#ifdef QCOM2 |
|
|
|
|
|
|
|
s->rgb_wide_width = s->cameras.wide.ci.frame_width; |
|
|
|
|
|
|
|
s->rgb_wide_height = s->cameras.wide.ci.frame_height; |
|
|
|
|
|
|
|
#endif |
|
|
|
|
|
|
|
|
|
|
|
for (int i=0; i<UI_BUF_COUNT; i++) { |
|
|
|
for (int i=0; i<UI_BUF_COUNT; i++) { |
|
|
|
VisionImg img = visionimg_alloc_rgb24(s->rgb_front_width, s->rgb_front_height, &s->rgb_front_bufs[i]); |
|
|
|
VisionImg img = visionimg_alloc_rgb24(s->rgb_front_width, s->rgb_front_height, &s->rgb_front_bufs[i]); |
|
|
|
@ -1081,6 +1352,17 @@ void init_buffers(VisionState *s) { |
|
|
|
} |
|
|
|
} |
|
|
|
} |
|
|
|
} |
|
|
|
tbuffer_init(&s->ui_front_tb, UI_BUF_COUNT, "frontrgb"); |
|
|
|
tbuffer_init(&s->ui_front_tb, UI_BUF_COUNT, "frontrgb"); |
|
|
|
|
|
|
|
#ifdef QCOM2 |
|
|
|
|
|
|
|
for (int i=0; i<UI_BUF_COUNT; i++) { |
|
|
|
|
|
|
|
VisionImg img = visionimg_alloc_rgb24(s->rgb_wide_width, s->rgb_wide_height, &s->rgb_wide_bufs[i]); |
|
|
|
|
|
|
|
s->rgb_wide_bufs_cl[i] = visionbuf_to_cl(&s->rgb_wide_bufs[i], s->device_id, s->context); |
|
|
|
|
|
|
|
if (i == 0){ |
|
|
|
|
|
|
|
s->rgb_wide_stride = img.stride; |
|
|
|
|
|
|
|
s->rgb_wide_buf_size = img.size; |
|
|
|
|
|
|
|
} |
|
|
|
|
|
|
|
} |
|
|
|
|
|
|
|
tbuffer_init(&s->ui_wide_tb, UI_BUF_COUNT, "widergb"); |
|
|
|
|
|
|
|
#endif |
|
|
|
|
|
|
|
|
|
|
|
// yuv back for recording and orbd
|
|
|
|
// yuv back for recording and orbd
|
|
|
|
pool_init(&s->yuv_pool, YUV_COUNT); |
|
|
|
pool_init(&s->yuv_pool, YUV_COUNT); |
|
|
|
@ -1112,6 +1394,23 @@ void init_buffers(VisionState *s) { |
|
|
|
s->yuv_front_bufs[i].v = s->yuv_front_bufs[i].u + (s->yuv_front_width/2 * s->yuv_front_height/2); |
|
|
|
s->yuv_front_bufs[i].v = s->yuv_front_bufs[i].u + (s->yuv_front_width/2 * s->yuv_front_height/2); |
|
|
|
} |
|
|
|
} |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
// yuv wide for recording
|
|
|
|
|
|
|
|
#ifdef QCOM2 |
|
|
|
|
|
|
|
pool_init(&s->yuv_wide_pool, YUV_COUNT); |
|
|
|
|
|
|
|
s->yuv_wide_tb = pool_get_tbuffer(&s->yuv_wide_pool); |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
s->yuv_wide_width = s->rgb_wide_width; |
|
|
|
|
|
|
|
s->yuv_wide_height = s->rgb_wide_height; |
|
|
|
|
|
|
|
s->yuv_wide_buf_size = s->rgb_wide_width * s->rgb_wide_height * 3 / 2; |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
for (int i=0; i<YUV_COUNT; i++) { |
|
|
|
|
|
|
|
s->yuv_wide_ion[i] = visionbuf_allocate_cl(s->yuv_wide_buf_size, s->device_id, s->context, &s->yuv_wide_cl[i]); |
|
|
|
|
|
|
|
s->yuv_wide_bufs[i].y = (uint8_t*)s->yuv_wide_ion[i].addr; |
|
|
|
|
|
|
|
s->yuv_wide_bufs[i].u = s->yuv_wide_bufs[i].y + (s->yuv_wide_width * s->yuv_wide_height); |
|
|
|
|
|
|
|
s->yuv_wide_bufs[i].v = s->yuv_wide_bufs[i].u + (s->yuv_wide_width/2 * s->yuv_wide_height/2); |
|
|
|
|
|
|
|
} |
|
|
|
|
|
|
|
#endif |
|
|
|
|
|
|
|
|
|
|
|
if (s->cameras.rear.ci.bayer) { |
|
|
|
if (s->cameras.rear.ci.bayer) { |
|
|
|
// debayering does a 2x downscale
|
|
|
|
// debayering does a 2x downscale
|
|
|
|
s->yuv_transform = transform_scale_buffer(s->cameras.rear.transform, 0.5); |
|
|
|
s->yuv_transform = transform_scale_buffer(s->cameras.rear.transform, 0.5); |
|
|
|
@ -1137,7 +1436,24 @@ void init_buffers(VisionState *s) { |
|
|
|
s->krnl_debayer_front = clCreateKernel(s->prg_debayer_front, "debayer10", &err); |
|
|
|
s->krnl_debayer_front = clCreateKernel(s->prg_debayer_front, "debayer10", &err); |
|
|
|
assert(err == 0); |
|
|
|
assert(err == 0); |
|
|
|
} |
|
|
|
} |
|
|
|
|
|
|
|
#ifdef QCOM2 |
|
|
|
|
|
|
|
if (s->cameras.wide.ci.bayer) { |
|
|
|
|
|
|
|
s->prg_debayer_wide = build_debayer_program(s, s->cameras.wide.ci.frame_width, s->cameras.wide.ci.frame_height, |
|
|
|
|
|
|
|
s->cameras.wide.ci.frame_stride, |
|
|
|
|
|
|
|
s->rgb_wide_width, s->rgb_wide_height, s->rgb_wide_stride, |
|
|
|
|
|
|
|
s->cameras.wide.ci.bayer_flip, s->cameras.wide.ci.hdr); |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
s->krnl_debayer_wide = clCreateKernel(s->prg_debayer_wide, "debayer10", &err); |
|
|
|
|
|
|
|
assert(err == 0); |
|
|
|
|
|
|
|
} |
|
|
|
|
|
|
|
#endif |
|
|
|
|
|
|
|
s->debayer_cl_localMemSize = (DEBAYER_LOCAL_WORKSIZE + 2 * (3 / 2)) * (DEBAYER_LOCAL_WORKSIZE + 2 * (3 / 2)) * sizeof(float); |
|
|
|
|
|
|
|
s->debayer_cl_globalWorkSize[0] = s->rgb_width; |
|
|
|
|
|
|
|
s->debayer_cl_globalWorkSize[1] = s->rgb_height; |
|
|
|
|
|
|
|
s->debayer_cl_localWorkSize[0] = DEBAYER_LOCAL_WORKSIZE; |
|
|
|
|
|
|
|
s->debayer_cl_localWorkSize[1] = DEBAYER_LOCAL_WORKSIZE; |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
#ifdef QCOM |
|
|
|
s->prg_rgb_laplacian = build_conv_program(s, s->rgb_width/NUM_SEGMENTS_X, s->rgb_height/NUM_SEGMENTS_Y, |
|
|
|
s->prg_rgb_laplacian = build_conv_program(s, s->rgb_width/NUM_SEGMENTS_X, s->rgb_height/NUM_SEGMENTS_Y, |
|
|
|
3); |
|
|
|
3); |
|
|
|
s->krnl_rgb_laplacian = clCreateKernel(s->prg_rgb_laplacian, "rgb2gray_conv2d", &err); |
|
|
|
s->krnl_rgb_laplacian = clCreateKernel(s->prg_rgb_laplacian, "rgb2gray_conv2d", &err); |
|
|
|
@ -1157,9 +1473,13 @@ void init_buffers(VisionState *s) { |
|
|
|
s->conv_cl_localWorkSize[1] = CONV_LOCAL_WORKSIZE; |
|
|
|
s->conv_cl_localWorkSize[1] = CONV_LOCAL_WORKSIZE; |
|
|
|
|
|
|
|
|
|
|
|
for (int i=0; i<(ROI_X_MAX-ROI_X_MIN+1)*(ROI_Y_MAX-ROI_Y_MIN+1); i++) {s->lapres[i] = 16160;} |
|
|
|
for (int i=0; i<(ROI_X_MAX-ROI_X_MIN+1)*(ROI_Y_MAX-ROI_Y_MIN+1); i++) {s->lapres[i] = 16160;} |
|
|
|
|
|
|
|
#endif |
|
|
|
|
|
|
|
|
|
|
|
rgb_to_yuv_init(&s->rgb_to_yuv_state, s->context, s->device_id, s->yuv_width, s->yuv_height, s->rgb_stride); |
|
|
|
rgb_to_yuv_init(&s->rgb_to_yuv_state, s->context, s->device_id, s->yuv_width, s->yuv_height, s->rgb_stride); |
|
|
|
rgb_to_yuv_init(&s->front_rgb_to_yuv_state, s->context, s->device_id, s->yuv_front_width, s->yuv_front_height, s->rgb_front_stride); |
|
|
|
rgb_to_yuv_init(&s->front_rgb_to_yuv_state, s->context, s->device_id, s->yuv_front_width, s->yuv_front_height, s->rgb_front_stride); |
|
|
|
|
|
|
|
#ifdef QCOM2 |
|
|
|
|
|
|
|
rgb_to_yuv_init(&s->wide_rgb_to_yuv_state, s->context, s->device_id, s->yuv_wide_width, s->yuv_wide_height, s->rgb_wide_stride); |
|
|
|
|
|
|
|
#endif |
|
|
|
} |
|
|
|
} |
|
|
|
|
|
|
|
|
|
|
|
void free_buffers(VisionState *s) { |
|
|
|
void free_buffers(VisionState *s) { |
|
|
|
@ -1167,18 +1487,28 @@ void free_buffers(VisionState *s) { |
|
|
|
for (int i=0; i<FRAME_BUF_COUNT; i++) { |
|
|
|
for (int i=0; i<FRAME_BUF_COUNT; i++) { |
|
|
|
visionbuf_free(&s->camera_bufs[i]); |
|
|
|
visionbuf_free(&s->camera_bufs[i]); |
|
|
|
visionbuf_free(&s->front_camera_bufs[i]); |
|
|
|
visionbuf_free(&s->front_camera_bufs[i]); |
|
|
|
|
|
|
|
#ifdef QCOM |
|
|
|
visionbuf_free(&s->focus_bufs[i]); |
|
|
|
visionbuf_free(&s->focus_bufs[i]); |
|
|
|
visionbuf_free(&s->stats_bufs[i]); |
|
|
|
visionbuf_free(&s->stats_bufs[i]); |
|
|
|
|
|
|
|
#elif defined(QCOM2) |
|
|
|
|
|
|
|
visionbuf_free(&s->wide_camera_bufs[i]); |
|
|
|
|
|
|
|
#endif |
|
|
|
} |
|
|
|
} |
|
|
|
|
|
|
|
|
|
|
|
for (int i=0; i<UI_BUF_COUNT; i++) { |
|
|
|
for (int i=0; i<UI_BUF_COUNT; i++) { |
|
|
|
visionbuf_free(&s->rgb_bufs[i]); |
|
|
|
visionbuf_free(&s->rgb_bufs[i]); |
|
|
|
visionbuf_free(&s->rgb_front_bufs[i]); |
|
|
|
visionbuf_free(&s->rgb_front_bufs[i]); |
|
|
|
|
|
|
|
#ifdef QCOM2 |
|
|
|
|
|
|
|
visionbuf_free(&s->rgb_wide_bufs[i]); |
|
|
|
|
|
|
|
#endif |
|
|
|
} |
|
|
|
} |
|
|
|
|
|
|
|
|
|
|
|
for (int i=0; i<YUV_COUNT; i++) { |
|
|
|
for (int i=0; i<YUV_COUNT; i++) { |
|
|
|
visionbuf_free(&s->yuv_ion[i]); |
|
|
|
visionbuf_free(&s->yuv_ion[i]); |
|
|
|
visionbuf_free(&s->yuv_front_ion[i]); |
|
|
|
visionbuf_free(&s->yuv_front_ion[i]); |
|
|
|
|
|
|
|
#ifdef QCOM2 |
|
|
|
|
|
|
|
visionbuf_free(&s->yuv_wide_ion[i]); |
|
|
|
|
|
|
|
#endif |
|
|
|
} |
|
|
|
} |
|
|
|
|
|
|
|
|
|
|
|
clReleaseMemObject(s->rgb_conv_roi_cl); |
|
|
|
clReleaseMemObject(s->rgb_conv_roi_cl); |
|
|
|
@ -1189,7 +1519,11 @@ void free_buffers(VisionState *s) { |
|
|
|
clReleaseProgram(s->prg_debayer_front); |
|
|
|
clReleaseProgram(s->prg_debayer_front); |
|
|
|
clReleaseKernel(s->krnl_debayer_rear); |
|
|
|
clReleaseKernel(s->krnl_debayer_rear); |
|
|
|
clReleaseKernel(s->krnl_debayer_front); |
|
|
|
clReleaseKernel(s->krnl_debayer_front); |
|
|
|
|
|
|
|
#ifdef QCOM2 |
|
|
|
|
|
|
|
clReleaseProgram(s->prg_debayer_wide); |
|
|
|
|
|
|
|
clReleaseKernel(s->krnl_debayer_wide); |
|
|
|
|
|
|
|
#endif |
|
|
|
|
|
|
|
|
|
|
|
clReleaseProgram(s->prg_rgb_laplacian); |
|
|
|
clReleaseProgram(s->prg_rgb_laplacian); |
|
|
|
clReleaseKernel(s->krnl_rgb_laplacian); |
|
|
|
clReleaseKernel(s->krnl_rgb_laplacian); |
|
|
|
|
|
|
|
|
|
|
|
@ -1211,13 +1545,18 @@ void party(VisionState *s) { |
|
|
|
processing_thread, s); |
|
|
|
processing_thread, s); |
|
|
|
assert(err == 0); |
|
|
|
assert(err == 0); |
|
|
|
|
|
|
|
|
|
|
|
#if !defined(QCOM2) && !defined(__APPLE__) |
|
|
|
#ifndef __APPLE__ |
|
|
|
// TODO: fix front camera on qcom2
|
|
|
|
|
|
|
|
pthread_t frontview_thread_handle; |
|
|
|
pthread_t frontview_thread_handle; |
|
|
|
err = pthread_create(&frontview_thread_handle, NULL, |
|
|
|
err = pthread_create(&frontview_thread_handle, NULL, |
|
|
|
frontview_thread, s); |
|
|
|
frontview_thread, s); |
|
|
|
assert(err == 0); |
|
|
|
assert(err == 0); |
|
|
|
#endif |
|
|
|
#endif |
|
|
|
|
|
|
|
#ifdef QCOM2 |
|
|
|
|
|
|
|
pthread_t wideview_thread_handle; |
|
|
|
|
|
|
|
err = pthread_create(&wideview_thread_handle, NULL, |
|
|
|
|
|
|
|
wideview_thread, s); |
|
|
|
|
|
|
|
assert(err == 0); |
|
|
|
|
|
|
|
#endif |
|
|
|
|
|
|
|
|
|
|
|
// priority for cameras
|
|
|
|
// priority for cameras
|
|
|
|
err = set_realtime_priority(51); |
|
|
|
err = set_realtime_priority(51); |
|
|
|
@ -1229,15 +1568,23 @@ void party(VisionState *s) { |
|
|
|
tbuffer_stop(&s->ui_front_tb); |
|
|
|
tbuffer_stop(&s->ui_front_tb); |
|
|
|
pool_stop(&s->yuv_pool); |
|
|
|
pool_stop(&s->yuv_pool); |
|
|
|
pool_stop(&s->yuv_front_pool); |
|
|
|
pool_stop(&s->yuv_front_pool); |
|
|
|
|
|
|
|
#ifdef QCOM2 |
|
|
|
|
|
|
|
tbuffer_stop(&s->ui_wide_tb); |
|
|
|
|
|
|
|
pool_stop(&s->yuv_wide_pool); |
|
|
|
|
|
|
|
#endif |
|
|
|
|
|
|
|
|
|
|
|
zsock_signal(s->terminate_pub, 0); |
|
|
|
zsock_signal(s->terminate_pub, 0); |
|
|
|
|
|
|
|
|
|
|
|
#if (defined(QCOM) && !defined(QCOM_REPLAY)) || defined(WEBCAM) |
|
|
|
#if (defined(QCOM) && !defined(QCOM_REPLAY)) || defined(WEBCAM) || defined(QCOM2) |
|
|
|
LOG("joining frontview_thread"); |
|
|
|
LOG("joining frontview_thread"); |
|
|
|
err = pthread_join(frontview_thread_handle, NULL); |
|
|
|
err = pthread_join(frontview_thread_handle, NULL); |
|
|
|
assert(err == 0); |
|
|
|
assert(err == 0); |
|
|
|
#endif |
|
|
|
#endif |
|
|
|
|
|
|
|
#ifdef QCOM2 |
|
|
|
|
|
|
|
LOG("joining wideview_thread"); |
|
|
|
|
|
|
|
err = pthread_join(wideview_thread_handle, NULL); |
|
|
|
|
|
|
|
assert(err == 0); |
|
|
|
|
|
|
|
#endif |
|
|
|
LOG("joining visionserver_thread"); |
|
|
|
LOG("joining visionserver_thread"); |
|
|
|
err = pthread_join(visionserver_thread_handle, NULL); |
|
|
|
err = pthread_join(visionserver_thread_handle, NULL); |
|
|
|
assert(err == 0); |
|
|
|
assert(err == 0); |
|
|
|
@ -1271,11 +1618,17 @@ int main(int argc, char *argv[]) { |
|
|
|
|
|
|
|
|
|
|
|
init_buffers(s); |
|
|
|
init_buffers(s); |
|
|
|
|
|
|
|
|
|
|
|
#if (defined(QCOM) && !defined(QCOM_REPLAY)) || defined(QCOM2) |
|
|
|
#if defined(QCOM) && !defined(QCOM_REPLAY) |
|
|
|
s->pm = new PubMaster({"frame", "frontFrame", "thumbnail"}); |
|
|
|
s->pm = new PubMaster({"frame", "frontFrame", "thumbnail"}); |
|
|
|
|
|
|
|
#elif defined(QCOM2) |
|
|
|
|
|
|
|
s->pm = new PubMaster({"frame", "frontFrame", "wideFrame", "thumbnail"}); |
|
|
|
#endif |
|
|
|
#endif |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
#ifndef QCOM2 |
|
|
|
cameras_open(&s->cameras, &s->camera_bufs[0], &s->focus_bufs[0], &s->stats_bufs[0], &s->front_camera_bufs[0]); |
|
|
|
cameras_open(&s->cameras, &s->camera_bufs[0], &s->focus_bufs[0], &s->stats_bufs[0], &s->front_camera_bufs[0]); |
|
|
|
|
|
|
|
#else |
|
|
|
|
|
|
|
cameras_open(&s->cameras, &s->camera_bufs[0], &s->focus_bufs[0], &s->stats_bufs[0], &s->front_camera_bufs[0], &s->wide_camera_bufs[0]); |
|
|
|
|
|
|
|
#endif |
|
|
|
|
|
|
|
|
|
|
|
party(s); |
|
|
|
party(s); |
|
|
|
|
|
|
|
|
|
|
|
|
robbederks
5 years ago
committed by
GitHub