[AI/機械学習] CUDA入門と実践：GPU並列計算の基礎からAI/画像生成の高速化まで

はじめに

過去、Pythonで生成AI系のローカル環境を何度も作ってきましたが、その際にエラー情報を含めて度々みる「CUDA」。

NVIDIAのGPUで並列計算を可能にするための開発プラットフォーム（＋API）という簡単な情報しか知らなかったので、改めて情報をまとめてもらいました。

TL;DR

CUDAはNVIDIA製GPUで汎用計算を爆速にするための開発基盤。
仕事は「スレッド」の大群で並列に実行し、ワープ（32スレッド）単位で動く。
コツはメモリ階層を意識し、分岐を減らし、転送を隠すこと。

1) CUDAとは？

CUDA（Compute Unified Device Architecture）は、NVIDIAが提供するGPU並列計算プラットフォーム＋API。
これにより、グラフィック用途に限らずAI/画像・動画/科学計算など**汎用計算（GPGPU）**にGPUの力を使える。

必要条件：NVIDIA GPU（GeForce/RTX/Quadro/データセンター系）
主な利用言語：C/C++（nvcc）、Python（PyTorch/CuPy/Numba等）
代表的ユースケース：Stable DiffusionやComfyUI、Real-ESRGAN、物理シミュ、行列演算等

2) 実行モデル：グリッド / ブロック / スレッドとワープ

GPUは超多数スレッドを同時に走らせるのが得意。CUDAでは次の階層で仕事を配る：

スレッド（thread）：最小単位の作業者
ブロック（block）：スレッドのグループ。共有メモリを使って協調
グリッド（grid）：ブロックの集合（1カーネル呼び出しあたり1グリッド）

さらにGPU内部では、32スレッド＝1ワープとして同一命令を同時実行（SIMT）。
→ 分岐が多いと（ifで分かれる）ワープが足並みを崩し、性能が落ちる（ワープダイバージェンス）。

3) メモリ階層（速い順）

種類	スコープ	特徴
レジスタ	スレッド	最速・超小容量
共有メモリ	ブロック	低遅延・手動管理。行列タイルなどで威力発揮
L2/キャッシュ	デバイス	自動キャッシュ
グローバルメモリ（VRAM）	デバイス	大容量・高遅延。コアレッシング（連続アクセス）で高速化
定数/テクスチャ	デバイス	読み取り最適化
ホストメモリ（CPU側）	ホスト	PCIe越えの転送が律速。Pinnedメモリで転送最適化

キモ：

コアレッシング（隣り合うスレッドが隣り合うアドレスを読む）→ 帯域を最大化
共有メモリで再利用→ グローバルメモリ往復を減らす
分岐削減→ ワープ効率を上げる

4) 典型フロー（CPU→GPU）

前提ミニ知識

ホスト = CPU側, デバイス = GPU側。GPUは“共同処理装置”で、処理は非同期で進むのが基本。
何もしないと**転送（PCIe）**が律速になりやすい。メモリの使い方と非同期化が速度のカギ。

1. メモリ確保（cudaMalloc / Pythonはライブラリ管理）

何をしてる？ GPUの**グローバルメモリ（VRAM）**上にバッファを作る。

主な選択肢

cudaMalloc(&d_ptr, bytes): ふつうのデバイスメモリ。
cudaMallocPitch: 2D配列向け。行頭アラインでコアレッシング（連続アクセス）しやすくなる。
Unified Memory（cudaMallocManaged）: CPU/GPUで同じポインタを共有。ページ移動はランタイム任せ。
メモリプール / 非同期確保: cudaMallocAsync（CUDA 11.2+）で確保・解放のオーバーヘッドを低減。

速くするコツ

使い回せるバッファは再利用（cudaMallocは重い）。
多数の小確保より大きめに1回確保して切り分け。
画像はpitch確保でアクセス効率UP。

落とし穴

まだ走っているカーネルが使っているバッファを解放しない（同期が必要）。

2. 転送（cudaMemcpy / Unified Memory なら自動）

何をしてる？
CPUメモリ⇔GPUメモリ間のデータ移動。ここが遅いと全体が詰まる。

主な選択肢

cudaMemcpy(HtoD/DtoH/DtoD)
: 同期転送（呼び出しがブロックされる）。
cudaMemcpyAsync(…, stream)
: 非同期転送（要：ホスト側がPinned（ページロック）メモリ）。
Pinnedメモリ
:cudaMallocHost/cudaHostAlloc。PyTorchはDataLoader(pin_memory=True)でOK。
Unified Memory
: 自動ページング。cudaMemPrefetchAsyncで明示プリフェッチすると速くて安定。
ゼロコピー（mapped host）
: cudaHostAllocMapped → GPUがホストRAMを直接読む（帯域制限大。特殊用途）。
GPU間
: cudaMemcpyPeer（NVLink/PCIeでP2P）。

速くするコツ

Pinned + cudaMemcpyAsync + streamsで計算と転送を重ね合わせる。
Unified Memoryは**cudaMemPrefetchAsync(ptr, size, device)**を癖にする（スラッシング回避）。

落とし穴

Pinnedじゃないのに MemcpyAsync しても重なり合わない。
Unified Memoryを大量・頻繁にチラ見するとページング地獄。

3. カーネル起動（kernel«<blocks, threads, sharedMem, stream»>(…)）

何をしてる？
GPU上で並列に走る**関数（カーネル）**を実行。
grid = blocks の集まり, block = threads の集まり。ワープ（32スレッド）単位でSIMT動作。

設計ポイント

インデックス計算：i = blockIdx.x*blockDim.x + threadIdx.x;
スレッド数：だいたい128〜1024/ブロックから実測で詰める（cudaOccupancyMaxPotentialBlockSizeも活用）。
共有メモリ：タイル化やデータ再利用に使うと激速になる。

速くするコツ

連続アドレスへ連続スレッド（コアレッシング）。
分岐を減らす（ワープダイバージェンス回避）。
sharedMemでグローバルメモリ往復を削減。

落とし穴

起動直後は非同期。エラーはcudaGetLastError() or 後段の同期で露呈。
インデックス範囲チェック忘れによるillegal memory access。

4. 同期（cudaDeviceSynchronize だけが答えじゃない）

何をしてる？
「この時点までにGPUの処理が終わっている」ことを保証する。

主な手段

全体同期：cudaDeviceSynchronize()（重い。最後の最後だけにしたい）
ストリーム同期：cudaStreamSynchronize(stream)（局所同期）
イベント：cudaEventRecord/ElapsedTime（計測・順序制御に◎）

速くするコツ

全体同期は最小限。イベント/ストリーム同期で必要部分だけ待つ。
cudaMemcpy（同期版）自体が暗黙の同期になることも理解しておく。

落とし穴

毎ステップDeviceSynchronizeすると非同期化の旨味が消える。

5. 結果転送 → ホスト

何をしてる？ GPUで計算した結果をCPU側へ戻す。2)の逆。

速くするコツ

ここも**cudaMemcpyAsync(..., stream) + Pinnedで計算と重ねる**。
Unified Memoryならそのまま読むだけでもOKだが、prefetch + 同期をセットで。

落とし穴

結果バッファがまだ計算中なのに転送開始して破綻（→ イベントで順序保証）。

6. 解放

何をしてる？ 確保したメモリを返す。GPU/CPUの両方。

API

cudaFree(d_ptr), cudaFreeHost(h_ptr)
非同期解放：cudaFreeAsync(d_ptr, stream)（対応環境のみ）

落とし穴

未完了の非同期操作が使っているメモリを解放しない（イベントで完了確認 → 解放が安全）。

まとめ図（最小同期 vs 高速版）

A. 最小の同期版（分かりやすいが遅い）

cudaMalloc(&dA, bytes);
cudaMemcpy(dA, hA, bytes, cudaMemcpyHostToDevice); // 同期
kernel<<<blocks, threads>>>(dA);
cudaDeviceSynchronize();                           // 全体待ち
cudaMemcpy(hA, dA, bytes, cudaMemcpyDeviceToHost); // 同期
cudaFree(dA);

B. 高速版（重ね合わせ：Pinned + Async + Streams + Events）

// 準備
cudaStream_t sH2D, sK, sD2H;
cudaStreamCreate(&sH2D); cudaStreamCreate(&sK); cudaStreamCreate(&sD2H);
cudaEvent_t ready; cudaEventCreate(&ready);
cudaMallocHost(&hA, bytes);                 // ← Pinned
cudaMalloc(&dA, bytes);

// 1) 転送（非同期）
cudaMemcpyAsync(dA, hA, bytes, cudaMemcpyHostToDevice, sH2D);

// 2) H2D完了を合図 → カーネルへ
cudaEventRecord(ready, sH2D);
cudaStreamWaitEvent(sK, ready, 0);
kernel<<<blocks, threads, 0, sK>>>(dA);

// 3) カーネル完了後にD2H（別ストリームで重ねる）
cudaEventRecord(ready, sK);
cudaStreamWaitEvent(sD2H, ready, 0);
cudaMemcpyAsync(hA, dA, bytes, cudaMemcpyDeviceToHost, sD2H);

// 必要になった時だけ待つ
cudaStreamSynchronize(sD2H);

cudaFree(dA); cudaFreeHost(hA);
cudaEventDestroy(ready);
cudaStreamDestroy(sH2D); cudaStreamDestroy(sK); cudaStreamDestroy(sD2H);

C. Unified Memory（管理メモリ + プリフェッチ）

cudaMallocManaged(&u, bytes);                // CPU/GPUで同じポインタ
// CPUで初期化...
cudaMemPrefetchAsync(u, bytes, device, 0);   // 実行前にGPUへ寄せる
kernel<<<blocks, threads>>>(u);
cudaDeviceSynchronize();                     // or イベント/ストリーム同期
cudaMemPrefetchAsync(u, bytes, cudaCpuDeviceId, 0); // 使う前にCPUへ寄せる
// CPUで結果参照...
cudaFree(u);

Python勢の「ここだけは」

PyTorch
- DataLoader(pin_memory=True) + to(device, non_blocking=True)
- 前処理とGPU計算を別スレッド/ワーカーで重ねる（num_workers）。
- 計測はtorch.cuda.Event。
CuPy
- 配列はcp.asarray/cp.emptyでGPU上に直置き。ホスト⇔デバイス往復を減らす。
Numba CUDA
- cuda.to_device(a) / device_array_like(a) + copy_to_host()
- グリッド次元はthreads_per_block=256あたりから実測で調整。

5) 最小サンプル

C++（ベクトル和）

// nvcc vector_add.cu -o vector_add
#include <cuda_runtime.h>
#include <cstdio>

__global__ void vectorAdd(const float* A, const float* B, float* C, int N) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < N) C[i] = A[i] + B[i];
}

int main() {
    const int N = 1<<20; // 約100万
    size_t bytes = N * sizeof(float);
    float *hA = (float*)malloc(bytes), *hB = (float*)malloc(bytes), *hC = (float*)malloc(bytes);
    for (int i=0;i<N;i++){ hA[i]=i*1.0f; hB[i]=2.0f; }

    float *dA,*dB,*dC;
    cudaMalloc(&dA, bytes); cudaMalloc(&dB, bytes); cudaMalloc(&dC, bytes);
    cudaMemcpy(dA, hA, bytes, cudaMemcpyHostToDevice);
    cudaMemcpy(dB, hB, bytes, cudaMemcpyHostToDevice);

    int threads = 256;
    int blocks  = (N + threads - 1)/threads;
    vectorAdd<<<blocks, threads>>>(dA,dB,dC,N);
    cudaDeviceSynchronize();

    cudaMemcpy(hC, dC, bytes, cudaMemcpyDeviceToHost);
    printf("C[0]=%.1f  C[N-1]=%.1f\n", hC[0], hC[N-1]);

    cudaFree(dA); cudaFree(dB); cudaFree(dC);
    free(hA); free(hB); free(hC);
}

Python（CuPy：NumPy互換APIでGPU）

import cupy as cp
N = 1<<20
a = cp.arange(N, dtype=cp.float32)
b = cp.full(N, 2, dtype=cp.float32)
c = a + b
print(c[0].item(), c[-1].item())

Python（Numba CUDA：自作カーネル）

import numpy as np
from numba import cuda

@cuda.jit
def vec_add(a,b,c):
    i = cuda.grid(1)
    if i < a.size:
        c[i] = a[i] + b[i]

N = 1<<20
a = np.arange(N, dtype=np.float32)
b = np.full(N, 2, dtype=np.float32)
c = np.empty_like(a)

d_a = cuda.to_device(a); d_b = cuda.to_device(b); d_c = cuda.device_array_like(a)
threads = 256; blocks = (N + threads - 1)//threads
vec_add[blocks, threads](d_a, d_b, d_c)
d_c.copy_to_host(c)
print(c[0], c[-1])

PyTorch（GPU確認＆移動）

import torch
device = torch.device("cuda" if torch.cuda.is_available() else "cpu")
x = torch.randn(8, 3, 224, 224, device=device)
print(device, x.shape)

6) 速度の正体：3つの柱

並列度：スレッドを大量に投げる（Occupancyを稼ぐ）
メモリ効率：コアレッシング＋共有メモリ＋再利用
非同期化：ストリームで計算と転送を重ねる（HtoD/DtHの隠蔽、Pinnedメモリ）

例：大きいバッチを複数チャンクに分け、Stream Aで計算中にStream Bで次チャンクを転送。

7) Tensor Coresと混合精度

近年のRTX系はTensor Coresを搭載。FP16/BF16/INT8など低精度で行列演算を爆速化。

推論：自動で半精度に最適化されるフレームワークが多い
学習：PyTorchのautocast + GradScalerで簡単導入（数値安定も確保）

# 学習時の例
scaler = torch.cuda.amp.GradScaler()
for x, y in loader:
    optimizer.zero_grad()
    with torch.cuda.amp.autocast():
        out = model(x.to(device))
        loss = criterion(out, y.to(device))
    scaler.scale(loss).backward()
    scaler.step(optimizer)
    scaler.update()

8) ライブラリエコシステム（使えるものは使う！）

ライブラリ	用途
cuBLAS	行列/ベクトル（BLAS）
cuDNN	深層学習（畳み込み/活性化/プーリング）
NCCL	複数GPU/複数ノード通信
CUTLASS/CUB/Thrust	自作カーネル補助（テンプレ/アルゴリズム）
TensorRT	推論最適化（量子化/融合/最適実行計画）
nvJPEG / NVENC/NVDEC	画像/動画のエンコ・デコ

原則：「まずライブラリ → 足りなければ自作」。性能・保守の両面で得。

9) インストールと互換性の要点

GPUドライバとCUDA Toolkitの互換が肝（PyTorch/各ライブラリが対応するCUDA版に揃える）
チェック：
- nvidia-smi（ドライバ＆GPU認識）
- nvcc --version（Toolchainの有無）
サンプル：deviceQuery（CUDA Samples）で環境確認が安心

迷ったら：フレームワークの公式「対応CUDA版」を優先。Toolkitは無理に最新へ上げない。

10) よくあるエラーと対処

Out of Memory：解像度/バッチ/モデルサイズを下げる、混合精度、不要テンソル del 、torch.cuda.empty_cache() （断片化の根本解決ではない）
driver/runtime mismatch：「ドライバが古くてランタイムが動かない」→ ドライバ更新
illegal memory access：境界外アクセス、未同期のまま参照等 → インデックス/同期を見直す
invalid device ordinal：GPU番号の指定ミス
nvccが見つからない：PATH設定

11) パフォーマンス・チートシート

配列アクセスは連続に（コアレッシング）
可能なら共有メモリでタイル化（行列演算など）
分岐は減らしてワープの足並みを揃える
大きめのブロック（例えば256/512）でGPUを埋める（要実測）
ストリーム + Pinnedメモリで転送と計算の重ね合わせ
計測はNsight Systems / Nsight Compute / cudaEventで可視化

12) 画像生成AI（Stable Diffusion/ComfyUI等）での実践Tips

VRAMは正義：解像度・バッチ・モデルサイズで一気に消費
半精度/メモリ節約：FP16/BF16、重ね合わせ、（可能なら）注意メモリ最適化
順序：まずバッチ→ 解像度の順に攻める
I/O律速も侮らない：前処理・画像読込・保存を並列/バッファリング

13) まとめ

CUDAは**「並列度 × メモリ効率 × 非同期」**の三位一体。
困ったらライブラリを使い、可視化ツールで実測し、設計を少しずつ正す。
画像生成や動画処理でも、ちょっとした設定（精度/バッチ/転送の隠蔽）で体感が激変する。

近道は「最適化の勘」を育てること。測って→直して→また測るのループが最強です。

[AI/機械学習] CUDA入門と実践：GPU並列計算の基礎からAI/画像生成の高速化まで

はじめに

TL;DR

1) CUDAとは？

2) 実行モデル：グリッド / ブロック / スレッドとワープ

3) メモリ階層（速い順）

4) 典型フロー（CPU→GPU）

前提ミニ知識

1. メモリ確保（cudaMalloc / Pythonはライブラリ管理）

2. 転送（cudaMemcpy / Unified Memory なら自動）

3. カーネル起動（kernel«<blocks, threads, sharedMem, stream»>(…)）

4. 同期（cudaDeviceSynchronize だけが答えじゃない）

5. 結果転送 → ホスト

6. 解放

まとめ図（最小同期 vs 高速版）

Python勢の「ここだけは」

5) 最小サンプル

C++（ベクトル和）

Python（CuPy：NumPy互換APIでGPU）

Python（Numba CUDA：自作カーネル）

PyTorch（GPU確認＆移動）

6) 速度の正体：3つの柱

7) Tensor Coresと混合精度

8) ライブラリエコシステム（使えるものは使う！）

9) インストールと互換性の要点

10) よくあるエラーと対処

11) パフォーマンス・チートシート

12) 画像生成AI（Stable Diffusion/ComfyUI等）での実践Tips

13) まとめ

関連動画

💬 コメント

はじめに#

TL;DR#

1) CUDAとは？#

2) 実行モデル：グリッド / ブロック / スレッド と ワープ#

3) メモリ階層（速い順）#

4) 典型フロー（CPU→GPU）#

前提ミニ知識#

1. メモリ確保（cudaMalloc / Pythonはライブラリ管理）#

2. 転送（cudaMemcpy / Unified Memory なら自動）#

3. カーネル起動（kernel«<blocks, threads, sharedMem, stream»>(…)）#

4. 同期（cudaDeviceSynchronize だけが答えじゃない）#

5. 結果転送 → ホスト#

6. 解放#

まとめ図（最小同期 vs 高速版）#

Python勢の「ここだけは」#

5) 最小サンプル#

C++（ベクトル和）#

Python（CuPy：NumPy互換APIでGPU）#

Python（Numba CUDA：自作カーネル）#

PyTorch（GPU確認＆移動）#

6) 速度の正体：3つの柱#

7) Tensor Coresと混合精度#

8) ライブラリエコシステム（使えるものは使う！）#

9) インストールと互換性の要点#

10) よくあるエラーと対処#

11) パフォーマンス・チートシート#

12) 画像生成AI（Stable Diffusion/ComfyUI等）での実践Tips#

13) まとめ#

関連動画#

💬 コメント

はじめに

TL;DR

1) CUDAとは？

2) 実行モデル：グリッド / ブロック / スレッドとワープ

3) メモリ階層（速い順）

4) 典型フロー（CPU→GPU）

前提ミニ知識

1. メモリ確保（cudaMalloc / Pythonはライブラリ管理）

2. 転送（cudaMemcpy / Unified Memory なら自動）

3. カーネル起動（kernel«<blocks, threads, sharedMem, stream»>(…)）

4. 同期（cudaDeviceSynchronize だけが答えじゃない）

5. 結果転送 → ホスト

6. 解放

まとめ図（最小同期 vs 高速版）

Python勢の「ここだけは」

5) 最小サンプル

C++（ベクトル和）

Python（CuPy：NumPy互換APIでGPU）

Python（Numba CUDA：自作カーネル）

PyTorch（GPU確認＆移動）

6) 速度の正体：3つの柱

7) Tensor Coresと混合精度

8) ライブラリエコシステム（使えるものは使う！）

9) インストールと互換性の要点

10) よくあるエラーと対処

11) パフォーマンス・チートシート

12) 画像生成AI（Stable Diffusion/ComfyUI等）での実践Tips

13) まとめ

関連動画