Matmul Inspector

See what your GPU or AI-accelerator compiler actually produced.

Matmul Inspector turns saved NVIDIA PTX, NVIDIA SASS, and AWS Neuron NKI compiler artifacts into readable JSON and Markdown reports. This matters because a small source change can disappear in an intermediate representation, produce different machine instructions, change numerical results, or affect performance—and none of those outcomes can safely be inferred from another.

Install the Python toolkit from PyPI:

python -m pip install matmul-inspector==0.3.0
matmul-inspector-python --version

Version 0.3.0 supports Python 3.10–3.13. Saved-artifact analysis has no mandatory CUDA, AWS Neuron, GPU, or native-library dependency. The optional experiment commands require an NVIDIA GPU, a compatible driver, CUDA-enabled PyTorch, and Triton; installing Matmul Inspector does not install or configure a driver and does not guarantee that a particular GPU/software combination is compatible.

Why Matmul Inspector?

Kernel work crosses several evidence layers:

source -- compilation --> compiler IR / PTX -- lowering --> SASS / NEFF
SASS / NEFF -- execution -------------------------------> numerical outputs
            `-- timed/profiled execution ---------------> performance measurements

A source-level optimization may be canonicalized away. A PTX instruction need not map one-to-one to a hardware instruction. Fewer static instructions do not guarantee fewer dynamically executed instructions or lower latency. A faster run is not successful if it violates the required numerical tolerance.

Matmul Inspector keeps these layers separate while making them easy to compare.

Quick start: compare two PTX modules

This CPU-only example compares one fused multiply-add instruction with separate multiply and add instructions. It demonstrates a static PTX difference; it does not show the final machine code or predict performance.

The inline PTX fixtures are synthetic inputs for static analysis, not executable numerical kernels. They contain no initialized inputs or observable stored output.

Create the fused module:

cat > fused.ptx <<'PTX'
.version 8.0
.target sm_80
.address_size 64

.visible .entry update_value()
{
    .reg .f32 %f<4>;
    fma.rn.f32 %f3, %f1, %f2, %f0;
    ret;
}
PTX

Create the separate-operation module:

cat > separate.ptx <<'PTX'
.version 8.0
.target sm_80
.address_size 64

.visible .entry update_value()
{
    .reg .f32 %f<5>;
    mul.rn.f32 %f3, %f1, %f2;
    add.rn.f32 %f4, %f3, %f0;
    ret;
}
PTX

Compare them with the explicit Python entry point:

matmul-inspector-python ptx compare \
  --baseline fused.ptx \
  --candidate separate.ptx \
  --kernel update_value \
  --output comparison

The command writes comparison/comparison.json, comparison/report.md, two normalized PTX files, and a bounded unified diff. With the published 0.1.0 wheel, the generated Markdown and JSON report:

PTX registers are virtual declarations, not final physical register allocation, hardware resource use, or occupancy. The output directory must not already exist; this protects previous reports from accidental overwrite.

The package also installs matmul-inspector. The native C++ source build uses the same name, so matmul-inspector-python is the unambiguous Python command when both are on PATH.

See it in action: FMA versus No-FMA

A controlled CUDA 13.3 experiment on an NVIDIA GeForce RTX 4060 Ti compared an explicit fused multiply-add kernel with one using separate round-to-nearest multiply and add operations. Inspection of the captured sm_89 SASS found:

This is hardware-specific final-machine-code evidence for that build. These are static instruction counts, not dynamic execution counts. The inspector also kept embedded division helpers separate from the selected kernel.

The same study recorded CUDA-event kernel timings and numerical comparisons:

The timing direction changes with size, measurements used fixed reference-before-candidate order, and the candidate fails the configured tolerance at larger sizes. The static instruction difference therefore does not establish the cause of any timing difference or demonstrate a successful optimization. See the PTX/SASS FMA case study and floating-point experiment report for the capture conditions, methodology, numerical evidence, and limitations.

Capabilities

The source tree also includes a CPU-only TensorGame research environment for exact matrix multiplication tensor decomposition. Its standard and Strassen trajectories are known correctness references, not newly discovered algorithms. A bounded sampled MCTS planner supports deterministic reference-guided validation and small unguided search experiments without neural networks or hardware dependencies. The independent discovery study compares uniform, residual-informed, and hybrid proposals under matched CPU search budgets. The synthetic demonstration engine creates exact, leakage-checked policy/value training examples without requiring an ML framework. The optional policy/value learning baseline trains a small supervised PyTorch MLP and compares held-out proposal quality without changing MCTS or adding reinforcement learning. Its versioned evaluation audit reports explicit deterministic and seeded denominators, paired outcomes, grouped uncertainty, and proposal-generation cost. The optional neural-guided MCTS experiment separates learned proposals from learned leaf values and reports five matched ablations without updating the trained model. The neural value study compares trained, untrained, constant, and handcrafted leaf values under fixed residual proposals. The value-scale study traces PUCT mechanics and separates optimistic value magnitude from learned state ranking. The experience replay guide records actual MCTS-driven environment decisions and validates exact episode replay. The optional offline learning guide trains policy and direct-return value targets on validated sampled MCTS supports. Its TG-4A run is a pipeline smoke test, not evidence of improved discovery. The diverse experience guide defines a leakage-safe 200-target curriculum and staged MCTS collection audit for future offline learning. The controlled offline ablation compares strict policy-only, value-only, and joint updates using validation-only selection and a frozen held-out search plan. The held-out protocol addendum resolves pre-execution path, comparison-matrix, bootstrap, and wall-time ambiguities while leaving the original frozen plan unchanged. The offline-learning held-out case study reports the completed 1,800-trial discovery evaluation and its exact evidence audit. The policy-regression investigation traces the policy-only failure to critical one-step actions removed by bounded decoding. The candidate-coverage investigation tests bounded recall, experiment-only proposal unions, and training support. The out-of-sample policy evaluation compares frozen, trained, and equal-budget hybrid proposals on fresh targets. The search calibration study validates reference trajectories and separates candidate coverage from search budget. The held-out policy and search evaluation compares all frozen candidate providers under the preregistered TG-4G protocol. The residual-supervision feasibility study audits teacher compatibility, supervision overlap, multi-step limits, and cost. The coverage-aware policy study tests verified-positive supervision against a reproduced MCTS-only control. The coverage-retention investigation compares individual-positive supervision and selective frozen-policy retention. The prefix-coverage investigation traces the exact beam stages and tests an experimental prefix-survival objective. The final-ranking investigation tests a decoder-cutoff hinge under matched tiny-set training and validation. The prefix-generalization investigation tests broader verified-positive coverage and audits beam-aware supervision cost. The policy-interference investigation separates local gradient effects from independent-token combinatorial dilution. The conditional-policy study implements and evaluates an isolated u -> v -> w policy and conditional beam decoder. The adaptive-routing study compares frozen, exhaustive residual, fixed-mix, and train-selected proposal routing cost. The bounded-residual study evaluates a continuous-bound coarse-to-fine approximation to exhaustive residual ranking. The held-out bounded-search study compares bounded and exhaustive residual providers in paired end-to-end search. The provider-integration study documents opt-in residual-provider selection and the expanded regression corpus. The bounded-memory study validates an exact-output-equivalent streaming heap selector for the opt-in bounded provider. The streaming integration report records its integration and permanent differential/search regressions. The default-readiness study evaluates the integrated provider on an independent structurally targeted corpus. The optimization history records the TG-4C through TG-4V decision trail, including negative results. The rollout decision applies the final release gates on a 150-target corpus while preserving exhaustive as the default. The cross-machine verification guide records the completed integrity-pinned, resumable TG-4X run on a separate AWS CPU host and the current clean-checkout reproduction procedure. The default-rollout gate report records the repaired TG-4Y release gates and the rollout of bounded_128 as the public default. Explicit exhaustive selection remains available as a fallback. The release-preparation summary collects usage, evidence, limitations, and the remaining publication checklist.

Direct script entry points remain available from a source checkout, but the installed dispatcher is normally simpler:

matmul-inspector-python ptx --help
matmul-inspector-python sass --help
matmul-inspector-python nki --help
matmul-inspector-python report --help

The installed package includes the end-to-end CUDA/Triton experiment runner. It validates the established packaged FC1 candidate, compares it with PyTorch/cuBLAS, alternates CUDA-event measurement order, and preserves successful and failed records under the ignored artifacts/experiments/ directory:

matmul-inspector-python experiment

The runner supports the compatibility-default allocating mode and an explicit preallocated mode with reusable PyTorch and Triton outputs. Results keep CUDA event, host submission, and synchronized host-wall timings separate; modes are not combined in rankings. It also accepts explicitly opted-in, trusted local Triton source candidates in preallocated mode. Their bytes are hashed, preserved with the run, and loaded only in the worker process.

See the experiment runner guide for configuration, result fields, failure behavior, and the callable Python interface.

Run a reproducibly shuffled, sequential batch of predefined FC1 candidates with fresh matched baselines. From a source checkout, the checked-in example is:

matmul-inspector-python experiment-batch \
  --config examples/fc1_experiment_batch.json \
  --output-root artifacts/experiment-batches

Persistent search sessions add immutable experiment settings, structured hypotheses and history, canonical duplicate detection, locking, crash recovery, and an enforced experiment budget without adding a model API:

matmul-inspector-python session create \
  --config examples/fc1_experiment_session.json \
  --output-root artifacts/experiment-sessions

See the experiment runner guide for proposal, status, context, and recovery commands.

How artifact comparison works

1. Read existing artifacts without modifying them.
2. Parse identities, instructions, declarations, metadata, and available resource records without guessing across ambiguous boundaries.
3. Preserve raw hashes and apply only documented conservative normalization.
4. Compare full-file text separately from selected-kernel structure.
5. Write deterministic structured JSON and a concise Markdown evidence report.
6. Keep compiler observations, numerical correctness, benchmark completion, and performance measurements independently labeled.

Output directories are never overwritten. Unknown or unavailable information is reported rather than silently treated as equal, empty, or successful.

Native C++/CUDA laboratory

The native executable is not bundled in the Python wheel. Build it from a source checkout with CMake 3.17+ and a C++17 compiler:

cmake -S . -B build-cpu \
  -DCMAKE_BUILD_TYPE=Release \
  -DMATMUL_INSPECTOR_ENABLE_CUDA=OFF
cmake --build build-cpu --parallel
ctest --test-dir build-cpu -L cpu --output-on-failure
./build-cpu/matmul-inspector compare \
  --reference cpu --candidate cpu --sizes 4,17 --seed 42 \
  --output results/cpu-check

The CUDA source build also provides an FP32 comparison between a one-thread-per- system batched Thomas baseline and a one-thread-per-equation PCR baseline. It retains raw CUDA-event samples and correctness evidence. See the CUDA Thomas versus PCR case study.

Enable CUDA with -DMATMUL_INSPECTOR_ENABLE_CUDA=ON when a supported toolkit is available. The laboratory includes naive and tiled matmul, controlled FMA and accumulation-order variants, correctness comparisons, benchmark distributions, CPU/GPU crossover analysis, reductions, and reproducible artifact capture. It is an experimental laboratory, not a replacement for a tuned BLAS.

Documentation and reproducible experiments

• Experimental findings — concise results and qualifications
• Benchmark methodology — timing boundaries, trials, bootstrap intervals, and reliability flags
• Floating-point experiments — contraction, reordering, correctness, and replay
• PTX compiler analysis
• CUDA PTX/SASS FMA case study
• NKI compiler-artifact comparison
• Compiler-to-performance reporting
• Trainium1 NKI GEMM case study
• NKI GEMM benchmarking
• Roofline methodology
• Cross-hardware comparison and reproducibility
• Summation methods, operations, and tridiagonal benchmarking
• Transformer optimization study
• CUDA/Triton experiment runner
• AI-assisted FC1 kernel search case study
• Python toolkit release guide, 0.3.0 release notes, and 0.2.0 release notes
• Checked-in RTX 4060 Ti floating-point captures

Requirements and limitations

• Saved-artifact analysis supports Linux and Python 3.10–3.13 and uses only the Python standard library.
• GPU experiments are optional. Follow PyTorch's platform-specific installation instructions for a CUDA-enabled build, then install matmul-inspector[gpu]==0.3.0. The extra requests PyTorch and Triton but does not configure NVIDIA drivers, select a system CUDA driver, or guarantee GPU compatibility. Verify torch.cuda.is_available() before running experiments.
• Creating PTX/CUBIN/SASS requires NVIDIA compiler tools; running CUDA benchmarks additionally requires compatible hardware and drivers.
• Creating NKI/NEFF/profiler captures requires an AWS Neuron environment and, for hardware measurements, Trainium. Comparing saved captures does not.
• PTX is an intermediate representation; SASS is hardware-specific machine code.
• Static counts do not establish dynamic execution, scheduling, stalls, traffic, occupancy, or performance causality.
• Conservative normalized equality is not proof of binary or semantic equality.
• Benchmark conclusions apply only to their recorded hardware, software, inputs, timing boundaries, and numerical acceptance criteria.

Contributing

Bug reports and focused contributions are welcome through GitHub Issues. Preserve reproducible inputs and outputs, report correctness independently from timing, and avoid optimization claims without controlled measurements.

License

Matmul Inspector is licensed under the Apache License 2.0. The project copyright notice is recorded in NOTICE.