CUDA代碼中的常量內存使用情況

我自己弄不明白，確保內核中使用的內存不變的最佳方法是什麼？有一個類似的問題在http://stackoverflow...r-pleasant-way。我正在使用GTX580並僅編譯2.0功能。我的內核看起來像CUDA代碼中的常量內存使用情況

__global__ Foo(const int *src, float *result) {...}

我執行在主機下面的代碼：

cudaMalloc(src, size); 
cudaMemcpy(src, hostSrc, size, cudaMemcpyHostToDevice); 
Foo<<<...>>>(src, result);

的另一種方法是添加

__constant__ src[size];

到.CU文件，從刪除SRC指針內核並執行

cudaMemcpyToSymbol("src", hostSrc, size, 0, cudaMemcpyHostToDevice); 
Foo<<<...>>>(result);

這兩種方式是等價的還是第一種不能保證使用常量內存而不是全局內存？大小動態變化，所以第二種方法在我的情況下不方便。

來源

2012-01-30 markhor

第二種方法是確保將數組編譯爲CUDA常量內存並通過常量內存緩存正確訪問的唯一方法。但是你應該問自己如何在一個線程塊內訪問該數組的內容。如果每個線程均勻地訪問陣列，那麼在使用常量內存時會有性能優勢，因爲常量內存緩存具有廣播機制（它還可以節省全局內存帶寬，因爲常量內存存儲在片外DRAM和緩存中減少DRAM交易次數）。但是如果訪問是隨機的，那麼可以訪問本地存儲器的序列化，這會對性能產生負面影響。

可能適合__constant__內存的典型東西是模型係數，權重和其他需要在運行時設置的常量值。例如，在Fermi GPU上，內核參數列表存儲在常量內存中。但是，如果內容訪問非一致，並且成員的類型或大小在調用時不是恆定的，那麼正常的全局內存是可取的。

另外請記住，每個GPU上下文有64kb的常量內存限制，因此在常量內存中存儲大量數據是不現實的。如果你需要大量帶緩存的只讀存儲器，可能值得嘗試將數據綁定到紋理，並查看性能如何。在費米卡之前，它通常會帶來一個方便的性能增益，費米的結果與全局存儲器相比，可預測性較差，因爲該架構中緩存佈局的改進。

來源

2012-01-30 06:27:15 talonmies

第一種方法將保證函數Foo內的內存不變。這兩個不相等，第二個保證它在初始化之後在那裏存在。如果你需要動態的比你想使用類似於第一種方式的東西。

來源

2012-01-30 04:41:05 rudolph9

CUDA代碼中的常量內存使用情況

回答

相關問題