pub fn cublas_bf16_proj_dense(
act: DevicePtr,
weight_bf16: DevicePtr,
out: DevicePtr,
m: u32,
n: u32,
k: u32,
stream: u64,
) -> Result<()>Expand description
Route a projection out[M,N] = act[M,K] @ weightᵀ through cuBLASLt BF16 for
a weight that is already BF16 [N,K]. Two kinds of caller: models whose
attention/shared-expert weights ship unquantized (e.g. Laguna), and the
row-wise GDN prefill arms, which hand it the ledgered BF16 dequant
qwen3_ssm/rowwise_bf16.rs writes once per layer.
There is deliberately NO cublas_bf16_proj beside it any more — the
dequant-and-cache variant that used to own the FP8→BF16 expansion is the
#917 off-ledger allocation (see dequant_fp8_bf16_into). Splitting
“who owns the BF16 bytes” from “multiply them” is what keeps the ledger
honest: this function cannot allocate.