PTX Lint

ptxlint reads the .ptx your GPU kernels compile to and reports local memory, FP64 use, register pressure and occupancy. No GPU, no CUDA install, no dependencies.

NVIDIA PTX 静态分析工具 —— 读取 kernel 编译出的 .ptx，报告 local memory、FP64、寄存器压力与占用率。不需要显卡和 CUDA。

---

⚡ Quick Start

install via cargo

cargo install ptxlint

install via brew

brew install rustq/tap/ptxlint

run

ptxlint -h

---

What It Is And Why It Exists

Rust can now compile kernels to PTX, but nothing tells you that the array you just wrote landed in DRAM instead of registers. That kind of mistake is invisible in the source: whether a scratch array stays in registers depends on how it is indexed and on whether the loop got unrolled, not on how the Rust reads. It usually surfaces much later, on a machine with a GPU, long after the change that caused it.

ptxlint moves that feedback earlier. PTX is the first point in the pipeline where the compiler has committed to its decisions — memory space, vector width, instruction types are all fixed — and it is still typed, readable text. Analysing it needs nothing from NVIDIA, so the same check runs on a laptop, in review, and on a CI runner with no GPU attached.

Note

Every Rust GPU toolchain ends at PTX — the in-tree nvptx64-nvidia-cuda target, rust-cuda, and NVIDIA's cuda-oxide alike. A tool that reads the common output works regardless of which one you picked.

---

What You Get

Capability	What It Detects	Why It Matters
Lints	Local memory, spills, FP64, integer division, register pressure, shared memory, narrow accesses, launch bounds, stray calls	Ten specific mistakes with a fix for each, not a wall of statistics
Occupancy model	Register, shared-memory, warp and block limits per architecture	Tells you which resource is capping residency, so you tune the right one
Baseline diff	Metric deltas between two builds of the same kernels	Answers "did my change make it worse", which is the question review actually asks
ptxas integration	Exact register counts and spill bytes	Turns the estimates into measurements when a CUDA toolkit is available
CI gating	Per-lint and per-regression exit codes	Fails the build on the findings you chose, and only those

---

Lints

Every lint has a runnable case in cases/, compiled from the Rust kernel linked beside it. Try one with ptxlint cases/ptx001_local_memory.ptx.

	What It Detects	Case	Kernel
PTX001	Local memory in use — an array indexed by a runtime value, backed by DRAM	ptx001	rs
PTX002	Register spills, needs --ptxas	ptx002	rs
PTX003	FP64 instructions — in Rust a bare 0.5 is f64	ptx003	rs
PTX004	Integer div/rem — GPUs have no integer divider	ptx004	rs
PTX005	High register pressure	ptx005	rs
PTX006	Low estimated occupancy	ptx006	rs
PTX007	Shared memory over budget, or capping residency	ptx007	hand-written
PTX008	Narrow, non-vectorised global accesses	ptx008	rs
PTX009	No .maxntid/.reqntid launch bounds	ptx009	rs
PTX010	Calls that were not inlined	ptx010	rs

cases/ also holds clean_saxpy, the control that must report nothing, modern_tensor_cores for wmma and cp.async, and nanoid_regression, a kernel this tool found a real bug in.

---

Baseline Diff

A report tells you whether a kernel is bad. Review usually asks something else: did this change make it worse? --baseline matches kernels by name across two builds and prints only what moved.

ptxlint --baseline cases/diff_before.ptx cases/diff_after.ptx

  mix16  improved
    ↓ local memory (B)     64 → 0 (-64)
    ↓ registers/thread     180 → 69 (-111)
    ↑ occupancy (%)        13 → 38 (+25)
    ↓ instructions         156 → 63 (-93)
    fixed    PTX001

Only the exact metrics can fail a build: local memory, spills, shared memory, a kernel that disappeared, a new error, and registers when both sides came from ptxas. Instruction counts and the virtual-register estimate move around too much to gate on, so they are reported and never block.

Tip

Keep the previous build's .ptx as a CI artifact and point --baseline at the directory. Kernels are paired by file name, then by kernel name.

---

CI Integration

- run: cargo build --release --target nvptx64-nvidia-cuda
- run: ptxlint --deny error target/nvptx64-nvidia-cuda/release/
- run: ptxlint --deny regression --baseline baseline/ target/nvptx64-nvidia-cuda/release/

Exit codes are 0 for a clean run, 1 when a denied lint or regression fired, and 2 when ptxlint itself could not run. Keeping 1 and 2 apart lets a check assert that a lint still fires rather than silently passing on a missing file.

---

Accuracy And Limits

PTX only carries virtual registers, which ptxas coalesces on the way to SASS. Local and shared memory, instruction mix, FP64 use and vector widths are therefore exact; registers and occupancy are an upper bound and are labelled as such in the report. Pass --ptxas or --ptxas-report to replace the estimate with the real number.

The scanner is deliberately not a full PTX grammar. An unknown opcode is recorded and matches no lint, rather than failing CI over an instruction NVIDIA shipped last month — wmma, cp.async and friends parse fine without the tool knowing what they mean.

Important

This is not a profiler. For real tuning use Nsight Compute, and for races use compute-sanitizer. ptxlint is the smoke alarm, not the fire brigade.

---

Development

cargo test
cargo clippy --all-targets -- -D warnings

Each case in fixtures/examples/ is a real Rust kernel that compiles to its own .ptx, so a fixture only ever triggers the lint it demonstrates. Regenerating them needs a nightly toolchain with the nvptx64-nvidia-cuda target.

./fixtures/generate.sh

showcase.sh walks through every feature against the cases, and is what the Showcase step in CI runs — the log is a live demo on a runner with no GPU.

cargo build --release && ./showcase.sh

Prior art: cuda-sage is a Python static PTX analyser covering similar ground, and the baseline diff idea came from it.

License

MIT