Struct QuantizedLinearPlan 

pub struct QuantizedLinearPlan<TIn: Element, TWQ: IntElement> { /* private fields */ }

quantized_linear plan (W8A8 fused).

Composes two passes internally:

1. Activation quantize — per-token symmetric dynamic-range quantization, fused max-abs reduce + scale compute + quantize. Implemented by the same dynamic_range_quantize_per_token_sym kernel that backs super::DynamicRangeQuantizePlan.
2. Quantized matmul — fused int8 GEMM + per-row/per-col dequantize + FP store. Implemented by the bespoke quantized_linear_w8a8 kernel.

Both passes share the same stream and execute back-to-back; the Plan does NOT own an internal DynamicRangeQuantizePlan instance — it invokes the FFI directly to keep the launch ordering explicit.

When to use: W8A8 inference matmul (SmoothQuant / AWQ-runtime style). Inference-only — no BW; for QAT use FakeQuantizePlan + normal FP matmul.

Dtypes (trailblazer): TIn (act/out) ∈ {f32, f64}; TWQ = S8. f16 / bf16 activations and u8 weight not yet wired.

Shape limits: activation [M, K]; weight_q [C_out, K]; weight_scale [C_out]; output [M, C_out]. The W4 layout [C_out, K] matches y = x · W^T (PyTorch nn.Linear.weight).

Workspace: zero in Workspace. Caller supplies act_q_scratch [M, K] (int8) and act_scale_scratch [M] (FP) in QuantizedLinearArgs for the fused activation-quant pass.

Precision guarantee: deterministic, bit-stable. Naive kernel (one thread per output cell, register-only int32 acc) for correctness; tiled-smem / mma.sync optimizations land in a perf milestone — current variant is correctness-scaffold, not throughput-optimized.

Implementations

impl<TIn: Element, TWQ: IntElement> QuantizedLinearPlan<TIn, TWQ>

pub fn select( _stream: &Stream, desc: &QuantizedLinearDescriptor, _pref: PlanPreference, ) -> Result<Self>

Pick a kernel for desc.

pub fn can_implement( &self, args: &QuantizedLinearArgs<'_, TIn, TWQ>, ) -> Result<()>

Validate args at run time.

pub fn workspace_size(&self) -> usize

Workspace bytes — none. Scratch buffers are caller-owned via the args bundle (act_q_scratch + act_scale_scratch), allowing reuse across launches.

pub fn sku(&self) -> KernelSku

Identity of the selected kernel.

pub fn precision_guarantee(&self) -> PrecisionGuarantee

Numerical guarantees.

pub fn run( &self, stream: &Stream, _workspace: Workspace<'_>, args: QuantizedLinearArgs<'_, TIn, TWQ>, ) -> Result<()>

Launch.