From dba466d5f1b82fde7c25d1223fdfae0f0e7030b5 Mon Sep 17 00:00:00 2001
From: Comma Device <device@comma.ai>
Date: Fri, 15 Mar 2024 18:01:18 -0500
Subject: [PATCH] take

---
 system/camerad/cameras/camera_common.cc | 5 +++--
 system/camerad/cameras/real_debayer.cl  | 2 +-
 2 files changed, 4 insertions(+), 3 deletions(-)

diff --git a/system/camerad/cameras/camera_common.cc b/system/camerad/cameras/camera_common.cc
index d2e3f6739a..ab9d9529a5 100644
--- a/system/camerad/cameras/camera_common.cc
+++ b/system/camerad/cameras/camera_common.cc
@@ -41,7 +41,7 @@ public:
 
   }
 
-  void queue(cl_command_queue q, VisionBuf cam_buff, cl_mem cam_buf_cl, cl_mem buf_cl, int width, int height, cl_event *debayer_event, bool todump) {
+  void queue(cl_command_queue q, VisionBuf cam_buff, cl_mem cam_buf_cl, cl_mem buf_cl, int width, int height, cl_event *debayer_event, bool todump, int expo_time) {
     if (todump && 0) {
       uint8_t *cam_buf = (uint8_t *)cam_buff.addr;
       printf("queueing buf %zu \n", cam_buff.len);
@@ -52,6 +52,7 @@ public:
 
     CL_CHECK(clSetKernelArg(krnl_, 0, sizeof(cl_mem), &cam_buf_cl));
     CL_CHECK(clSetKernelArg(krnl_, 1, sizeof(cl_mem), &buf_cl));
+    CL_CHECK(clSetKernelArg(krnl_, 2, sizeof(cl_int), &expo_time));
 
     const size_t globalWorkSize[] = {size_t(width / 2), size_t(height / 2)};
     const int debayer_local_worksize = 16;
@@ -128,7 +129,7 @@ bool CameraBuf::acquire() {
 
   double start_time = millis_since_boot();
   cl_event event;
-  debayer->queue(q, camera_bufs[cur_buf_idx], camera_bufs[cur_buf_idx].buf_cl, cur_yuv_buf->buf_cl, rgb_width, rgb_height, &event, (stream_type==0));
+  debayer->queue(q, camera_bufs[cur_buf_idx], camera_bufs[cur_buf_idx].buf_cl, cur_yuv_buf->buf_cl, rgb_width, rgb_height, &event, (stream_type==0), cur_frame_data.integ_lines);
   clWaitForEvents(1, &event);
   CL_CHECK(clReleaseEvent(event));
   cur_frame_data.processing_time = (millis_since_boot() - start_time) / 1000.0;
diff --git a/system/camerad/cameras/real_debayer.cl b/system/camerad/cameras/real_debayer.cl
index 72aff61f84..dce6d9eb0e 100644
--- a/system/camerad/cameras/real_debayer.cl
+++ b/system/camerad/cameras/real_debayer.cl
@@ -170,7 +170,7 @@ float get_k(float a, float b, float c, float d) {
   return 2.0 - (fabs(a - b) + fabs(c - d));
 }
 
-__kernel void debayer10(const __global uchar * in, __global uchar * out)
+__kernel void debayer10(const __global uchar * in, __global uchar * out, int expo_time)
 {
   const int gid_x = get_global_id(0);
   const int gid_y = get_global_id(1);