Chaos and Order

💡 왼쪽 원문을 읽으면서 오른쪽에 따라 써보세요. Tab 키로 힌트를 받을 수 있습니다.

원문 렌더가 준비되기 전까지 텍스트 가이드로 표시합니다.

はじめに

GPU プログラミングに初めて触れると、もっとも分かりにくいのが「自分が書いた一行のカーネルコードが、実際のハードウェア上でどのように数千個へと展開されるのか」という点です。CPU は数個のコアで逐次的な処理を高速にこなすよう設計されているのに対し、GPU は数千個の単純な演算ユニットで同じ演算を同時に処理するよう設計されています。この違いを理解しないと、コードは動くものの性能は CPU 以下、という結果になりがちです。

本記事では、CUDA の論理的な実行モデル(グリッド・ブロック・ワープ・スレッド)から、実際のハードウェア(SM、メモリ階層、テンソルコア)までを、図を中心に整理します。抽象的な概念を図に置き換えておくと、後で性能チューニングをするときに「なぜこのコードは遅いのか」を頭の中でシミュレーションできるようになります。

本記事で扱う内容は次のとおりです。

- 実行モデル:グリッド、ブロック、ワープ、スレッドの階層構造

- ハードウェア:SM(Streaming Multiprocessor)の内部構造

- メモリ階層:レジスタ、共有メモリ、L1/L2 キャッシュ、グローバル HBM

- コアレス(coalesced)メモリアクセスのパターン

- ワープスケジューリングと占有率(occupancy)

- テンソルコア(Tensor Core)と混合精度演算

- ストリーム(stream)と非同期実行

- シンプルなカーネル例とよくある落とし穴

> アーキテクチャごと(例:Hopper、Blackwell)の具体的な数値は世代やチップによって異なりますので、正確な値は公式ドキュメントと各チップの仕様書をご確認ください。本記事では、概念的に一貫している部分を中心に説明します。

実行モデル:論理的な階層構造

CUDA の実行モデルは徹底して階層的です。カーネル(kernel)を実行すると多数のスレッドが生成されますが、これらのスレッドは無秩序に散らばるのではなく、次のような階層に組織されます。

┌──────────────────────────────┐

│ GRID │

│ (カーネル1回実行 = 1グリッド) │

│ │

│ ┌────────┐ ┌────────┐ │

│ │ BLOCK │ │ BLOCK │ │

│ │ (0,0) │ │ (1,0) │ │

│ └────────┘ └────────┘ │

│ ┌────────┐ ┌────────┐ │

│ │ BLOCK │ │ BLOCK │ │

│ │ (0,1) │ │ (1,1) │ │

│ └────────┘ └────────┘ │

└──────────────────────────────┘

│

▼ (ブロック1つを拡大)

┌──────────────────────────────┐

│ BLOCK │

│ 最大1024スレッド (例) │

│ │

│ Thread Thread Thread ... │

│ (0) (1) (2) │

│ │

│ -- 32個ずつ束ねてWARP形成 -- │

│ [warp 0] [warp 1] ... │

└──────────────────────────────┘

│

▼ (ワープ1つを拡大)

┌──────────────────────────────┐

│ WARP │

│ 32スレッドが同じ命令を │

│ 同時に(SIMT)実行 │

│ │

│ T0 T1 T2 ... T30 T31 │

│ └──── lockstep実行 ────┘ │

└──────────────────────────────┘

グリッド(Grid)

カーネルを一度実行したときに生成される全スレッド集合が、1 つのグリッドです。グリッドは 1 次元、2 次元、3 次元で構成でき、行列やボリュームデータに自然に対応づけられます。

ブロック(Block)

グリッドは複数のブロックに分割されます。同じブロック内のスレッドは、次の 2 つの特権を持ちます。

1. **共有メモリ(shared memory)** を通じて高速にデータをやり取りできます。

2. **同期(`__syncthreads()`)** を通じて互いの進行状況をそろえられます。

異なるブロック同士は、原則として直接同期できません。この制約こそが GPU のスケーラビリティを保証します。ブロックは互いに独立しているため、ハードウェアは空いている SM に自由にブロックを配置できます。

ワープ(Warp)

ハードウェアの立場で見ると、実際のスケジューリング単位はスレッドではなく **ワープ** です。ワープは 32 個のスレッドの束で、同じ命令を同時に実行します。この方式を SIMT(Single Instruction, Multiple Threads)と呼びます。

SIMT実行方式 (1命令 -> 32スレッド)

PC ──▶ [ ADD r1, r2, r3 ]

│

┌────────┼─────────────────────────────────┐

▼ ▼ ▼ ... ▼ ▼

T0 T1 T2 T30 T31

(各スレッドは自分のレジスタで同じ命令を実行)

ここで重要な落とし穴が登場します。同じワープ内のスレッドが異なる分岐をたどると(例:`if`-`else`)、ハードウェアは両方の経路を逐次的に実行します。これを **ワープダイバージェンス(warp divergence)** と呼び、性能低下の主要な原因です。

ワープダイバージェンス (分岐による直列化)

if (threadIdx.x < 16) { A() } else { B() }

サイクル1..k : T0..T15 -> A()実行, T16..T31 -> 待機(マスクoff)

サイクルk+1..: T0..T15 -> 待機, T16..T31 -> B()実行

結果: A経路の時間 + B経路の時間 (並列ではなく直列)

インデックスの計算

各スレッドは自分のグローバルな位置を次のように計算します。この式は、ほぼすべての CUDA カーネルの先頭に登場します。

int idx = blockIdx.x * blockDim.x + threadIdx.x;

blockDim.x = 4 の場合のグローバルインデックス

blockIdx.x: 0 1 2

┌────────┐ ┌────────┐ ┌────────┐

threadIdx.x: │0 1 2 3 │ │0 1 2 3 │ │0 1 2 3 │

global idx : │0 1 2 3 │ │4 5 6 7 │ │8 9 10 11│

└────────┘ └────────┘ └────────┘

ハードウェア構造:SM(Streaming Multiprocessor)

論理的なモデルを実際に実行するハードウェア単位が **SM(Streaming Multiprocessor)** です。GPU 1 枚は数十個の SM で構成され、ブロックは SM に割り当てられて実行されます。1 つの SM は通常、複数のブロックを同時に収容できます。

┌─────────────────────────────────────────────────────────────┐

│ SM (Streaming Multiprocessor) │

│ │

│ ┌───────────────┐ ┌───────────────┐ │

│ │ Warp Scheduler│ │ Warp Scheduler│ (通常4パーティション)│

│ │ + Dispatch │ │ + Dispatch │ │

│ └───────┬───────┘ └───────┬───────┘ │

│ ▼ ▼ │

│ ┌──────────────┐ ┌──────────────┐ │

│ │ CUDA Cores │ │ CUDA Cores │ (FP32 / INT 演算) │

│ │ ████████████│ │ ████████████│ │

│ └──────────────┘ └──────────────┘ │

│ ┌──────────────┐ ┌──────────────┐ │

│ │ Tensor Cores │ │ Tensor Cores │ (行列積の加速) │

│ └──────────────┘ └──────────────┘ │

│ ┌──────────────┐ ┌──────────────┐ │

│ │ LD/ST │ │ SFU │ (メモリ / 超越関数) │

│ └──────────────┘ └──────────────┘ │

│ │

│ ┌───────────────────────────────────────────────────────┐ │

│ │ Register File (数万個の32ビットレジスタ) │ │

│ └───────────────────────────────────────────────────────┘ │

│ ┌───────────────────────────────────────────────────────┐ │

│ │ Shared Memory / L1 Cache (ブロック単位で分割) │ │

│ └───────────────────────────────────────────────────────┘ │

└─────────────────────────────────────────────────────────────┘

SM 内部の主な構成要素は次のとおりです。

- **ワープスケジューラ(Warp Scheduler)**:実行準備ができたワープを選び、実行ユニットへ送ります。通常、1 つの SM は複数のスケジューラを持つパーティションに分かれています。

- **CUDA コア**:FP32/INT などの基本的な算術演算を行う ALU です。

- **テンソルコア(Tensor Core)**:小さな行列の積和(MMA)を一度に処理する専用ユニットです。

- **LD/ST ユニット**:メモリのロード/ストアを担当します。

- **SFU(Special Function Unit)**:`sin`、`sqrt`、`exp` などの超越関数を高速に計算します。

- **レジスタファイル**:スレッドごとのローカル変数が格納される、もっとも高速なストレージです。

- **共有メモリ/L1**:SM 内でブロックが使用する高速なオンチップメモリです。

ここで重要な洞察は、レジスタファイルと共有メモリが **有限の資源** だという点です。1 つの SM が同時に収容できるブロックとワープの数は、これらの資源をどれだけ使うかで決まります。これこそが、次に説明する **占有率** の本質です。

メモリ階層

GPU 性能の 8 割はメモリで決まると言っても過言ではありません。演算ユニットは非常に高速ですが、データが間に合わなければ遊んでしまいます。CUDA のメモリは、速度と容量が反比例する階層構造をなしています。

メモリ階層 (上に行くほど速くて小さい)

▲ 速い / 小さい

│ ┌──────────────────────────┐

│ │ Registers (スレッド専用) │ 数サイクル, 数十~数百KB/SM

│ ├──────────────────────────┤

│ │ Shared Memory / L1 (ブロック)│ ~数十サイクル, 数十~数百KB/SM

│ ├──────────────────────────┤

│ │ L2 Cache (全SM共有) │ ~数百サイクル, 数~数十MB

│ ├──────────────────────────┤

│ │ Global Memory (HBM, DRAM)│ ~数百サイクル, 数十GB

│ └──────────────────────────┘

│ 遅い / 大きい

▼

各メモリ空間の特性を表にまとめると次のとおりです。(具体的な数値はアーキテクチャごとに異なりますので、相対的な傾向としてご理解ください。)

| --- | --- | --- | --- | --- | --- |

| 共有メモリ | ブロック | 低い | 高い | 小さい | 明示的に管理、バンク構造 |

| L1 キャッシュ | SM | 低い | 高い | 小さい | 共有メモリと資源を共有 |

| L2 キャッシュ | GPU 全体 | 中 | 中 | 中 | 自動キャッシュ |

中核となる戦略は明確です。**グローバルメモリへのアクセスを最小化し、一度読んだデータを共有メモリやレジスタに載せて再利用する** ことです。タイリング(tiling)技法は、まさにこの原理を実装したものです。

コアレス(Coalesced)メモリアクセス

グローバルメモリは、一定サイズのメモリセグメント(例:32、64、128 バイト)を 1 トランザクションとして取得します。1 つのワープの 32 スレッドが **連続したアドレス** にアクセスすると、ハードウェアはこれを少数のトランザクションにまとめて処理します。これがコアレスです。

コアレスドアクセス (良い): ワープが連続アドレスにアクセス

Thread: T0 T1 T2 T3 ... T31

Addr : 0 4 8 12 ... 124

└──────────────────────┘

-> 1回のメモリトランザクションで処理

アンコアレスドアクセス (悪い): スレッドが散らばったアドレスにアクセス

Thread: T0 T1 T2 T3 ...

Addr : 0 512 1024 1536 ...

└─┐ └─┐ └─┐ └─┐

▼ ▼ ▼ ▼

-> スレッドごとに別トランザクション -> 帯域の浪費

コアレスを壊す代表的な原因は次のとおりです。

- 誤ったデータレイアウト(例:行優先/列優先の不一致)

- 大きなストライド(stride)を持つアクセス

- 構造体配列(AoS)の代わりに配列構造体(SoA)を使っていない場合

AoS vs SoA

AoS (Array of Structs): [x0 y0 z0][x1 y1 z1][x2 y2 z2]...

-> x だけ読みたくても y,z が割り込み非連続アクセス

SoA (Struct of Arrays): [x0 x1 x2 ...][y0 y1 y2 ...][z0 z1 z2 ...]

-> x 配列が連続 -> コアレスに有利

ワープスケジューリングと占有率(Occupancy)

GPU が速い本当の秘訣は、コアの個数ではなく **レイテンシ隠蔽(latency hiding)** にあります。あるワープがグローバルメモリを待って止まると、スケジューラはすぐに実行準備ができた別のワープへ切り替えます。十分に多くのワープが待機していれば、メモリ遅延がほとんど見えなくなります。

ワープスケジューリングのタイムライン (レイテンシ隠蔽)

時間 ───────────────────────────────────────▶

Warp0: [実行]──[メモリ待ち ........]──[実行]

Warp1: [実行]──[メモリ待ち ......]──[実行]

Warp2: [実行]──[メモリ待ち ....]──

Warp3: [実行]──[メモリ待ち]

-> どれか1つのワープが止まっても、他のワープが実行ユニットを埋める

-> 実行ユニットが遊ばなくなる (高い使用率)

占有率(Occupancy)

**占有率** は、1 つの SM で同時に活性化されたワープ数を、その SM が対応できる最大ワープ数で割った比率です。占有率が高いほど、遅延を隠すワープが多くなります。ただし、占有率が 100% でなければ最高性能が出ないわけではありません。適度な占有率でも十分に遅延を隠せますし、むしろスレッドあたりのレジスタを多めに使ってより速くなる場合もあります。

占有率を制限する要因を表にまとめると次のとおりです。

| 制限要因 | 説明 | 緩和方法 |

| --- | --- | --- |

| スレッドあたりレジスタ数 | レジスタが多いと同時ワープ数が減る | 変数を単純化、コンパイラで上限指定 |

| ブロックあたり共有メモリ | 共有メモリが多いと同時ブロック数が減る | タイルサイズの調整 |

| ブロックあたりスレッド数 | 少なすぎるとワープ不足、多すぎると資源不足 | 128、256 など適切に選択 |

| SM あたり最大ブロック/ワープ数 | ハードウェアの上限 | アーキテクチャ仕様を確認 |

占有率の直感

レジスタ/共有メモリを少なく使うカーネル:

SM資源 ████████████████ -> 多くのワープを同時収容 -> 高い占有率

レジスタ/共有メモリを多く使うカーネル:

SM資源 ████ -> 少しのワープしか収容できず -> 低い占有率

実務では、NVIDIA の Occupancy Calculator や `cudaOccupancyMaxPotentialBlockSize` API で適切なブロックサイズを見つけ、Nsight Compute で実際のボトルネックを確認するのが定石です。

テンソルコア(Tensor Core)

ディープラーニングの中核となる演算は行列積(GEMM)です。通常の CUDA コアがスカラ単位で積和を処理するのに対し、**テンソルコア** は小さな行列ブロックの積和(MMA, Matrix Multiply-Accumulate)を 1 命令で処理します。そのおかげで、スループットが数倍から数十倍まで向上します。

テンソルコアのMMA演算 (概念図)

A (M×K) B (K×N) C (M×N)

┌──────────┐ ┌──────────┐ ┌──────────┐

│ タイル │ × │ タイル │ + │ 累算器 │ ──▶ D = A·B + C

└──────────┘ └──────────┘ └──────────┘

(低精度) (低精度) (高精度で累算)

1つのワープが協力して1つのタイルMMAを実行

入力はFP16/BF16/INT8など、累算はFP32で精度を維持

テンソルコアの特徴は **混合精度(mixed precision)** です。入力は FP16、BF16、TF32、INT8、そして最新世代では FP8 のような低い精度で受けつつ、累算は FP32 のような高い精度で行うことで、精度の損失を抑えます。

| 入力精度 | 特徴 | 主な用途 |

| --- | --- | --- |

| FP16 | 半精度、広いサポート | 学習/推論 |

| BF16 | FP32 と同じ指数範囲 | 学習の安定性 |

| TF32 | FP32 入力を内部的に処理 | コード変更を最小化 |

| INT8 | 整数、量子化推論 | 低遅延推論 |

| FP8 | 最新世代でサポート | 大規模な学習/推論 |

多くの場合、テンソルコアを直接扱うよりも、cuBLAS、cuDNN、CUTLASS といった高水準ライブラリを通じて活用します。自分で書く場合は、`mma` 系のワープレベル API を使用します。

> Hopper、Blackwell などの最新世代は、テンソルコアのスループットやサポート精度(例:FP8)、非同期データ移動機能を世代ごとに拡張してきました。具体的なサポート精度や性能数値は世代やチップによって異なりますので、公式仕様をご確認ください。

ストリームと非同期実行

デフォルトでは、CUDA の処理は 1 つのデフォルトストリーム(default stream)上で逐次的に実行されます。しかし **ストリーム(stream)** を複数使うと、互いに独立した処理を重ねて(overlap)実行できます。代表的には、データ転送とカーネル演算を同時に進めて全体時間を短縮します。

単一ストリーム (直列): 転送と演算が順番に

H2D[==] Kernel[======] D2H[==]

─────────────────────────────────▶ 時間

複数ストリーム (オーバーラップ): チャンク単位で重ねて

S1: H2D[==] Kernel[====] D2H[==]

S2: H2D[==] Kernel[====] D2H[==]

S3: H2D[==] Kernel[====] D2H[==]

──────────────────────────────────────▶ 時間

-> 転送と演算が重なり総時間が短縮

オーバーラップを正しく活用するには、次の条件が必要です。

- ホストメモリが **ページ固定(pinned)** されている必要があります(非同期コピーのため)。

- `cudaMemcpyAsync` とカーネル実行に、明示的にストリームを指定する必要があります。

- イベント(`cudaEvent`)でストリーム間の依存関係を管理できます。

cudaStream_t stream;

cudaStreamCreate(&stream);

cudaMemcpyAsync(d_in, h_in, bytes, cudaMemcpyHostToDevice, stream);

myKernel<<<grid, block, 0, stream>>>(d_in, d_out, n);

cudaMemcpyAsync(h_out, d_out, bytes, cudaMemcpyDeviceToHost, stream);

cudaStreamSynchronize(stream);

cudaStreamDestroy(stream);

最近の世代では、繰り返される処理グラフをあらかじめキャプチャして実行オーバーヘッドを減らす CUDA Graphs や、非同期メモリコピー機能で転送と演算をより緻密に重ねる技法がよく使われています。

カーネル例:タイルド行列積

ここまでの概念を一つにまとめた例として、共有メモリを活用したタイルド行列積を見ていきます。核心となるアイデアは、グローバルメモリから小さなタイルを一度読んで共有メモリに載せたうえで、そのタイルを何度も再利用してグローバルアクセスを減らすことです。

まず、もっとも単純なベクトル加算から見ます。

__global__ void vectorAdd(const float* a, const float* b, float* c, int n) {

int idx = blockIdx.x * blockDim.x + threadIdx.x;

if (idx < n) {

c[idx] = a[idx] + b[idx];

}

// ホスト側の実行例

// int threads = 256;

// int blocks = (n + threads - 1) / threads;

// vectorAdd<<<blocks, threads>>>(d_a, d_b, d_c, n);

次に、共有メモリのタイリングを適用した行列積です。`C = A * B` を計算し、各ブロックは 16x16 のタイルを担当します。

#define TILE 16

__global__ void matMulTiled(const float* A, const float* B, float* C, int N) {

__shared__ float As[TILE][TILE];

__shared__ float Bs[TILE][TILE];

int row = blockIdx.y * TILE + threadIdx.y;

int col = blockIdx.x * TILE + threadIdx.x;

float acc = 0.0f;

// K 次元をタイル単位で巡回

for (int t = 0; t < (N + TILE - 1) / TILE; ++t) {

// 1つのタイルを共有メモリにロード (コアレスドアクセス)

int aCol = t * TILE + threadIdx.x;

int bRow = t * TILE + threadIdx.y;

As[threadIdx.y][threadIdx.x] = (row < N && aCol < N) ? A[row * N + aCol] : 0.0f;

Bs[threadIdx.y][threadIdx.x] = (bRow < N && col < N) ? B[bRow * N + col] : 0.0f;

__syncthreads(); // タイルのロード完了を保証

// 共有メモリから部分和を計算 (グローバルアクセスなし)

for (int k = 0; k < TILE; ++k) {

acc += As[threadIdx.y][k] * Bs[k][threadIdx.x];

}

__syncthreads(); // 次のタイルをロードする前に同期

}

if (row < N && col < N) {

C[row * N + col] = acc;

}

このカーネルが速い理由を図にまとめると次のとおりです。

タイリングでグローバルアクセスを再利用

素朴な版: C[i][j] の計算ごとに A行N個 + B列N個を毎回グローバルから読む

-> 同じデータを何度も重複ロード

タイルド版: TILE×TILE ブロックを共有メモリに1回ロード ->

ブロック内の全スレッドがそのタイルを再利用

-> グローバルトラフィックが約 TILE 倍減少

┌─────────┐ load ┌──────────────┐ reuse ┌──────────┐

│ Global │ ───────▶ │ Shared (タイル)│ ──────▶ │ 256スレッド│

│ (HBM) │ 1回 │ 16×16 │ 何度も │ 演算 │

└─────────┘ └──────────────┘ └──────────┘

`__syncthreads()` が 2 回登場する理由に注目してください。1 回目はすべてのスレッドがタイルのロードを終えるまで待ってデータ競合を防ぎ、2 回目は次のタイルを上書きする前に現在の計算が終わったことを保証します。この同期を抜かすと、結果が非決定的に壊れます。

よくある落とし穴とチェックリスト

最後に、実務で頻繁に出会う落とし穴をまとめます。

1. アンコアレスドメモリアクセス

もっとも一般的で致命的です。データレイアウトを SoA に変えるか、インデックスを連続アクセスになるよう調整してください。プロファイラのメモリ効率指標をまず確認する習慣が重要です。

2. ワープダイバージェンス

ワープ内部でデータに応じて分岐が分かれると直列化されます。可能であれば分岐条件をワープ境界に揃えるか、分岐の代わりに算術演算(predication)へ置き換えることを検討してください。

3. 共有メモリのバンクコンフリクト

共有メモリは複数のバンクに分かれており、1 つのワープのスレッドが同じバンクに同時にアクセスするとコンフリクトが起きて直列化されます。1 列分のパディングを追加して(例:`[TILE][TILE+1]`)コンフリクトを避ける技法がよく使われます。

4. 過度なレジスタ使用

レジスタを使いすぎると占有率が下がり、上限を超えると遅いローカルメモリへスピルします。カーネルを細かく分割したり、変数を減らしたりしてみてください。

5. 同期の漏れや誤用

`__syncthreads()` を分岐の中で一部のスレッドだけが呼ぶと、デッドロックが発生します。すべてのスレッドが同じ同期地点を通過するよう記述してください。

6. エラーチェックの省略

CUDA 呼び出しの戻り値や `cudaGetLastError()` を確認しないと、静かに誤った結果を得てしまいます。デバッグビルドでは呼び出しごとにエラーを検査するのが安全です。

| 症状 | 疑わしい原因 | 確認ツール |

| --- | --- | --- |

| 遅いのにコアは遊んでいるよう | メモリバウンド、アンコアレス | Nsight Compute |

| 占有率が低い | レジスタ/共有メモリの過多 | Occupancy Calculator |

| 分岐の多いカーネルが遅い | ワープダイバージェンス | Nsight Compute |

| 結果がたまに誤る | 同期漏れ、競合状態 | compute-sanitizer |

| 転送が演算と重ならない | pinned メモリ未使用 | Nsight Systems |

おわりに

CUDA をうまく使う核心は、結局のところ 2 つに集約されます。第 1 に、**論理的な実行モデル(グリッド・ブロック・ワープ)が実際のハードウェア(SM・メモリ階層)にどうマッピングされるか** を頭の中に描けることです。第 2 に、**データをメモリ階層にどう上手く配置して再利用するか** が性能のほとんどを決めます。

本記事で扱った図を思い浮かべながらコードを書けば、「なぜこのカーネルは遅いのか」を推測ではなく構造で説明できるようになります。その次は、Nsight のようなプロファイラで仮説を検証し、反復的に磨き上げる過程です。

最後に、アーキテクチャは世代ごとに急速に進化します。テンソルコアの精度サポート、非同期メモリ移動、SM あたりの資源といった具体的な数値は、チップやバージョンによって異なりますので、実際のチューニング時には常に該当アーキテクチャの公式ドキュメントをご確認ください。

参考資料

- [CUDA Toolkit Documentation](https://docs.nvidia.com/cuda/)

- [CUDA C++ Programming Guide](https://docs.nvidia.com/cuda/cuda-c-programming-guide/)

- [CUDA C++ Best Practices Guide](https://docs.nvidia.com/cuda/cuda-c-best-practices-guide/)

- [Parallel Thread Execution (PTX) ISA](https://docs.nvidia.com/cuda/parallel-thread-execution/)

- [cuBLAS Library Documentation](https://docs.nvidia.com/cuda/cublas/)

- [Matrix Multiplication Performance Guide](https://docs.nvidia.com/deeplearning/performance/dl-performance-matrix-multiplication/index.html)

- [NVIDIA Developer Blog](https://developer.nvidia.com/blog/)

- [Nsight Compute Documentation](https://docs.nvidia.com/nsight-compute/)

- [CUTLASS GitHub Repository](https://github.com/NVIDIA/cutlass)