cuda全局和共享內存訪問時間

最近我在研究CUDA。我想知道CUDA內存訪問時間。cuda全局和共享內存訪問時間

在，CUDA編程指南寫入存儲器存取時間：

全局存儲器存取時間爲400〜600循環
共用存儲器（L1高速緩存）訪問時間爲20〜40週期

我認爲Cycle和Clock是一樣的。它是否正確？如果這是正確的，所以我檢查了內存訪問時間。主機是固定的，但內核代碼有3個版本。這是我的代碼：

主機代碼

float* H1 = (float*)malloc(sizeof(float)*100000); 
float* D1; 

for(int i = 0 ; i < 100000 ; i++){ 
    H1[i] = i; 
} 

cudaMalloc((void**)&D1, sizeof(float)*100000); 
cudaMemcpy(D1, H1, sizeof(float)*100000, cudaMemcpyHostToDevice); 


cudaPrintfInit(); 
test<<<1,1>>>(D1); 
cudaPrintfDisplay(stdout, true); 

cudaPrintfEnd();

內核版本1

float Global1; 
float Global2; 
float Global3; 

int Clock; 

Clock = clock(); 
Global1 = Dev_In1[1]; 
Clock = clock() - Clock; 
cuPrintf("Global Memory Access #1 : %d\n", Clock); 

Clock = clock(); 
Global2 = Dev_In1[2]; 
Clock = clock() - Clock; 
cuPrintf("Global Memory Access #2 : %d\n", Clock); 

Clock = clock(); 
Global3 = Dev_In1[3]; 
Clock = clock() - Clock; 
cuPrintf("Global Memory Access #3 : %d\n", Clock);

把它`導致

全球內存訪問＃1：882次
全球內存訪問＃2：312次
全球內存訪問＃3：312

我覺得首先訪問不緩存，以便把800週期但第二訪問第三訪問了312週期，因爲，Dev_In [2]，Dev_In [3]被緩存..

內核版本2

int Global1, Global2, Global3;    
int Clock;    

Clock = clock();     
Global1 = Dev_In1[1];    
Clock = clock() - Clock;     
cuPrintf("Global Memory Access #1 : %d\n", Clock);    

Clock = clock();     
Global2 = Dev_In1[50000];    
Clock = clock() - Clock;     
cuPrintf("Global Memory Access #2 : %d\n", Clock);    

Clock = clock();     
Global3 = Dev_In1[99999];    
Clock = clock() - Clock;     
cuPrintf("Global Memory Access #3 : %d\n", Clock);

它單曲導致

全球內存訪問＃1：872次
全球內存訪問＃2：776次
全球內存訪問＃3：782

我認爲不緩存在第一次訪問時Dev_In1 [50000]和Dev_In2 [99999]

so ...＃1，＃2，＃3遲到...

內核版本3

int Global1, Global2, Global3;     
int Clock;     

Clock = clock();      
Global1 = Dev_In1[1];     
Clock = clock() - Clock;      
cuPrintf("Global Memory Access #1 : %d\n", Clock);     

Clock = clock();      
Global1 = Dev_In1[50000];     
Clock = clock() - Clock;      
cuPrintf("Global Memory Access #2 : %d\n", Clock);     

Clock = clock();      
Global1 = Dev_In1[99999];     
Clock = clock() - Clock;      
cuPrintf("Global Memory Access #3 : %d\n", Clock);

結果

全球內存訪問＃1：168
全球內存訪問＃2：168次
全球內存訪問＃3： 168

我不理解這個結果

Dev_In [50000]，Dev_In [99999]沒有被緩存，但訪問時間非常快！只是，我用1變量....

SO ..我的問題是，GPU週期== GPU時鐘？

and result1，result2，result3爲什麼內存訪問時間在result3中非常快？

來源

2012-09-18 오승택

您的編譯器可能會優化代碼。也就是說，因爲你沒有在任何地方使用Global1，所以你的編譯器可能不會閱讀全部內容。嘗試在讀取內存中的下一個值之前使用Global1。 – phoad

clock（）返回的值應該是由設備屬性CU_DEVICE_ATTRIBUTE_CLOCK_RATE指定的值。時鐘（）的值是循環的，它不是固定的頻率。在執行內核期間，頻率固定在2.x和更舊的設備上。在3.x設備上，時鐘頻率將與電源管理和散熱管理密切相關。 –

您是否在同一個進程的同一執行過程中在同一個CUDA上下文中運行這三個內核，或者您是否啓動該進程3次？如果前者運行kernel2，然後kernel3使用相同的設備內存，則可能導致該值在L2中被緩存，這將佔用顯着較低的流逝時間。使用clock（）時，您應該驗證SASS（彙編代碼，不是PTX）是否具有正確的指令序列。編譯器移動或消除對clock（）的調用是很常見的。我建議在cuPrintf之後插入一個threadfence（）來刷新LSU單元。 –

由於@phoad聲明的原因，您的評估無效。在內存訪問之後和時鐘停止之前，應該重新使用內存讀取值來對未完成的加載進行指令依賴。否則，GPU依次發出獨立指令，並在時鐘啓動和加載後立即執行時鐘結束。我建議你嘗試Henry Wong在here準備的microbenchmarking西裝。使用這個套裝，您可以檢索各種微體系結構詳細信息，包括內存訪問延遲。如果您只需要內存延遲，那麼嘗試使用由Sylvain Collange開發的CUDA latency更容易。

來源

2012-09-19 10:57:45 ahmad

cuda全局和共享內存訪問時間

回答

相關問題