[JayceSu98]

Generated CUDA for per-token cast benchmarks can reduce pre-quantized FP8/FP4 values through half2 absmax operations. The generated loads cast packed fp8_e4_2_t/fp4_e2_2_t storage with from_uint1<__half2>, but common.h only exposes from_uint1 for uint1 storage, so NVCC rejects overload resolution for the packed FP8/FP4 arguments.

This is needed for TileKernels workloads that exercise the packed FP8/FP4 per-token cast path to compile and run on the H100 CUDA backend.

Previous compile errors looked like:

<build-temp>/tvm_kernels.cu(45): error: no instance of function template "tl::from_uint1" matches the argument list
            argument types are: (fp8_e4_2_t)
        stage1_amax_fragment_pack[i_1] = tl::to_uint1(tl::max2(tl::from_uint1<__half2>(stage1_amax_fragment_pack[i_1]), tl::from_uint1<__half2>(tl::to_uint1(tl::abs2(tl::from_uint1<__half2>(*(fp8_e4_2_t*)(x_stage1_fragment_reshaped + ((i_1 * 16) + (rv * 2)))))))));
                                                                                                                                                                      ^
<tilelang>/src/tl_templates/cuda/common.h(696): note #3326-D: function template "tl::from_uint1" does not match because argument #1 does not match parameter
  template <typename T> TL_DEVICE T from_uint1(uint1 v) {

and:

<build-temp>/tvm_kernels.cu(46): error: no instance of function template "tl::from_uint1" matches the argument list
            argument types are: (fp4_e2_2_t)
        stage1_amax_fragment_pack[i_1] = tl::to_uint1(tl::max2(tl::from_uint1<__half2>(stage1_amax_fragment_pack[i_1]), tl::from_uint1<__half2>(tl::to_uint1(tl::abs2(tl::from_uint1<__half2>(*(fp4_e2_2_t*)(x_stage1_fragment_reshaped_packed + ((i_1 * 8) + rv))))))));
                                                                                                                                                                       ^
<tilelang>/src/tl_templates/cuda/common.h(696): note #3326-D: function template "tl::from_uint1" does not match because argument #1 does not match parameter
  template <typename T> TL_DEVICE T from_uint1(uint1 v) {

Add dtype-local conversion overloads in the CUDA FP8 and FP4 template headers. The FP8 bridge converts fp8_e4_2_t through __nv_fp8x2_storage_t and __nv_cvt_fp8x2_to_halfraw2(..., __NV_E4M3), while the FP4 bridge converts fp4_e2_2_t through the existing __tl_cvt_fp4x2_to_half2 helper. Keeping these bridges next to the packed dtype definitions avoids making common.h depend on dtype-specific packed storage.

Verified on H100 (sm_90, 114 SMs) with CUDA 13.2 and PyTorch 2.12.0+cu132:

• representative packed/TMA FP8 and FP4 per_token_cast benchmark cases: 2 passed
• full per_token_cast benchmark: 448 passed
• representative functional validation for engram_grad_w_reduce and packed/TMA per_token_cast cases: 8 passed

The original NVCC errors for tl::from_uint1 with fp8_e4_2_t and fp4_e2_2_t arguments no longer appear.

Summary by CodeRabbit

• New Features
  • Device-side support to convert packed 4-bit and 8-bit floating formats to half-precision on CUDA for using compressed FP4/FP8 values in half-precision computations.
  • Added a handle byte-offset utility to create tensors at byte offsets, plus compiler backend support to emit the corresponding call.
• Refactor
  • Metal module creation now delegates to a fallback-capable constructor to improve robustness.

[github-actions]

👋 Hi! Thank you for contributing to the TileLang project.

Please remember to run pre-commit run --all-files in the root directory of the project to ensure your changes are properly linted and formatted. This will help ensure your contribution passes the format check.

We appreciate you taking this step! Our team will review your contribution, and we look forward to your awesome work! 🚀

[coderabbitai]

Note

Reviews paused

It looks like this branch is under active development. To avoid overwhelming you with review comments due to an influx of new commits, CodeRabbit has automatically paused this review. You can configure this behavior by changing the reviews.auto_review.auto_pause_after_reviewed_commits setting.

Use the following commands to manage reviews:

• @coderabbitai resume to resume automatic reviews.
• @coderabbitai review to trigger a single review.

Use the checkboxes below for quick actions:

• ▶️ Resume reviews
• 🔍 Trigger review

📝 Walkthrough

Walkthrough

Adds CUDA device helpers tl::from_uint1<T> for packed FP4/FP8→__half2 conversions, introduces CutDSL handle_add_byte_offset and its CUDA codegen emission, and refactors Metal module creation to use the target fallback helper.

Changes

Precision Conversion Helpers

Layer / File(s)	Summary
FP4 to target type conversion src/tl_templates/cuda/cuda_fp4.h	Adds from_uint1<T>(fp4_e2_2_t) that converts packed fp4_e2_2_t to __half2 via __tl_cvt_fp4x2_to_half2, then reinterprets the result as T.
FP8 to target type conversion src/tl_templates/cuda/cuda_fp8.h	Adds from_uint1<T>(fp8_e4_2_t) that memcpy's bytes into __nv_fp8x2_storage_t, converts via __nv_cvt_fp8x2_to_halfraw2 with __NV_E4M3, and reinterprets the resulting __half2 as T.

CutDSL handle_add_byte_offset

Layer / File(s)	Summary
handle_add_byte_offset utility and codegen tilelang/contrib/cutedsl/utils.py, src/backend/cuda/codegen/codegen_cutedsl.cc	Exports handle_add_byte_offset(handle, byte_offset) (updates __all__), derives a byte-pointer from handle (or handle.iterator), constructs a 1-D tensor at the byte offset via make_tensor_at_offset, and adds codegen to emit tl.handle_add_byte_offset(arg0, arg1).

Metal module creation

Layer / File(s)	Summary
MetalModuleCreate fallback delegation src/runtime/metal/metal_module.h	Replaces direct FFI factory lookups with a call to target::MetalModuleCreateWithFallback(...) and adds the Metal fallback module header; removes prior LOG(FATAL) error path.

Sequence Diagram(s)

sequenceDiagram
  participant Caller
  participant memcpy_storage
  participant __nv_cvt_fp8x2_to_halfraw2
  Caller->>memcpy_storage: memcpy input into __nv_fp8x2_storage_t
  memcpy_storage->>__nv_cvt_fp8x2_to_halfraw2: call with __NV_E4M3
  __nv_cvt_fp8x2_to_halfraw2->>memcpy_storage: return __half2
  memcpy_storage->>Caller: return as T via reinterpret_cast

Loading

Estimated code review effort

🎯 4 (Complex) | ⏱️ ~45 minutes

Possibly related PRs

• tile-ai/tilelang#1438: Related FP8 helper additions and FP8x2→float2 conversion utilities.
• tile-ai/tilelang#1880: Related changes to FP4→half conversion helpers used by the new FP4 helper.
• tile-ai/tilelang#1724: Related packed FP4 storage and codegen changes for fp4_e2_2_t buffers.

Suggested reviewers

• LJC00118
• SiriusNEO
• LeiWang1999

Poem

🐰 I nibble bits and memcpy tunes,
Tiny floats dance beneath the moons,
FP4 and FP8 hop into half2's light,
Handles skip bytes, codegen takes flight,
A rabbit claps — conversions done right.

🚥 Pre-merge checks | ✅ 4 | ❌ 1

❌ Failed checks (1 warning)

Check name	Status	Explanation	Resolution
Docstring Coverage	⚠️ Warning	Docstring coverage is 33.33% which is insufficient. The required threshold is 80.00%.	Write docstrings for the functions missing them to satisfy the coverage threshold.

✅ Passed checks (4 passed)

Check name	Status	Explanation
Description Check	✅ Passed	Check skipped - CodeRabbit’s high-level summary is enabled.
Title check	✅ Passed	The title '[CUDA] Support packed fp8/fp4 pair reduction casts' directly and accurately summarizes the main changes: adding support for packed fp8/fp4 pair reductions with new from_uint1 conversion overloads in CUDA headers.
Linked Issues check	✅ Passed	Check skipped because no linked issues were found for this pull request.
Out of Scope Changes check	✅ Passed	Check skipped because no linked issues were found for this pull request.

_{✏️ Tip: You can configure your own custom pre-merge checks in the settings.}

✨ Finishing Touches

🧪 Generate unit tests (beta)

• Create PR with unit tests

---

Thanks for using CodeRabbit! It's free for OSS, and your support helps us grow. If you like it, consider giving us a shout-out.

❤️ Share

• X
• Mastodon
• Reddit
• LinkedIn

_{Comment @coderabbitai help to get the list of available commands and usage tips.}