新人メイドさくらがお届け！NVIDIA Cutlassで学ぶCUDAテンプレートの極意

NVIDIA/cutlass

今日はね、ご主人様がとっても興味を持ってくれた「NVIDIA/cutlass」について、さくらが分かりやすく丁寧にご説明しちゃうね！

えへへ、ちょっと難しいお話になっちゃうかもしれないけど、さくらが頑張ってご主人様を楽しませるコントにしちゃうから、最後までお付き合いしてくれたら嬉しいな！

登場人物

さくら
元気いっぱいの新人メイド。GPUとディープラーニングが大好き！

ご主人様
ちょっと困り顔のソフトウェアエンジニアさん。

(舞台
ぴゅあはーと店内。さくらがニコニコしながらご主人様にお茶を出している。)

さくら
ご主人様、お待たせしました！さくら特製、愛情たっぷりのロイヤルミルクティーですっ！

ご主人様
ありがとう、さくらちゃん。んー、美味しいね。

さくら
えへへ、嬉しいな！ところでご主人様、さっきから「NVIDIA/cutlass」って単語が気になってるみたいだけど…何かお困りですか？

ご主人様
ああ、実はね、仕事でディープラーニングのモデルを開発してるんだけど、もっと速く計算させたいんだ。それで、Cutlassっていうのが良いって聞いたんだけど、いまいちピンとこなくて…。

さくら
(キラキラした目で) そうなんですね！ご主人様、それはさくらにとってまさに「お呼びがかかった」って感じですよ！Cutlassはね、ご主人様の悩みを解決してくれる、とっても頼りになる「勇者さま」みたいなものなんです！

ご主人様
勇者さま…？

さくら
はいっ！じゃあ、ご主人様、さくらと一緒にCutlassの冒険に出発しませんか？

さくら
ご主人様、「Cutlass」って、一言で言うとね、「CUDAを使った、すっごく速い行列計算のライブラリ」なんです！

ご主人様
CUDA…行列計算…？ちょっと難しいな。

さくら
そうですよね！じゃあね、想像してみてください！

(さくらが身振り手振りで説明する)

さくら
ご主人様のお仕事って、ディープラーニングのモデルで「計算」をいっぱいするんですよね？特に、「行列の掛け算」（数学で習った、あれです！）っていうのが、たっくさん出てくるんです。

ご主人様
うん、出てくるね。

さくら
その計算、もしご主人様が自分でイチからコードを書いたら…「よいしょ、よいしょ」って、すごく時間がかかっちゃうんです。まるで、ご主人様が一人で大きな荷物を運んでるみたいに！

ご主人様
確かに、自分で書くと結構大変だ…。

さくら
そこで登場するのが、「Cutlass」っていう勇者さまなんです！この勇者さまはね、CUDAっていう「魔法の力」（NVIDIAのGPUを動かすための技術です！）をめちゃくちゃ使いこなしてて、行列の掛け算を「ビューン！」って、あっという間に終わらせちゃうんです！

ご主人様
へえ、すごいんだね！

さくら
しかもね、Cutlassはただ速いだけじゃないんです！色々な種類の行列の掛け算に対応できるように、たくさんの「必殺技」（テンプレートって言います！）を用意してくれているんですよ！だから、ご主人様がどんな計算をしたいかによって、ぴったりの必殺技を選べるんです！

ご主人様
なるほど、汎用性も高いんだね。

さくら
はいっ！まとめると、Cutlassは、ディープラーニングで必要な行列計算を、GPUの力を最大限に引き出して、誰よりも速く、そして柔軟に実行するための、スーパー便利なツールなんです！ご主人様がソフトウェアエンジニアとして、ディープラーニングの性能を極めたいなら、まさに「導入必須」の勇者さまってわけですね！

さくら
じゃあ次に、この勇者さまがご主人様にどんな「良いこと」をもたらしてくれるのか、ご説明しますね！

(さくらがボードを指差す)

さくら
主に3つのポイントがありますっ！

爆速計算！ご主人様の時間を節約！

一番の魅力はこれです！ディープラーニングの計算は、ちょっとしたモデルでもとーっても時間がかかりますよね。Cutlassを使えば、その計算時間を劇的に短縮できます。まるで、ご主人様が「ワープ」できるようになったみたいに！

計算が速くなると、新しいアイデアをどんどん試せるようになるので、開発のスピードが格段に上がりますよ！

GPUの力を最大限に！無駄をなくすエコエンジニア！

ご主人様がお持ちのGPU、せっかくならその力を100%引き出したいですよね？Cutlassは、GPUの内部構造を熟知したプロフェッショナルが作ったライブラリなので、GPUの計算リソースを余すことなく使い切ってくれます。

ご主人様が自分でCUDAコードを書くよりも、はるかに効率的な計算ができますよ！

自分でゴリゴリ書く必要なし！ご主人様はアイデアに集中！

行列計算って、実はGPU上で効率的に動かすのがすごく難しいんです。メモリのアクセスとか、スレッドの同期とか、考えることがたくさんあって、とっても大変…。

でも、Cutlassを使えば、そういった面倒な部分を全部勇者さまが肩代わりしてくれます！ご主人様は、「どんな計算をしたいか」というアイデアに集中できるようになりますよ！まるで、さくらがご主人様の代わりに雑用を全部引き受けるみたいに！

ご主人様
なるほど、それは魅力的だね。自分でCUDAのコードを書かなくても、最高のパフォーマンスが出せるっていうのは大きいな。

さくら
はいっ！まさに、ご主人様の「エンジニア魂」を解放してくれる、そんな勇者さまなんです！

さくら
じゃあ次に、このCutlass勇者さまを、ご主人様のパソコンにお招きする方法をご説明しますね！

(さくらが笑顔で説明する)

さくら
ちょっと専門的な言葉が出てくるけど、さくらが分かりやすく説明するから安心してね！

ステップ1
必要なものを用意しよう！

NVIDIA製のGPU
これがないと始まりません！GeForceとかQuadroとか、ご主人様のPCに入ってるかな？

CUDA Toolkit
GPUで計算するための「魔法の道具セット」みたいなものです。NVIDIAの公式サイトからダウンロードしてインストールしておいてくださいね。バージョンは、Cutlassのドキュメントで推奨されているものを選ぶのがおすすめです。

Git
Cutlassのソースコードを手に入れるためのツールです。

CMake
Cutlassをコンパイル（プログラムを作ること）するためのツールです。

Visual Studio (Windowsの場合) または g++ / clang (Linuxの場合)
C++のプログラムをコンパイルするためのコンパイラです。

ステップ2
Cutlassのソースコードを手に入れよう！

さくら
まずは、Cutlassの「設計図」を手に入れるところから始めましょう！ご主人様のターミナル（真っ黒な画面で文字を打ち込むところ）を開いて、こんなコマンドを打ってみてください！

git clone https://github.com/NVIDIA/cutlass.git
cd cutlass

さくら
これで、Cutlassの「秘密基地」にご主人様が辿り着きました！

ステップ3
Cutlassを組み立てよう！

さくら
次に、手に入れた設計図をもとに、Cutlassを「使える状態」に組み立てていきます！ちょっと専門的だけど、大丈夫！

# ビルド用のディレクトリを作成
mkdir build
cd build

# CMakeを使ってビルド設定を行う
# --config Release は、リリースビルド（最適化されたもの）を作るという意味です
# -DCUTLASS_NVCC_ARCHITECTURES="<ご自身のGPUのcompute capability>" は、
# ご自身のGPUに合わせた最適化を行うための設定です。
# 例：RTX 3080なら 80 とか、RTX 2070なら 75 とかになります。
# 詳しくはお使いのGPUのスペックを見て調べてみてくださいね！
# 調べ方が分からなかったら、いつでもさくらに聞いてね！
cmake .. -DCUTLASS_NVCC_ARCHITECTURES="<ご自身のGPUのcompute capability>" --config Release

# ビルドを実行！
# 並列でビルドすると速いので、-j オプションでCPUのコア数を指定すると良いです
cmake --build . --config Release -j <CPUのコア数>

さくら
これで、ご主人様のパソコンにCutlass勇者さまが「召喚」されました！えへへ、ちょっと時間がかかるかもしれないけど、頑張って待ってあげてね！

ご主人様
compute capabilityか、それは調べておく必要があるね。

さくら
はいっ！もし分からなかったら「NVIDIA compute capability」で検索すると、すぐに調べられますよ！

さくら
さあ、いよいよCutlass勇者さまの「必殺技」を見てみましょう！今回は、一番よく使う「行列の掛け算（GEMM
General Matrix Multiply）」の例をご紹介しますね！

(さくらが笑顔で説明する)

さくら
Cutlassにはたくさんのサンプルコードが用意されているんだけど、今回はとってもシンプルな例を、さくらがカスタマイズしてみたよ！

#include <iostream>
#include <vector>
#include <numeric>
#include <algorithm>

// Cutlassのヘッダーファイル
// インストールディレクトリやビルドディレクトリからのパスに応じて調整してください
#include "cutlass/gemm/device/gemm.h" // GEMMのデバイスカーネル
#include "cutlass/half.h"             // float16 (半精度浮動小数点数) を使う場合
#include "cutlass/numeric_types.h"    // 数値型定義

// ホスト側のヘルパー関数 (CPU側で動く関数)
template <typename T>
void print_matrix(const std::string& name, const T* matrix, int rows, int cols) {
    std::cout << name << ":" << std::endl;
    for (int i = 0; i < rows; ++i) {
        for (int j = 0; j < cols; ++j) {
            std::cout << matrix[i * cols + j] << "\t";
        }
        std::cout << std::endl;
    }
    std::cout << std::endl;
}

// 行列の掛け算をCPUで実行する参考関数
template <typename T>
void cpu_gemm(const T* A, const T* B, T* C, int M, int N, int K) {
    for (int m = 0; m < M; ++m) {
        for (int n = 0; n < N; ++n) {
            C[m * N + n] = 0;
            for (int k = 0; k < K; ++k) {
                C[m * N + n] += A[m * K + k] * B[k * N + n];
            }
        }
    }
}

int main() {
    // 行列のサイズを定義
    // M: 行列Aの行数、行列Cの行数
    // N: 行列Bの列数、行列Cの列数
    // K: 行列Aの列数、行列Bの行数
    int M = 4;
    int N = 4;
    int K = 4;

    // 行列の要素数を計算
    size_t size_A = M * K;
    size_t size_B = K * N;
    size_t size_C = M * N;

    // 計算に使う型 (例: float)
    using ElementType = float; // float32

    // ホスト（CPU）側でメモリを確保し、行列を初期化
    std::vector<ElementType> h_A(size_A);
    std::vector<ElementType> h_B(size_B);
    std::vector<ElementType> h_C_cutlass(size_C); // Cutlassの結果を格納
    std::vector<ElementType> h_C_cpu(size_C);     // CPUの結果を格納（比較用）

    // 行列AとBを適当な値で初期化
    for (int i = 0; i < M; ++i) {
        for (int j = 0; j < K; ++j) {
            h_A[i * K + j] = static_cast<ElementType>(i * K + j + 1);
        }
    }
    for (int i = 0; i < K; ++i) {
        for (int j = 0; j < N; ++j) {
            h_B[i * N + j] = static_cast<ElementType>(i * N + j + 10);
        }
    }

    // ホスト側の行列を表示
    print_matrix("Matrix A", h_A.data(), M, K);
    print_matrix("Matrix B", h_B.data(), K, N);

    // デバイス（GPU）側でメモリを確保
    ElementType* d_A;
    ElementType* d_B;
    ElementType* d_C;
    cudaMalloc(&d_A, size_A * sizeof(ElementType));
    cudaMalloc(&d_B, size_B * sizeof(ElementType));
    cudaMalloc(&d_C, size_C * sizeof(ElementType));

    // ホストからデバイスへデータをコピー
    cudaMemcpy(d_A, h_A.data(), size_A * sizeof(ElementType), cudaMemcpyHostToDevice);
    cudaMemcpy(d_B, h_B.data(), size_B * sizeof(ElementType), cudaMemcpyHostToDevice);
    
    //----------------------------------------------------------------------
    // Cutlassを使ってGEMMを実行！
    //----------------------------------------------------------------------

    // Cutlass GEMMのインスタンスを作成
    // gemm::device::Gemm<
    //     ElementType, // Aの要素型
    //     cutlass::layout::RowMajor, // Aのレイアウト (RowMajor/ColumnMajor)
    //     ElementType, // Bの要素型
    //     cutlass::layout::RowMajor, // Bのレイアウト
    //     ElementType, // Cの要素型
    //     cutlass::layout::RowMajor, // Cのレイアウト
    //     ElementType, // 計算の内部精度 (Accumlatorの型)
    //     cutlass::arch::OpClassSimt, // 演算クラス (Simt:汎用的な演算、TensorOp:Tensor Coreを使う場合)
    //     cutlass::arch::Sm80, // Target Sm (GPUのCompute Capability) 例: Sm75, Sm80, Sm86
    //     cutlass::gemm::GemmShape<128, 128, 8>, // Threadblock Shape (スレッドブロックごとの計算量)
    //     cutlass::gemm::GemmShape<64, 64, 8>,  // Warp Shape (Warpごとの計算量)
    //     cutlass::gemm::GemmShape<16, 8, 8>,   // Instruction Shape (Instructionごとの計算量)
    //     cutlass::epilogue::thread::LinearCombination<
    //         ElementType, // ElementOutput
    //         128,         // ElementsPerAccess
    //         ElementType, // ElementAccumulator
    //         ElementType  // ElementVector
    //     >
    // >
    // 上記のテンプレート引数は非常に多岐にわたりますが、Cutlassのドキュメントやサンプルを参照すると、
    // 用途に合わせた推奨設定が見つかります。今回は簡単のため、基本的な設定を使用します。
    // 詳細なテンプレート引数の説明は割愛しますが、
    // 簡単に言えば「どういう種類の行列を、どういう効率で計算するか」を指定しています。
    using Gemm = cutlass::gemm::device::Gemm<
        ElementType, cutlass::layout::RowMajor,
        ElementType, cutlass::layout::RowMajor,
        ElementType, cutlass::layout::RowMajor,
        ElementType
    >;

    Gemm gemm_op;

    // GEMMの引数を設定
    typename Gemm::Arguments arguments{
        {M, N, K},          // Problem size (M, N, K)
        {d_A, K},           // Tensor A (data pointer, stride)
        {d_B, N},           // Tensor B
        {d_C, N},           // Tensor C (出力)
        {d_C, N},           // Tensor D (出力, Cと同じでOK)
        {1.0f, 0.0f}        // Alpha, Beta (C = alpha * A * B + beta * C)
    };

    // GEMMを実行！
    cutlass::Status status = gemm_op(arguments);
    if (status != cutlass::Status::kSuccess) {
        std::cerr << "Cutlass GEMM failed: " << cutlassGetStatusString(status) << std::endl;
        return 1;
    }

    // デバイスからホストへ結果をコピー
    cudaMemcpy(h_C_cutlass.data(), d_C, size_C * sizeof(ElementType), cudaMemcpyDeviceToHost);

    // Cutlassの結果を表示
    print_matrix("Result (Cutlass)", h_C_cutlass.data(), M, N);

    //----------------------------------------------------------------------
    // CPUで同じ計算を実行（結果の比較用）
    //----------------------------------------------------------------------
    cpu_gemm(h_A.data(), h_B.data(), h_C_cpu.data(), M, N, K);
    print_matrix("Result (CPU)", h_C_cpu.data(), M, N);

    // 結果の比較（簡易的なチェック）
    bool ok = true;
    for (size_t i = 0; i < size_C; ++i) {
        if (std::abs(h_C_cutlass[i] - h_C_cpu[i]) > 1e-4) { // 浮動小数点数の比較は誤差を考慮
            ok = false;
            break;
        }
    }
    if (ok) {
        std::cout << "Results match between Cutlass and CPU!" << std::endl;
    } else {
        std::cout << "Results MISMATCH between Cutlass and CPU!" << std::endl;
    }


    // メモリ解放
    cudaFree(d_A);
    cudaFree(d_B);
    cudaFree(d_C);

    return 0;
}

さくら
わー、ご主人様、ちょっとコードが長くなっちゃったけど、大丈夫かな？

ご主人様
うーん、確かにテンプレート引数の部分が複雑そうだね。

さくら
そうなんです！Cutlassは、とっても柔軟に設定できる代わりに、この「テンプレート引数」で、どんなGPUのアーキテクチャで、どんなデータ形式で、どんな計算のブロックサイズで…っていうのを細かく指定するんです！

さくら
でもね、安心してください！これはあくまで一例で、Cutlassの公式サンプルコードには、もっとたくさんのパターンや、ご自身のGPUに合わせた最適な設定方法が、たっくさん載っているんです！

ご主人様
なるほど、まずは簡単な例から試して、少しずつカスタマイズしていくのが良さそうだね。

さくら
はいっ！その通りです！このコードをコンパイルするには、CUDAのコンパイラである nvcc を使う必要があります。

コンパイル方法の例

ご主人様のmain.cppファイルに上記のコードを保存して、Cutlassのビルドディレクトリ内で、以下のようにコンパイルしてみてください。（cutlass/includeディレクトリへのパスは適宜調整してくださいね）

# Cutlassのビルドディレクトリにいることを前提
nvcc -std=c++17 -I../../include -I. main.cpp -o cutlass_gemm_example -L./lib -lcutlass -lcuda -lcudart

さくら
えへへ、-Iとか-Lとか、おまじないみたいでしょ？これはね、「Cutlassの頭脳（ヘッダーファイル）はここだよ！」とか、「Cutlassの筋肉（ライブラリ）はここにあるよ！」って、コンパイラさんに教えてあげるおまじないなんです！

ご主人様
ふむふむ、これで動くのか。

さくら
はい！実行すると、Cutlassを使ったGPUでの計算結果と、CPUでの計算結果が表示されて、ちゃんと一致しているか確認できるはずです！

さくら
ご主人様、Cutlassの冒険、どうだったかな？

ご主人様
うん、すごくよく分かったよ、さくらちゃん。最初は難しそうだったけど、君の説明のおかげで、導入してみる気になったよ。

さくら
えへへ、嬉しいな！最後に、さくらからご主人様に、ちょっとだけアドバイスがあります！

最初は公式ドキュメントとサンプルから！ Cutlassは奥が深いけど、公式ドキュメントやGithubのサンプルコードがとっても充実しています。まずはそれらを動かしてみて、基本的な使い方に慣れるのが一番の近道ですよ！

GPUのアーキテクチャを意識しよう！ Cutlassは、GPUのアーキテクチャ（Sm75, Sm80など）に特化した最適化を行うので、ご自身のGPUのCompute Capabilityに合わせて設定を調整することが、最高のパフォーマンスを引き出す秘訣です！

テンプレートの沼は楽しい沼！ C++のテンプレートをガッツリ使うので、最初は戸惑うかもしれません。でも、このテンプレートのおかげで、Cutlassはどんな計算にも柔軟に対応できるんです！ちょっとずつ慣れていくと、きっと「沼」にはまっちゃうかも？えへへ！

さくら
ご主人様がCutlassを使って、ディープラーニングの計算を「爆速」にできたら、さくらもとっても嬉しいな！何か困ったことがあったら、いつでも「ぴゅあはーと」に来て、さくらを呼んでくださいね！さくらが全力でサポートしますっ！

ご主人様
ありがとう、さくらちゃん！君のおかげで、一歩前に進めそうだ。また会いに来るよ。

さくら
はいっ！ご主人様、またのご来店を心よりお待ちしております！いってらっしゃーい！

さくら
(ご主人様を見送った後、ニコニコしながら) よーし、これでご主人様のディープラーニングがもっと速くなるはず！さくら、今日もご奉仕できて満足ですっ！えへへ！