3

最近CUDAについて勉強しています。CUDAのメモリアクセス時間について知りたいです。

では、CUDA プログラミング ガイドに書かれているメモリ アクセス時間:

  • グローバル メモリ アクセス時間は 400 ~ 600 サイクル
  • 共有メモリ ( L1 キャッシュ ) アクセス時間は 20 ~ 40 サイクル

サイクルは時計と同じだと思います。これは正しいです ？ということで、メモリアクセス時間を調べてみました。ホストは固定されていますが、カーネル コードには 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

Dev_In[2]、Dev_In[3]がキャッシュされているので、1回目のアクセスはキャッシュではないので800サイクルかかりましたが、2回目のアクセス3回目のアクセスは312サイクルかかったと思います。

カーネル バージョン 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] がキャッシュされていないと思います

だから… #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 クロックですか?

また、result1、result2、result3 では、result3 のメモリ アクセス時間が非常に速いのはなぜですか?

4

1 に答える 1

1

@phoad が述べた理由により、あなたの評価は無効です。メモリ アクセスの後、クロック ストップの前に、メモリの読み取り値を再利用して、未処理の負荷に対する命令の依存関係を作成する必要があります。それ以外の場合、GPU は独立した命令を次々に発行し、クロック開始とロードの直後にクロック終了が実行されます。Henry Wong が用意したマイクロベンチマーク スーツを試すことをお勧めします。このスーツを使用すると、メモリ アクセス レイテンシを含むさまざまなマイクロアーキテクチャの詳細を取得できます。メモリ レイテンシだけが必要な場合は、Sylvain Collange によって開発されたCUDA レイテンシを試す方が簡単です。

于 2012-09-19T10:57:45.537 に答える