[Common] Experimental CuTeDSL MXFP4 backend

[janekb04]

Description

This PR adds a CuTe DSL NVFP4 quantization backend. It is based on the work in #3137.

In this PR only a CuTe DSL version of quantize_transpose_tuned_1D is implemented. It handles the case of 1D bfloat16 to NVFP4 quantization, mirroring the capabilities of the original CUDA implementation in quantize_transpose_nvfp4_tuned_1D.cuh.

To be rebased once #3137 merges.

Overview

Below is the overview of #3137, which introduces the CuTe DSL infrastructure that this PR is based on.

Kernels using CuTe DSL are implemented in quantize_mxfp8.py.
These kernels are implemented as @cute.jit functions, which are wrapped in Kernel classes with a @cute.kernel __call__ method.
get_mxfp8_quantization_function is registered from the Python side with TVM FFI. It returns a Python function callable from the C++ side that runs the CuTE DSL kernel for a particular configuration.

The C++ side of the TVM FFI bridge is implemented in quantize_mxfp8_cutedsl.cuh. The mxfp8_quantize_cutedsl function has the same interface as the original mxfp8::quantize. It tries to call the Python side implementation.
The MXFP8QuantConfig is what gets carried over between the sides. The kernel is instantiated once per every config. It is like the kernel's template parameters.

quantize.cuh simply tries to call the CuTe DSL version if it is available instead of the regular implementation.

This PR mirrors the structure of #3137:

• The CuTe DSL kernel lives in quantize_transpose_nvfp4.py (vs. quantize_mxfp8.py), registering get_nvfp4_quantization_function (vs. get_mxfp8_quantization_function) with TVM FFI.
• The C++ side of the bridge is quantize_transpose_nvfp4_cutedsl.cuh (vs. quantize_mxfp8_cutedsl.cuh), exposing nvfp4_quantize_transpose_cutedsl (vs. mxfp8_quantize_cutedsl) with the same interface as nvfp4::quantize_transpose.
• The per-config instantiation parameters are carried in NVFP4QuantConfig (vs. MXFP8QuantConfig).
• quantize.cuh tries cutedsl_backend::nvfp4_quantize_transpose_cutedsl first and falls back to the CUDA nvfp4::quantize_transpose kernel if it declines (same fallback pattern as the MXFP8 path).

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

Support for the CuTe DSL backend is implemented in a top-down manner, one commit at a time:

• Add the CuTe DSL NVFP4 backend dispatch path to quantize.cuh (done)
  • Route the optimized 1D NVFP4 quantize-transpose path through the new CuTe DSL backend, falling back to the existing CUDA kernel when the backend declines to handle the case.
  • Add the FP4 dtype mapping to the TVM FFI bridge (tvm_ffi_bridge.h) and an empty stub header for the backend.
• Add the TVM FFI bridge scaffolding (done)
  • Establish the general structure and file layout of the backend, copied from the MXFP8 backend.
  • The C++ side is scaffolded in quantize_transpose_nvfp4_cutedsl.cuh, mirroring quantize_mxfp8_cutedsl.cuh.
  • The Python side is scaffolded in quantize_transpose_nvfp4.py, mirroring quantize_mxfp8.py.
  • Kernel instantiation parameters and validation are left as todos.
• Adapt the CuTe DSL scaffolding to the NVFP4 case (done)
  • Populate NVFP4QuantizeConfig with the kernel instantiation parameters (stochastic rounding, fast math, row-scaled NVFP4 and transpose flags) on both the C++ and Python sides, mirroring quantize_transpose_nvfp4.cuh:1441. The config values are derived from the quantization config and output tensor and forwarded across the TVM FFI boundary.
  • Add the input/output validation checks mirroring the CUDA implementation, guarding the backend behind FP4_TYPE_SUPPORTED.
• Pass tensors from the C++ side (done)
  • Wrap every kernel argument (input, rowwise/columnwise output + scales, rowwise/columnwise amax, noop flag, RNG state, stream) in a DLTensorWrapper and actually invoke the Python-side entrypoint, instead of just resolving it.
• Implement CuTe DSL kernel compilation (done)
  • Build the fake/symbolic tensor signatures (carrying the TMA alignment/divisibility constraints and the config-dependent optional arguments) that fix the AOT-compiled ABI.
  • Compile the kernel with cute.compile(..., options="--enable-tvm-ffi") and register it under the config-derived key via tvm_ffi.register_global_func, with graceful fallback to the CUDA kernel on unsupported configs or compilation failure.
• Implement the CuTe DSL kernel's device-side logic (todo)
  • The @cute.kernel device function and the host-side launch (tile/thread/stage sizing, grid/block computation, TMA atom setup) in __call__ are still stubs.
  • Add a dedicated backend test (mirroring test_mxfp8_cutedsl_backend.py) once the kernel is functional.

Checklist:

• I have read and followed the contributing guidelines
• 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