内積をタイル化する
フェーズごとにAとBのサブタイルを読み込みます。
「内積をタイル化する」はCoddyKit上の無料CUDA Academyレッスンです。 これはレッスン2/4です。 下記で完全なレッスンを無料で読むことができます。その後、ブラウザ内の組み込みコードエディタと24時間対応のAIチューターでハンズオン演習できます。 これはCUDA Academy学習パスの一部であり、ウェブとCoddyKitアプリ全体で進捗が同期されます。 CUDA Academyコースには全4レッスンが含まれています。
タイル化の考え方
タイル化では、行列を高速なオンチップメモリに収まる小さな正方形のタイルに分割します。スレッドが協調してタイルを一度だけ読み込み、それを何度も再利用します。
共有メモリを使う理由
タイルはブロック内のすべてのスレッドから見える__shared__メモリに置きます。何度もglobal memoryへアクセスするより、そこから読み取るほうがはるかに高速です。⚡
__shared__ float As[TILE][TILE];
__shared__ float Bs[TILE][TILE];1ブロック、1出力タイル
各blockは、出力CのTILE×TILEの領域を1つ担当します。そのスレッドたちが協力して、その領域全体を計算します。
タイルサイズをブロックに合わせる
TILEにはブロックの幅と同じ値を選びます。多くの場合は16または32です。これにより、各スレッドが各tileの要素をちょうど1つずつ読み込めます。
#define TILE 16
dim3 threads(TILE, TILE);スレッドをタイル内の位置に対応させる
タイル内でのスレッドの位置は、単にthreadIdxです。グローバルな行と列は、これまでどおりブロックとスレッドのインデックスから求めます。
int ty = threadIdx.y, tx = threadIdx.x;
int row = blockIdx.y*TILE + ty;
int col = blockIdx.x*TILE + tx;Aのタイルを読み込む
各スレッドは、Aの現在のタイルから要素を1つコピーしてshared memoryに格納します。ブロック全体でAのTILE×TILEの領域を読み込みます。
As[ty][tx] = A[row*N + (phase*TILE + tx)];Bのタイルを読み込む
同時に、各スレッドがBのタイルから要素を1つ読み込みます。これで両方のtilesがチップ上に置かれ、高速に繰り返し読み取れる状態になります。
Bs[ty][tx] = B[(phase*TILE + ty)*N + col];計算前に同期する
すべてのスレッドが読み込みを終えてから誰もタイルを読み取らないよう、__syncthreads()を呼び出します。これを省くと不正な結果になります。
__syncthreads();タイル上で計算する
次に各スレッドがtileに対する短いループを実行し、shared memoryだけを読み取ります。これらの読み取りはglobal memoryより大幅に低コストです。
for (int k = 0; k < TILE; ++k)
sum += As[ty][k] * Bs[k][tx];再利用による効果
読み込んだ各値は、1スレッドではなくTILE個のスレッドによって使われます。このデータ再利用こそ、タイル化したmatmulが高速になる理由です。
1つのタイルでは不十分
1つのタイルで扱えるのは、内積の一部だけです。次のレッスンで扱うように、複数のフェーズにわたって読み込み・同期・計算を繰り返します。
確認問題
共有タイルを使うときの手順の順番を思い出してください。
まとめ
AとBのタイルをshared memoryに読み込み、同期してから、低コストなオンチップ読み取りで計算しました。勝負を分けるのは再利用です。次はフェーズを扱います。🧱
AI チューターと学ぶ C++ — 無料
ブラウザでリアルコードを書いて実行し、24/7 の AI チューターから瞬時にサポートを受け、ウェブまたはアプリで続きから学習できます。
- コース
- 30
- レッスン
- 120
よくある質問
「内積をタイル化する」レッスンは無料ですか？
はい。「内積をタイル化する」の完全なテキストはこのウェブで無料で読めます。インタラクティブに演習し（組み込みコードエディタと24時間対応のAIチューター）、CUDA Academyコースの残りをアンロックするには、CoddyKit PROにアップグレードしてください。 CUDA Academyコースには全4レッスンが含まれています。
「内積をタイル化する」で何を学びますか？
フェーズごとにAとBのサブタイルを読み込みます。 ブラウザで直接実行するハンズオンコードでCUDA Academyを演習し、24時間対応のAIチューターがレッスンを進める中での質問に答えます。
CUDA Academyを始めるのに経験は必要ですか？
事前経験は必要ありません。CoddyKitのCUDA Academyは初級者から上級者向けに構成されているため、ここから始めるか最初から始めて、自分のペースで進むことができます。 これはレッスン2/4です。
「内積をタイル化する」レッスンにはどのくらい時間がかかりますか？
ほとんどのCoddyKitレッスンは約5～10分かかります。各レッスンはコンパクトでインタラクティブなので、着実に進歩し、ウェブとアプリ全体で正確に前回の場所から再開できます。
このCUDA Academyレッスンでコードを書いて実行できますか？
はい。すべてのCUDA Academyレッスンに組み込みコードエディタが含まれているため、ブラウザでリアルコードを書いて実行し、即座のAIフィードバックを取得できます。ローカル設定は不要です。
このコースのすべてのレッスン
- 単純な行列乗算カーネル
- 内積をタイル化する
- タイルのフェーズをループする
- 高速化の効果を測定する