1

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 ?

4

1 に答える 1

2

使用するバリアGLOBAL_MEM_FENCEは、意図したとおりにすべての作業項目を同期するわけではありません。アクセスを単一のワークグループと同期するだけです。

通常、すべてのワークグループが同時に実行されることはありません。これは、少数の物理コアでのみスケジュールされ、カーネル内でグローバル同期ができないためです。

解決策は、出力を別のバッファーに書き込むことです。

于 2012-09-15T17:50:43.337 に答える