2012-09-15 61 views
1

我有一個4點模板OpenCL代碼的問題。代碼運行良好,但我沒有得到預期的symetrics最終2D值。問題與OpenCL模板代碼

我懷疑這是內核代碼中更新值的問題。以下是內核代碼:

// 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)"後更新它。

但是,似乎不同工作項目之間存在衝突。我怎麼能解決這個問題?

回答

2

您使用的屏障GLOBAL_MEM_FENCE將而不是按預期同步所有工作項目。它只與一個工作組同步訪問。

通常,所有工作組不會同時執行,因爲它們僅在少量物理內核上進行調度,並且在內核中不可能進行全局同步。

解決方法是將輸出寫入不同的緩衝區。