Expand description
2026-09-25: Host-side driver for InnerQ, the per-channel K equalization of the WHT-rotated (turbo) KV dtypes.
TURBO_INNERQ=N (N > 0 calibration tokens) turns it on;
TURBO_INNERQ_STRENGTH in (0, 1] sets the exponent, 0.5 otherwise. The
state is __device__ globals in namespace tq_plus, defined in
kernels/gb10/common/tq_plus_innerq_apply.cu. The driver reaches them
through the registry: device_symbol for the address, async copies to read
and write.
start zeroes the accumulators and sets d_innerq_calibrating = 1.
maybe_finalize reads d_innerq_count; once it reaches target_tokens it
clears d_innerq_calibrating, derives per-channel scale/scale_inv from
d_innerq_sq_accum, and either stops (every channel ratio within 1.2x) or
uploads them and sets d_innerq_active = 1.
Owner: model-layers (attention).
Invariants: the d_innerq_active = 1 write is issued only after the scale
uploads were synchronised on the driver’s stream.