Repository navigation
[SYCL] Make sycl::ext::oneapi::bfloat16 operations constexpr - #23569
Open
inteltdelisle wants to merge 2 commits into
Open
inteltdelisle wants to merge 2 commits into
inteltdelisle wants to merge 2 commits into
Conversation
Mark bfloat16 operations as constexpr, where applicable. Implementation is reliant on sycl::bit_cast and therefore are constexpr only when __SYCL_BITCAST_IS_CONSTEXPR is defined. Conversions to and from sycl::half are only constexpr when sycl::half itself is. - Replace union type punning in to_float/from_float_to_uint16_t with sycl::bit_cast. - On device, use __builtin_is_constant_evaluated() to fall back to the portable implementation during constant evaluation instead of calling __devicelib_* functions or NVPTX inline asm. - Move NVPTX inline asm into separate non-constexpr helpers, since C++17 forbids asm and uninitialized variables in constexpr functions. - Make explicit operator bool() const. This was not a requirement but I don't believe there are any reasons not to make it const. Add a regression test covering constexpr evaluation on host, SPIR-V and NVPTX sm_80 devices. Signed-off-by: Delisle, Thierry <thierry.delisle@intel.com>
inteltdelisle
requested review from
0x12CC
and
a balanced review from Copilot
October 20, 2026 20:26
Contributor
There was a problem hiding this comment.
🟢 Approval recommended
No blocking issues were identified in the conditional constexpr paths, preserved runtime behavior, or regression coverage.
No open issues 🎉
What changed in this PR
Adds conditional constexpr support to SYCL’s bfloat16 operations while preserving device runtime implementations.
Changes:
- Replaces union type punning with
sycl::bit_cast. - Uses portable conversions during constant evaluation and isolates NVPTX assembly in runtime helpers.
- Adds compile-time regression coverage for host, SPIR-V, and NVPTX sm_80.
| File | Description |
|---|---|
| sycl/test/regression/constexpr-bfloat16.cpp | Tests constexpr operations, conversions, and rounding. |
| sycl/test/regression/constexpr-bfloat16-nvptx.cpp | Runs constexpr coverage for NVPTX sm_80. |
| sycl/include/sycl/ext/oneapi/bfloat16.hpp | Enables conditional constexpr operations and makes boolean conversion const. |
Review effort: Balanced
This branch has not been deployed
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Mark bfloat16 operations as constexpr, where applicable. Implementation is reliant on sycl::bit_cast and therefore are constexpr only when __SYCL_BITCAST_IS_CONSTEXPR is defined. Conversions to and from sycl::half are only constexpr when sycl::half itself is.
Add a regression test covering constexpr evaluation on host, SPIR-V and NVPTX sm_80 devices.