加載從全局內存

假設簡單的內核是這樣的：加載從全局內存

__global__ void fg(struct s_tp tp, struct s_param p) 

{ 

    const uint bid = blockIdx.y * gridDim.x + blockIdx.x; 
    const uint tid = threadIdx.x; 
    const uint idx = bid * blockDim.x + tid; 

    if(idx >= p.ntp) return; 

    double3 r = tp.rh[idx]; 

    double d = sqrt(r.x*r.x + r.y*r.y + r.z*r.z); 

    tp.d[idx] = d; 

}

這是真的：

double3 r = tp.rh[idx];

數據從全局內存加載到R參數。
r存儲在寄存器中，或者在本地存儲器中存在多個變量。
r未存儲在共享內存中。
d是計算出來的，然後寫回全局內存。
寄存器比其他存儲器更快。
如果寄存器的空間已滿（一些大的內核），局部存儲器使用，並且訪問速度較慢
當我需要的雙打，是有什麼辦法可以加快步伐？例如，首先將數據加載到共享內存中，然後對其進行操作？

謝謝大家。

來源

2014-11-22 Henry

'double3 v'用於什麼？它被賦予一個從未使用過的值。你的表述看起來很準確由於每個線程都從'tp.rh'中讀取自己的值，讀入共享內存沒有任何好處。您可以通過在一個線程中處理多個數組元素來加快內核速度。 – 2014-11-22 09:22:20

我忘了刪除它。現在沒問題。 – Henry 2014-11-22 09:24:00

是的，這幾乎都是真的。

•當我需要雙打時，有什麼方法可以加速嗎？例如，首先將數據加載到共享內存中，然後對其進行操作？

當存在數據重用（通常由線程塊中的多個線程加載相同的數據項）時，或者可能在專門使用共享內存時，使用共享內存很有用以幫助全球合併，例如在optimized matrix transpose期間。

數據重用意味着您不止一次地使用（加載）數據，並且爲了共享內存的有用性，這意味着您正在通過多個線程多次加載它。如果您在單個線程中多次使用它，則只需將單個負載加上將其存儲在寄存器中的編譯器（自動）「優化」即可。

編輯 @Jez給出的答案對於優化加載有一些好的想法。我建議另一個想法是將您的AoS數據存儲方案轉換爲SoA方案。數據存儲轉換是提高CUDA代碼速度的常用方法。

您的s_tp結構，您沒有顯示，似乎有每個項目/結構的數量double數量的存儲。如果您爲這些數量中的每一個創建單獨的數組，您將有機會獲得最佳的加載/存儲。事情是這樣的：

__global__ void fg(struct s_tp tp, double* s_tp_rx, double* s_tp_ry, double* s_tp_rz, double* s_tp_d, struct s_param p) 

{ 

    const uint bid = blockIdx.y * gridDim.x + blockIdx.x; 
    const uint tid = threadIdx.x; 
    const uint idx = bid * blockDim.x + tid; 

    if(idx >= p.ntp) return; 

    double rx = s_tp_rx[idx]; 
    double ry = s_tp_ry[idx]; 
    double rz = s_tp_rz[idx]; 

    double d = sqrt(rx*rx + ry*ry + rz*rz); 

    s_tp_d[idx] = d; 

}

這種做法很可能會在你的設備代碼的其他地方有優惠也爲相似類型的使用模式。

來源

2014-11-22 16:24:11

這都是事實。

當我需要雙打時，有什麼方法可以加速嗎？例如首先將數據加載到共享內存中，然後操作它們？

對於您給出的示例，您的實現可能不是最優的。你應該做的第一件事是比較一個參考內核的帶寬，例如一個cudaMemcpy。如果差距很大，並且從縮小這一差距中獲得的提速是非常重要的，那麼優化也許是可能的。在你的內核

尋找有一些罷工我可能不理想的兩件事情：

這裏沒有每個線程太多的工作。如果可能的話，每個線程處理多個元素可以提高性能。這部分是因爲它避免了線程初始化/刪除開銷。
從double3加載不如從其他類型的加載效率。加載數據的最佳方式通常是每個線程使用128位加載。加載三個連續的64位值將會變慢，可能不會很多，但速度會更慢。

編輯：下面羅伯特Crovella的回答給出了一個很好的解決這需要在你的數據類型改變第二點。出於某種原因，我原本以爲這不是一種選擇，所以如果你只是改變你的數據類型，那麼下面的解決方案可能是過頂的！

雖然爲每個線程添加更多工作是一件相當簡單的事情，但對於解決方案來說，優化內存訪問模式（不更改數據類型）並不那麼簡單。幸運的是，有些庫可以提供幫助。我認爲使用CUB，尤其是BlockLoad集體，應該允許您更有效地加載。通過使用轉置運算符爲每個線程加載6 double項，您可以爲每個線程處理兩個元素，將它們打包成double2，並將它們正常存儲。

來源

2014-11-22 18:23:04 Jez

加載從全局內存

回答

相關問題