將共享內存用於相鄰的數組元素？

Question

我想用CUDA處理圖像。每個像素的新值是基於一行中的兩個相鄰像素計算的。對於像素值使用__shared__內存是否有意義，因爲每個值只能使用兩次？不是瓷磚也是錯誤的做法，因爲它不適合問題結構？我的方法是在每個像素上運行一個線程，併爲每個線程每次加載相鄰的像素值。

Answer 1

目前支持的所有CUDA架構都有緩存。 從計算能力3.5開始，這些對於只讀數據特別有效（因爲讀寫數據只能在L2中緩存，L1緩存僅限於只讀數據）。如果您將輸入數據的指針標記爲const __restrict__，則編譯器很可能會加載它via the L1 texture cache。您也可以通過明確使用__ldg() builtin來強制執行此操作。

雖然可以通過共享內存顯式管理從鄰近像素重新使用數據，但您可能會發現這比只依賴緩存沒有什麼好處。

當然，無論您是否使用共享內存，您希望在x方向上最大化塊大小，並使用1的blockSize.y作爲最佳訪問局部性。

Answer 2

結合使用共享內存和利用合併內存訪問。你所需要做的就是確保圖像以行方式存儲。每個塊都會處理一個線性數組塊。由於數據重用（除了第一個和最後一個像素都會參與處理三次），如果在內核開始時您將複製將要處理的所有像素的值到共享內存，將是有益的。