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>
This commit is contained in:
Joost Wooning 2022-05-06 18:27:41 +02:00 committed by GitHub
parent 9ef17a877f
commit 49e10dc773
No known key found for this signature in database
GPG Key ID: 4AEE18F83AFDEB23
2 changed files with 14 additions and 16 deletions

View File

@ -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);

View File

@ -1 +1 @@
1f9907122a9bda9279d73f698addaa0e22796059
14f411de8085d1cc9e467592c90bcaf95447a467