I have a problem, I am trying to implement a median filter. The task to be solved is to calculate a median filter on a matrix, which has to be repeated x times. The problem I have is that when the matrix gets bigger than 2048x2048, the result is no longer correct.
__kernel void saxpy_kernel(float alpha, int size, __global float *in,
__global float *out, int loops) {
int i = get_global_id(0);
int index = 0;
for (int loop = 0; loop < loops; loop++) {
for (int j = 0; j < size; j++) {
index = j * size + i;
float center = in[index];
float centerTop = (index > size) ? in[index - size] : 0;
float centerLeft =
(index > 0 && (index - 1) % size < (size - 1)) ? in[index - 1] : 0;
float centerRight = ((index + 1) % size > 0) ? in[index + 1] : 0;
float centerDown =
(index + size < (size * size)) ? centerDown = in[index + size] : 0;
// write in buffer
out[index] =
alpha * (center + centerTop + centerLeft + centerRight + centerDown);
}
if (loop < loops - 1) {
barrier(CLK_GLOBAL_MEM_FENCE);
for (int j = 0; j < size; j++) {
index = j * size + i;
*(in + index) = *(out + index);
}
barrier(CLK_GLOBAL_MEM_FENCE);
}
}
}
The kernel is loaded as follows.
size_t global_size[] = {VECTOR_SIZE, VECTOR_SIZE, 0}; //My matrix pattern
size_t group_pattern[] = {1, VECTOR_SIZE, 0}; //My workgroups pattern
size_t local_size[] = {global_size[0] / group_pattern[0], global_size[1] / group_pattern[1]}; //My workgroups pattern
clStatus = clEnqueueNDRangeKernel(command_queue, kernel, 2, NULL, global_size, local_size, 0, NULL, NULL);
I hope someone can help me. I would be happy to receive any other feedback on what can be optimised.