CUDA在哪裏申明共享內存分配常量

Question

我正在爲1024個矩陣運行適應度函數，每個矩陣都有自己的塊並且大小相同。每個塊都有n*n個線程（矩陣的維度），並且需要共享內存，這樣我可以輕鬆地減少總和。但是，所有矩陣的維數n在運行時間之前是可變的（即它可以手動更改，但總是2的乘方因此總和很簡單）。這裏的問題是共享內存必須使用常量分配，但我也需要從主機傳遞給內核的值。我在哪裏聲明維度n，以便它對CPU可見（用於傳遞給內核）並可用於聲明共享內存的大小（在內核中）？

我的代碼的結構是這樣的：

從main.cu我調用內核：

const int num_states = 1024 
const int dimension = 4 

fitness <<< num_states, dimension * dimension >>> (device_array_of_states, dimension, num_states, device_fitness_return); 

，然後在kernel.cu我有：

__global__ void fitness(
    int *numbers, 
    int dimension, 
    int num_states, 
    int *fitness_return) { 
    __shared__ int fitness[16]; <<-- needs to be dimension * dimension 
    //code 
} 

numbers是表示1024點矩陣的陣列，dimension是行和列的長度，num_states是1024,fitness_return我是一個長度爲1024的數組，它保存每個矩陣的適應值。在內核中，共享內存硬編碼爲dimension的正方形（因此在此示例中dimension爲4）。

在哪裏以及如何聲明dimension，以便它可以用來分配共享內存以及調用內核，這樣我只需要在一個地方更新dimension？謝謝你的幫助。

Answer 1

分配的共享內存量在所有塊上是均勻的。您可能會在每個塊中實際使用不同數量的共享內存，但它仍然可用。另外，無論如何，共享內存的數量都非常有限，因此n * n個元素不能超過maximum amount of space（通常爲48KiB）;對於float類型的元素（每個4字節），這將意味着n大約< 340左右。

現在有兩種方法可以分配共享內存：靜態和動態。

靜態分配是你給什麼作爲一個例子，這是行不通的：

__shared__ int fitness[16]; 

在這些情況下

，大小必須在編譯時（被稱爲在設備端代碼編譯時間） - 這不適合你。

隨着動態共享內存分配，你不在內核代碼中指定的大小 - 你leave it empty和前面加上extern：

extern __shared__ int fitness[]; 

相反，在啓動內核時指定的金額，並且不同塊的線程不一定知道它是什麼。

但是在你的情況下，線程做需要知道n是什麼。那麼，只需將它作爲內核參數傳遞即可。所以，

__global__ void fitness(
    int *numbers, 
    int dimension, 
    int num_states, 
    int *fitness_return, 
    unsigned short fitness_matrix_order /* that's your n*/) 
{ 
    extern __shared__ int fitness[]; 
    /* ... etc ... */ 
} 

nVIDIA的Parallel-for-all博客有a nice post有更深入的介紹瞭如何使用共享內存，具體包括靜態和動態共享內存分配。