2023-03-31 05:54:47 +00:00
|
|
|
int twoToOne(int x, int y, int width, int height) {
|
2023-03-31 09:47:33 +00:00
|
|
|
while (x < 0) {
|
|
|
|
x += width;
|
|
|
|
}
|
|
|
|
while (y < 0) {
|
|
|
|
y += height;
|
|
|
|
}
|
|
|
|
x = x % width;
|
|
|
|
y = y % height;
|
|
|
|
return x + y * width;
|
2021-11-08 11:04:58 +00:00
|
|
|
}
|
|
|
|
|
2023-03-31 09:47:33 +00:00
|
|
|
// float gaussian(float x, float y) {
|
|
|
|
// return exp(-(x*x + y*y) / (1.5F * 1.5F * 2.0F));
|
|
|
|
// }
|
2021-11-08 11:04:58 +00:00
|
|
|
|
2023-03-31 09:47:33 +00:00
|
|
|
__kernel void do_filter(__global float *filter_out,
|
|
|
|
__global const float *precomputed,
|
|
|
|
__global const int *pbp, const int width,
|
|
|
|
const int height, const int filter_size) {
|
|
|
|
int i = get_global_id(0);
|
|
|
|
if (i < 0 || i >= width * height) {
|
|
|
|
return;
|
|
|
|
}
|
2021-11-10 09:12:04 +00:00
|
|
|
|
2023-03-31 09:47:33 +00:00
|
|
|
int x = i % width;
|
|
|
|
int y = i / width;
|
2021-01-23 07:42:45 +00:00
|
|
|
|
2023-03-31 09:47:33 +00:00
|
|
|
float sum = 0.0F;
|
|
|
|
for (int q = 0; q < filter_size; ++q) {
|
|
|
|
int q_prime = height - filter_size / 2 + y + q;
|
|
|
|
for (int p = 0; p < filter_size; ++p) {
|
|
|
|
int p_prime = width - filter_size / 2 + x + p;
|
|
|
|
if (pbp[twoToOne(p_prime, q_prime, width, height)] != 0) {
|
|
|
|
sum += precomputed[twoToOne(p, q, filter_size, filter_size)];
|
|
|
|
// sum += gaussian(p - filter_size / 2.0F + 0.5F, q -
|
|
|
|
// filter_size / 2.0F + 0.5F);
|
|
|
|
}
|
2021-01-23 07:42:45 +00:00
|
|
|
}
|
2023-03-31 09:47:33 +00:00
|
|
|
}
|
2021-01-23 07:42:45 +00:00
|
|
|
|
2023-03-31 09:47:33 +00:00
|
|
|
filter_out[i] = sum;
|
2021-01-23 07:42:45 +00:00
|
|
|
}
|
|
|
|
|
|
|
|
// vim: syntax=c
|