issue with OpenCL stencil code

Question

I have a problem with a 4-point stencil OpenCL code. The code runs fine but I don't get symetrics final 2D values which are expected.

I suspect it is a problem of updates values in the kernel code. Here's the kernel code :

// kernel code

const char *source ="__kernel void line_compute(const double diagx, const double diagy,\
const double weightx,  const double weighty,  const int size_x,\
 __global double* tab_new, __global double* r)\
{ int iy = get_global_id(0)+1;\
  int ix = get_global_id(1)+1;\
  double new_value, cell, cell_n, cell_s, cell_w, cell_e;\
  double rk;\
  cell_s = tab_new[(iy+1)*(size_x+2)+ix];\
  cell_n = tab_new[(iy-1)*(size_x+2)+ix];\
  cell_e = tab_new[iy*(size_x+2)+(ix+1)];\
  cell_w = tab_new[iy*(size_x+2)+(ix-1)];\
  cell     = tab_new[iy*(size_x+2)+ix];\
  new_value = weighty *( cell_n + cell_s + cell*diagy)+\
                      weightx *( cell_e + cell_w + cell*diagx);\
  rk = cell - new_value;\
  r[iy*(size_x+2)+ix] = rk *rk;\
  barrier(CLK_GLOBAL_MEM_FENCE);\
  tab_new[iy*(size_x+2)+ix] = new_value;\
}";

cell_s, cell_n, cell_e, cell_w represents the 4 values for the 2D stencil. I compute the new_value and update it after a "barrier(CLK_GLOBAL_MEM_FENCE)".

However, it seems there are conflicts between differents work-items. How could I fix this ?

score 1 · Answer 1 · answered Sep 15 '12 at 17:50

The barrier GLOBAL_MEM_FENCE you use will not synchronize all work-items as intended. It does only synchronize access with one single workgroup.

Usually all workgroups won't be executed at the same time, because they are scheduled on only a small number of physical cores, and global synchronization is not possible within a kernel.

The solution is to write the output to a different buffer.

issue with OpenCL stencil code

1 Answers1