何の話か
例えばGeForce RTX 3080 (Shared memory/L1 Cache: 128KB)で走らせることを想定した以下のコードがあります。
このコードは64KiB分のShared memoryのデータをGlobal memoryに書き出すだけのコードです。
このコードをnvccでコンパイルすると、アクセスするShared memoryのアドレスは搭載されているShared memoryの大きさを超えていないですが、エラーとなります。
ptxas error : Entry function '_Z6kernelPf' uses too much shared data (0x10000 bytes, 0xc000 max)
0xc000 bytesは48KiBです。
ではどう書いたら48KiB以上のSharedメモリを使えるようになるかというのがこの記事です。
48KiBを超えるSharedメモリの確保の仕方
48KiB以上のShared memoryを確保するために行うことは3つです。
- カーネル関数内でのShared memoryの宣言を修正
- cudaFuncSetAttribute関数で、立ち上げるカーネル関数が必要とするShared memoryの大きさを設定
- カーネル関数の立ち上げ時に確保するShared memoryのサイズを指定
この3つを行うよう上記のコードを書き換えると以下のようになります。
終わりに
はじめのコードでは、コンパイルの時点ではどのアーキで実行されるかは判定できないため、サポートしているGPUの最小Shared memoryサイズを上限としてエラーを出しているということですかね。
ptxasでエラーが出ていることからも分かるとおり、ptxからカーネルイメージに落とす際にエラーが出るのですが、cuからptxへの変換はエラーなく行われます。
ですので、nvcc -ptx main.error.cuでptxを見てみると、64KiBのShared memoryをとろうとしていることが確認できます。
カテゴリー:CUDA
記事作成日:2021-02-03