mawei/StarPilot
1
0
Fork 0
mirror of https://github.com/firestar5683/StarPilot.git synced 2026-10-01 03:43:46 +08:00
Code Issues Packages Projects Releases Wiki Activity

nv12: encoderd avoids a full frame copy (#24519)

Browse Source 
* rgb to nv12

* nv12 works (w memcpy)

* correct now

* no copy

* fix nv12 with fast debayer

* reverts of unused stuff

* ui use nv12

* comment out thumbnails for now

* rebase fix

* dm read nv12

* model read nv12

* fix ffmpeg encoder

* thumbnails from nv12

* replay to nv12

* python framereader support nv12

* remove hardcoded frame/buffer sizes

* fix build

* ffmpeg encoder fix buffers

* small cleanup

* reduce power usage test

* fix cpu usage test

* fix snapshot

* fix loggerd test

* bump cereal

Co-authored-by: Comma Device <device@comma.ai>
Co-authored-by: Joost Wooning <jwooning@gmail.com>
old-commit-hash: ea5b8cdfb10686c972d8d2bc1778e6e6567e45a9
This commit is contained in:
George Hotz
2022-06-01 08:18:28 -07:00
committed by GitHub
parent 82f9f302ed
commit 2a6ffaf71d
28 changed files with 235 additions and 186 deletions
Download Patch File Download Diff File Expand all files Collapse all files
selfdrive/modeld/transforms/transform.cc
+30 -28
View File
@@ -25,7 +25,7 @@ void transform_destroy(Transform* s) {
void transform_queue(Transform* s,
cl_command_queue q,
cl_mem in_yuv, int in_width, int in_height,
cl_mem in_yuv, int in_width, int in_height, int in_stride, int in_uv_offset,
cl_mem out_y, cl_mem out_u, cl_mem out_v,
int out_width, int out_height,
const mat3& projection) {
@@ -44,28 +44,30 @@ void transform_queue(Transform* s,
const int in_y_width = in_width;
const int in_y_height = in_height;
const int in_y_px_stride = 1;
const int in_uv_width = in_width/2;
const int in_uv_height = in_height/2;
const int in_y_offset = 0;
const int in_u_offset = in_y_offset + in_y_width*in_y_height;
const int in_v_offset = in_u_offset + in_uv_width*in_uv_height;
const int in_uv_px_stride = 2;
const int in_u_offset = in_uv_offset;
const int in_v_offset = in_uv_offset + 1;
const int out_y_width = out_width;
const int out_y_height = out_height;
const int out_uv_width = out_width/2;
const int out_uv_height = out_height/2;
CL_CHECK(clSetKernelArg(s->krnl, 0, sizeof(cl_mem), &in_yuv));
CL_CHECK(clSetKernelArg(s->krnl, 1, sizeof(cl_int), &in_y_width));
CL_CHECK(clSetKernelArg(s->krnl, 2, sizeof(cl_int), &in_y_offset));
CL_CHECK(clSetKernelArg(s->krnl, 3, sizeof(cl_int), &in_y_height));
CL_CHECK(clSetKernelArg(s->krnl, 4, sizeof(cl_int), &in_y_width));
CL_CHECK(clSetKernelArg(s->krnl, 5, sizeof(cl_mem), &out_y));
CL_CHECK(clSetKernelArg(s->krnl, 6, sizeof(cl_int), &out_y_width));
CL_CHECK(clSetKernelArg(s->krnl, 7, sizeof(cl_int), &zero));
CL_CHECK(clSetKernelArg(s->krnl, 8, sizeof(cl_int), &out_y_height));
CL_CHECK(clSetKernelArg(s->krnl, 9, sizeof(cl_int), &out_y_width));
CL_CHECK(clSetKernelArg(s->krnl, 10, sizeof(cl_mem), &s->m_y_cl));
CL_CHECK(clSetKernelArg(s->krnl, 0, sizeof(cl_mem), &in_yuv)); // src
CL_CHECK(clSetKernelArg(s->krnl, 1, sizeof(cl_int), &in_stride)); // src_row_stride
CL_CHECK(clSetKernelArg(s->krnl, 2, sizeof(cl_int), &in_y_px_stride)); // src_px_stride
CL_CHECK(clSetKernelArg(s->krnl, 3, sizeof(cl_int), &zero)); // src_offset
CL_CHECK(clSetKernelArg(s->krnl, 4, sizeof(cl_int), &in_y_height)); // src_rows
CL_CHECK(clSetKernelArg(s->krnl, 5, sizeof(cl_int), &in_y_width)); // src_cols
CL_CHECK(clSetKernelArg(s->krnl, 6, sizeof(cl_mem), &out_y)); // dst
CL_CHECK(clSetKernelArg(s->krnl, 7, sizeof(cl_int), &out_y_width)); // dst_row_stride
CL_CHECK(clSetKernelArg(s->krnl, 8, sizeof(cl_int), &zero)); // dst_offset
CL_CHECK(clSetKernelArg(s->krnl, 9, sizeof(cl_int), &out_y_height)); // dst_rows
CL_CHECK(clSetKernelArg(s->krnl, 10, sizeof(cl_int), &out_y_width)); // dst_cols
CL_CHECK(clSetKernelArg(s->krnl, 11, sizeof(cl_mem), &s->m_y_cl)); // M
const size_t work_size_y[2] = {(size_t)out_y_width, (size_t)out_y_height};
@@ -74,21 +76,21 @@ void transform_queue(Transform* s,
const size_t work_size_uv[2] = {(size_t)out_uv_width, (size_t)out_uv_height};
CL_CHECK(clSetKernelArg(s->krnl, 1, sizeof(cl_int), &in_uv_width));
CL_CHECK(clSetKernelArg(s->krnl, 2, sizeof(cl_int), &in_u_offset));
CL_CHECK(clSetKernelArg(s->krnl, 3, sizeof(cl_int), &in_uv_height));
CL_CHECK(clSetKernelArg(s->krnl, 4, sizeof(cl_int), &in_uv_width));
CL_CHECK(clSetKernelArg(s->krnl, 5, sizeof(cl_mem), &out_u));
CL_CHECK(clSetKernelArg(s->krnl, 6, sizeof(cl_int), &out_uv_width));
CL_CHECK(clSetKernelArg(s->krnl, 7, sizeof(cl_int), &zero));
CL_CHECK(clSetKernelArg(s->krnl, 8, sizeof(cl_int), &out_uv_height));
CL_CHECK(clSetKernelArg(s->krnl, 9, sizeof(cl_int), &out_uv_width));
CL_CHECK(clSetKernelArg(s->krnl, 10, sizeof(cl_mem), &s->m_uv_cl));
CL_CHECK(clSetKernelArg(s->krnl, 2, sizeof(cl_int), &in_uv_px_stride)); // src_px_stride
CL_CHECK(clSetKernelArg(s->krnl, 3, sizeof(cl_int), &in_u_offset)); // src_offset
CL_CHECK(clSetKernelArg(s->krnl, 4, sizeof(cl_int), &in_uv_height)); // src_rows
CL_CHECK(clSetKernelArg(s->krnl, 5, sizeof(cl_int), &in_uv_width)); // src_cols
CL_CHECK(clSetKernelArg(s->krnl, 6, sizeof(cl_mem), &out_u)); // dst
CL_CHECK(clSetKernelArg(s->krnl, 7, sizeof(cl_int), &out_uv_width)); // dst_row_stride
CL_CHECK(clSetKernelArg(s->krnl, 8, sizeof(cl_int), &zero)); // dst_offset
CL_CHECK(clSetKernelArg(s->krnl, 9, sizeof(cl_int), &out_uv_height)); // dst_rows
CL_CHECK(clSetKernelArg(s->krnl, 10, sizeof(cl_int), &out_uv_width)); // dst_cols
CL_CHECK(clSetKernelArg(s->krnl, 11, sizeof(cl_mem), &s->m_uv_cl)); // M
CL_CHECK(clEnqueueNDRangeKernel(q, s->krnl, 2, NULL,
(const size_t*)&work_size_uv, NULL, 0, 0, NULL));
CL_CHECK(clSetKernelArg(s->krnl, 2, sizeof(cl_int), &in_v_offset));
CL_CHECK(clSetKernelArg(s->krnl, 5, sizeof(cl_mem), &out_v));
CL_CHECK(clSetKernelArg(s->krnl, 3, sizeof(cl_int), &in_v_offset)); // src_ofset
CL_CHECK(clSetKernelArg(s->krnl, 6, sizeof(cl_mem), &out_v)); // dst
CL_CHECK(clEnqueueNDRangeKernel(q, s->krnl, 2, NULL,
(const size_t*)&work_size_uv, NULL, 0, 0, NULL));