この記事の対象読者
- CUDAの基本的なコンパイルフローを理解している方
- nvccの
-arch/-codeオプションを使ったことがある方 - 「PTXって何？」と聞かれて答えられるようになりたい方
- GPUコードの最適化や低レベル処理に興味がある方
この記事で得られること
- PTX（Parallel Thread Execution）の役割と設計思想の理解
- PTXの基本構文と命令セットの知識
- インラインPTXアセンブリの実践的な書き方
- PTXを活用したパフォーマンス分析手法
この記事で扱わないこと
- PTX ISAの全命令リファレンス（公式ドキュメント参照）
- 高度なPTX最適化テクニック
- cuda-gdbを使ったPTXレベルデバッグの詳細
1. PTXとの出会い
「この命令、なんでこんなに遅いんだ...？」
あるとき、CUDAカーネルのボトルネックを追っていて、Nsight Computeの出力を眺めていた。すると、見慣れない文字列が目に入った。
ld.global.f32 %f1, [%rd1];
add.f32 %f2, %f1, %f3;
st.global.f32 [%rd2], %f2;
これがPTXだった。
最初は「なんだこの呪文は」と思ったが、調べていくうちに、これがCUDAの「中間言語」であり、GPUプログラミングの根幹を支える重要な技術だとわかった。
PTX（Parallel Thread Execution）は、NVIDIAが設計した仮想GPUアセンブリ言語だ。CPUにとってのx86アセンブリのように、PTXはGPUにとっての低レベル言語...ただし、実際のハードウェアではなく「仮想的なGPU」のための言語という点が異なる。
例えるなら、PTXは「GPU世界のエスペラント語」。どんな世代のNVIDIA GPUでも理解できる共通語だ。
ここまでで、PTXがどんなものか、なんとなくイメージできただろうか。次は、PTXが生まれた背景を見ていこう。
2. 前提知識の確認
本題に入る前に、この記事で使う用語を整理しておく。
2.1 ISA（Instruction Set Architecture）とは
プロセッサが実行できる命令の仕様のこと。CPUならx86_64やARM64、GPUならNVIDIAの各世代で異なるISAがある。いわば「プロセッサの母国語」と考えるとわかりやすい。
2.2 仮想マシンとは
実際のハードウェアを抽象化したソフトウェア上の計算機。JavaのJVM（Java Virtual Machine）が有名。PTXも「仮想GPU」のための言語。
2.3 中間表現（IR: Intermediate Representation）とは
高級言語と機械語の間に位置する表現形式。LLVMのLLVM IRや、JavaのバイトコードがIRの例。PTXもGPU向けの中間表現の一種。
2.4 JIT（Just-In-Time）コンパイルとは
プログラムの実行時にコンパイルを行う方式。PTXから実際のGPUバイナリへの変換は、このJITで行われることがある。
これらの用語が押さえられたら、次に進もう。
3. なぜPTXが必要なのか
3.1 GPUの世代間互換性問題
NVIDIAはGPUを頻繁に更新する。Maxwell、Pascal、Volta、Turing、Ampere、Ada Lovelace、Hopper、Blackwell...各世代で命令セット（ISA）が異なる。
CPUの世界ではx86_64向けにコンパイルすればほぼすべてのPCで動くが、GPUでは「RTX 3090用バイナリがRTX 4090で動かない」という状況があり得た。
3.2 NVIDIAの解決策: 二段階コンパイル
この問題を解決するため、NVIDIAは「仮想GPU」という概念を導入した。
| 段階 | ターゲット | 出力形式 | 例 |
|---|---|---|---|
| 第1段階 | 仮想アーキテクチャ | PTX | compute_89 |
| 第2段階 | 実アーキテクチャ | cubin/SASS | sm_89 |
PTXは仮想アーキテクチャ向けに生成され、実際のGPUバイナリ（cubin）への変換はビルド時またはJIT実行時に行われる。
3.3 PTXの設計目標
NVIDIA公式ドキュメントによると、PTXには5つの設計目標がある。
- 安定したISAの提供: 複数のGPU世代をまたいで使える
- ネイティブ並みの性能: コンパイル後のコードが効率的
- マシン非依存: C/C++以外のコンパイラからもターゲットにできる
- 配布用フォーマット: アプリケーション配布時の共通形式
- 手書き可能: 最適化やテスト用にエキスパートが直接記述できる
PTXの設計思想が理解できたところで、次は実際のPTXの構文を見ていこう。
4. PTXの基本構文
4.1 PTXファイルの構造
PTXファイルはテキスト形式で、以下の要素で構成される。
// PTXファイルの基本構造
.version 8.5 // PTX ISAバージョン
.target sm_89 // ターゲットアーキテクチャ
.address_size 64 // アドレスサイズ（32 or 64）
// グローバル変数宣言
.global .align 4 .f32 my_global;
// カーネル定義
.visible .entry my_kernel(
.param .u64 param_ptr
)
{
// レジスタ宣言
.reg .f32 %f<10>; // %f0〜%f9
.reg .b64 %rd<5>; // %rd0〜%rd4
.reg .pred %p<2>; // 述語レジスタ
// 命令
ld.param.u64 %rd1, [param_ptr];
// ...
ret;
}
4.2 データ型
PTXは型付けされた言語で、以下の型をサポートする。
| 型 | 説明 | サイズ |
|---|---|---|
.pred |
述語（boolean） | 1ビット |
.b8, .b16, .b32, .b64
|
ビット列 | 8/16/32/64ビット |
.s8, .s16, .s32, .s64
|
符号付き整数 | 8/16/32/64ビット |
.u8, .u16, .u32, .u64
|
符号なし整数 | 8/16/32/64ビット |
.f16, .f32, .f64
|
浮動小数点 | 16/32/64ビット |
4.3 状態空間（State Space）
PTXでは、メモリの種類を「状態空間」として明示的に指定する。これがPTXの特徴的な点だ。
| 状態空間 | キーワード | 説明 | CUDAでの対応 |
|---|---|---|---|
| レジスタ | .reg |
最速。スレッドローカル | ローカル変数 |
| 共有メモリ | .shared |
ブロック内で共有 | __shared__ |
| グローバル | .global |
全スレッドからアクセス可能 | __device__ |
| 定数 | .const |
読み取り専用 | __constant__ |
| ローカル | .local |
スレッドローカル（スピル用） | 自動変数 |
| パラメータ | .param |
カーネル引数 | 関数引数 |
4.4 基本命令
PTX命令は以下の形式をとる。
命令.修飾子.型 出力, 入力1, 入力2;
代表的な命令を見てみよう。
// 算術演算
add.s32 %r1, %r2, %r3; // r1 = r2 + r3 (符号付き32bit)
mul.lo.s32 %r1, %r2, %r3; // r1 = (r2 * r3) の下位32bit
mad.f32 %f1, %f2, %f3, %f4; // f1 = f2 * f3 + f4 (FMA)
// メモリアクセス
ld.global.f32 %f1, [%rd1]; // グローバルメモリから読み込み
st.global.f32 [%rd1], %f1; // グローバルメモリへ書き込み
ld.shared.f32 %f1, [%rd1]; // 共有メモリから読み込み
// 比較と分岐
setp.lt.s32 %p1, %r1, %r2; // p1 = (r1 < r2)
@%p1 bra LABEL; // p1がtrueならLABELへジャンプ
@!%p1 bra LABEL; // p1がfalseならLABELへジャンプ
// 特殊レジスタ（スレッド情報取得）
mov.u32 %r1, %tid.x; // スレッドID取得
mov.u32 %r2, %ctaid.x; // ブロックID取得
mov.u32 %r3, %ntid.x; // ブロックサイズ取得
4.5 述語実行（Predicated Execution）
PTXの特徴的な機能が「述語実行」。条件分岐の代わりに、命令の実行自体を条件付きにできる。
// 述語実行（分岐なし）
setp.lt.s32 %p1, %r1, 10;
@%p1 add.s32 %r2, %r2, 1; // r1 < 10 のときだけ実行
@!%p1 add.s32 %r2, %r2, 2; // r1 >= 10 のときだけ実行
分岐を減らすことで、GPU特有の**ワープ発散（warp divergence）**を軽減できる。これは32スレッドが同時に同じ命令を実行するGPUアーキテクチャにおいて、パフォーマンスに大きく影響する。
PTXの基本構文が理解できたところで、次は実際にPTXを生成・確認する方法を見ていこう。
5. 実際に使ってみよう
5.1 環境設定ファイル
以下の3種類の設定ファイルを用意した。用途に応じて選択してほしい。
開発環境用（ptx_dev.sh）
#!/bin/bash
# ptx_dev.sh - PTX調査用開発環境（このままコピーして使える）
export CUDA_HOME=/usr/local/cuda
export PATH=$CUDA_HOME/bin:$PATH
# PTX生成用オプション
export NVCC_PTX_FLAGS="-ptx -arch=compute_89"
# 詳細出力用（中間ファイルを保持）
export NVCC_VERBOSE="-v --keep --keep-dir=./ptx_output"
echo "PTX開発環境を設定しました"
nvcc --version
本番環境用（ptx_prod.sh）
#!/bin/bash
# ptx_prod.sh - 本番ビルド環境（このままコピーして使える）
export CUDA_HOME=/usr/local/cuda
export PATH=$CUDA_HOME/bin:$PATH
# マルチアーキテクチャ対応（PTX埋め込みあり）
export NVCC_GENCODE="\
-gencode arch=compute_80,code=sm_80 \
-gencode arch=compute_86,code=sm_86 \
-gencode arch=compute_89,code=sm_89 \
-gencode arch=compute_89,code=compute_89"
# 最適化オプション
export NVCC_OPT="-O3 -use_fast_math"
echo "本番環境を設定しました"
テスト環境用（ptx_test.sh）
#!/bin/bash
# ptx_test.sh - CI/テスト環境（このままコピーして使える）
export CUDA_HOME=/usr/local/cuda
export PATH=$CUDA_HOME/bin:$PATH
# 高速コンパイル（単一アーキテクチャ）
export NVCC_FLAGS="-arch=sm_89 -O0"
# 警告を厳格に
export NVCC_WARNINGS="-Werror all-warnings"
echo "テスト環境を設定しました"
5.2 サンプルコード
// vector_add.cu - PTX生成サンプル（このままコピーして使える）
#include <cuda_runtime.h>
#include <stdio.h>
__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 main() {
const int N = 1024;
float *d_a, *d_b, *d_c;
cudaMalloc(&d_a, N * sizeof(float));
cudaMalloc(&d_b, N * sizeof(float));
cudaMalloc(&d_c, N * sizeof(float));
vectorAdd<<<(N+255)/256, 256>>>(d_a, d_b, d_c, N);
cudaDeviceSynchronize();
cudaFree(d_a);
cudaFree(d_b);
cudaFree(d_c);
printf("Kernel executed successfully\n");
return 0;
}
5.3 PTXを生成するビルドスクリプト
#!/bin/bash
# build_ptx.sh - PTX生成スクリプト（このままコピーして使える）
set -e
INPUT=${1:-"vector_add.cu"}
OUTPUT_DIR="./ptx_output"
mkdir -p $OUTPUT_DIR
echo "=== PTXを生成します ==="
nvcc -ptx -arch=compute_89 $INPUT -o $OUTPUT_DIR/kernel.ptx
echo "=== 生成されたPTX（先頭50行） ==="
head -50 $OUTPUT_DIR/kernel.ptx
echo ""
echo "=== PTX統計 ==="
echo "総行数: $(wc -l < $OUTPUT_DIR/kernel.ptx)"
echo "命令数: $(grep -c '^\s*[a-z]' $OUTPUT_DIR/kernel.ptx || echo 0)"
echo "レジスタ宣言: $(grep -c '\.reg' $OUTPUT_DIR/kernel.ptx || echo 0)"
5.4 実行結果
上記のスクリプトを実行すると、以下のようなPTXが生成される。
$ ./build_ptx.sh vector_add.cu
=== PTXを生成します ===
=== 生成されたPTX（先頭50行） ===
//
// Generated by NVIDIA NVVM Compiler
//
.version 8.5
.target sm_89
.address_size 64
.visible .entry _Z9vectorAddPKfS0_Pfi(
.param .u64 _Z9vectorAddPKfS0_Pfi_param_0,
.param .u64 _Z9vectorAddPKfS0_Pfi_param_1,
.param .u64 _Z9vectorAddPKfS0_Pfi_param_2,
.param .u32 _Z9vectorAddPKfS0_Pfi_param_3
)
{
.reg .pred %p<2>;
.reg .f32 %f<4>;
.reg .b32 %r<6>;
.reg .b64 %rd<11>;
ld.param.u64 %rd1, [_Z9vectorAddPKfS0_Pfi_param_0];
ld.param.u64 %rd2, [_Z9vectorAddPKfS0_Pfi_param_1];
ld.param.u64 %rd3, [_Z9vectorAddPKfS0_Pfi_param_2];
ld.param.u32 %r2, [_Z9vectorAddPKfS0_Pfi_param_3];
mov.u32 %r3, %ctaid.x;
mov.u32 %r4, %ntid.x;
mov.u32 %r5, %tid.x;
mad.lo.s32 %r1, %r3, %r4, %r5;
setp.ge.s32 %p1, %r1, %r2;
@%p1 bra $L__BB0_2;
// 配列アクセスとベクトル加算
cvta.to.global.u64 %rd4, %rd1;
mul.wide.s32 %rd5, %r1, 4;
add.s64 %rd6, %rd4, %rd5;
ld.global.f32 %f1, [%rd6];
ld.global.f32 %f2, [%rd8];
add.f32 %f3, %f1, %f2;
st.global.f32 [%rd10], %f3;
$L__BB0_2:
ret;
}
5.5 よくあるエラーと対処法
| エラー | 原因 | 対処法 |
|---|---|---|
ptxas fatal: Unsupported .version X.Y |
PTXバージョンがptxasでサポートされていない | CUDA Toolkitを更新する |
ptxas error: Entry function uses too much register |
レジスタ使用量オーバー |
--maxrregcountで制限 |
ptxas warning: Double is not supported |
sm_13未満でdouble使用 | sm_13以上を指定 |
Unknown symbol |
シンボル未定義 |
.extern宣言を追加 |
Instruction not supported |
アーキテクチャ非対応命令 | より新しいcompute_XXを指定 |
PTXの生成・確認方法がわかったところで、次はインラインPTXアセンブリについて見ていこう。
6. インラインPTXアセンブリ
CUDAでは、asm()文を使ってPTX命令を直接埋め込める。これにより、コンパイラが生成しない特殊な命令を使用できる。
6.1 基本構文
// 基本形
asm("命令" : 出力 : 入力);
// 例: add.s32命令
int a = 10, b = 20, c;
asm("add.s32 %0, %1, %2;" : "=r"(c) : "r"(a), "r"(b));
// c = 30
6.2 制約文字
| 制約 | 説明 | 用途 |
|---|---|---|
r |
32ビット整数レジスタ | int, unsigned int |
l |
64ビット整数レジスタ | long long, ポインタ |
f |
32ビット浮動小数点レジスタ | float |
d |
64ビット浮動小数点レジスタ | double |
n |
即値定数 | コンパイル時定数 |
= |
出力専用 | 書き込みのみ |
+ |
入出力両用 | 読み書き両方 |
6.3 実践例: 高速clock取得
// clock_sample.cu - インラインPTXサンプル（このままコピーして使える）
#include <cuda_runtime.h>
#include <stdio.h>
__device__ __forceinline__ unsigned int getClock() {
unsigned int clock;
asm volatile("mov.u32 %0, %%clock;" : "=r"(clock));
return clock;
}
__device__ __forceinline__ unsigned long long getClock64() {
unsigned long long clock;
asm volatile("mov.u64 %0, %%clock64;" : "=l"(clock));
return clock;
}
__global__ void measureLatency(unsigned int* output) {
unsigned int start = getClock();
// 何か処理（ここでは単純な計算）
float x = 1.0f;
#pragma unroll
for (int i = 0; i < 100; i++) {
x = x * 1.001f + 0.001f;
}
unsigned int end = getClock();
if (threadIdx.x == 0) {
output[0] = end - start;
printf("Elapsed cycles: %u\n", end - start);
}
}
int main() {
unsigned int* d_output;
cudaMalloc(&d_output, sizeof(unsigned int));
measureLatency<<<1, 32>>>(d_output);
cudaDeviceSynchronize();
cudaFree(d_output);
return 0;
}
6.4 Warp Shuffle（インラインPTX版）
// warp_shuffle.cu - Warp Shuffleサンプル（このままコピーして使える）
__device__ __forceinline__ float warpShuffleDown(float val, int offset) {
float result;
asm volatile(
"shfl.sync.down.b32 %0, %1, %2, 0x1f, 0xffffffff;"
: "=f"(result) : "f"(val), "r"(offset)
);
return result;
}
__device__ float warpReduceSum(float val) {
// Warp内の全スレッドの値を合計
val += warpShuffleDown(val, 16);
val += warpShuffleDown(val, 8);
val += warpShuffleDown(val, 4);
val += warpShuffleDown(val, 2);
val += warpShuffleDown(val, 1);
return val;
}
インラインPTXが理解できたところで、次はユースケース別のガイドを見ていこう。
7. ユースケース別ガイド
7.1 ユースケース1: パフォーマンス分析
想定読者: カーネルのボトルネックを特定したい方
推奨構成: nvcc -ptx + grep分析
サンプルコード:
#!/bin/bash
# analyze_ptx.sh - PTX分析スクリプト（このままコピーして使える）
INPUT=${1:-"kernel.cu"}
PTX_FILE="analysis.ptx"
# PTX生成
nvcc -ptx -arch=compute_89 -lineinfo $INPUT -o $PTX_FILE
echo "=== メモリアクセス命令 ==="
echo "グローバルロード: $(grep -c 'ld\.global' $PTX_FILE)"
echo "グローバルストア: $(grep -c 'st\.global' $PTX_FILE)"
echo "共有メモリロード: $(grep -c 'ld\.shared' $PTX_FILE)"
echo "共有メモリストア: $(grep -c 'st\.shared' $PTX_FILE)"
echo ""
echo "=== 演算命令 ==="
echo "加算: $(grep -c 'add\.' $PTX_FILE)"
echo "乗算: $(grep -c 'mul\.' $PTX_FILE)"
echo "FMA: $(grep -c 'mad\.\|fma\.' $PTX_FILE)"
echo ""
echo "=== 分岐命令 ==="
echo "条件分岐: $(grep -c 'bra' $PTX_FILE)"
echo "述語付き命令: $(grep -c '@%p' $PTX_FILE)"
echo ""
echo "=== レジスタ使用量 ==="
grep '\.reg' $PTX_FILE | sort | uniq -c
7.2 ユースケース2: 前方互換性の確保
想定読者: 将来のGPUでも動くバイナリを配布したい方
推奨構成: cubin + PTX埋め込み
サンプルコード:
#!/usr/bin/env python3
# check_ptx_embed.py - PTX埋め込み確認（このままコピーして使える）
"""
実行可能ファイルにPTXが埋め込まれているか確認
使い方: python check_ptx_embed.py ./my_cuda_app
"""
import subprocess
import sys
def check_embedded_code(binary_path):
"""バイナリに含まれるcubin/PTXを確認"""
# cubin確認
result = subprocess.run(
['cuobjdump', '-lelf', binary_path],
capture_output=True, text=True
)
print("=== 埋め込まれたcubin ===")
if result.stdout.strip():
print(result.stdout)
else:
print("(なし)")
# PTX確認
result = subprocess.run(
['cuobjdump', '-lptx', binary_path],
capture_output=True, text=True
)
print("\n=== 埋め込まれたPTX ===")
if 'PTX file' in result.stdout:
print(result.stdout)
print("\n[OK] 将来のGPUとの前方互換性があります")
return True
else:
print("(なし)")
print("\n[WARNING] PTXが埋め込まれていません")
print("将来のGPUで動作しない可能性があります")
return False
if __name__ == "__main__":
if len(sys.argv) < 2:
print("Usage: python check_ptx_embed.py <binary>")
sys.exit(1)
check_embedded_code(sys.argv[1])
7.3 ユースケース3: DSL（ドメイン特化言語）からのPTX生成
想定読者: 独自言語やコンパイラを開発している方
推奨構成: libNVVM または PTX直接生成
サンプルコード:
// ptx_generator.cpp - PTX生成サンプル（このままコピーして使える）
#include <cuda.h>
#include <nvrtc.h>
#include <iostream>
#include <vector>
int main() {
// CUDA初期化
cuInit(0);
CUdevice device;
cuDeviceGet(&device, 0);
CUcontext context;
cuCtxCreate(&context, 0, device);
// カーネルソース
const char* kernelSource = R"(
extern "C" __global__ void addKernel(float* a, float* b, float* c, int n) {
int idx = blockIdx.x * blockDim.x + threadIdx.x;
if (idx < n) {
c[idx] = a[idx] + b[idx];
}
}
)";
// NVRTCでPTXにコンパイル
nvrtcProgram prog;
nvrtcCreateProgram(&prog, kernelSource, "kernel.cu", 0, NULL, NULL);
const char* opts[] = {"--gpu-architecture=compute_89"};
nvrtcResult compileResult = nvrtcCompileProgram(prog, 1, opts);
if (compileResult == NVRTC_SUCCESS) {
// PTXを取得
size_t ptxSize;
nvrtcGetPTXSize(prog, &ptxSize);
std::vector<char> ptx(ptxSize);
nvrtcGetPTX(prog, ptx.data());
std::cout << "=== Generated PTX ===" << std::endl;
std::cout << ptx.data() << std::endl;
}
nvrtcDestroyProgram(&prog);
cuCtxDestroy(context);
return 0;
}
ユースケースが把握できたところで、この記事を読んだ後の学習パスを確認しよう。
8. 学習ロードマップ
この記事を読んだ後、次のステップとして以下をおすすめする。
初級者向け（まずはここから）
-
自分のカーネルのPTXを見てみる
-
nvcc -ptx -arch=compute_89 your_kernel.cuで生成 - grep等で命令を数えてみる
-
-
cuobjdumpで既存バイナリを調査
cuobjdump -lptx -lelf your_binary- CUDA Binary Utilities
中級者向け（実践に進む）
-
PTX ISA仕様書を読む
- PTX ISA Documentation
- 特に状態空間とメモリモデルの章
-
インラインPTXで特殊命令を使う
- Inline PTX Assembly
- warp shuffle, atomic, memory barrierなど
上級者向け（さらに深く）
-
NVVM IRとlibNVVMを使ったカスタムコンパイラ開発
- libNVVM API
- LLVM IRからPTXへの変換
-
OpenAI Triton等のDSLの内部実装を読む
- PTX生成のベストプラクティスを学ぶ
9. まとめ
この記事では、PTXについて以下を解説した。
- 役割: PTXはNVIDIA GPUの仮想アセンブリ言語で、世代を超えた互換性を実現
- 構文: 型付き言語で、状態空間や述語実行といった特徴を持つ
-
生成方法:
nvcc -ptxで生成、cuobjdump -lptxで確認 - 活用法: インラインPTXで特殊命令を使う、パフォーマンス分析に活用
私の所感
正直、PTXを初めて見たときは「これを読めるようになる日が来るのか」と思った。しかし、基本的な構文さえ覚えてしまえば、意外と読める。
PTXを理解する最大のメリットは**「コンパイラが何をしているか見える」**こと。カーネルが遅いとき、PTXを見れば「ああ、ここで余計なメモリアクセスが入ってるな」と気づける。
もう一つのメリットは**「将来への保険」**。PTXを埋め込んでおけば、10年後のGPUでもJITコンパイルで動く（たぶん）。これは配布用バイナリを作る際に、かなり重要な設計判断だと思う。
参考文献
- PTX ISA Documentation - NVIDIA公式PTX仕様書
- Understanding PTX, the Assembly Language of CUDA GPU Computing - NVIDIA Developer Blog
- Inline PTX Assembly in CUDA - インラインPTX公式ガイド
- CUDA Binary Utilities - cuobjdump/nvdisasm公式ガイド