feat: support build on gfx1250 (Part2)

[RuibinCheung]

Description

This PR is Part 2 of enabling Primus-Turbo on gfx1250 (CDNA5). Part 1 made the project build for gfx1250; this part makes the MXFP4 / MXFP8 quantization and de-quantization kernels actually run correctly on gfx1250, and unblocks the turbo (MFMA) FP8 GEMM host binding so it links uniformly.

Fixes # (issue)

Type of change

• Documentation change (change only to the documentation, either a fix or a new content)
• Bug fix (non-breaking change which fixes an issue)
• New feature (non-breaking change which adds functionality)
• Breaking change (fix or feature that would cause existing functionality to not work as expected)
• Infra/Build change
• Code refactoring

Changes

Please list the changes introduced in this PR:

• Parameterized the wavefront size in the MXFP4 / MXFP8 quantization kernels.
  WARP_SIZE now resolves to the per-arch THREADS_PER_WARP (64 on gfx942/gfx950, 32 on gfx1250), WARPS_PER_BLOCK is fixed at 4 so each warp still owns exactly one 32x32 chunk, and THREADS_PER_BLOCK scales to WARP_SIZE * WARPS_PER_BLOCK (256 on gfx950, 128 on gfx1250).
  This keeps a 1-warp-per-chunk mapping, which avoids both the extra __syncthreads deadlock and the LDS blow-up that a fixed 256-thread block would cause with 32-lane wavefronts.

• Fixed the host launch geometry for the fat-binary case.
  Because the host translation unit is compiled once but device code is compiled per --offload-arch, the block dimension is now computed from the runtime warp_size() (dim3 block(warp_size() * WARPS_PER_BLOCK)) instead of the compile-time THREADS_PER_BLOCK.

• Generalized the coalesced colwise FP4 / FP8 write-out and the LDS zero-init into grid-stride loops.
  These previously assumed a one-item-per-thread mapping tied to a 256-thread block; they now cover the buffer for both 256-thread (gfx950) and 128-thread (gfx1250) blocks, and the corresponding static_asserts were relaxed from strict equality to divisibility.

• Added gfx1250 conversion paths using compiler builtins instead of inline asm.
  Quantization uses __builtin_amdgcn_cvt_scalef32_pk8_fp4_f32, __builtin_amdgcn_cvt_scalef32_sr_pk8_fp4_f32 (stochastic rounding), and __builtin_amdgcn_cvt_scalef32_pk8_{fp8,bf8}_f32, padding the unused upper lanes and keeping the low bits so the element-to-nibble/byte order and the divide-by-scale semantics match the gfx950 path exactly.
  De-quantization uses __builtin_amdgcn_cvt_scale_pk8_f32_fp4 purely as an FP4-to-F32 unpack (scale byte 0x7f, opsel 0) and applies the per-thread scale in registers, sidestepping the cross-lane scale selection of the PK8 de-quant converter.

• MXFP8 de-quantization needed no intrinsic change (it is a pure C++ cast plus scale multiply); only the gfx1250 guard was removed.

• Removed all the temporary is_gfx1250() "not implemented" guards from the MXFP4 / MXFP8 quantize / de-quantize entry points added in Part 1.

• Unblocked the turbo (MFMA) FP8 GEMM binding.
  Dropped the #ifdef BUILD_TURBO_BACKEND stub branch in turbo_gemm.cu and moved the gfx1250 runtime guard down into turbo_gemm_mxfp8_impl, so the host binding in turbo_gemm.cpp links uniformly and fails clearly at call time on gfx1250 rather than at build configuration.

Verification

• Both gfx1250 and gfx950 device compilations pass (gfx950 checked for regressions).
• Validated on a real gfx1250 GPU with quantize -> de-quantize round trips using the same tolerances as the unit tests:
  • 416 cases across MXFP4 / MXFP8, rowwise / colwise, 2D-block on/off, fused dual (with-transpose), multiple shapes and dtypes.
  • 194 cases across e4m3 / e5m2 / fp4, batched (B = 1, 4) and 2D-block; stochastic-rounding outputs differ across consecutive calls with a sane round-trip error.
  • Grouped MXFP8 rowwise (single and dual) round trips match, and grouped MXFP4 single and dual kernels run cleanly.

Note: tests/conftest.py still carries a blanket "not yet supported on gfx1250" skip from Part 1, so the above was validated with standalone scripts using the same ops and tolerances. Relaxing that skip is left out of this PR because other ops (e.g. GEMM) are still unsupported on gfx1250.

Checklist:

• The functionality is complete
• I have commented my code, particularly in hard-to-understand areas
• I have made corresponding changes to the documentation
• My changes generate no new warnings
• I have added tests that prove my fix is effective or that my feature works
• New and existing unit tests pass locally with my changes

[Copilot]

Pull request overview

This PR is part 2 of adding gfx1250 (CDNA5) support by making MXFP4/MXFP8 quantize/dequantize kernels run correctly with 32-lane wavefronts, and by adjusting the turbo (MFMA) FP8 GEMM binding to link more uniformly across builds.

Changes:

• Parameterizes quantization kernels by wavefront size (THREADS_PER_WARP), updating block sizing and converting some fixed thread↔work mappings into grid-stride loops.
• Adds gfx1250-specific conversion paths for MXFP4/MXFP8 (quantize/dequantize) using AMDGCN compiler builtins and removes “not implemented on gfx1250” runtime guards.
• Moves the gfx1250 “turbo GEMM unsupported” runtime guard from the PyTorch binding into turbo_gemm_mxfp8_impl and removes the previous stub-branch approach.

Reviewed changes

Copilot reviewed 6 out of 6 changed files in this pull request and generated 4 comments.

Show a summary per file

File	Description
csrc/pytorch/gemm/turbo_gemm.cpp	Removes the entry-point gfx1250 guard so the binding relies on the lower-level impl to reject gfx1250.
csrc/kernels/quantization/quantization_mxfp8.cu	Makes MXFP8 quantization wavefront-size aware; adds gfx1250 builtin conversion path; updates colwise write-out mapping.
csrc/kernels/quantization/quantization_mxfp4.cu	Makes MXFP4 quantization wavefront-size aware; adds gfx1250 builtin conversion paths; updates LDS init + colwise write-out mapping.
csrc/kernels/quantization/dequantization_mxfp8.cu	Removes gfx1250 “not implemented” guards for MXFP8 dequantization paths.
csrc/kernels/quantization/dequantization_mxfp4.cu	Adds gfx1250 builtin-based FP4 unpack path and removes gfx1250 “not implemented” guard.
csrc/kernels/gemm/turbo_gemm.cu	Removes BUILD_TURBO_BACKEND stub branch and moves gfx1250 rejection to runtime check inside turbo_gemm_mxfp8_impl.

---

💡 Add Copilot custom instructions for smarter, more guided reviews. Learn how to get started.