1

これはパフォーマンス関連の質問です。「CUDA By Example」サンプルコードに基づいて、次の単純な CUDA カーネルを作成しました。

#define N 37426 /* the (arbitrary) number of hashes we want to calculate */
#define THREAD_COUNT 128

__device__ const unsigned char *m = "Goodbye, cruel world!";

__global__ void kernel_sha1(unsigned char *hval) {
  sha1_ctx ctx[1];
  unsigned int tid = threadIdx.x + blockIdx.x * blockDim.x;

  while(tid < N) {
    sha1_begin(ctx);
    sha1_hash(m, 21UL, ctx);
    sha1_end(hval+tid*SHA1_DIGEST_SIZE, ctx);
    tid += blockDim.x * gridDim.x;
  }
}

コードは私には正しいように見え、実際に同じハッシュの 37,426 個のコピーを吐き出します (予想どおりです。第 5 章のセクション 5.3 を読んだことに基づいて、「hval」として渡されたグローバル メモリに書き込む各スレッドは、非常に非効率になります。

次に、共有メモリを使用してパフォーマンスを向上させるキャッシュと思われるものを実装しました。コードは次のように変更されました。

#define N 37426 /* the (arbitrary) number of hashes we want to calculate */
#define THREAD_COUNT 128

__device__ const unsigned char *m = "Goodbye, cruel world!";

__global__ void kernel_sha1(unsigned char *hval) {
  sha1_ctx ctx[1];
  unsigned int tid = threadIdx.x + blockIdx.x * blockDim.x;
  __shared__ unsigned char cache[THREAD_COUNT*SHA1_DIGEST_SIZE];

  while(tid < N) {
    sha1_begin(ctx);
    sha1_hash(m, 21UL, ctx);
    sha1_end(cache+threadIdx.x*SHA1_DIGEST_SIZE, ctx);

    __syncthreads();
    if( threadIdx.x == 0) {
      memcpy(hval+tid*SHA1_DIGEST_SIZE, cache, sizeof(cache));
    }
    __syncthreads();
    tid += blockDim.x * gridDim.x;
  }
}

2 番目のバージョンも正しく動作しているように見えますが、最初のバージョンより数倍遅くなります。後者のコードは約 8.95 ミリ秒で完了し、前者は約 1.64 ミリ秒で実行されます。スタック オーバーフロー コミュニティへの私の質問は単純です。なぜですか?

4

1 に答える 1

2

CUDA by Example を調べたところ、これに似たものは見つかりませんでした。はい、付録に GPU ハッシュ テーブルの説明がありますが、このようには見えません。したがって、あなたの関数、特にsha1_endが何をしているのか本当にわかりません。このコードがその本の何かに似ている場合は、それを指摘してください。

ただし、sha1_end が (スレッドごとに) 1 回グローバル メモリに書き込み、結合された方法で行う場合、非常に効率的でない理由はありません。おそらく、各スレッドは異なる場所に書き込みを行っているため、それらが多かれ少なかれ隣接している場合、確実に合体する可能性があります。合体の詳細には触れずに、複数のスレッドが単一のトランザクションでデータをメモリに書き込むことができると言えば十分です。また、データをグローバル メモリに書き込む場合は、このペナルティを少なくとも 1 回、どこかで支払う必要があります。

あなたの修正のために、あなたはこの概念を完全に殺しました。これで、単一のスレッドからすべてのデータのコピーを実行しました。memcpy は、後続のデータ書き込み (int または char など) が個別のトランザクションで発生していることを意味します。はい、これに役立つキャッシュがありますが、GPU でそれを行う方法は完全に間違っています。各スレッドがグローバル メモリを更新できるようにし、それを並行して行う機会を利用します。ただし、単一のスレッドですべての更新を強制すると、そのスレッドはデータを順番にコピーする必要があります。これはおそらく、タイミングの違いにおける最大のコスト要因です。

__syncthreads() を使用すると、追加のコストもかかります。

CUDA by Examples のセクション 12.2.7 では、ビジュアル プロファイラーについて言及されています (また、結合されたアクセスに関する情報を収集できることも言及されています)。ビジュアル プロファイラーは、このような質問に答えるのに役立つ優れたツールです。

効率的なメモリ技術と合体について詳しく知りたい場合は、「CUDA C を使用した GPU コンピューティング – Advanced 1 (2010)」というタイトルの NVIDIA GPU コンピューティングウェビナーをお勧めします。それへの直接リンクはスライド付きのhereです。

于 2012-11-18T22:20:43.753 に答える