Skip to content

Keep checked conversions out of the exception runtime - #568

Merged
maleadt merged 2 commits into
mainfrom
tb/consistent-runtime
Oct 21, 2026
Merged

maleadt merged 2 commits into
mainfrom
tb/consistent-runtime

Conversation

@maleadt

@maleadt maleadt commented Oct 19, 2026 •

Copy link
Copy Markdown
Member

Same issue as JuliaGPU/Metal.jl#567, fixed there in JuliaGPU/Metal.jl#1070.

On Julia 1.12, x < π in a kernel fails to compile: MPFR calls end up in device code because the BigFloat-based Float32(π, RoundDown) is not evaluated at compile time. The reason is an inference cycle through the exception runtime:

throw_inexacterror (overlay) → record_exception! → lock_output! → claim_output!
  → atomic_claim_device! → Int32(::Int64) → checked_trunc_sint → throw_inexacterror
  → get_local_id(1)     → UInt32(::Int64) (the `1u32` literal) → … → throw_inexacterror
  → UInt64(0)           → UInt64(::Int64) → … → throw_inexacterror

Julia 1.12 merges the effects of all frames in a cycle. The runtime's atomics reach UnsafeAtomics' is_native(), which is (correctly) not a consistent overlay, so that merge taints Int32(::Int64), UInt32(::Int64) and UInt64(::Int64) and disables concrete evaluation of any foldable code that uses them.

This PR keeps throwing calls out of the runtime: the constants are now written with %, and SPIRVIntrinsics' 1u32 is interpolated so the value is built when the method is defined. With that, no cycle passes through an overlay any more. I checked on 1.12 and 1.13 with PoCL.

The new concrete evaluation test fails on 1.12 without the fix.

On 1.10, which lacks consistent overlays, the conversion can't be evaluated at compile time at all. SPIRVIntrinsics now overrides the π/ℯ conversions with their precomputed values there, as Metal.jl does.

oneAPI.jl isn't affected today, because its exception runtime is a stub. It does use the same get_local_id code, though, so the SPIRVIntrinsics change is worth releasing.

@maleadt maleadt added the kernels Writing and compiling kernels: device intrinsics, codegen, KernelAbstractions back-end. label Oct 19, 2026
Without consistent overlays, the overlays on Base's throw paths keep the
BigFloat-based `Float32(π, RoundDown)` from being evaluated at compile time,
so comparing against `π` or `ℯ` in a kernel pulled MPFR calls into device
code. Override these conversions with their precomputed values on 1.10, as
Metal.jl does.
The exception runtime reached from the throw overlays performed checked
integer conversions (`Int32(1)`, `UInt64(0)`, and `UInt32(1)` through the
`1u32` literal in `get_local_id`/`get_group_id`), whose error paths lead back
into `throw_inexacterror`. On Julia 1.12, inference merges the effects of all
frames in such a cycle, so the inconsistent overlay reached through the
runtime's atomics tainted Base's conversion methods and disabled concrete
evaluation of any foldable code using them. For example, `x < π` pulled
MPFR calls into the kernel.

Same issue as JuliaGPU/Metal.jl#567.
@maleadt
maleadt force-pushed the tb/consistent-runtime branch from d1428ec to 68865aa Compare October 21, 2026 08:52
@codecov

codecov Bot commented Oct 21, 2026 •

Copy link
Copy Markdown

Codecov Report

✅ All modified and coverable lines are covered by tests.
✅ Project coverage is 87.73%. Comparing base (9c27048) to head (68865aa).
⚠️ Report is 2 commits behind head on main.

Additional details and impacted files
@@           Coverage Diff           @@
##             main     #568   +/-   ##
=======================================
  Coverage   87.73%   87.73%           
=======================================
  Files          18       18           
  Lines        1737     1737           
=======================================
  Hits         1524     1524           
  Misses        213      213           

☔ View full report in Codecov by Harness.
📢 Have feedback on the report? Share it here.

🚀 New features to boost your workflow:
  • ❄️ Test Analytics: Detect flaky tests, report on failures, and find test suite problems.

@maleadt
maleadt merged commit 38aeb64 into main Oct 21, 2026
17 checks passed
@maleadt
maleadt deleted the tb/consistent-runtime branch October 21, 2026 19:59
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

kernels Writing and compiling kernels: device intrinsics, codegen, KernelAbstractions back-end.

Projects

None yet

Development

Successfully merging this pull request may close these issues.

1 participant