From 1f1bcb246e1560e63386d9828ccc8fe9cf1a1949 Mon Sep 17 00:00:00 2001
From: Joost Wooning <jwooning@gmail.com>
Date: Thu, 5 May 2022 13:06:13 +0200
Subject: [PATCH] make debayering consistent at edges (#24437)

old-commit-hash: 1f9907122a9bda9279d73f698addaa0e22796059
---
 selfdrive/camerad/cameras/real_debayer.cl | 18 +++++++++---------
 1 file changed, 9 insertions(+), 9 deletions(-)

diff --git a/selfdrive/camerad/cameras/real_debayer.cl b/selfdrive/camerad/cameras/real_debayer.cl
index 6452e44ebb..8cb0aeb166 100644
--- a/selfdrive/camerad/cameras/real_debayer.cl
+++ b/selfdrive/camerad/cameras/real_debayer.cl
@@ -104,38 +104,38 @@ __kernel void debayer10(const __global uchar * in,
   half pv = val_from_10(in, x_global, y_global, black_level);
   cached[localOffset] = pv;
 
-  // don't care
-  if (x_global < 1 || x_global >= RGB_WIDTH - 1 || y_global < 1 || y_global >= RGB_HEIGHT - 1) {
-    return;
-  }
-
   // cache padding
   int localColOffset = -1;
   int globalColOffset = -1;
 
   // cache padding
-  if (x_local < 1) {
+  if (x_global >= 1 && x_local < 1) {
     localColOffset = x_local;
     globalColOffset = -1;
     cached[(y_local + 1) * localRowLen + x_local] = val_from_10(in, x_global-1, y_global, black_level);
-  } else if (x_local >= get_local_size(0) - 1) {
+  } else if (x_global < RGB_WIDTH - 1 && x_local >= get_local_size(0) - 1) {
     localColOffset = x_local + 2;
     globalColOffset = 1;
     cached[localOffset + 1] = val_from_10(in, x_global+1, y_global, black_level);
   }
 
-  if (y_local < 1) {
+  if (y_global >= 1 && y_local < 1) {
     cached[y_local * localRowLen + x_local + 1] = val_from_10(in, x_global, y_global-1, black_level);
     if (localColOffset != -1) {
       cached[y_local * localRowLen + localColOffset] = val_from_10(in, x_global+globalColOffset, y_global-1, black_level);
     }
-  } else if (y_local >= get_local_size(1) - 1) {
+  } else if (y_global < RGB_HEIGHT - 1 && y_local >= get_local_size(1) - 1) {
     cached[(y_local + 2) * localRowLen + x_local + 1] = val_from_10(in, x_global, y_global+1, black_level);
     if (localColOffset != -1) {
       cached[(y_local + 2) * localRowLen + localColOffset] = val_from_10(in, x_global+globalColOffset, y_global+1, black_level);
     }
   }
 
+  // don't care
+  if (x_global < 1 || x_global >= RGB_WIDTH - 1 || y_global < 1 || y_global >= RGB_HEIGHT - 1) {
+    return;
+  }
+
   // sync
   barrier(CLK_LOCAL_MEM_FENCE);