[maleadt]

Kernels that throw now raise a KernelException. It is thrown when the kernel's stream is synchronized, or when the host waits for memory the stream owns. Before, such kernels printed a message (for @gputhrow quirks) or failed silently:

julia> function kernel(a)
           a[get_global_id() + 1] = 1
           return
       end

julia> a = oneArray{Int}(undef, 1);

julia> @oneapi kernel(a)

julia> synchronize()
ERROR: KernelException: A BoundsError was thrown on Intel(R) Arc(TM) A770 Graphics: Out-of-bounds array access
For more details, run Julia with `-g2`, or launch the kernel with `@oneapi debug_level=2`.

At -g2, the exception also names the work-item and work-group that threw, followed by a device backtrace. At -g0, it only reports that an exception was thrown. The new debug_level keyword to @oneapi sets the level for a single kernel.

The design follows OpenCL.jl's exception reporting (JuliaGPU/OpenCL.jl#478), but is simpler because oneAPI.jl already has a stream per task:

• Report and claim word. Each oneStream has a report in host USM and a claim word in device USM.
  • The first work-item to throw claims the report by storing its id in the claim word with a CAS. It is then the only writer of the report, using plain stores, so host USM needs no atomics. DG2 and PVC do not advertise them.
  • Reports and claim words come from slabs shared by a device's streams. Kernels declare indirect access, so every separate allocation is made resident on every launch. With one allocation per stream, 100 live task streams raised the cost of an empty launch from 2.4 to 5.4 µs. With the slabs it stays at 2.4 µs.
• Host side. The host reads the report while synchronizing the stream (or draining it, for memory another task waits for), under the stream's lock. Checking the report adds no submission, wait or lock; only reading the strings of a reported exception copies and waits. The claim word is reset by a fill appended to the stream, ahead of any later kernel.
• Strings by address. Names, reasons, and the functions and files of backtraces are reported by the address of the constant holding them. The host reads them, but only at addresses that zeMemGetAddressRange places in an allocation, such as a module's constants, and only up to its end, so a corrupted report cannot fault the device. CUDA.jl does the same in Add hostcalls for device-to-host services CUDA.jl#3243. An earlier version copied the strings on the device, and GPUCompiler replicates those copies at every throw site and, at -g2, every frame. IGC then generated up to 20 times more code: building a GPUArrays getindex kernel took 0.17 s instead of 0.05 s at -g1, and 0.55 s at -g2 (10.5 s with a debug build of IGC, which timed out CI). It now takes 0.07 and 0.13 s (2.7 s with debug IGC).
• No per-work-item state. Throwing only costs code on the cold path. An earlier version kept @gputhrow's name and reason in private memory. That added 192 bytes of private memory to every bounds-checked kernel and slowed a vector add by 11%.
• @gputhrow. It now records its name and reason (e.g. "BoundsError", "Out-of-bounds array access") in the report. This replaces the device printf, which makes bounds-checked kernels faster than on main. On an A770, a 2^24-element bounds-checked vector add goes from 597 to 498 µs at -g1, the same as at -g0. At -g2, reporting is inlined, which IGC compiles faster, and the vector add is as fast as on main. Neither level uses private memory or spills registers to memory.

Exceptions in kernels launched on an explicitly created ZeCommandQueue are not reported, as there is no stream to hold the report.

This supersedes #621. Thanks to @michel2323 for the earlier work there.

Closes #65

[github-actions]

Documentation preview: https://juliagpu.github.io/oneAPI.jl/previews/PR731/

[codecov]

Codecov Report

❌ Patch coverage is 94.82759% with 9 lines in your changes missing coverage. Please review.
✅ Project coverage is 84.54%. Comparing base (705aeb0) to head (24ec33b).
⚠️ Report is 7 commits behind head on main.

Files with missing lines	Patch %	Lines
src/compiler/exceptions.jl	90.90%	8 Missing ⚠️
src/compiler/execution.jl	75.00%	1 Missing ⚠️

Additional details and impacted files

@@            Coverage Diff             @@
##             main     #731      +/-   ##
==========================================
- Coverage   84.67%   84.54%   -0.14%     
==========================================
  Files          52       52              
  Lines        4543     4554      +11     
==========================================
+ Hits         3847     3850       +3     
- Misses        696      704       +8     

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