2011-05-03 113 views
6

我試圖寫在OpenCL的直方圖內核來計算256倉R,G,和B的直方圖。我的內核是這樣的:OpenCL的圖像直方圖

const sampler_t mSampler = CLK_NORMALIZED_COORDS_FALSE | 
          CLK_ADDRESS_CLAMP| 
          CLK_FILTER_NEAREST; 


__kernel void computeHistogram(read_only image2d_t input, __global int* rOutput, 
           __global int* gOutput, __global int* bOutput) 
{ 

    int2 coords = {get_global_id(0), get_global_id(1)}; 

    float4 sample = read_imagef(input, mSampler, coords); 

    uchar rbin = floor(sample.x * 255.0f); 
    uchar gbin = floor(sample.y * 255.0f); 
    uchar bbin = floor(sample.z * 255.0f); 

    rOutput[rbin]++; 
    gOutput[gbin]++; 
    bOutput[bbin]++; 


} 

當我的2100 X 894的圖像上運行(1877400個像素)我往往只在看到或正在錄製的周圍1,870,000總價值時,我總結直方圖值對每個通道。每次都是不同的數字。我確實希望這樣做,因爲偶爾有兩個內核可能會從輸出數組中獲取相同的值並增加它,從而有效地取消一個增量操作(我假設?)。

的1,870,000輸出爲{1,1}工作組規模(這是什麼似乎得到默認設置,如果我沒有另行指定)。如果我強制像{10,6}這樣的較大的工作組大小,我會在直方圖中得到一個非常小的總和(與工作組大小的變化成正比)。這似乎很奇怪,我,但我猜會發生什麼情況是,所有的組中的工作項目的同時增加輸出數組值,所以它只是被計爲一次增量?

不管怎麼說,我讀過在OpenCL的沒有全局內存同步化,只有使用他們的__local內存本地工作組內同步化規範。由NVIDIA直方圖示例打破了直方圖工作量成一束特定大小的子問題,計算它們的局部直方圖,然後將結果合併到後一個單一的直方圖。對於任意大小的圖像來說,這看起來不會很好。我想我可以用虛擬值填充圖像數據...

作爲OpenCL的新手,我想我想知道是否有更直接的方法來做到這一點(因爲它看起來應該是一個相對簡單的方法GPGPU問題)。

謝謝!

回答

5

如前所述,你寫入到共享存儲器的非同步和非原子。這導致錯誤。如果畫面足夠大,我有一個建議:

拆分工作組爲一維之一的cols或行。使用每個內核總結col或row的直方圖,然後使用atomic atom_inc全局求和它。這帶來了私人記憶中最快速和最快的原子操作。

如果你在兩個維度上工作,你可以在圖片的某些部分進行。

[編輯:]

我想,我有更好的回答:;-)

,看一下:http://developer.download.nvidia.com/compute/opencl/sdk/website/samples.html#oclHistogram

他們有一個有趣的實施有...

5

是的,你寫的,從在同一時間許多工作項的共享內存,因此你將失去的元素,如果你不這樣做在一個安全的方式(或更新更糟?乾脆不要做它)。組大小的增加實際上增加了計算設備的利用率,反過來又增加了衝突的可能性。所以你最終會失去更多的更新。但是,您似乎混淆了同步(訂購線程執行順序)和共享內存更新(通常需要原子操作或代碼同步內存屏障,以確保內存更新對其他線程可見這是同步的)。

對於您的情況,同步+屏障並不是特別有用(因爲您注意到無法用於全局同步,原因是,2個線程組可能永遠不會同時運行,因此嘗試同步它們是無意義的)。它通常用於所有線程開始生成公共數據集的工作,然後所有線程開始以不同的訪問模式使用該數據集。

在你的情況,你可以使用原子操作(例如atom_inc,見http://www.cmsoft.com.br/index.php?option=com_content&view=category&layout=blog&id=113&Itemid=168)。但是,請注意更新高度爭內存地址(比如說,因爲你有數千個線程嘗試所有寫僅256整數的)有可能產生業績不佳。所有箍式典型直方圖代碼都經過以減少對直方圖數據的爭用。