llvm_passes: expand llvm.llround and llvm.llrint - #1406
Merged
Conversation
pvelesko
force-pushed
the
2026-07-29-github-1404-llround-clean
branch
from
August 20, 2026 12:08
0fe6c3b to
f74d3a7
Compare
Collaborator
Author
|
/run-aurora-ci |
pvelesko
force-pushed
the
2026-07-29-github-1404-llround-clean
branch
2 times, most recently
from
August 24, 2026 09:55
b4e9a23 to
ba3357b
Compare
pvelesko
force-pushed
the
2026-07-29-github-1404-llround-clean
branch
from
August 24, 2026 12:23
ba3357b to
9fbf37b
Compare
std::llround and std::llrint on a float expand to llvm.llround and llvm.llrint, and std::lround on a float expands to llvm.lround. None of these has a 1:1 OpenCL.std ExtInst, because every OpenCL.std rounding instruction is float to float, so a producer that does not expand them itself rejects the module. llvm.lrint is not reachable through std::lrint: chipStar declares lrint(float) in sp_math.hh and pulls it into namespace std, so that overload wins and routes to the devicelib entry point. The test reaches the intrinsic through __builtin_lrintf instead. llvm.llround, llvm.llrint and llvm.lrint have no translation today, so the test does not link. llvm.lround does lower natively and is covered here to pin that it keeps producing the same values.
llvm.lround, llvm.llround, llvm.lrint and llvm.llrint take a floating point value and return an integer. OpenCL.std has no instruction of that shape, since ceil, floor, rint, round and trunc are all float to float, so these four need a two instruction expansion rather than a 1:1 ExtInst and a producer that does not do it rejects the module: InvalidFunctionCall: Unexpected llvm intrinsic: llvm.llround.i64.f32 or, with the in-tree SPIR-V backend: LLVM ERROR: unable to legalize instruction: %8:iid(s64) = G_INTRINSIC_LRINT The failure surfaces at hipspv-link with no source location at all, and it cannot be headed off in the headers: libstdc++ declares these for float as constexpr, which HIP turns into implicitly __host__ __device__ functions, so they win overload resolution over anything chipStar can declare (a plain __device__ overload with the same signature is rejected outright). So expand them here: round in floating point with llvm.round or llvm.rint, which both producers map to the OpenCL round and rint ExtInsts, then convert. Calling the devicelib entry points instead does not work, because devicelib is linked in before the post-link passes run and only the symbols referenced at that point are pulled in. Measured against llvm-spirv and the in-tree SPIR-V backend of LLVM 22.1.0: intrinsic translator backend lround ok ok llround fails ok lrint fails fails llrint fails fails All four are handled rather than only the three that fail today, so the set is defined by the shape of the intrinsic instead of by which producer version happens to be in use, and the pass can be removed in one go once the producers cover them. lround already lowers natively to exactly this expansion (OpExtInst round, then OpConvertFToS) on both producers, so covering it changes no emitted code. Fixes #1404
pvelesko
force-pushed
the
2026-07-29-github-1404-llround-clean
branch
from
August 25, 2026 21:58
9fbf37b to
8f7adfb
Compare
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
1 participant
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.
Fixes #1404
libstdc++ declares std::llround and std::llrint for float as constexpr, which HIP turns into implicitly host device functions. They win overload resolution in device code over anything chipStar can declare (a plain device overload with the same signature is rejected outright, see the analysis in #1404) and expand to __builtin_llroundf, so the module ends up with an llvm.llround intrinsic that llvm-spirv has no translation for. The failure surfaces at hipspv-link with no source location.
The new pass rounds in floating point with llvm.round or llvm.rint, which the translator maps to the OpenCL round and rint ExtInsts, then converts.
Calling the devicelib entry points (__chip_llround_f32 and friends) was the first attempt and does not work: devicelib is linked in before the post-link passes run, so a symbol first referenced by a post-link pass is never pulled in and the program build fails with
unresolved external symbol __chip_llround_f32.New test tests/devicelib/TestStdLlroundFloat.cpp, verified on Intel Arc B570 (OpenCL): does not link before, PASSED after. It pins down both rounding modes (llround(2.5f) is 3, llrint(2.5f) is 2) for float and double.
Found while getting the Kokkos HIP backend running on chipStar; this is what blocks Kokkos_CoreUnitTest_HIP from linking.
This pass is a bridge, not the real fix. The real fix belongs in the SPIR-V producers, and both are open upstream: llvm/llvm-project#217618 llvm/llvm-project#217618 for the in tree backend, and KhronosGroup/SPIRV-LLVM-Translator#3987 KhronosGroup/SPIRV-LLVM-Translator#3987 for the translator. Neither reaches chipStar automatically, and the LLVM 21 lanes never receive either, so the pass is still required. Removal is tracked in #1476.