OpenCL Matrix Median Filter

Viewed 96

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.

1 Answers

The problem is that global_size[] in your code is 3D and there is a maximum size limit for a 3D range, for example 2048x2048x64 (Max work item sizes in clinfo). Make global_size 1-dimensional and use linear indexing (n=size_x*y+x) to access your 2D matrix coordinates.

Related