動的共有メモリ
起動時に共有メモリのサイズを指定します。
「動的共有メモリ」はCoddyKit上の無料CUDA Academyレッスンです。 これはレッスン4/4です。 下記で完全なレッスンを無料で読むことができます。その後、ブラウザ内の組み込みコードエディタと24時間対応のAIチューターでハンズオン演習できます。 これはCUDA Academy学習パスの一部であり、ウェブとCoddyKitアプリ全体で進捗が同期されます。 CUDA Academyコースには全4レッスンが含まれています。
サイズがまだ決まっていない場合
shared arrayのサイズが実行時の値によって決まり、ハードコードできない場合があります。CUDAではdynamic shared memoryを使ってこの問題を解決します。
extern宣言
extern __shared__と空の角括弧を使って宣言します。コード内にサイズを指定しないため、コンパイラはサイズを未確定のままにします。
extern __shared__ float sdata[];起動時にサイズを設定
バイト数を3番目の起動パラメータとして、スレッド数と閉じ角括弧の間に渡します。これで実行時にサイズが設定されます。
kern<<<blocks, threads, n * sizeof(float)>>>();常にバイト単位
よくあるバグは、3番目の引数が要素数ではなくバイト数であることを見落とすことです。正しく指定するには、要素数にsizeofを掛けてください。
デフォルトは0バイト
3番目の引数を省略すると、CUDAは0バイトのdynamic shared memoryしか確保しません。その状態でsdataを読み取ると未定義動作になるため、絶対に忘れないでください。
静的と動的を併用
1つのkernelで両方の種類を使えます。Static arrayは固定サイズを維持し、dynamic blockは起動時の引数に応じてサイズが変わります。
1つのバッファ、多数の配列
dynamic shared memoryが提供するのは1つのフラットなバッファです。複数の配列を格納するには、ポインタのオフセットを使って自分で分割します。🔪
オフセットで分割
2つ目の配列をバッファの途中から指すようにします。ただし、floatがワードの途中から始まらないよう、それぞれの領域をアラインメントしてください。
extern __shared__ float buf[];
float* a = buf;
float* b = &buf[blockDim.x];dynamicを使う理由
1つのコンパイル済みkernelでさまざまなタイルサイズに対応できます。ケースごとに再コンパイルせず、起動ごとにshared memoryのサイズを調整できます。
ハードウェアの上限に注意
それでも、ブロックごとの最大値を超えることはできません。起動時に多すぎるバイト数を要求すると、kernelは起動に失敗します。
大きな割り当てを有効にする
新しいGPUでデフォルトの上限を超えるには、起動前にcudaFuncSetAttributeを使って明示的に有効化する必要があります。そうしないと、大きな要求は拒否されます。
クイックチェック
dynamic shared memoryのサイズ指定方法を確認しましょう。
まとめ
extern __shared__ arrayは起動時にバイト数が決まり、複数の配列に分割でき、ブロックごとの上限に従うことを学びました。コース完了です！🎉
よくある質問
「動的共有メモリ」レッスンは無料ですか？
はい。「動的共有メモリ」の完全なテキストはこのウェブで無料で読めます。インタラクティブに演習し（組み込みコードエディタと24時間対応のAIチューター）、CUDA Academyコースの残りをアンロックするには、CoddyKit PROにアップグレードしてください。 CUDA Academyコースには全4レッスンが含まれています。
「動的共有メモリ」で何を学びますか？
起動時に共有メモリのサイズを指定します。 ブラウザで直接実行するハンズオンコードでCUDA Academyを演習し、24時間対応のAIチューターがレッスンを進める中での質問に答えます。
CUDA Academyを始めるのに経験は必要ですか？
事前経験は必要ありません。CoddyKitのCUDA Academyは初級者から上級者向けに構成されているため、ここから始めるか最初から始めて、自分のペースで進むことができます。 これはレッスン4/4です。
「動的共有メモリ」レッスンにはどのくらい時間がかかりますか？
ほとんどのCoddyKitレッスンは約5～10分かかります。各レッスンはコンパクトでインタラクティブなので、着実に進歩し、ウェブとアプリ全体で正確に前回の場所から再開できます。
このCUDA Academyレッスンでコードを書いて実行できますか？
はい。すべてのCUDA Academyレッスンに組み込みコードエディタが含まれているため、ブラウザでリアルコードを書いて実行し、即座のAIフィードバックを取得できます。ローカル設定は不要です。