針對CUDA中的圖像行/列的線索索引

Question

因此，我正在研究CUDA程序，並且在索引塊和線程時遇到了一些問題。基本上，我試圖在CUDA中實現Pixel Sort算法。

我被可視化它是簡單地與每個1級的線程中運行的N個塊（對於行數的方式（有一個修改中，我們只，而不是兩個同時處理要麼行或列），或列），並讓每個塊獨​​立處理該行/列。

所以，如果我們要進行排序列，我們推出這樣的內核（也有一對額外的參數僅僅是我們的具體處理相關的，所以我離開了出來爲簡單起見）

pixel_sort<<<cols, 1>>>(d_image, d_imageSort, rows, cols, ...); 

然後在內核中，我訪問與

int tid = blockIdx.x; 

塊索引這使我每塊一個行/列中的數據的工作，但仍存在一些問題。它的運行速度比我們對較小圖像的算法的串行執行速度要慢，並且在圖像尺寸變得過大時直接崩潰。

我正在考慮的替代線程方案是將每個圖像的像素映射到一個線程，但是我有幾個關於這個問題的問題。

1. 如果我們要用M個線程（用M行表示N個col）來啓動N個塊，我們應該如何避免每個塊的線程數512（或1024？）。在這個例子中，我們可以讓每個線程在列中處理多個像素嗎？索引在內核中的外觀如何？
2. 該算法基本上要求我們在整個列上工作，因此每個線程都不能只是對該像素做了一些工作，他們必須進行通信，可能使用共享內存。如果每塊有一個「主」線程，實際的排序計算，然後讓所有其他線程參與共享內存，這是否是一種有效的策略？

其他註釋：

• 我們的圖像數據通過OpenCV的讀入，並具有存儲在uchar4陣列

Answer 1

如果每個塊的單個線程的RGBA值，則很快就會遇到線程佔用問題。如果您的目標是進行全行排序（對於列，您可以在發送給GPU之前轉換圖像以利用全局合併），但獲得相當好結果的最快方法可能是進行基數或合併排序以每行爲基礎，基本上從http://mgarland.org/files/papers/nvr-2008-001.pdf複製步驟。您可以爲每行分配k個m線程，使得km> =圖像寬度。然後你會啓動k *（圖像高度）塊。那麼你的網格將會是大小（k，高度，1）。

至於你具體的問題：

1. 不能圍繞512/1024線每塊限制得到，你必須調整你的算法。
2. 「主」線程通常設計不佳，造成停頓，開銷，並沒有充分利用多核。您有時可能需要使用單個線程，比如輸出/廣播結果，但大多數情況下您希望避免它。請參閱鏈接的文章以瞭解大多數避免這種情況的示例算法