1

我对 4 点模板 OpenCL 代码有疑问。代码运行良好,但我没有得到预期的 symetrics 最终 2D 值。

我怀疑这是内核代码中更新值的问题。这是内核代码:

// 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 表示 2D 模板的 4 个值。我计算 new_value 并在"barrier(CLK_GLOBAL_MEM_FENCE)".

但是,不同的工作项之间似乎存在冲突。我该如何解决这个问题?

4

1 回答 1

2

您使用的屏障 GLOBAL_MEM_FENCE不会按预期同步所有工作项。它只与一个工作组同步访问。

通常不会同时执行所有工作组,因为它们仅在少数物理内核上调度,并且在内核内不可能进行全局同步。

解决方案是将输出写入不同的缓冲区。

于 2012-09-15T17:50:43.337 回答