OpenCL atom_inc分離的好主意？

我想使用OpenCL對圖像中的非零點數進行計數。OpenCL atom_inc分離的好主意？

由於它是一項添加工作，我使用了atom_inc。

此處顯示內核代碼。

__kernel void points_count(__global unsigned char* image_data, __global int* total_number, __global int image_width) 
{ 
    size_t gidx = get_global_id(0); 
    size_t gidy = get_global_id(1); 
    if(0!=*(image_data+gidy*image_width+gidx)) 
    { 
     atom_inc(total_number); 
    } 
}

我的問題是，通過使用atom_inc它會是多餘的權利？

當我們遇到一個非零點時，我們應該等待atom_inc。

我有這樣的想法，我們可以將整行分成數百個組，我們在不同的組中找到數字並最後添加它們。

如果我們可以做這樣的事情：

__kernel void points_count(__global unsigned char* image_data, __global int* total_number_array, __global int image_width) 
{ 
    size_t gidx = get_global_id(0); 
    size_t gidy = get_global_id(1); 
    if(0!=*(image_data+gidy*image_width+gidx)) 
    { 
     int stepy=gidy%10; 
     atom_inc(total_number_array+stepy); 
    }  
}

我們將整個問題分成多個組。在這種情況下，我們可以將數字逐個添加到total_number_array中。

從理論上講，它會有很大的性能提升嗎？

那麼，有沒有人有關於求和問題的一些建議嗎？

謝謝！

來源

2014-02-11 David Ding

減少問題。解釋數千次。 Atomic_inc（）效率不高，需要在O（log2（））模式下執行。 – DarkZeros

就像評論中提到的那樣，這是一個減少問題。

這個想法是保持單獨的計數，然後在最後把它們放回去。

考慮使用本地內存來存儲值。

聲明本地緩衝區以供每個工作組使用。
通過使用local_id作爲索引來跟蹤此緩衝區中出現的次數。
在執行結束時將這些值相加。

來源

2014-02-11 14:18:51 OlivierLi

一個很好的介紹瞭如何使用OpenCL的還原問題如下所示： http://developer.amd.com/resources/documentation-articles/articles-whitepapers/opencl-optimization-case-study-simple-reductions/

減少內核可能看起來像這樣（從上面的鏈接所）：

__kernel 
void reduce(
      __global float* buffer, 
      __local float* scratch, 
      __const int length, 
      __global float* result) { 

    int global_index = get_global_id(0); 
    int local_index = get_local_id(0); 
    // Load data into local memory 
    if (global_index < length) { 
    scratch[local_index] = buffer[global_index]; 
    } else { 
    // Infinity is the identity element for the min operation 
    scratch[local_index] = INFINITY; 
    } 
    barrier(CLK_LOCAL_MEM_FENCE); 
    for(int offset = get_local_size(0)/2; 
     offset > 0; 
     offset >>= 1) { 
    if (local_index < offset) { 
     float other = scratch[local_index + offset]; 
     float mine = scratch[local_index]; 
     scratch[local_index] = (mine < other) ? mine : other; 
    } 
    barrier(CLK_LOCAL_MEM_FENCE); 
    } 
    if (local_index == 0) { 
    result[get_group_id(0)] = scratch[0]; 
    } 
}

爲了進一步解釋請參閱建議的鏈接。

來源

2014-02-13 15:48:30 cguenther

OpenCL atom_inc分離的好主意？

回答

相關問題