[CUDA] Add pack functions for FP8 types

[LJC00118]

Summary by CodeRabbit

• New Features
  • Added helper utilities to assemble packed FP8 vectors in E4 and E5 formats for sizes 2, 4, 8, 16, and 32. These constructors simplify creating and passing larger FP8 vector values from individual elements, improving FP8 data handling, readability, and consistency in GPU-related code paths.

[github-actions]

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

Please remember to run bash format.sh 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]

Walkthrough

Adds TL_DEVICE helper functions in src/tl_templates/cuda/cuda_fp8.h to construct packed FP8 vector types for both e4 and e5 formats at sizes 2, 4, 8, 16, and 32 elements, using recursive composition. No existing types modified; only new free functions returning constructed packed values.

Changes

Cohort / File(s)	Summary
FP8 packer helpers src/tl_templates/cuda/cuda_fp8.h	Introduces make_fp8_e4_{2,4,8,16,32}_t and make_fp8_e5_{2,4,8,16,32}_t functions that assemble packed vector structs from individual fp8 elements, composed hierarchically (e.g., 32 from two 16s).

Sequence Diagram(s)

sequenceDiagram
  autonumber
  participant U as Caller
  participant H as cuda_fp8.h
  rect rgba(230,240,255,0.5)
    note over U,H: Build fp8_e4_32_t via hierarchical packers
    U->>H: make_fp8_e4_32_t(x0..x15, y0..y15)
    activate H
    H->>H: make_fp8_e4_16_t(x0..x7, x8..x15)
    H->>H: make_fp8_e4_16_t(y0..y7, y8..y15)
    H-->>U: fp8_e4_32_t packed
    deactivate H
  end

Loading

Estimated code review effort

🎯 2 (Simple) | ⏱️ ~10 minutes

Poem

A nibble, a byte, FP8 in a row,
I stack little pebbles where bitstreams flow.
Two into four, then eight into more—
Sixteen, thirty-two, neat bundles I store.
Thump goes my paw: pack, pack, hooray! 🥕🐇

Pre-merge checks and finishing touches

❌ Failed checks (1 warning)

Check name	Status	Explanation	Resolution
Docstring Coverage	⚠️ Warning	Docstring coverage is 0.00% which is insufficient. The required threshold is 80.00%.	You can run @coderabbitai generate docstrings to improve docstring coverage.

✅ Passed checks (2 passed)

Check name	Status	Explanation
Description Check	✅ Passed	Check skipped - CodeRabbit’s high-level summary is enabled.
Title Check	✅ Passed	The title clearly and concisely summarizes the main change of adding pack functions for FP8 types in the CUDA context, which directly matches the new helper functions introduced in the pull request.

✨ Finishing touches

• 📝 Generate docstrings

🧪 Generate unit tests (beta)

• Create PR with unit tests
• Post copyable unit tests in a comment

---

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.}

[coderabbitai]

Actionable comments posted: 2

📜 Review details

Configuration used: CodeRabbit UI

Review profile: CHILL

Plan: Pro

📥 Commits

Reviewing files that changed from the base of the PR and between 7cd0da9 and fe15b64.

📒 Files selected for processing (1)

• src/tl_templates/cuda/common.h (1 hunks)

🧰 Additional context used

🧬 Code graph analysis (1)

src/tl_templates/cuda/common.h (1)

src/tl_templates/cuda/cuda_fp8.h (9)

• fp8_e4_2_t (9-12)
• fp8_e4_4_t (14-19)
• fp8_e4_8_t (21-24)
• fp8_e5_16_t (61-64)
• fp8_e4_32_t (31-42)
• fp8_e5_2_t (44-47)
• fp8_e5_4_t (49-54)
• fp8_e5_8_t (56-59)
• fp8_e5_32_t (66-77)

🪛 GitHub Actions: CI Test on Metal

src/tl_templates/cuda/common.h

[error] 1-1: Code formatted by format.sh. Changes not staged for commit. Please review and commit the formatted changes.

🔇 Additional comments (9)

src/tl_templates/cuda/common.h (9)

117-123: LGTM!

The function correctly packs two fp8_e4_t values into fp8_e4_2_t by direct field assignment.

---

125-132: LGTM!

The function correctly packs four fp8_e4_t values into fp8_e4_4_t by direct field assignment.

---

134-141: LGTM!

The function correctly packs eight fp8_e4_t values into fp8_e4_8_t by hierarchically composing two fp8_e4_4_t structures.

---

154-169: LGTM after the previous fix!

The function correctly packs thirty-two fp8_e4_t values into fp8_e4_32_t by hierarchically composing two fp8_e4_16_t structures. This will compile successfully once the make_fp8_e4_16_t function name is corrected.

---

171-177: LGTM!

The function correctly packs two fp8_e5_t values into fp8_e5_2_t by direct field assignment.

---

179-186: LGTM!

The function correctly packs four fp8_e5_t values into fp8_e5_4_t by direct field assignment.

---

188-195: LGTM!

The function correctly packs eight fp8_e5_t values into fp8_e5_8_t by hierarchically composing two fp8_e5_4_t structures.

---

197-206: LGTM!

The function correctly packs sixteen fp8_e5_t values into fp8_e5_16_t by hierarchically composing two fp8_e5_8_t structures.

---

208-223: LGTM!

The function correctly packs thirty-two fp8_e5_t values into fp8_e5_32_t by hierarchically composing two fp8_e5_16_t structures.

[coderabbitai]

Actionable comments posted: 1

♻️ Duplicate comments (1)

src/tl_templates/cuda/common.h (1)

148-159: Critical: Fix return type, parameter types, and function name (duplicate issue).

This function has multiple type consistency issues that were previously flagged but remain unfixed:

1. Return type is fp8_e5_16_t but should be fp8_e4_16_t
2. Parameters y6 and y7 at Lines 153-154 are declared as fp8_e5_t but should be fp8_e4_t
3. Function name is make_fp8_e5_16_t but should be make_fp8_e4_16_t

The function body correctly uses fp8_e4_16_t and calls make_fp8_e4_8_t, indicating this is meant to pack e4 values, not e5. Additionally, make_fp8_e4_32_t at lines 171-174 calls make_fp8_e4_16_t, which does not exist due to the incorrect naming.

Apply this diff to fix the type mismatches and function name:

-// Pack sixteen fp8_e4_t values.
-TL_DEVICE fp8_e5_16_t make_fp8_e5_16_t(fp8_e4_t x0, fp8_e4_t x1, fp8_e4_t x2,
+// Pack sixteen fp8_e4_t values.
+TL_DEVICE fp8_e4_16_t make_fp8_e4_16_t(fp8_e4_t x0, fp8_e4_t x1, fp8_e4_t x2,
                                        fp8_e4_t x3, fp8_e4_t x4, fp8_e4_t x5,
                                        fp8_e4_t x6, fp8_e4_t x7, fp8_e4_t y0,
                                        fp8_e4_t y1, fp8_e4_t y2, fp8_e4_t y3,
-                                       fp8_e4_t y4, fp8_e4_t y5, fp8_e5_t y6,
-                                       fp8_e5_t y7) {
+                                       fp8_e4_t y4, fp8_e4_t y5, fp8_e4_t y6,
+                                       fp8_e4_t y7) {
   fp8_e4_16_t result;
   result.x = make_fp8_e4_8_t(x0, x1, x2, x3, x4, x5, x6, x7);
   result.y = make_fp8_e4_8_t(y0, y1, y2, y3, y4, y5, y6, y7);
   return result;
 }

📜 Review details

Configuration used: CodeRabbit UI

Review profile: CHILL

Plan: Pro

📥 Commits

Reviewing files that changed from the base of the PR and between 6a47b0e and e2a1fee.

📒 Files selected for processing (1)

• src/tl_templates/cuda/common.h (2 hunks)

🧰 Additional context used

🧬 Code graph analysis (1)

src/tl_templates/cuda/common.h (1)

src/tl_templates/cuda/cuda_fp8.h (9)

• fp8_e4_2_t (9-12)
• fp8_e4_4_t (14-19)
• fp8_e4_8_t (21-24)
• fp8_e5_16_t (61-64)
• fp8_e4_32_t (31-42)
• fp8_e5_2_t (44-47)
• fp8_e5_4_t (49-54)
• fp8_e5_8_t (56-59)
• fp8_e5_32_t (66-77)

🔇 Additional comments (3)

src/tl_templates/cuda/common.h (3)

20-21: LGTM! Clear type aliases for FP8 types.

The type aliases provide convenient shorthand for the FP8 types and follow clear naming conventions.

---

119-146: LGTM! Pack functions follow consistent patterns.

The pack functions for 2, 4, and 8 element vectors are correctly implemented and follow a clear hierarchical pattern.

---

178-235: LGTM! All e5 pack functions are correctly implemented.

The e5 pack functions (2, 4, 8, 16, and 32 element variants) follow consistent patterns and have correct type signatures. They properly mirror the structure of the e4 pack functions.

[coderabbitai]

Actionable comments posted: 0

🧹 Nitpick comments (2)

src/tl_templates/cuda/cuda_fp8.h (2)

79-195: Consider adding __forceinline__ for consistency.

The existing assignment operators (lines 35-41, 70-76) use __device__ __forceinline__. For consistency and potential performance benefits, consider adding __forceinline__ to these packing functions, especially for the base cases (2 and 4 element packs) that are likely to be called frequently.

Example for make_fp8_e4_2_t:

-TL_DEVICE fp8_e4_2_t make_fp8_e4_2_t(fp8_e4_t x, fp8_e4_t y) {
+TL_DEVICE __forceinline__ fp8_e4_2_t make_fp8_e4_2_t(fp8_e4_t x, fp8_e4_t y) {

Apply similar changes to other packing functions if desired.

---

79-195: Consider adding tests for the new packing functions.

These new packing functions would benefit from unit tests to ensure correctness, especially for the larger packed types where the recursive composition could be prone to ordering errors.

Would you like me to help generate unit tests for these packing functions, or open a new issue to track this task?

📜 Review details

Configuration used: CodeRabbit UI

Review profile: CHILL

Plan: Pro

📥 Commits

Reviewing files that changed from the base of the PR and between 2be0c5e and 22f1b06.

📒 Files selected for processing (1)

• src/tl_templates/cuda/cuda_fp8.h (1 hunks)

🔇 Additional comments (5)

src/tl_templates/cuda/cuda_fp8.h (5)

79-96: LGTM! Base packing functions are correctly implemented.

The make_fp8_e4_2_t and make_fp8_e4_4_t functions correctly pack individual FP8 values into their respective struct types, with field assignments matching the struct definitions.

---

98-136: LGTM! Recursive packing strategy is well-designed.

The functions make_fp8_e4_8_t, make_fp8_e4_16_t, and make_fp8_e4_32_t correctly use recursive composition to build larger packed types from smaller ones, matching the hierarchical structure of the underlying structs. The parameter distribution is correct.

---

138-155: LGTM! Base packing functions for e5 format are correctly implemented.

The make_fp8_e5_2_t and make_fp8_e5_4_t functions correctly pack individual FP8 e5m2 values, mirroring the structure of the e4m3 counterparts.

---

157-195: LGTM! Recursive packing for e5 format is correctly implemented.

The functions make_fp8_e5_8_t, make_fp8_e5_16_t, and make_fp8_e5_32_t correctly implement the same recursive composition strategy as their e4 counterparts, maintaining consistency across both FP8 formats.

---

79-195: No changes needed for TL_DEVICE macro
TL_DEVICE already expands to __forceinline__ __device__, satisfying both device and inlining requirements.

[coderabbitai]

Actionable comments posted: 0

🧹 Nitpick comments (2)

src/tl_templates/cuda/cuda_fp8.h (2)

79-138: LGTM! FP8 E4 pack functions are correctly implemented.

The recursive composition pattern is clean and the logic correctly maps input parameters to struct fields. However, note a minor style inconsistency and an AI summary discrepancy:

1. The attribute ordering __forceinline__ __device__ differs from the existing code style __device__ __forceinline__ (see lines 35, 70). Consider maintaining consistency.
2. The AI summary mentions __DEVICE__ decorator, but the actual code uses lowercase __device__.

Apply this diff to align with the existing code style:

-__forceinline__ __device__ fp8_e4_2_t make_fp8_e4_2_t(fp8_e4_t x, fp8_e4_t y) {
+__device__ __forceinline__ fp8_e4_2_t make_fp8_e4_2_t(fp8_e4_t x, fp8_e4_t y) {

-__forceinline__ __device__ fp8_e4_4_t make_fp8_e4_4_t(fp8_e4_t x0, fp8_e4_t x1,
+__device__ __forceinline__ fp8_e4_4_t make_fp8_e4_4_t(fp8_e4_t x0, fp8_e4_t x1,

-__forceinline__ __device__ fp8_e4_8_t make_fp8_e4_8_t(fp8_e4_t x0, fp8_e4_t x1,
+__device__ __forceinline__ fp8_e4_8_t make_fp8_e4_8_t(fp8_e4_t x0, fp8_e4_t x1,

-__forceinline__ __device__ fp8_e4_16_t
+__device__ __forceinline__ fp8_e4_16_t

-__forceinline__ __device__ fp8_e4_32_t make_fp8_e4_32_t(
+__device__ __forceinline__ fp8_e4_32_t make_fp8_e4_32_t(

---

140-199: LGTM! FP8 E5 pack functions correctly mirror the E4 implementation.

The e5 variant functions follow the same correct recursive composition pattern. The same style note applies regarding __forceinline__ __device__ vs __device__ __forceinline__ ordering for consistency with existing code.

Apply this diff to align with the existing code style:

-__forceinline__ __device__ fp8_e5_2_t make_fp8_e5_2_t(fp8_e5_t x, fp8_e5_t y) {
+__device__ __forceinline__ fp8_e5_2_t make_fp8_e5_2_t(fp8_e5_t x, fp8_e5_t y) {

-__forceinline__ __device__ fp8_e5_4_t make_fp8_e5_4_t(fp8_e5_t x0, fp8_e5_t x1,
+__device__ __forceinline__ fp8_e5_4_t make_fp8_e5_4_t(fp8_e5_t x0, fp8_e5_t x1,

-__forceinline__ __device__ fp8_e5_8_t make_fp8_e5_8_t(fp8_e5_t x0, fp8_e5_t x1,
+__device__ __forceinline__ fp8_e5_8_t make_fp8_e5_8_t(fp8_e5_t x0, fp8_e5_t x1,

-__forceinline__ __device__ fp8_e5_16_t
+__device__ __forceinline__ fp8_e5_16_t

-__forceinline__ __device__ fp8_e5_32_t make_fp8_e5_32_t(
+__device__ __forceinline__ fp8_e5_32_t make_fp8_e5_32_t(

📜 Review details

Configuration used: CodeRabbit UI

Review profile: CHILL

Plan: Pro

📥 Commits

Reviewing files that changed from the base of the PR and between 22f1b06 and 7b18464.

📒 Files selected for processing (1)

• src/tl_templates/cuda/cuda_fp8.h (1 hunks)

⏰ Context from checks skipped due to timeout of 90000ms. You can increase the timeout in your CodeRabbit configuration to a maximum of 15 minutes (900000ms). (2)

• GitHub Check: build-test-metal
• GitHub Check: format-check

[coderabbitai]

Actionable comments posted: 0

🧹 Nitpick comments (2)

src/tl_templates/cuda/cuda_fp8.h (2)

79-138: LGTM! FP8 E4 pack functions are correctly implemented.

The pack functions correctly construct FP8 E4 vectors using recursive composition that matches the struct definitions. The logic is sound:

• Basic 2 and 4-element functions directly assign fields
• Larger functions (8, 16, 32) recursively compose from smaller packers
• Parameter ordering (x0-xN, y0-yN) correctly maps to struct fields {x, y}

For consistency with existing code (lines 35, 70), consider swapping the qualifier order:

-__forceinline__ __device__ fp8_e4_2_t make_fp8_e4_2_t(fp8_e4_t x, fp8_e4_t y) {
+__device__ __forceinline__ fp8_e4_2_t make_fp8_e4_2_t(fp8_e4_t x, fp8_e4_t y) {

Apply the same pattern to all new pack functions.

---

140-199: LGTM! FP8 E5 pack functions mirror the E4 implementation correctly.

The E5 pack functions are correctly implemented with identical structure to the E4 variants, ensuring consistency across both FP8 formats.

Same qualifier order suggestion applies here for consistency with the existing codebase.

📜 Review details

Configuration used: CodeRabbit UI

Review profile: CHILL

Plan: Pro

📥 Commits

Reviewing files that changed from the base of the PR and between 7b18464 and d3d643a.

📒 Files selected for processing (1)

• src/tl_templates/cuda/cuda_fp8.h (1 hunks)

⏰ Context from checks skipped due to timeout of 90000ms. You can increase the timeout in your CodeRabbit configuration to a maximum of 15 minutes (900000ms). (2)

• GitHub Check: build-test-nvidia
• GitHub Check: build-test-amd