Nvfp4 Scale

namespace trt_edgellm

Argument bundles for the CuTe DSL FMHA launchers.

Each public run* entry point fills one of these once from its own arguments and members, and the launcher then spreads it across the generated descriptors. One struct per wrapper shape rather than a single superset: a superset would let a field that matters on one path (kvCacheCapacity on the dense path, tokensPerPage on the paged one) sit silently at zero on another.

These carry no generated types, so this header is safe to include from translation units that use the required FMHA-v2 runner, the optional CUTE_DSL_FMHA_BLACKWELL_ENABLED runner, or both. The descriptor-filling machinery itself lives in cuteDslTensorDescriptors.h, which stays free of any FMHA concept.

Helpers for populating the tensor descriptors emitted by the CuTe DSL C exporter.

Every AOT variant exports its own nominally distinct but layout-identical descriptor structs (fmha_d64_Tensor_q_tensor_t vs fmha_d128_Tensor_q_tensor_t) plus a cute_dsl_<variant>_wrapper entry point. There is no umbrella C type, so descriptor types are recovered here from the signature of the wrapper that consumes them: a call site names only the wrapper and the kernel module, and pairing a descriptor with the wrong variant is not expressible.

The exporter (cutlass/cute/export/c_header_generator.py) always names the members data, dynamic_shapes and dynamic_strides, but emits each array only when its dynamic mask is non-empty. A rank-1 descriptor therefore has no dynamic_strides member at all and needs makeCuSeqLenTensor() rather than the strided builders below.

namespace kernel

Per-position sum across codec groups (Qwen3.5-Omni speaker codec embedding fold).

Mirrors HF _get_codec_input_embeddings: concatenate per-group embeddings then sum across the group axis. Implemented here as a single-row reduction so the caller only invokes one kernel per row of the speaker codec template.

namespace nvfp4

Functions

__device__ __forceinline__ uint16_t e4m3ToFloat16Times128 (uint32_t const raw)

__device__ __forceinline__ uint32_t e4m3x2ToFloat16x2Times128 (uint16_t const raw)

__device__ __forceinline__ uint16_t e4m3ToBfloat16Times128 (uint32_t const raw)

__device__ __forceinline__ uint32_t e4m3x2ToBfloat16x2Times128 (uint16_t const raw)