如何使用常量內存將結構數組傳遞給CUDA內核

Question

我的內核需要配置的列表/數組作爲輸入參數。我甚至有這樣的列表/數組的列表/數組，其中之一是傳遞給內核。這些配置在主機上準備好，不會更改。所以這對於不斷的記憶是一個完美的使用。但說實話，我真的不知道該怎麼做。

我嘗試在下面的代碼草案中給出我的想法。基本上，我看到兩種方法如何定義/通過列表：

• 將它們定義爲具有固定lenghts陣列和按值傳遞它們的內核
• 將它們定義爲指針，只是傳遞一個指針內核（當然必須先複製到設備上）

我應該採取哪種方法，以及如何修改下面的代碼以確保使用常量內存？

我希望每個列表的大小通常小於200-300字節。如果我要製作所有相同大小的列表，我可能會選擇512 Bytes或1 kB的大小。

class Configuration{ 
    // some constants 
} 

// We need a list of lists Configurations, these could be implemented either as... 
Configuration a[10][100]; // fixed-length array or... 
Configuration ** b; // as a dynamic array to pointers of arrays 

// Parameter will take an array of Configuration, either as a pointer or directly as an array 
__global__ kernel(Configuration * config){ 

} 

// According to the above example, we use the pointer-version. Could also be a call directly using a[i] 
kernel<<...>>(b[i], lengthOfB[i]); 

Answer 1

如果希望數據是在__constant__存儲器（其可以不是明智之舉，這取決於你如何訪問內核中的數據），那麼第一種方法（固定長度陣列）是唯一明智的。同樣爲了簡單起見，我將平面二維數組轉換爲一維數組，以便於使用/複製。

除了爲只讀，__constant__內存旨在被訪問的效率，使每個線程在一個warp請求相同的值。你的問題沒有提到這一點，所以你可能想參考this question/answer的解釋/例子。

如果你使用指針方法，只有指針將在常量內存中（大概），所以這不是你想要的（推測）。

如果使用__constant__內存，則不需要將該指針作爲內核參數傳遞。數據聲明具有全局範圍。

像這樣的東西可能會奏效：

class Configuration{ 
    // some constants 
    int cdata; 
} 

__constant__ Configuration const_data[10*100]; 

// ***setup in host code 
Configuration h_data[10*100]; 
// fill in h_data ... 
// then copy to device 
cudaMemcpyToSymbol(const_data, h_data, sizeof(h_data)); 
// *** 

//use in kernel code 
__global__ void mykernel(){ 

    int my_data = const_data[5].cdata; 

} 

注意，在總，__constant__內存被限制爲64K字節。