CUDA Academy · レッスン

WMMA フラグメント API

フラグメントのロード、mma_sync、ストア

レッスン 3/413 ステップ

「WMMA フラグメント API」はCoddyKit上の無料CUDA Academyレッスンです。 これはレッスン3/4です。 下記で完全なレッスンを無料で読むことができます。その後、ブラウザ内の組み込みコードエディタと24時間対応のAIチューターでハンズオン演習できます。 これはCUDA Academy学習パスの一部であり、ウェブとCoddyKitアプリ全体で進捗が同期されます。 CUDA Academyコースには全4レッスンが含まれています。

WMMA API

Tensor Coreを直接プログラミングするには、nvcuda::wmma名前空間にあるWMMA(warp matrix multiply-accumulate)APIを使用します。🧩

ワープ全体で協調する

WMMAはワープ単位で動作します。1つのタイルに対して、ワープ内の32スレッドすべてがワープとして協調します。単一スレッドではなく、タイル単位で考えます。

フラグメントに出会う

フラグメントは、行列タイルのうち1つのワープが担当する部分を保持するWMMAのデータ型です。各スレッドがその要素の一部を担当します。

3種類のフラグメントの役割

matrix_a、matrix_b、またはaccumulatorというタグを付けてフラグメントを宣言します。このタグによって、WMMAはそのタイルを積和演算にどう使うかを判断します。

タイル形状は固定されている

フラグメントでは、16 by 16 by 16のようにタイルサイズが決められています。対応している形状を選びます。ハードウェアが受け付ける組み合わせは限られています。

ステップ1:ロード

まずload_matrix_syncを呼び出し、メモリからタイルをフラグメントに読み込みます。このロードによって、データはワープ全体に自動的に分配されます。

wmma::load_matrix_sync(a_frag, ptr, ldm);

ステップ2:アキュムレーターをクリアする

加算する前に、fill_fragmentで結果タイルをゼロにします。きれいなアキュムレーターを用意することで、既知の値から累積を開始できます。

wmma::fill_fragment(c_frag, 0.0f);

ステップ3:積和演算

中心となる呼び出しはmma_syncです。読み込んだフラグメントに対してD = A times B plus Cを実行し、Tensor Coreを直接使用します。

wmma::mma_sync(c_frag, a_frag, b_frag, c_frag);

ステップ4:ストア

最後にstore_matrix_syncでアキュムレーターフラグメントをメモリに書き戻します。このストアによって、各スレッドの担当部分が1つのタイルにまとめられます。

wmma::store_matrix_sync(out, c_frag, ldm, layout);

Syncはワープ同期を意味する

_syncという接尾辞は、すべてのWMMA呼び出しがワープ同期であることを示します。32個のレーンすべてが一緒に到達しなければ、動作は未定義になります。

レイアウトと先行次元

ロードとストアでは、行間のストライドである先行次元に加えて、メモリを正しく読み取るための行優先または列優先レイアウトが必要です。

クイックチェック

Tensor Core上で行列積和演算を実際に実行するWMMA呼び出しはどれですか。

まとめ

WMMAの流れを確認しました。フラグメントを宣言し、タイルをロードし、アキュムレーターをクリアし、mma_syncを呼び出してからストアします。すべての手順はワープ同期で実行されます。🙌

無料で開始

AI チューターと学ぶ C++ — 無料

ブラウザでリアルコードを書いて実行し、24/7 の AI チューターから瞬時にサポートを受け、ウェブまたはアプリで続きから学習できます。

コース
30
レッスン
120

よくある質問

「WMMA フラグメント API」レッスンは無料ですか?

はい。「WMMA フラグメント API」の完全なテキストはこのウェブで無料で読めます。インタラクティブに演習し(組み込みコードエディタと24時間対応のAIチューター)、CUDA Academyコースの残りをアンロックするには、CoddyKit PROにアップグレードしてください。 CUDA Academyコースには全4レッスンが含まれています。

「WMMA フラグメント API」で何を学びますか?

フラグメントのロード、mma_sync、ストア ブラウザで直接実行するハンズオンコードでCUDA Academyを演習し、24時間対応のAIチューターがレッスンを進める中での質問に答えます。

CUDA Academyを始めるのに経験は必要ですか?

事前経験は必要ありません。CoddyKitのCUDA Academyは初級者から上級者向けに構成されているため、ここから始めるか最初から始めて、自分のペースで進むことができます。 これはレッスン3/4です。

「WMMA フラグメント API」レッスンにはどのくらい時間がかかりますか?

ほとんどのCoddyKitレッスンは約5~10分かかります。各レッスンはコンパクトでインタラクティブなので、着実に進歩し、ウェブとアプリ全体で正確に前回の場所から再開できます。

このCUDA Academyレッスンでコードを書いて実行できますか?

はい。すべてのCUDA Academyレッスンに組み込みコードエディタが含まれているため、ブラウザでリアルコードを書いて実行し、即座のAIフィードバックを取得できます。ローカル設定は不要です。

このコースのすべてのレッスン

  1. Tensor Core が計算するもの
  2. 混合精度: FP16、BF16、TF32
  3. WMMA フラグメント API
  4. 数値安定性のトレードオフ
← CUDA Academyに戻る