Support Functions Typed with __nv_bfloat16_raw

[isVoid]

Currently Numbast's type system does not recognize __nv_bfloat16_raw data type. Certain CUDA C++ functions takes the raw type as input, so we also provide those support here.

Summary by CodeRabbit

Release Notes

• New Features

  • Added support for bfloat16 raw type conversions, enabling bidirectional conversion between bfloat16 values and their raw representations in CUDA kernels.
  • Extended type system to recognize and properly handle the new bfloat16 raw variant.

• Tests

  • Added comprehensive test coverage for bfloat16 raw conversion functionality.

_{✏️ Tip: You can customize this high-level summary in your review settings.}

[coderabbitai]

Walkthrough

Adds support for a new Numba CUDA type __nv_bfloat16_raw with type mappings and imports, introduces two device conversion helpers in test CUDA headers, and adds a test that verifies round-trip conversion between nv_bfloat16 and __nv_bfloat16_raw.

Changes

Cohort / File(s)	Summary
Type mapping support numbast/src/numbast/static/renderer.py, numbast/src/numbast/static/types.py, numbast/src/numbast/types.py	Added handling for __nv_bfloat16_raw including import of the internal Numba BF16 type and mapping to the renderer/type-string logic.
Test CUDA helpers numbast/src/numbast/static/tests/data/bf16.cuh	Added two inline device functions: bf16_to_raw and bf16_from_raw for conversions between nv_bfloat16 and __nv_bfloat16_raw.
Test coverage numbast/src/numbast/static/tests/test_bf16_support.py	Added test_bindings_from_bf16raw to exercise the new raw bfloat16 conversion bindings and assert round-trip value correctness.

Estimated code review effort

🎯 3 (Moderate) | ⏱️ ~20 minutes

• Verify the import path and symbol name _type_unnamed1405307 from numba.cuda._internal.cuda_bf16 is correct and stable.
• Confirm consistent handling of __nv_bfloat16_raw across renderer and type-mapping layers.
• Review the new device functions in bf16.cuh for correct conversion semantics and ABI compatibility.
• Ensure the test's numerical tolerance and kernel setup are appropriate and reliable on CI hardware.

Poem

🐰 I hopped on CUDA paths anew,
A raw bf16 I brought to view.
To raw, from raw, the numbers play,
Round-trip checked — they leap and stay.
Tiny hops, big bytes — hooray! ✨

Pre-merge checks and finishing touches

❌ Failed checks (1 warning)

Check name	Status	Explanation	Resolution
Docstring Coverage	⚠️ Warning	Docstring coverage is 14.29% 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 describes the main change: adding support for functions typed with __nv_bfloat16_raw, which aligns with the PR's core objective of supporting the raw bfloat16 type in Numbast's type system.

✨ Finishing touches

• 📝 Generate docstrings

🧪 Generate unit tests (beta)

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

---

📜 Recent review details

Configuration used: Path: .coderabbit.yaml

Review profile: ASSERTIVE

Plan: Pro

📥 Commits

Reviewing files that changed from the base of the PR and between 0a3e183 and dcabfe2.

📒 Files selected for processing (1)

• numbast/src/numbast/static/renderer.py (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). (8)

• GitHub Check: Build Conda Python / build (arm64, 3.13, 12.9.1, rockylinux8)
• GitHub Check: Build Conda Python / build (arm64, 3.12, 12.9.1, rockylinux8)
• GitHub Check: Build Conda Python / build (arm64, 3.11, 12.9.1, rockylinux8)
• GitHub Check: Build Conda Python / build (amd64, 3.12, 12.9.1, rockylinux8)
• GitHub Check: Build Conda Python / build (amd64, 3.13, 12.9.1, rockylinux8)
• GitHub Check: Build Conda Python / build (arm64, 3.10, 12.9.1, rockylinux8)
• GitHub Check: Build Conda Python / build (amd64, 3.10, 12.9.1, rockylinux8)
• GitHub Check: Build Conda Python / build (amd64, 3.11, 12.9.1, rockylinux8)

🔇 Additional comments (1)

numbast/src/numbast/static/renderer.py (1)

112-117: Remove the verification request—the import pattern is correct for Numbast's CUDA bfloat16 bridging.

The past review comment regarding control flow has been addressed correctly—line 112 now uses elif.

The import from numba.cuda._internal.cuda_bf16 and the type name _type_unnamed1405307 are part of Numbast's documented approach to bridging CUDA's bfloat16 support to Python/Numba. The type name is compiler-generated from CUDA headers and is expected. For a bridging project like Numbast, relying on internal Numba APIs and CUDA-derived types is the appropriate implementation pattern and poses no stability concern beyond normal library maintenance practices.

---

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: 1

📜 Review details

Configuration used: Path: .coderabbit.yaml

Review profile: ASSERTIVE

Plan: Pro

📥 Commits

Reviewing files that changed from the base of the PR and between 22f06ec and 0a3e183.

📒 Files selected for processing (5)

• numbast/src/numbast/static/renderer.py (1 hunks)
• numbast/src/numbast/static/tests/data/bf16.cuh (1 hunks)
• numbast/src/numbast/static/tests/test_bf16_support.py (1 hunks)
• numbast/src/numbast/static/types.py (1 hunks)
• numbast/src/numbast/types.py (2 hunks)

🧰 Additional context used

🧬 Code graph analysis (2)

numbast/src/numbast/static/types.py (1)

numbast/src/numbast/static/renderer.py (2)

• BaseRenderer (15-137)
• _try_import_numba_type (104-129)

numbast/src/numbast/static/tests/test_bf16_support.py (1)

numbast/src/numbast/static/tests/conftest.py (1)

• make_binding (25-58)

🔇 Additional comments (6)

numbast/src/numbast/static/types.py (1)

62-64: LGTM - Consistent with existing bfloat16 handling.

The implementation correctly mirrors the pattern used for __nv_bfloat16 at lines 58-60 and properly registers the type import.

numbast/src/numbast/types.py (2)

12-12: Import from internal module may be fragile.

Same concern as in renderer.py: importing _type_unnamed1405307 from numba.cuda._internal.cuda_bf16 relies on an internal API with an auto-generated name. Ensure this is stable across the target numba versions.

---

88-89: LGTM - Early return pattern is correct.

The early return for __nv_bfloat16_raw follows the established pattern and correctly returns the imported type.

numbast/src/numbast/static/tests/test_bf16_support.py (1)

28-49: LGTM - Well-structured test for raw bfloat16 conversion.

The test appropriately exercises the round-trip conversion between bfloat16 and its raw representation, verifying the bindings work correctly.

numbast/src/numbast/static/tests/data/bf16.cuh (1)

13-19: LGTM - Conversion functions use CUDA BF16 API correctly.

The implementation correctly uses CUDA constructors for BF16 conversions. The __nv_bfloat16_raw constructor accepts the nv_bfloat16 value, and the __nv_bfloat16 constructor accepts the raw representation via its constexpr constructor signature. Both inline device function declarations are correct.

Note: __nv_bfloat16_raw may be deprecated in future CUDA releases; consider this for long-term maintenance.

numbast/src/numbast/static/renderer.py (1)

113-115: _type_unnamed1405307 is an auto-generated internal type name that is fragile across numba/numba-cuda versions.

The import relies on _type_unnamed1405307 from numba.cuda._internal.cuda_bf16, which is not a standard numba export but rather an auto-generated internal type from the Numbast bfloat16 bindings. The auto-generated naming convention (_type_unnamed{id}) makes this fragile:

• Different numba-cuda or Numbast builds may generate different type IDs
• No guarantee this type name remains stable across patch or minor version updates
• Not a documented public API, so changes are not tracked in release notes

Consider documenting the required numba/numba-cuda version constraints or adding a runtime check that validates this type's availability at import time with a clear error message if the expected type cannot be found.