2013-04-20 68 views
9

在CUDA中使用常量的最佳方式是哪種?在CUDA中使用常量

一種方式是在不斷的內存來定義常量,如:

// CUDA global constants 
__constant__ int M; 

int main(void) 
{ 
    ... 
    cudaMemcpyToSymbol("M", &M, sizeof(M)); 
    ... 
} 

的備選方法是使用C預處理:

#define M = ... 

我認爲定義與C預處理器常數速度要快得多。那麼在CUDA設備上使用恆定內存的好處是什麼?被在編譯時已知

+1

編譯時已知的常量應該使用預處理器宏(即'#define')來定義。在其他情況下,'__constant__' [變量](http://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#constant)可能是CUDA程序員用來優化訪問代碼的一個選項不改變的計算變量。請注意,使用''M''來引用一個符號在cuda 5中不再有效。 – 2013-04-20 12:43:28

+0

知道這兩種可能性之間的運行時間差異是多麼有趣。我是currenlty在一些cfd代碼上工作,我想將參數作爲選項傳遞給程序,因此有必要使用第一種方法。另一方面,如果我使用預處理器宏,這是不可能的。 – jrsm 2013-04-20 14:13:50

+0

既然你的第二個例子不生成任何類型的機器代碼,這不是一個明智的問題。您需要提出實際的運行時使用情況,以便對此問題有任何意義。對於將單個標量立即值初始加載到變量或寄存器中,第二種方法將始終更快。 – 2013-04-20 14:14:27

回答

9
  1. 常數應該使用 預處理宏(例如#define)或經由全球/文件範圍內的C/C++ const變量來定義。
  2. __constant__ memory的使用對於那些使用某些在內核持續時間內不發生變化並且存在某些訪問模式(例如,所有線程同時訪問相同值)的程序的應用程序可能是有益的。這並不比滿足上述第1項要求的常數更好或更快。
  3. 如果通過程序進行選擇的數量在數量上相對較少,而這些選擇會影響內核執行額外的編譯時優化一個可能的方法是使用templated code/kernels
+2

由於多種原因,使用C/C++樣式常量,預處理器宏或C++模板可能比使用__constant__內存更快:1.編譯器可以應用其他優化,以及2.常量可以作爲直接內嵌在指令中。常量高速緩存訪​​問可能會錯過常量高速緩存,增加額外的延遲 – 2014-06-10 06:07:17

+0

由於每次調用每個線程中的構造函數,不會定義非重要類型的'#definitiond常量? – 2016-09-12 08:15:37

4

常規的C/C++風格常量:在CUDA C(本身是C99的修改)中,常量是絕對編譯時間實體。考慮到GPU處理的性質,考慮到在NVCC中發生的優化量非常大,這並不奇怪。

#define:宏一如既往非常不雅但在捏中有用。

__constant__變量說明符是,但是一個全新的動物,在我看來是一個用詞不當的東西。我會放下什麼Nvidia已經在下面的空白處here

__constant__限定符可選擇與 __device__一起使用,聲明的變量:

  • 寓於常量內存空間,
  • 具有應用程序的生命週期,
  • 可以從網格內的所有線程和主機通過運行時庫訪問(cudaGetSymbolAddress()/ cudaGetSymbolSize()/ cudaMemcpyToSymbol()/ cudaMemcpyFromSymbol())。

Nvidia的文檔指定__constant__可在寄存器水平速度(接近零延遲)只要是由一個經紗的所有線程被訪問的相同的恆定。

它們在CUDA代碼的全局範圍內聲明。但是,基於個人(當前正在進行的)經驗,在分離編譯時必須小心使用該說明符,例如分離CUDA代碼(.cu和.c)。cuh文件)從你的C/C++代碼中加入封裝函數到C風格的頭文件中。

與傳統的「常量」指定變量不同,它們是在運行時從分配設備內存並最終啓動內核的主機代碼初始化的。我再說一遍,我目前正在使用的代碼演示這些可以在內核執行前使用cudaMemcpyToSymbol()在運行時上設置。

鑑於保證訪問的L1緩存級別的速度,至少可以說很方便。