fix: cast fp32 block scales to UE8M0 before packing on SM100

[ipiszy]

Summary

• Replace UE8M0 packing device asserts with an explicit power-of-2 cast in smxx_layout.cuh.
• Match the torch fallback path by rounding fp32 scales before exponent packing.

Why

DeepGEMM 2.6.1 asserts when masked MoE scale buffers contain non-UE8M0 fp32 values (including padded rows). The SM100 transform path with disable_ue8m0_cast=false should cast inputs to UE8M0, not require callers to pre-sanitize buffers.

Test plan

• Text CUDA B200 TestDeepseek3.test_dsv2 on fireworks PR after bumping to this commit

Made with Cursor

[cursor]

Stale comment

Risk: medium. Cursor Bugbot completed with status skipping and no Bugbot review comment was found, so auto-approval is withheld. Requesting human review of the UE8M0 casting change in the SM100 pack kernels.

  

_{Sent by Cursor Approval Agent: Pull Request Approver}

[vipulSharma18]

Please check fireworks code for parity. Deepgemm or flash MLA was breaking in the past and leading to quality issue due to quantization logic difference between fireworks triton kernels and their helpers. Clamp and saturation go a long way to end quality.

[cursor]

Cursor Bugbot has reviewed your changes and found 1 potential issue.

^{❌ Bugbot Autofix is OFF. To automatically fix reported issues with cloud agents, have a team admin enable autofix in the Cursor dashboard.}

^{Reviewed by Cursor Bugbot for commit b95c96e. Configure here.}

[cursor]

Missing UE8M0 cast before pack

High Severity

The UE8M0 packing asserts were removed, but fp32 block scales are still packed by shifting raw IEEE bits (>> 23) with no power-of-two rounding. Non-UE8M0 scales (nonzero mantissa) therefore encode a floored exponent instead of the rounded UE8M0 code used by ceil_to_ue8m0, tilelang get_sf_and_inv, and fast_log2_ceil/fast_pow2, which can skew SM100 GEMM dequantization and model quality.

Additional Locations (1)

• deep_gemm/include/deep_gemm/impls/smxx_layout.cuh#L227-L236

 

^{Reviewed by Cursor Bugbot for commit b95c96e. Configure here.}

[cursor]

Risk: medium. Not approving: Cursor Bugbot finished with status skipping and reported 1 unresolved finding (missing UE8M0 cast before exponent packing in smxx_layout.cuh). Human review is required; vipulSharma18 is already requested and zheanxu was added as a second reviewer.

  

_{Sent by Cursor Approval Agent: Pull Request Approver}