GitHub

Kernel Programming

For maximum performance or custom operations not covered by high-level array abstractions, you can write custom kernels in Julia that execute on the GPU.

The @oneapi Macro

The @oneapi macro is used to launch a kernel on the device. It takes configuration arguments like the number of items (threads) and groups (blocks).

using oneAPIfunction kernel(a, b)    i = get_global_id()    if i <= length(a)        @inbounds a[i] += b[i]    end    returnenda = oneArray(rand(Float32, 100))b = oneArray(rand(Float32, 100))# Launch configurationitems = 100groups = 1@oneapi items=items groups=groups kernel(a, b)

KernelAbstractions.jl

For portable kernel programming, it is highly recommended to use KernelAbstractions.jl. This allows you to write kernels that work on CPU, CUDA, ROCm, and oneAPI.

using KernelAbstractions, oneAPI@kernel function my_kernel!(a, b)    i = @index(Global, Linear)    @inbounds a[i] += b[i]end# Get the backendbackend = get_backend(a)# Instantiate the kernelk = my_kernel!(backend)# Launch with configurationk(a, b; ndrange=length(a))

Device Intrinsics

Inside a kernel, you can use various intrinsics to interact with the hardware:

  • get_global_id(): Get the global thread ID.
  • get_local_id(): Get the local thread ID within a workgroup.
  • get_group_id(): Get the workgroup ID.
  • barrier(flags): Synchronize threads within a workgroup.

These are provided by SPIRVIntrinsics.jl and correspond to the standard OpenCL built-in functions. Note that the indices they return are 1-based, so they can be used to index Julia arrays directly. See Device Intrinsics for the full list.

Dynamic Memory Allocation

Kernels can allocate Julia objects, such as a Ref passed to a @noinline function or a boxed value in an Any field. Allocations that survive optimization use a 1 KiB heap private to each work-item. Each allocation is rounded up to 16 bytes, and memory is only reclaimed when the work-item exits. Allocated objects must not be shared with other work-items or retained across kernel launches.

When the heap is exhausted, the work-item prints an error and throws an exception (see Exceptions), so kernel output may be incomplete:

ERROR: Out of dynamic GPU memory (trying to allocate 4 bytes)

Kernels without remaining allocations do not reserve an arena. The device compiler may optimize away some heap storage, but allocations can increase private-memory use and reduce performance. Avoid repeated allocations in loops: even short-lived objects consume heap space until the work-item exits.

Exceptions

A work-item that throws an exception, e.g. on an out-of-bounds access, exits without completing its work. The other work-items keep running. The exception is reported as a KernelException when the stream the kernel was launched on is synchronized, e.g. by synchronize() or by copying an array to the host:

julia> function kernel(a)           a[get_global_id() + 1] = 1           return       endjulia> 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 accessFor more details, run Julia with `-g2`, or launch the kernel with `@oneapi debug_level=2`.

The amount of detail depends on the debug level the kernel is compiled with, which defaults to Julia's -g level. At level 0 only the fact that an exception was thrown is reported; at level 1 (the default) also its type and reason; at level 2 also the work-item that threw it and a backtrace. Only one exception is reported per synchronization. If the task that launched the kernel finishes without synchronizing, the exception is reported by the next synchronize() of another task, unless something waited for the kernel's results first. Exceptions in kernels launched on an explicitly created ZeCommandQueue are not reported.