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

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.

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