是否值得通過共享內存傳遞內核參數？

Question

假設我們有一個數組int * data，每個線程將訪問這個數組的一個元素。由於這個數組將在所有線程之間共享，它將被保存在全局內存中。

讓我們創建一個測試內核：

__global__ void test(int *data, int a, int b, int c){ ... } 

我知道肯定的data陣列將在全局內存，因爲我使用cudaMalloc這個數組分配內存。現在至於其他變量，我已經看到了一些不通過分配內存就將整數傳遞給內核函數的例子。在我的情況下，這些變量是ab和c。

如果我沒有記錯的話，即使我們不直接調用cudaMalloc分配4個字節爲每三個整數，CUDA會自動爲我們做，所以最終的變量ab和c將在分配全球記憶。

現在這些變量只是輔助的，線程只能讀取它們，沒有別的。

我的問題是，將這些變量轉移到共享內存不是更好嗎？

我可以想象，如果我們有例如10塊與1024線程，我們需要10*3 = 30以便存儲在每個塊中的共享存儲器中的號碼讀取4字節。

沒有共享內存，並且如果每個線程都必須讀取所有這三個變量一次，則全局內存讀取的總數將是1024*10*3 = 30720，這非常低效。

現在，這裏的問題是，我有點新的CUDA，我不知道，如果有可能，而不必每個線程讀取這些變量傳遞變量ab和c每個塊的共享內存的內存從全局內存中讀取並將它們加載到共享內存中，因此最終全局內存讀取的總數將是1024*10*3 = 30720而不是10*3 = 30。

在以下website有這個例子：

__global__ void staticReverse(int *d, int n) 
{ 
    __shared__ int s[64]; 
    int t = threadIdx.x; 
    int tr = n-t-1; 
    s[t] = d[t]; 
    __syncthreads(); 
    d[t] = s[tr]; 
} 

這裏每個線程的負載共享變量s內不同的數據。因此，每個線程根據其索引將指定的數據加載到共享內存中。

在我的情況下，我想只加載變量ab和c到共享內存。這些變量總是相同的，它們不會改變，所以它們與線程本身沒有任何關係，它們是輔助的，並且被每個線程用來運行一些算法。

我該如何解決這個問題？全局內存讀取只能做到total_amount_of_blocks*3是否可以實現這一點？

Answer 1

GPU運行時已經在不需要做任何事情的情況下（以及關於CUDA中的參數傳遞如何工作的假設是錯誤的）做到了這一點。這是目前發生的事情：

• 在計算能力1.0/1.1/1.2/1.3設備中，內核參數由運行時在共享內存中傳遞。
• 在計算能力2.x/3.x/4.x/5.x/6.x設備中，內核參數由運行時在保留的常量內存庫（具有帶廣播的專用高速緩存）中傳遞。

因此，在假想的內核

__global__ void test(int *data, int a, int b, int c){ ... } 

data，a，b，和c都由值在共用存儲器或常數存儲器（取決於GPU架構）傳遞給每個塊自動。做你的建議沒有任何好處。