diff --git a/selfdrive/camerad/cameras/real_debayer.cl b/selfdrive/camerad/cameras/real_debayer.cl index 783abaabb4..6be53c2144 100644 --- a/selfdrive/camerad/cameras/real_debayer.cl +++ b/selfdrive/camerad/cameras/real_debayer.cl @@ -111,34 +111,32 @@ __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 barrier(CLK_LOCAL_MEM_FENCE); diff --git a/selfdrive/test/process_replay/debayer_replay_ref_commit b/selfdrive/test/process_replay/debayer_replay_ref_commit index 44fe33001c..2f37c33669 100644 --- a/selfdrive/test/process_replay/debayer_replay_ref_commit +++ b/selfdrive/test/process_replay/debayer_replay_ref_commit @@ -1 +1 @@ -1f9907122a9bda9279d73f698addaa0e22796059 \ No newline at end of file +14f411de8085d1cc9e467592c90bcaf95447a467 \ No newline at end of file