Module occupancy_register_count 

PTX-parseable SM occupancy from register declarations.

The CUDA driver’s cuOccupancyMaxActiveBlocksPerMultiprocessorWithFlags needs the registers-per-thread figure to give an accurate occupancy answer. On a live GPU that number comes from ptxas; offline, the only source of truth is the PTX text itself — every kernel begins with a set of .reg directives that declare its virtual register file.

This module parses those .reg directives per .entry and feeds the resulting register count into the CPU-side OccupancyCalculator from crate::occupancy_ext. It is entirely host-side — no GPU is required — which makes it useful for build-time kernel auto-tuning and for occupancy reports in CI.

PTX register declaration grammar

Inside a kernel body, virtual registers are declared with lines such as:

.reg .f32   %f<10>;   // 10 single-precision registers
.reg .b32   %r<5>;    //  5 untyped 32-bit registers
.reg .pred  %p<3>;    //  3 predicate registers
.reg .f64   %fd<4>;   //  4 double-precision registers (each 2x 32-bit)

The bracketed <N> is the vectorised count: it declares %f0 .. %f9. A bare declaration without brackets (e.g. .reg .u32 %tid;) declares a single register.

The hardware register file is 32-bit-word addressed, so wide types (.f64, .b64, .s64, .u64) occupy two 32-bit slots each, while .pred registers do not consume general-purpose registers at all (they live in a separate condition-code file). PtxRegisterUsage models all of this so the occupancy estimate matches what ptxas would compute for the -O0 (no register coalescing) lower bound.

Example

use oxicuda_driver::occupancy_register_count::OccupancyFromPtx;

let ptx = r#"
.version 8.0
.target sm_90
.visible .entry saxpy() {
    .reg .f32   %f<8>;
    .reg .b32   %r<4>;
    .reg .pred  %p<2>;
    ret;
}
"#;

let from_ptx = OccupancyFromPtx::parse(ptx).expect("valid PTX");
// saxpy uses 8 f32 + 4 b32 = 12 general-purpose 32-bit registers.
assert_eq!(from_ptx.kernel("saxpy").unwrap().registers_per_thread(), 12);

// Feed straight into the CPU occupancy model for an sm_90 device.
let est = from_ptx.estimate_for("saxpy", 9, 0, 256, 0).unwrap();
assert!(est.occupancy_ratio > 0.0);

Structs

OccupancyFromPtx

Parses register usage out of a PTX string and wires it into the CPU-side occupancy model.

PtxKernel

One parsed kernel: its mangled entry name plus its register usage.

PtxRegisterUsage

Register-file usage of a single PTX kernel, derived from its .reg directives.