當按參數值傳遞給內核函數時，參數在哪裏複製？

Question

我是CUDA編程的初學者，有一個問題。

當我通過值傳遞參數，如下所示：

__global__ void add(int a, int b, int *c) { 
    // some operations 
} 

由於可變一個和b被傳遞到核函數添加如在函數調用棧複製的值，我猜到一些存儲器空間將需要複製。

如果我是正確的，那些參數在GPU或Host的主內存中複製 的額外內存空間？

我想知道這個問題的原因是我應該將一個大結構傳遞給內核函數。

我也想過傳遞結構的指針，但這些方式似乎需要調用cudamalloc爲結構和每個成員變量。

Answer 1

很簡短的答案是所有的CUDA內核參數都通過值傳遞，這些參數被主機通過API複製到GPU上的專用內存參數緩衝區中。目前，此緩衝區存儲在常量內存中，每次啓動內核時參數限制爲4kb - 請參閱here。

---

在更多的細節，所述PTX標準（技術上因爲計算能力2.0硬件和CUDA ABI出現）定義的專用邏輯狀態空間呼叫.param持有內核和設備參數的參數。請參閱here。從該文檔引用：

每個內核函數定義包括一個可選列表 參數。這些參數是可尋址的，在.param狀態空間中聲明的只讀變量 。從主機傳遞到 內核的值可通過使用ld.param 指令的這些參數變量進行訪問。內核參數變量在網格內的所有 CTA中共享。

委員會還指出，：

注：的參數空間中的位置是特定的實現。例如，在一些實現中，內核參數駐留在全局內存中。在這種情況下，參數和全局空間之間不提供訪問保護。類似地，基於應用程序二進制接口 （ABI）的函數調用約定，函數參數被映射到參數傳遞寄存器和/或堆棧位置 。

所以參數狀態空間的精確位置是實現特定的。在CUDA硬件的第一次迭代中，它實際上映射到用於內核參數的共享內存和用於設備函數參數的寄存器。但是，自從計算2.0硬件和PTX 2.2標準以來，它在大多數情況下映射到內核的常量內存。該文件說，following on the matter：

常數（.const）狀態空間是一個只讀存儲器由主機進行初始化 。使用ld.const 指令訪問常量內存。恆定內存的大小受到限制，目前限制爲 至64 KB，可用於保存靜態大小的常量 變量。 還有另外640 KB的常量內存， 組織爲10個獨立的64 KB區域。驅動程序可以分配 並初始化這些區域中的常量緩衝區，並將緩衝區的指針作爲內核函數參數傳遞給 。由於十個區域是 不連續，則驅動程序必須確保恆定的緩衝劑是 分配，使得每個緩衝器大小爲64KB的區域內完全配合和 不跨越區域邊界。

靜態大小的常量變量有一個可選的變量初始值設定項;沒有顯式初始化器的常量變量默認爲 ，初始化爲零。 由主機初始化 驅動程序分配的常量緩衝區，指向此類緩衝區的指針爲 作爲參數傳遞給內核。

[着重礦]

因此，儘管內核參數被存儲在常數存儲器中，這是不一樣的恆定存儲器，其通過在CUDA C或定義一個變量作爲__constant__映射到訪問的.const狀態空間在Fortran或Python中是等效的。相反，它是由驅動程序管理的內部設備內存池，不能被程序員直接訪問。