レジスターとローカルメモリ
スレッドごとに使える最速の保存領域と、あふれた場合の動作を学びます。
「レジスターとローカルメモリ」はCoddyKit上の無料CUDA Academyレッスンです。 これはレッスン1/4です。 下記で完全なレッスンを無料で読むことができます。その後、ブラウザ内の組み込みコードエディタと24時間対応のAIチューターでハンズオン演習できます。 これはCUDA Academy学習パスの一部であり、ウェブとCoddyKitアプリ全体で進捗が同期されます。 CUDA Academyコースには全4レッスンが含まれています。
最速のメモリ
すべてのスレッドは、チップ上で最も高速なストレージである専用のレジスタを持ちます。レジスタは演算ユニットのすぐ隣にあるため、読み取りのコストはほとんどありません。
通常の変数はレジスタになる
カーネル内で通常のローカル変数を宣言すると、コンパイラは通常それをレジスタに保持します。特別な構文はまったく必要ありません。🙂
__global__ void k() {
int x = 5; // x lives in a register
float y = x * 2.0f;
}1つのスレッドだけが使用する
レジスタの値は、厳密に1つのスレッドに属します。スレッド7からスレッド8のレジスタは見えないため、スレッドごとの処理をきれいに分離できます。
レジスタは共有プール
各ストリーミングマルチプロセッサには、容量が固定されたレジスタファイルがあります。常駐するすべてのスレッドでこれを分けて使うため、1スレッドあたりのレジスタ使用量を減らすと、同時に実行できるスレッド数が増えます。
使い切った場合
カーネルが使用可能な数を超えるレジスタを必要とすると、余分な値は別の場所に移されます。この処理をレジスタスピルと呼びます。
スピル先: ローカルメモリ
スピルした値はローカルメモリに格納されます。名前とは異なり、これはチップ上にはありません。実際には、低速なチップ外のDRAM上にあります。
ローカルでもスレッドごと
ローカルメモリはレジスタと同じように各スレッド専用です。違いは速度だけで、ローカルメモリへのアクセスははるかに低速です。
大きなローカル配列はスピルする
スレッドごとに大きな配列を宣言すると、すべての要素を保持できるほどレジスタが存在しないため、多くの場合ローカルメモリに置かれます。
__global__ void k() {
float buf[64]; // likely lives in local memory
}スピルは性能を損なう
スピルが発生するたびに、無料同然のレジスタアクセスが低速なDRAMへのアクセスに変わります。レジスタの使用圧力を下げることは、最も簡単にできる最適化の1つです。
レジスタ数を確認する
コンパイラに、スレッドあたりのレジスタ数を報告させることができます。nvccに--ptxas-options=-vを渡すと、各カーネルの使用量が出力されます。
nvcc --ptxas-options=-v kernel.cu意図的にレジスタ数を制限する
__launch_bounds__修飾子を使うと、コンパイラにレジスタ数の制限を促せます。スレッドごとの速度を少し犠牲にして、同時に実行するスレッド数を増やします。
__global__ void __launch_bounds__(256)
myKernel() { /* ... */ }理解度チェック
スレッドごとに大きな配列を宣言したところ、性能が低下しました。何が起きたと考えられますか。
まとめ: 高速かつ専用
レジスタはスレッドごとに使える最速のストレージであり、その容量には限りがあることを学びました。あふれた値は低速なローカルメモリにスピルします。レジスタの使用圧力を低く保ちましょう。🎯
AI チューターと学ぶ C++ — 無料
ブラウザでリアルコードを書いて実行し、24/7 の AI チューターから瞬時にサポートを受け、ウェブまたはアプリで続きから学習できます。
- コース
- 30
- レッスン
- 120
よくある質問
「レジスターとローカルメモリ」レッスンは無料ですか？
はい。「レジスターとローカルメモリ」の完全なテキストはこのウェブで無料で読めます。インタラクティブに演習し（組み込みコードエディタと24時間対応のAIチューター）、CUDA Academyコースの残りをアンロックするには、CoddyKit PROにアップグレードしてください。 CUDA Academyコースには全4レッスンが含まれています。
「レジスターとローカルメモリ」で何を学びますか？
スレッドごとに使える最速の保存領域と、あふれた場合の動作を学びます。 ブラウザで直接実行するハンズオンコードでCUDA Academyを演習し、24時間対応のAIチューターがレッスンを進める中での質問に答えます。
CUDA Academyを始めるのに経験は必要ですか？
事前経験は必要ありません。CoddyKitのCUDA Academyは初級者から上級者向けに構成されているため、ここから始めるか最初から始めて、自分のペースで進むことができます。 これはレッスン1/4です。
「レジスターとローカルメモリ」レッスンにはどのくらい時間がかかりますか？
ほとんどのCoddyKitレッスンは約5～10分かかります。各レッスンはコンパクトでインタラクティブなので、着実に進歩し、ウェブとアプリ全体で正確に前回の場所から再開できます。
このCUDA Academyレッスンでコードを書いて実行できますか？
はい。すべてのCUDA Academyレッスンに組み込みコードエディタが含まれているため、ブラウザでリアルコードを書いて実行し、即座のAIフィードバックを取得できます。ローカル設定は不要です。
このコースのすべてのレッスン
- レジスターとローカルメモリ
- グローバルメモリのトレードオフ
- コンスタントメモリとそのキャッシュ
- メモリ階層のメンタルモデル