從CUDA內核中隨機訪問小常量數組的工作原理

我的內核使用float大小爲8乘8的數組，其隨機訪問模式如下所示。從CUDA內核中隨機訪問小常量數組的工作原理

// inds - big array of indices in range 0,...,7 
// flts - 8 by 8 array of floats 

// kernel essentially processes large 2D array by looping through slow coordinate 
// and having block/thread parallelization of fast coordinate. 

__global__ void kernel (int const* inds, float const* flt, ...) 
{ 
    int idx = threadIdx.x + blockDim.x * blockIdx.x; // Global fast coordinate 
    int idy;           // Global slow coordinate 
    int sx = gridDim.x * blockDim.x;     // Stride of fast coordinate 

    for (idy = 0; idy < 10000; idy++)  // Slow coordinate loop 
    { 
     int id = idx + idy * sx;    // Global coordinate in 2D array 

     int ind = inds[id];     // Index of random access to small array 

     float f0 = flt[ind * 8 + 0]; 
     ... 
     float f7 = flt[ind * 8 + 7]; 

     NEXT I HAVE SOME ALGEBRAIC FORMULAS THAT USE f0, ..., f7 
    } 
}

什麼是訪問flt數組的最佳方式？

請勿通過flt，請使用__const__內存。我不知道當不同的線程訪問不同的數據時，const內存有多快。
如上使用。因爲線程訪問不同的數據，所以不使用負載均勻。它會因爲緩存而變快嗎？
複製到共享內存並使用共享內存數組。
使用紋理。從未使用過的紋理......這種方法可以快嗎？

對於共享存儲器，它可能是更好的轉置flt陣列，即訪問這種方式，以避免存儲體衝突：

float fj = flt_shared[j * 8 + ind]; // where j = 0, ..., 7

PS：目標架構是費米和開普勒。

來源

2013-03-02 user2052436

那麼，你有沒有嘗試過使用'__constant__'或'__shared__'內存和計時內核？在內核中多次訪問'flt'。我認爲'__shared__'記憶會給你最好的表現。 – sgarizvi 2013-03-02 05:41:46

'__constant__'只有在warp中的所有線程都可以訪問相同的值時纔是一個不錯的選擇。恕我直言，這裏最好的選擇是使用包含轉置數據的共享內存。請指定您的目標計算架構，它可能會產生影響。 – RoBiK 2013-03-02 10:19:50

到目前爲止，我只實現了'inds'數組中所有值都爲0的情況，所以我沒有隨機訪問，我使用'__const__'而不是'flt'參數。順便說一句，在這種情況下，如果使用'-dlcm = cg'選項編譯我的內核，它的工作速度會更快。現在我需要將我的內核擴展到一般情況。 – user2052436 2013-03-03 02:43:14

「最佳」方式還取決於您正在處理的架構。我個人對隨機訪問的經驗（你的訪問似乎是一種隨機的，因爲使用映射inds[id]）在費米和開普勒是現在L1是如此之快，在許多情況下，最好是繼續使用全局內存，而不是共享內存或紋理內存。

加速全球內存隨機訪問：無效的L1高速緩存行

費米和開普勒架構支持兩種類型從全局內存負載。 全緩存是的默認模式，它試圖在L1中命中，然後L2，然後GMEM和負載粒度爲128字節的行。 L2-only嘗試在L2中命中，然後GMEM和負載粒度爲32字節。對於某些隨機訪問模式，可通過使L1無效並利用L2的較低粒度來提高內存效率。這可以通過編譯–Xptxas –dlcm=cg選項到nvcc來完成。

加速全球內存訪問的一般準則：禁止ECC支持

費米和開普勒的GPU支持糾錯碼（ECC）和ECC默認情況下啓用。 ECC降低峯值存儲器帶寬，並被要求增強醫療成像和大規模集羣計算等應用程序中的數據完整性。如果不需要，可以使用Linux上的nvidia-smi實用程序（請參閱link）或通過Microsoft Windows系統上的控制面板禁用以提高性能。請注意，打開或關閉ECC需要重新啓動才能生效。使用只讀數據緩存

開普勒的功能是已知只讀功能的持續時間數據的48KB緩存：

爲開普勒加速全局內存訪問的一般準則。使用只讀路徑是有益的，因爲它可以卸載共享/ L1緩存路徑，並且支持全速未對齊的內存訪問。使用只讀路徑可由編譯器自動管理（使用const __restrict關鍵字）或編程器明確（使用__ldg()固有）。

來源

2013-03-02 07:18:51 JackOLantern

是的，我的目標是費米和開普勒架構。 – user2052436 2013-03-03 02:21:55

您的建議的問題是，如果使用-dlcm = cg選項編譯，我的內核的其餘部分（某些代數操作）使用對其他幾個大型輸入和輸出數據的並集訪問，實際上運行得更快。 – user2052436 2013-03-03 02:30:13

我編輯了我的答案，並通過禁用ECC（Fermi和Kepler）以及使用只讀數據緩存（Kepler）來加速內存訪問。 – JackOLantern 2013-03-04 20:51:43

從CUDA內核中隨機訪問小常量數組的工作原理

回答

相關問題