この記事の対象読者
- PythonでGPUプログラミングを始めたい方
- CUDAの基礎は知っているが、コンパイルの仕組みを深く理解したい方
- 「なぜ自分のGPUで動かないんだ」というエラーに悩まされた経験がある方
この記事で得られること
- cubinとは何かを理解し、人に説明できるようになる
- PTX、fatbin、SASSとの関係を正確に把握できる
- nvccのコンパイルオプションを適切に設定できるようになる
- GPU互換性問題をトラブルシューティングできるようになる
この記事で扱わないこと
- CUDAカーネルの書き方（入門レベル）
- GPUアーキテクチャの詳細な設計論
- NVIDIA Driver APIの詳細な使い方
1. cubinとの出会い
「え、このGPUで動かないの...？」
PyTorchでモデルを動かそうとしたとき、こんなエラーに遭遇したことはないだろうか。
CUDA error: no kernel image is available for execution on the device
私は何度もこのエラーに泣かされた。同じCUDAプログラムなのに、あるGPUでは動いて別のGPUでは動かない。その謎を解く鍵が「cubin」だった。
cubinは、NVIDIA GPUで実行される実際のバイナリコードのこと。人間でいえば「母国語」のようなもので、GPUはこのcubinだけを直接理解できる。
最初にこの仕組みを理解したとき、「なるほど、だからGPUごとにビルドが必要なのか」と腑に落ちた。今回は、その感動を共有したい。
ここまでで、cubinがGPUの「ネイティブ言語」だということがなんとなくイメージできただろうか。次は、この技術が生まれた背景を見ていこう。
2. 前提知識の確認
本題に入る前に、この記事で使う用語を整理しておく。
2.1 CUDA（Compute Unified Device Architecture）とは
NVIDIAが開発したGPU上で汎用計算を行うためのプラットフォーム。C/C++に拡張を加えた言語でGPUプログラムを記述できる。深層学習フレームワーク（PyTorch、TensorFlowなど）の基盤技術でもある。
2.2 nvcc（NVIDIA CUDA Compiler）とは
CUDAプログラムをコンパイルするためのツールチェーン。ホスト（CPU）コードとデバイス（GPU）コードの両方を処理し、実行可能なバイナリを生成する。
2.3 Compute Capability（コンピュート能力）とは
各NVIDIA GPUが持つ機能セットを示す番号。例えば、RTX 3090はCompute Capability 8.6、RTX 4090は8.9を持つ。この番号がGPUの「世代」と「機能」を表す。
2.4 カーネル（Kernel）とは
GPUで実行される関数のこと。__global__キーワードで定義され、数千〜数百万のスレッドで並列実行される。
これらの用語が押さえられたら、次に進もう。
3. cubinが生まれた背景
3.1 GPUの多様化という課題
NVIDIAは毎年のように新しいGPUアーキテクチャをリリースしている。
| 世代 | アーキテクチャ | Compute Capability | 代表的なGPU |
|---|---|---|---|
| 2018 | Turing | 7.5 | RTX 2080 |
| 2020 | Ampere | 8.0, 8.6 | RTX 3090, A100 |
| 2022 | Ada Lovelace | 8.9 | RTX 4090 |
| 2022 | Hopper | 9.0 | H100 |
| 2024 | Blackwell | 10.0 | B200 |
問題は、各世代で命令セットが異なること。同じCUDAコードでも、実行するGPUに合わせたバイナリが必要になる。
3.2 NVIDIAの解決策：2段階コンパイル
NVIDIAはこの問題を、コンパイルを2段階に分けることで解決した。
CUDA C/C++ → PTX（仮想的な中間コード） → cubin（実際のGPUバイナリ）
PTX（Parallel Thread Execution）は「仮想的なGPU」向けのアセンブリ言語。これを間に挟むことで、将来のGPUとの互換性を確保できる仕組みになっている。
背景がわかったところで、抽象的な概念から順に、具体的な仕組みを見ていこう。
4. cubinの基本概念
4.1 cubinの正体：ELF形式のバイナリ
cubinは「CUDA Binary」の略で、実際にはELF（Executable and Linkable Format）形式のファイル。Linux の実行ファイルと同じフォーマットだが、中身はGPU用のネイティブ命令（SASS: Shader ASSembly）が詰まっている。
┌──────────────────────────────────────┐
│ cubin ファイル │
├──────────────────────────────────────┤
│ ELF ヘッダー │
├──────────────────────────────────────┤
│ .text セクション (SASS命令) │
├──────────────────────────────────────┤
│ .nv.info (CUDA固有のメタデータ) │
├──────────────────────────────────────┤
│ シンボルテーブル │
├──────────────────────────────────────┤
│ リロケーション情報 │
└──────────────────────────────────────┘
公式ドキュメントより:
A CUDA binary (also referred to as cubin) file is an ELF-formatted file which consists of CUDA executable code sections as well as other sections containing symbols, relocators, debug info, etc.
「CUDAバイナリ（cubinとも呼ばれる）ファイルは、ELF形式のファイルで、CUDA実行可能コードセクションと、シンボル、リロケータ、デバッグ情報などを含むその他のセクションで構成される。」
出典: NVIDIA CUDA Binary Utilities Documentation
4.2 PTXとcubinの関係
PTXは「仮想GPU」向けの中間表現、cubinは「実際のGPU」向けのネイティブコード。この関係を理解することが、GPU互換性を理解する鍵になる。
| 特性 | PTX | cubin |
|---|---|---|
| 形式 | テキスト（アセンブリ言語） | バイナリ（ELF形式） |
| ターゲット | 仮想アーキテクチャ（compute_XX） | 実アーキテクチャ（sm_XX） |
| 互換性 | 将来のGPUとも互換 | 同一メジャーバージョン内のみ |
| 実行時処理 | JITコンパイル必要 | そのまま実行可能 |
| 性能 | JITのオーバーヘッドあり | 最適化済みで高速 |
公式ドキュメントより:
CUDA applications and libraries are usually written in a higher-level language like C++. That higher-level language is compiled to PTX, and then the PTX is compiled into real binary for a physical GPU, called a CUDA binary, or cubin for short.
「CUDAアプリケーションやライブラリは通常、C++のような高水準言語で書かれる。その高水準言語はPTXにコンパイルされ、その後PTXは物理GPUのための実際のバイナリ、すなわちCUDAバイナリ（略してcubin）にコンパイルされる。」
出典: NVIDIA CUDA Programming Guide
4.3 fatbinという「コンテナ」
実際のアプリケーションでは、複数のcubinとPTXを「fatbin」というコンテナにまとめて配布する。
┌─────────────────────────────────────────────────┐
│ fatbin │
├─────────────────────────────────────────────────┤
│ cubin (sm_70) - Volta/Turing用 │
├─────────────────────────────────────────────────┤
│ cubin (sm_80) - Ampere用 │
├─────────────────────────────────────────────────┤
│ cubin (sm_89) - Ada Lovelace用 │
├─────────────────────────────────────────────────┤
│ PTX (compute_70) - 将来のGPU用（JIT対応） │
└─────────────────────────────────────────────────┘
実行時、CUDAランタイムは以下の順序で適切なコードを選択する。
- 実行中のGPUに対応するcubinがあれば、それを使用
- 互換性のあるcubin（同一メジャーバージョン、同等以下のマイナーバージョン）があれば使用
- cubinがなければ、PTXをJITコンパイルしてcubinを生成
基本概念が理解できたところで、これらの抽象的な概念を具体的なコードで実装していこう。
5. 実際に使ってみよう
5.1 環境構築
# CUDA Toolkitのインストール確認
nvcc --version
# 必要なツールの確認
which cuobjdump # cubinを調べるツール
which nvdisasm # cubinを逆アセンブルするツール
5.2 設定ファイルの準備
以下の3種類のMakefileテンプレートを用意した。用途に応じて選択してほしい。
開発環境用（Makefile.dev）
# Makefile.dev - 開発環境用（このままコピーして使える）
# デバッグ情報付き、単一アーキテクチャ向け
NVCC = nvcc
CUDA_ARCH = sm_86 # 開発マシンのGPUに合わせて変更
# デバッグフラグ
NVCC_FLAGS = -g -G -lineinfo
NVCC_FLAGS += -arch=$(CUDA_ARCH)
# ソースファイル
SOURCES = kernel.cu
TARGET = app_dev
all: $(TARGET)
$(TARGET): $(SOURCES)
$(NVCC) $(NVCC_FLAGS) -o $@ $^
# cubinのみを生成（デバッグ用）
cubin: $(SOURCES)
$(NVCC) $(NVCC_FLAGS) --cubin -o kernel.cubin $^
# PTXを確認（デバッグ用）
ptx: $(SOURCES)
$(NVCC) $(NVCC_FLAGS) --ptx -o kernel.ptx $^
# cubinの中身を確認
inspect: cubin
cuobjdump -sass kernel.cubin
clean:
rm -f $(TARGET) *.cubin *.ptx
.PHONY: all cubin ptx inspect clean
本番環境用（Makefile.prod）
# Makefile.prod - 本番環境用（このままコピーして使える）
# 複数アーキテクチャ対応、最適化済み
NVCC = nvcc
# 最適化フラグ
NVCC_FLAGS = -O3 -use_fast_math
# 複数GPUアーキテクチャをサポート（主要なものをカバー）
GENCODE_FLAGS = \
-gencode arch=compute_70,code=sm_70 \
-gencode arch=compute_75,code=sm_75 \
-gencode arch=compute_80,code=sm_80 \
-gencode arch=compute_86,code=sm_86 \
-gencode arch=compute_89,code=sm_89 \
-gencode arch=compute_90,code=sm_90 \
-gencode arch=compute_90,code=compute_90 # PTXも含める（将来のGPU用）
SOURCES = kernel.cu
TARGET = app_prod
all: $(TARGET)
$(TARGET): $(SOURCES)
$(NVCC) $(NVCC_FLAGS) $(GENCODE_FLAGS) -o $@ $^
# fatbinの中身を確認
inspect:
cuobjdump -lelf $(TARGET)
# fatbinのサイズを確認
size:
@echo "Binary size breakdown:"
@cuobjdump -lelf $(TARGET) | grep -E "(sm_|compute_)"
@ls -lh $(TARGET)
clean:
rm -f $(TARGET)
.PHONY: all inspect size clean
CI/テスト環境用（Makefile.test）
# Makefile.test - CI/テスト環境用（このままコピーして使える）
# JITコンパイル前提、最小バイナリサイズ
NVCC = nvcc
# PTXのみを含める（JITコンパイルで全GPUに対応）
# バイナリサイズを最小化
NVCC_FLAGS = -O2
GENCODE_FLAGS = -gencode arch=compute_70,code=compute_70
SOURCES = kernel.cu
TARGET = app_test
all: $(TARGET)
$(TARGET): $(SOURCES)
$(NVCC) $(NVCC_FLAGS) $(GENCODE_FLAGS) -o $@ $^
# JITコンパイルを強制してテスト
test-jit:
CUDA_FORCE_PTX_JIT=1 ./$(TARGET)
# バイナリに含まれるコードを確認
inspect:
cuobjdump -lelf $(TARGET)
clean:
rm -f $(TARGET)
.PHONY: all test-jit inspect clean
5.3 基本的な使い方
/**
* kernel.cu - cubinの動作確認用サンプル
*
* コンパイル方法:
* 開発環境: make -f Makefile.dev
* 本番環境: make -f Makefile.prod
* テスト: make -f Makefile.test
*/
#include <stdio.h>
#include <cuda_runtime.h>
// エラーチェックマクロ
#define CUDA_CHECK(call) \
do { \
cudaError_t err = call; \
if (err != cudaSuccess) { \
fprintf(stderr, "CUDA Error at %s:%d - %s\n", \
__FILE__, __LINE__, cudaGetErrorString(err)); \
exit(EXIT_FAILURE); \
} \
} while(0)
// シンプルなベクトル加算カーネル
__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];
}
}
// GPU情報を表示
void printGPUInfo() {
int deviceCount;
CUDA_CHECK(cudaGetDeviceCount(&deviceCount));
for (int i = 0; i < deviceCount; i++) {
cudaDeviceProp prop;
CUDA_CHECK(cudaGetDeviceProperties(&prop, i));
printf("GPU %d: %s\n", i, prop.name);
printf(" Compute Capability: %d.%d\n", prop.major, prop.minor);
printf(" SM Count: %d\n", prop.multiProcessorCount);
printf(" Global Memory: %.2f GB\n",
prop.totalGlobalMem / (1024.0 * 1024.0 * 1024.0));
}
}
int main() {
const int N = 1024;
const size_t size = N * sizeof(float);
// GPU情報を表示
printGPUInfo();
printf("\n");
// ホストメモリ確保
float *h_a = (float*)malloc(size);
float *h_b = (float*)malloc(size);
float *h_c = (float*)malloc(size);
// 初期化
for (int i = 0; i < N; i++) {
h_a[i] = i * 1.0f;
h_b[i] = i * 2.0f;
}
// デバイスメモリ確保
float *d_a, *d_b, *d_c;
CUDA_CHECK(cudaMalloc(&d_a, size));
CUDA_CHECK(cudaMalloc(&d_b, size));
CUDA_CHECK(cudaMalloc(&d_c, size));
// ホスト→デバイス転送
CUDA_CHECK(cudaMemcpy(d_a, h_a, size, cudaMemcpyHostToDevice));
CUDA_CHECK(cudaMemcpy(d_b, h_b, size, cudaMemcpyHostToDevice));
// カーネル実行
int threadsPerBlock = 256;
int blocksPerGrid = (N + threadsPerBlock - 1) / threadsPerBlock;
printf("Launching kernel: %d blocks x %d threads\n",
blocksPerGrid, threadsPerBlock);
vectorAdd<<<blocksPerGrid, threadsPerBlock>>>(d_a, d_b, d_c, N);
CUDA_CHECK(cudaGetLastError());
CUDA_CHECK(cudaDeviceSynchronize());
// デバイス→ホスト転送
CUDA_CHECK(cudaMemcpy(h_c, d_c, size, cudaMemcpyDeviceToHost));
// 結果確認
printf("Result check: c[0]=%.1f, c[100]=%.1f, c[1023]=%.1f\n",
h_c[0], h_c[100], h_c[1023]);
printf("Expected: c[0]=0.0, c[100]=300.0, c[1023]=3069.0\n");
// クリーンアップ
CUDA_CHECK(cudaFree(d_a));
CUDA_CHECK(cudaFree(d_b));
CUDA_CHECK(cudaFree(d_c));
free(h_a);
free(h_b);
free(h_c);
printf("\nSuccess!\n");
return 0;
}
5.4 実行結果
上記のコードを実行すると、以下のような出力が得られる。
$ make -f Makefile.dev && ./app_dev
nvcc -g -G -lineinfo -arch=sm_86 -o app_dev kernel.cu
GPU 0: NVIDIA GeForce RTX 3090
Compute Capability: 8.6
SM Count: 82
Global Memory: 24.00 GB
Launching kernel: 4 blocks x 256 threads
Result check: c[0]=0.0, c[100]=300.0, c[1023]=3069.0
Expected: c[0]=0.0, c[100]=300.0, c[1023]=3069.0
Success!
5.5 cubinの中身を確認する
# cubinを生成
$ nvcc --cubin -arch=sm_86 -o kernel.cubin kernel.cu
# cubinの基本情報を確認
$ cuobjdump -elf kernel.cubin
Fatbin elf code:
================
arch = sm_86
code version = [1,7]
host = linux
compile_size = 64bit
# SASS（ネイティブ命令）を確認
$ nvdisasm kernel.cubin
.headerflags @"EF_CUDA_SM86 EF_CUDA_PTX_SM(EF_CUDA_SM86)"
.elftype @"ET_EXEC"
.global _Z9vectorAddPKfS0_Pfi
.text._Z9vectorAddPKfS0_Pfi:
/*0000*/ MOV R1, c[0x0][0x28] ;
/*0010*/ S2R R0, SR_CTAID.X ;
...
5.6 よくあるエラーと対処法
| エラー | 原因 | 対処法 |
|---|---|---|
no kernel image is available for execution on the device |
実行GPUに対応するcubinもPTXもない |
-gencodeオプションで対象GPUを追加、またはPTXを含める |
CUDA driver version is insufficient for CUDA runtime version |
ドライバが古い | NVIDIAドライバを更新 |
invalid device function |
カーネルがGPUに対応していない | Compute Capabilityを確認し、適切な-archを指定 |
PTX JIT compilation failed |
PTXがGPUと互換性がない | より古いcompute_XXでPTXを生成 |
out of memory |
GPUメモリ不足 | バッチサイズを小さくする、またはcudaMemGetInfoで残量確認 |
基本的な使い方をマスターしたので、次は応用例を見ていこう。
6. ユースケース別ガイド
6.1 ユースケース1: 開発中のデバッグ
想定読者: CUDAプログラムを開発中で、特定のGPUでのみ動作確認したい方
推奨構成: 単一アーキテクチャのcubinのみ（ビルド時間短縮）
サンプルコード:
#!/bin/bash
# debug_build.sh - デバッグ用ビルドスクリプト
# 開発マシンのGPUアーキテクチャを自動検出
GPU_ARCH=$(nvidia-smi --query-gpu=compute_cap --format=csv,noheader | head -1 | tr -d '.')
echo "Detected GPU architecture: sm_${GPU_ARCH}"
# デバッグ情報付きでビルド
nvcc -g -G -lineinfo \
-arch=sm_${GPU_ARCH} \
-o app_debug \
kernel.cu
# cubinの内容を確認
echo ""
echo "=== Generated cubin info ==="
cuobjdump -elf app_debug
# 実行
echo ""
echo "=== Running ==="
./app_debug
6.2 ユースケース2: 複数GPU環境での本番デプロイ
想定読者: 様々なGPUを持つユーザーに配布するアプリケーションを作る方
推奨構成: 複数cubin + PTX（互換性最大化）
サンプルコード:
#!/bin/bash
# production_build.sh - 本番用ビルドスクリプト
# 主要なGPUアーキテクチャをすべてカバー
# Volta(7.0), Turing(7.5), Ampere(8.0, 8.6), Ada(8.9), Hopper(9.0)
# 最後にPTXを含めて将来のGPU互換性を確保
nvcc -O3 -use_fast_math \
-gencode arch=compute_70,code=sm_70 \
-gencode arch=compute_75,code=sm_75 \
-gencode arch=compute_80,code=sm_80 \
-gencode arch=compute_86,code=sm_86 \
-gencode arch=compute_89,code=sm_89 \
-gencode arch=compute_90,code=sm_90 \
-gencode arch=compute_90,code=compute_90 \
-o app_prod \
kernel.cu
echo "=== Binary size ==="
ls -lh app_prod
echo ""
echo "=== Included architectures ==="
cuobjdump -lelf app_prod | grep -E "(arch|code)"
echo ""
echo "=== PTX verification ==="
cuobjdump -ptx app_prod | head -20
6.3 ユースケース3: Dockerコンテナでの配布
想定読者: Dockerイメージを小さく保ちたいが、GPU互換性も確保したい方
推奨構成: PTXのみ（JIT前提、最小サイズ）
サンプルコード:
# Dockerfile - 最小サイズのCUDAアプリケーション
FROM nvidia/cuda:12.3-devel-ubuntu22.04 AS builder
WORKDIR /app
COPY kernel.cu .
# PTXのみでビルド（最小サイズ）
# compute_70は十分古いので、ほぼすべてのGPUで動作
RUN nvcc -O3 \
-gencode arch=compute_70,code=compute_70 \
-o app \
kernel.cu
# ランタイムイメージ（より小さい）
FROM nvidia/cuda:12.3-runtime-ubuntu22.04
WORKDIR /app
COPY --from=builder /app/app .
# JITコンパイルのキャッシュディレクトリ
ENV CUDA_CACHE_PATH=/app/.cache
CMD ["./app"]
#!/bin/bash
# docker_build.sh - Dockerイメージのビルドと実行
# ビルド
docker build -t cuda-app:minimal .
# イメージサイズ確認
echo "=== Image size ==="
docker images cuda-app:minimal
# 実行（JITコンパイルが走る）
echo ""
echo "=== Running with JIT compilation ==="
docker run --gpus all cuda-app:minimal
ユースケースが把握できたところで、この記事を読んだ後の学習パスを確認しよう。
7. 学習ロードマップ
この記事を読んだ後、次のステップとして以下をおすすめする。
初級者向け（まずはここから）
- CUDA C++ Programming Guide - The CUDA platform - 公式のCUDAプラットフォーム解説
-
NVIDIA CUDA Samples - 公式サンプルコード集。
0_Introductionから始めよう
中級者向け（実践に進む）
- CUDA Binary Utilities - cuobjdump、nvdisasmの詳細な使い方
- nvcc Documentation - コンパイラオプションの完全リファレンス
上級者向け（さらに深く）
- PTX ISA Reference - PTX命令セットの完全仕様
- NVIDIA GPU Architecture Compatibility Guide - 各アーキテクチャ間の互換性詳細
8. まとめ
この記事では、cubinについて以下を解説した。
- cubinの正体: ELF形式のGPUネイティブバイナリで、特定のCompute Capability向けにコンパイルされる
- PTX・fatbinとの関係: PTX（仮想GPU向け）→ cubin（実GPU向け）の2段階コンパイル、fatbinは複数アーキテクチャをまとめるコンテナ
- 互換性の仕組み: 同一メジャーバージョン内で下位互換性あり、PTXを含めれば将来のGPUでもJIT対応可能
私がハマったポイント
最初、「PTXだけ含めておけば全GPU対応でしょ」と思っていた。しかし、本番環境でJITコンパイルが走ると、アプリの初回起動が数十秒遅くなることがあった。ユーザーから「フリーズした」と報告を受けて初めて、cubinも含める重要性に気づいた。
今では、主要なアーキテクチャのcubinを含めつつ、PTXも含めるのがベストプラクティスだと確信している。
私の見解
正直なところ、NVIDIAの「Compute Capability」という命名は紛らわしいと思っている。「GPU世代+機能番号」と言った方がわかりやすい。
また、毎年のようにアーキテクチャが変わり、古いcubinが使えなくなるのは配布側にとって悩ましい。PTXのJITコンパイルがもっと高速になれば、もはやcubinを含める必要がなくなるかもしれない。
とはいえ、この2段階コンパイルという設計は、10年以上にわたってGPUの急速な進化を支えてきた。互換性と性能のバランスを取る優れたアーキテクチャだと評価している。
参考文献
- NVIDIA CUDA Binary Utilities Documentation - cubinに関する公式ドキュメント
- NVIDIA CUDA Programming Guide - The CUDA platform - CUDAプラットフォームの概要
- NVIDIA CUDA Compiler Driver NVCC - nvccの公式ドキュメント
- Understanding PTX, the Assembly Language of CUDA GPU Computing - NVIDIAブログによるPTX解説
- NVIDIA Ampere GPU Architecture Compatibility Guide - GPU互換性ガイド