CUDA 內核函數參數通過恒定存儲器傳遞給設備，并且被限制為 4096 字節。 CUDA 12.1 將此參數限制從 4096 字節增加到 32764 字節，在所有設備架構上都有效，包括 NVIDIA Volta 及以上。
以前，傳遞超過 4096 字節的內核參數需要通過將多余的參數復制到常量內存中來繞過內核參數限制cudaMemcpyToSymbol或cudaMemcpyToSymbolAsync，如下面的片段所示。
#define TOTAL_PARAMS (8000) // ints
#define KERNEL_PARAM_LIMIT (1024) // ints
#define CONST_COPIED_PARAMS (TOTAL_PARAMS - KERNEL_PARAM_LIMIT)
__constant__ int excess_params[CONST_COPIED_PARAMS];
typedef struct {
int param[KERNEL_PARAM_LIMIT];
} param_t;
__global__ void kernelDefault(__grid_constant__ const param_t p,...) {
// access <= 4,096 parameters from p
// access excess parameters from __constant__ memory
}
int main() {
param_t p;
int *copied_params = (int*)malloc(CONST_COPIED_PARAMS * sizeof(int));
cudaMemcpyToSymbol(excess_params,
copied_params,
CONST_COPIED_PARAMS * sizeof(int),
0,
cudaMemcpyHostToDevice);
kernelDefault<<<GRIDDIM,BLOCKDIM>>>(p,...);
cudaDeviceSynchronize();
}
這種方法限制了可用性，因為您必須顯式管理常量內存分配和副本。復制操作還增加了顯著的延遲，降低了接受大于 4096 字節參數的延遲綁定內核的性能。
從 CUDA 12 . 1 開始，您現在可以在 NVIDIA Volta 及更高版本上傳遞多達 32764 個字節作為內核參數，從而得到下面第二個片段中所示的簡化實現。
#define TOTAL_PARAMS (8000) // ints
typedef struct {
int param[TOTAL_PARAMS];
} param_large_t;
__global__ void kernelLargeParam(__grid_constant__ const param_large_t p,...) {
// access all parameters from p
}
int main() {
param_large_t p_large;
kernelLargeParam<<<GRIDDIM,BLOCKDIM>>>(p_large,...);
cudaDeviceSynchronize();
}
請注意，在前面的兩個示例中，內核參數都用__grid_constant__限定符，以指示它們是只讀的。
工具包和驅動程序兼容性
注意，使用 CUDA Toolkit 12.1 和 R530 驅動程序或更高版本編譯、啟動和調試具有大內核參數的內核是必需的。如果在較舊的驅動程序上嘗試啟動，CUDA 將發布 CUDA_ERROR_NOT_SUPPORTED 錯誤。
支持的體系結構
更高的參數限制適用于所有架構，包括 NVIDIA Volta 及以上版本。在低于 NVIDIA Volta 的體系結構上，參數限制保持在 4096 字節
CUDA 工具包修訂版之間的鏈接兼容性
當鏈接設備對象時，如果至少有一個設備對象包含具有更高參數限制的內核，則必須重新編譯設備源中的所有對象，并使用 CUDA Toolkit 12.1 將它們鏈接在一起。否則將導致鏈接器錯誤。
例如，考慮兩個設備對象 a.o 和 b.o 鏈接在一起的場景。如果 a.o 或 b.o 至少包含一個具有較高參數限制的內核，則必須重新編譯相應的源并將生成的對象鏈接在一起。
使用大內核參數可節省性能
圖 1 比較了上面提供的兩個代碼片段在單個代碼上的性能，在 NVIDIA H100 系統上測量了超過 1000 次迭代。在本例中，通過避免常量拷貝，使應用程序運行時總體節省了 28% 。圖 2 顯示，對于相同的代碼段，使用 NVIDIA Nsight Systems 后，內核執行時間提高了 9% 。
對于這兩個圖像，灰色條顯示了內核的執行時間，其中 1024 個整數作為內核參數傳遞，其余整數使用恒定內存傳遞（代碼片段 1 ）。綠條顯示了內核的執行時間，其中 8000 個整數作為內核參數傳遞（代碼片段 2 ）。兩個內核都累積了 8000 個整數。
請注意，如果省略 __grid_constant__限定符，然后從內核對其執行后續寫入操作，自動復制到thread-local-memory被觸發。這可能會抵消任何性能提升。
圖 3 顯示了使用 QUDA（一個用于在晶格量子色動力學中進行計算的 HPC 庫）的結果。
本例中的參考內核執行批量矩陣乘法 X * a + Y ，其中 a 、 X 和 Y 是矩陣。內核參數存儲 A 的系數。在 CUDA 12.1 之前，當系數超過 4096 字節的參數限制時，它們被顯式復制到恒定內存中，大大增加了內核延遲。刪除該副本后，可以觀察到顯著的性能改進（圖 3 ）。
總結
CUDA 12.1 為您提供了使用內核參數傳遞多達 32764 個字節的選項，可以利用這些參數簡化應用程序并提升性能。要查看本文中引用的完整代碼示例，請訪問 NVIDIA/cuda-samples GitHub 。
?
