0

非常に大きな文字列の1000バイトフラグメント内のAの数をカウントする単純なCUDAカーネルがあります。データベースは、メモリアクセスが合体するように配置されています。カーネルから戻った後、私のメイン関数はresults、さらに分析するためにデバイス配列をホスト上の配列にコピーします。

__global__ void kernel(unsigned int jobs_todo, char* database, float* results ) {

  unsigned int id = threadIdx.x + blockIdx.x * blockDim.x;
  float A = 0; int i; char ch;
  if(id < jobs_todo) {
    for(i = 0; i < 1000; i += 1){
     ch = database[jobs_todo*i + id];
     if(ch == 'A') A++;
   }
  results[id] = A;
}

カーネルは正常に動作します。results[id]=Aただし、そのような些細なもの に置き換えresults[id]=10たり、その行をコメントアウトしたりすると、実行速度がはるかに速く（10倍）、使用するレジスタがはるかに少なくなります--ptxas-options=-v。その行をコメントアウトすると、カーネルは役に立ちません。CUDAコンパイラは渡されたパラメータを見てこれを知っていますか？そして、それは何もしないことを選択しますか？

4

1 に答える 1

3

あなたが見ているのは、コンパイラの最適化の結果です。コンパイルにより、「デッド」コード、つまりメモリ書き込みが直接発生しないコードが削除されます。だからあなたのカーネル

__global__ void kernel(unsigned int jobs_todo, char* database, float* results ) {

  unsigned int id = threadIdx.x + blockIdx.x * blockDim.x;
  float A = 0; int i; char ch;
  if(id < jobs_todo) {
    for(i = 0; i < 1000; i += 1){
     ch = database[jobs_todo*i + id];
     if(ch == 'A') A++;
   }
   results[id]=10;
}

に効果的に最適化されています

__global__ void kernel(unsigned int jobs_todo, char* database, float* results ) {

  unsigned int id = threadIdx.x + blockIdx.x * blockDim.x;
  results[id]=10;
}

明らかに、削減されたコードのレジスタフットプリントと実行時間は、完全なコードよりもはるかに短くなります。これは、コードをPTXにコンパイルし、発行されたコードを調べることで確認できます。

于 2013-01-01T20:32:14.693 に答える