2

構造体の配列を使用していて、各ブロックが共有メモリに配列の1つのセルをロードするようにします。例：ブロック0は共有メモリにarray [0]をロードし、ブロック1はarray[1]をロードします。

そのために、構造体の配列をfloat *にキャストして、メモリアクセスを統合しようとしました。

私は2つのバージョンのコードを持っています

バージョン1

__global__ 
void load_structure(float * label){

  __shared__ float shared_label[48*16];
  __shared__ struct LABEL_2D* self_label;


  shared_label[threadIdx.x*16+threadIdx.y] = 
          label[blockIdx.x*sizeof(struct LABEL_2D)/sizeof(float) +threadIdx.x*16+threadIdx.y];
  shared_label[(threadIdx.x+16)*16+threadIdx.y] = 
          label[blockIdx.x*sizeof(struct LABEL_2D)/sizeof(float) + (threadIdx.x+16)*16+threadIdx.y];
  if((threadIdx.x+32)*16+threadIdx.y < sizeof(struct LABEL_2D)/sizeof(float))  {
    shared_label[(threadIdx.x+32)*16+threadIdx.y] = 
          label[blockIdx.x*sizeof(struct LABEL_2D)/sizeof(float) +(threadIdx.x+32)*16+threadIdx.y];
   }

  if(threadIdx.x == 0){
    self_label = (struct LABEL_2D *) shared_label;
  }
  __syncthreads();
  return;
}

...

dim3 dimBlock(16,16);
load_structure<<<2000,dimBlock>>>((float*)d_Label;

計算時間：0.740032ミリ秒

バージョン2

__global__ 
void load_structure(float * label){

  __shared__ float shared_label[32*32];
  __shared__ struct LABEL_2D* self_label;

  if(threadIdx.x*32+threadIdx.y < *sizeof(struct LABEL_2D)/sizeof(float))
    shared_label[threadIdx.x*32+threadIdx.y] = 
              label[blockIdx.x*sizeof(struct LABEL_2D)/sizeof(float)+threadIdx.x*32+threadIdx.y+];


  if(threadIdx.x == 0){
      self_label = (struct LABEL_2D *) shared_label;
    }
  __syncthreads();
  return;
}

dim3 dimBlock(32,32);
load_structure<<<2000,dimBlock>>>((float*)d_Label);

計算時間：2.559264ミリ秒

どちらのバージョンでも、nvidiaプロファイラーを使用しました。グローバルな負荷効率は8％です。

私には2つの問題があります：1-タイミングの違いがある理由がわかりません。2-私の通話は合体していますか？

2.1の計算機能（32スレッド/ラップ）を備えたビデオカードを使用しています

4

2 に答える 2

2

あなたのグローバルな負荷は合体していません。8％はかなり低く、あなたができる最悪の事態は3％です。

これの主な理由は、threadIdx.xとthreadIdx.yに基づいてインデックスを作成する方法にあると思います。2番目のカーネルからのこのコード行を考えてみましょう（最初のカーネルにも同様の問題があります）。

shared_label[threadIdx.x*32+threadIdx.y] =  label[blockIdx.x*sizeof(struct LABEL_2D)/sizeof(float)+threadIdx.x*32+threadIdx.y];

特に、次のインデックス付けを検討してください。

threadIdx.x*32+threadIdx.y

CUDAワープは、X、Y、Zの順序でグループ化されます。つまり、ワープ内の急速に変化するインデックスは、最初にXインデックス、次にY、次にZになる傾向があります。たとえば、16x16スレッドブロックがある場合、最初のワープにはthreadIdx.xスパンを持つスレッドが含まれます。 0から15までで、threadIdx.yは0から1のみにまたがっています。この場合、隣接するスレッドには、ほとんどの場合、隣接するthreadIdx.xインデックスがあります。

コードの結果は、インデックス作成のために合体が壊れたことです。このタイプのインデックスを使用するようにロードとストレージを再構築できる場合：

threadIdx.y*32+threadIdx.x

グローバルな負荷効率が突然大幅に向上します。（共有メモリの使用量も改善される可能性があります。）

質問が2つあると思いますが、最初の質問について考えると戸惑います。「計算時間」は約1時間です。2番目の実装では4倍長くなりますが、おそらくcompute_interpolationカーネルを参照しているため、詳細はまったく示されていません。ただし、2番目の場合は4倍のスレッドを起動しています。おそらくここには謎はありません。コードは表示されていません。また、カーネルを使用して共有メモリに大量のデータをロードしてから終了することも意味がありません。共有メモリの内容は、あるカーネル呼び出しから次の呼び出しまで持続しません。

于 2013-03-25T22:14:12.317 に答える
0

以前のバージョンではアクセスメモリパターンが正しくなかったため、問題を解決しました。cudaベストプラクティスガイドのパラグラフ6.2.1を読んだ後、それらが整列されているとアクセスが速くなることがわかりました。

アクセスパターンを調整するために、構造体に「偽の」変数を追加して、構造体のサイズを128（現金サイズの線）で割ることができるようにしました。

この戦略により、優れたパフォーマンスが得られます。2000構造を2000ブロックにロードするのに、わずか0.16msしかかかりませんでした。

コードのバージョンは次のとおりです。

struct TEST_ALIGNED{
  float data[745];
  float aligned[23];
}; 


__global__
void load_structure_v4(float * structure){

  // Shared structure within a block
  __shared__ float s_structure[768];
  __shared__ struct TEST_ALIGNED * shared_structure;

  s_structure[threadIdx.x] = 
    structure[blockIdx.x*sizeof(struct TEST_ALIGNED)/sizeof(float) + threadIdx.x];
  s_structure[threadIdx.x + 256] = 
    structure[blockIdx.x*sizeof(struct TEST_ALIGNED)/sizeof(float) + threadIdx.x + 256];
  if(threadIdx.x < 745)
        s_structure[threadIdx.x + 512] = 
            structure[blockIdx.x*sizeof(struct TEST_ALIGNED)/sizeof(float) +    threadIdx.x + 512];
  if(threadIdx.x == 0)
       shared_structure = (struct TEST_ALIGNED*) s_structure;

  __syncthreads();

    return;
}

dim3 dimBlock(256);
load_structure_v4<<<2000,dimBlock>>>((float*)d_test_aligned);

私はまだ最適化を探しています、そして私がそれを見つけたらここに投稿します。

于 2013-03-26T23:12:39.143 に答える