2

デバイスのグローバルメモリに大きな文字配列があり、スレッドによって合体的にアクセスされます。スレッドごとに1つのメモリトランザクションで4文字または16文字を読み取ることで、メモリアクセスを高速化できることをどこかで読んだことがあります。テクスチャとchar4またはint4構造体を使用する必要があると思います。ただし、これに関するドキュメントや例は見つかりません。ここにいる誰かが、私がこれについてもっと学ぶことができる場所への簡単な例またはポインタを提供できますか?

私のコードでは、char配列を次のように定義しています。

char *database = NULL;
cudaMalloc( (void**) &database, SIZE * sizeof(char) );

テクスチャとchar4(またはint4)を使用したい場合、定義はどうなりますか?

どうもありがとう。

4

1 に答える 1

1

私はついに自分の質問に対する答えを見つけました。char4での定義は次のようになります

char4 *database = NULL;
cudaMalloc( (void**) &database, SIZE * sizeof(char4)/4 );

このためにテクスチャは必要ありません。カーネルはchar4で3倍高速化しますが、ループ展開を行うと2倍に減少します。完全を期すために、私のカーネルは

__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;
  }
}

そしてchar4でそれは

__global__ void kernel4(unsigned int jobs_todo, char4* database, float* results ) {

  unsigned int id = threadIdx.x + blockIdx.x * blockDim.x;
  float A = 0; int i; char4 ch4;
  if(id < jobs_todo) {
    for(i = 0; i < 1000/4; i += 1){
     ch4 = database[jobs_todo*i + id];
     if(ch4.x == 'A') A++;
     if(ch4.y == 'A') A++;
     if(ch4.z == 'A') A++;
     if(ch4.w == 'A') A++;
    }
    results[id] = A;
  }
}

int4も試しましたが、char4の時間よりもわずか.0002秒速くなっています。

于 2013-01-01T21:11:15.163 に答える