CUDAでShared memoryを48KiB以上使うには

何の話か

例えば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つです。

1. カーネル関数内でのShared memoryの宣言を修正
2. cudaFuncSetAttribute関数で、立ち上げるカーネル関数が必要とするShared memoryの大きさを設定
3. カーネル関数の立ち上げ時に確保するShared memoryのサイズを指定

この3つを行うよう上記のコードを書き換えると以下のようになります。

---

終わりに

はじめのコードでは、コンパイルの時点ではどのアーキで実行されるかは判定できないため、サポートしているGPUの最小Shared memoryサイズを上限としてエラーを出しているということですかね。
ptxasでエラーが出ていることからも分かるとおり、ptxからカーネルイメージに落とす際にエラーが出るのですが、cuからptxへの変換はエラーなく行われます。
ですので、nvcc -ptx main.error.cuでptxを見てみると、64KiBのShared memoryをとろうとしていることが確認できます。