3

共有メモリを使用してOpenACCでキャッシュしようとしています。

基本的に私が取り組んでいるのは行列の乗算であり、私が持っているのはこれです：

typedef float ff; 

// Multiplies two square row-major matrices a and b, puts the result in c. 
void mmul(const restrict ff* a, 
          const restrict ff* b, 
          restrict ff* c, 
          const int n) { 
#pragma acc data copyin(a[0:n*n], b[0:n*n]) copy(c[0:n*n]) 
{ 

#pragma acc region 
{ 

#pragma acc loop independent vector(16) 
  for (int i = 0; i < n; ++i) { 
#pragma acc loop independent vector(16) 
    for (int j = 0; j < n; ++j) { 
      ff sum = 0; 
      for (int k = 0; k < n; ++k) { 
        sum += a[i + n * k] * b[k + n * j]; 
      } 
      c[i + n * j] = sum; 
    } 
  } 

} 
}
}

私がやりたいのは、CUDA mmulアルゴリズムと同様の方法で、共有メモリを使用して行列「a」と「b」のタイルをキャッシュし、「c」の計算に使用することです。

基本的にCUDAでは、ブロックの正確なサイズがわかり、次のことができるようになります。

  • ブロックのサイズで共有メモリを宣言する
  • データの「関連する」部分をブロックにコピーします
  • このデータを使用する

私は使用できることを理解しています

#pragma acc cached

ディレクティブであり、vectorおよびgangオプションを使用してブロックサイズを指定できますが、それがCUDAアーキテクチャにどのようにマッピングされるかを理解するのに問題があります。

OpenACCで同様のことを達成する方法はありますか？キャッシュされたディレクティブの使用や、共有メモリの能力の一部をCUDAからOpenACCにマッピングする方法に関する優れたチュートリアル/リソースはありますか？

4

1 に答える 1

4

PGIアクセラレータコンパイラを使用している場合は、生成されたPTXファイルをダンプして、実行の実行で何が起こっているかを確認できます。

pgcc -acc -fast -Minfo -ta=nvidia,cc13,keepptx matrixMult.c -o matrixMult

生成されたPTXは現在のディレクトリに保存されます。

編集：高水準コード（CまたはFortranの場合はCUDA）を見たいと思うかもしれません。したがって、次を使用します-ta=nvidia,cc13,keepptx,keepgpu。

于 2012-10-17T08:34:42.427 に答える