debayering: fill outer image borders (#24441)

* debayering: fill outer image borders

* correct conditions

* try something else

* remove global check

* update ref

Co-authored-by: Comma Device <device@comma.ai>
old-commit-hash: 49e10dc773

selfdrive/camerad/cameras/real_debayer.cl

@@ -111,32 +111,30 @@ __kernel void debayer10(const __global uchar * in,
   int localColOffset = -1;
   int globalColOffset = -1;
 
+  const int x_global_mod = (x_global == 0 || x_global == RGB_WIDTH - 1) ? -1: 1;
+  const int y_global_mod = (y_global == 0 || y_global == RGB_HEIGHT - 1) ? -1: 1;
+
   // cache padding
-  if (x_global >= 1 && x_local < 1) {
+  if (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_global < RGB_WIDTH - 1 && x_local >= get_local_size(0) - 1) {
+    cached[(y_local + 1) * localRowLen + x_local] = val_from_10(in, x_global-x_global_mod, y_global, black_level);
+  } else if (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);
+    cached[localOffset + 1] = val_from_10(in, x_global+x_global_mod, y_global, black_level);
   }
 
-  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 (y_local < 1) {
+    cached[y_local * localRowLen + x_local + 1] = val_from_10(in, x_global, y_global-y_global_mod, black_level);
     if (localColOffset != -1) {
-      cached[y_local * localRowLen + localColOffset] = val_from_10(in, x_global+globalColOffset, y_global-1, black_level);
+      cached[y_local * localRowLen + localColOffset] = val_from_10(in, x_global+(x_global_mod*globalColOffset), y_global-y_global_mod, black_level);
     }
-  } 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);
+  } else if (y_local >= get_local_size(1) - 1) {
+    cached[(y_local + 2) * localRowLen + x_local + 1] = val_from_10(in, x_global, y_global+y_global_mod, black_level);
     if (localColOffset != -1) {
-      cached[(y_local + 2) * localRowLen + localColOffset] = val_from_10(in, x_global+globalColOffset, y_global+1, black_level);
-    }
+      cached[(y_local + 2) * localRowLen + localColOffset] = val_from_10(in, x_global+(x_global_mod*globalColOffset), y_global+y_global_mod, black_level);
     }
-
-  // don't care
-  if (x_global < 1 || x_global >= RGB_WIDTH - 1 || y_global < 1 || y_global >= RGB_HEIGHT - 1) {
-    return;
   }
 
   // sync

selfdrive/test/process_replay/debayer_replay_ref_commit

@@ -1 +1 @@
-1f9907122a9bda9279d73f698addaa0e22796059
+14f411de8085d1cc9e467592c90bcaf95447a467