GitHub

Kernel Programming

This page documents the kernel programming API for writing custom GPU kernels in oneAPI.jl.

Kernel Launch

@oneapi [kwargs...] kernel(args...)

High-level interface for launching Julia kernels on Intel GPUs using oneAPI.

This macro compiles a Julia function to SPIR-V, prepares the arguments, and optionally launches the kernel on the GPU.

Keyword Arguments:

Macro Keywords (compile-time):

  • launch::Bool=true: Whether to launch the kernel immediately

Compiler Keywords:

  • kernel::Bool=false: Whether to compile as a kernel or device function
  • name::Union{String,Nothing}=nothing: Explicit name for the kernel
  • always_inline::Bool=false: Whether to always inline device functions
  • debug_level::Int: How much to report about exceptions thrown by the kernel (defaults to Julia's -g level; see KernelException)

Launch Keywords (runtime):

  • groups: Number of workgroups (required). Can be an integer or tuple.
  • items: Number of work-items per workgroup (required). Can be an integer or tuple.
  • queue::Union{oneStream,ZeCommandQueue}: Submission target (defaults to the task's global_stream). An explicit ZeCommandQueue submits through a per-dispatch command list and is not ordered with respect to the task's stream. Exceptions thrown by such kernels are not reported.

KernelException

Thrown when synchronizing a stream on which a kernel threw an exception. It carries the exception's name and reason (from debug level 1), and the work_item and work_group that threw it along with a backtrace (from debug level 2).

zefunction(f, tt; kwargs...)

Compile a Julia function to a Level Zero kernel function. This is the lower-level interface used by @oneapi. Returns a callable kernel object.

kernel_convert(x)

Convert arguments for kernel execution. This function is called for every argument passed to a kernel, allowing customization of argument conversion. By default, it converts oneArray to oneDeviceArray, and ZePtr{T} to LLVMPtr{T,AS.CrossWorkgroup}.

Basic Kernel Example

using oneAPIfunction vadd_kernel!(a, b, c)    i = get_global_id()    if i <= length(a)        @inbounds c[i] = a[i] + b[i]    end    returnendN = 1024a = oneArray(rand(Float32, N))b = oneArray(rand(Float32, N))c = similar(a)# Launch with 4 workgroups of 256 work-items each@oneapi groups=4 items=256 vadd_kernel!(a, b, c)

Launch Configuration

Workgroups and Work-Items

The oneAPI execution model is based on:

  • Work-items: Individual threads of execution (analogous to CUDA threads)
  • Workgroups: Groups of work-items that can synchronize and share local memory (analogous to CUDA blocks)
# 1D configuration@oneapi groups=10 items=64 kernel(args...)        # 640 work-items total# 2D configuration@oneapi groups=(10, 10) items=(8, 8) kernel(args...)  # 6400 work-items total# 3D configuration@oneapi groups=(4, 4, 4) items=(4, 4, 4) kernel(args...)  # 4096 work-items total

Determining Launch Configuration

# For simple element-wise operationsN = length(array)items = 256  # Typical workgroup sizegroups = cld(N, items)  # Ceiling division@oneapi groups=groups items=items kernel(array)

Compile Without Launch

You can compile a kernel without launching it:

# Compile the kernelkernel = @oneapi launch=false vadd_kernel!(a, b, c)# Launch later with different configurationskernel(a, b, c; groups=4, items=256)kernel(a, b, c; groups=8, items=128)

Device Intrinsics

Inside GPU kernels, you can use various intrinsics to query execution context and synchronize work-items.

Thread Indexing

# Global ID (unique across all work-items)i = get_global_id()      # 1D linear indexi = get_global_id(0)     # X dimensionj = get_global_id(1)     # Y dimensionk = get_global_id(2)     # Z dimension# Local ID (within workgroup)local_i = get_local_id()   # 1D linear indexlocal_i = get_local_id(0)  # X dimensionlocal_j = get_local_id(1)  # Y dimensionlocal_k = get_local_id(2)  # Z dimension# Workgroup IDgroup_i = get_group_id(0)  # X dimensiongroup_j = get_group_id(1)  # Y dimensiongroup_k = get_group_id(2)  # Z dimension# Workgroup sizelocal_size = get_local_size()   # Total work-items in workgrouplocal_size_x = get_local_size(0)local_size_y = get_local_size(1)# Global sizeglobal_size = get_global_size()   # Total work-itemsglobal_size_x = get_global_size(0)

2D Matrix Example

function matmul_kernel!(C, A, B)    # Get 2D indices    row = get_global_id(0)    col = get_global_id(1)    if row <= size(C, 1) && col <= size(C, 2)        sum = 0.0f0        for k in 1:size(A, 2)            @inbounds sum += A[row, k] * B[k, col]        end        @inbounds C[row, col] = sum    end    returnendM, N, K = 256, 256, 256A = oneArray(rand(Float32, M, K))B = oneArray(rand(Float32, K, N))C = oneArray{Float32}(undef, M, N)# Launch with 2D configurationitems = (16, 16)  # 16x16 work-items per workgroupgroups = (cld(M, items[1]), cld(N, items[2]))@oneapi groups=groups items=items matmul_kernel!(C, A, B)

Synchronization

# Barrier: synchronize all work-items in a workgroupbarrier()# Memory fences (ensure memory operations are visible)mem_fence()     # Both local and global memorylocal_mem_fence()   # Local memory onlyglobal_mem_fence()  # Global memory only

Local Memory

Local memory (workgroup-shared memory) enables cooperation between work-items:

function optimized_reduction!(result, input)    local_id = get_local_id()    local_size = get_local_size()    # Allocate local memory (shared within workgroup)    local_data = oneLocalArray(Float32, 256)    # Load into local memory    @inbounds local_data[local_id] = input[get_global_id()]    barrier()    # Tree reduction in local memory    stride = local_size ÷ 2    while stride > 0        if local_id <= stride            @inbounds local_data[local_id] += local_data[local_id + stride]        end        barrier()        stride ÷= 2    end    # First work-item writes result    if local_id == 1        @inbounds result[get_group_id()] = local_data[1]    end    returnend

Atomic Operations

For thread-safe operations on shared data:

# Atomic addoneAPI.atomic_add!(ptr, value)# Atomic exchangeold_value = oneAPI.atomic_xchg!(ptr, new_value)# Atomic compare-and-swapold_value = oneAPI.atomic_cas!(ptr, compare, new_value)# Atomic min/maxoneAPI.atomic_min!(ptr, value)oneAPI.atomic_max!(ptr, value)

Atomic operations on 8- and 16-bit values (e.g., Int8 or Float16, through UnsafeAtomics.jl, Atomix.jl or KernelAbstractions.jl) are generally implemented by the compiler as 32-bit compare-and-swap loops on the aligned 4-byte word that contains the value (a half-precision addition can use a native instruction instead, if the device supports it). Arrays allocated by oneAPI.jl are padded to a multiple of 4 bytes so that this word is always part of the allocation. When using such atomics on memory that oneAPI.jl did not allocate, e.g. a pointer obtained from another library, the whole aligned 4-byte word containing every element you access atomically must be accessible to the device. Other elements in that word, such as adjacent elements of the same array, must not be modified with plain stores while an atomic operation on the word may be in flight; access them atomically too.

Example histogram kernel:

function histogram_kernel!(hist, data, bins)    i = get_global_id()    if i <= length(data)        @inbounds val = data[i]        bin = clamp(floor(Int, val * bins) + 1, 1, bins)        oneAPI.atomic_add!(pointer(hist, bin), 1)    end    returnend

Kernel Restrictions

GPU kernels have certain restrictions:

  1. Must return nothing: Kernels cannot return values directly. Use output arrays instead.
  2. No dynamic memory allocation: Cannot allocate arrays inside kernels
  3. No I/O operations: Cannot print or write to files (use printf-style debugging with care)
  4. Limited recursion: Avoid or minimize recursive calls
  5. Type stability: Ensure type-stable code for best performance
# ❌ Bad: Returns a valuefunction bad_kernel(a)    return a[1] + 1end# ✅ Good: Returns nothing, uses output parameterfunction good_kernel!(result, a)    @inbounds result[1] = a[1] + 1    returnend

KernelAbstractions.jl

For portable GPU programming across CUDA, AMD, and Intel GPUs, use KernelAbstractions.jl:

using KernelAbstractionsusing oneAPI@kernel function generic_kernel!(a, b)    i = @index(Global)    @inbounds a[i] = a[i] + b[i]enda = oneArray(rand(Float32, 100))b = oneArray(rand(Float32, 100))backend = get_backend(a)  # oneAPIBackend()kernel! = generic_kernel!(backend)kernel!(a, b, ndrange=length(a))

See the KernelAbstractions.jl documentation for more details.

Debugging Kernels

See the Compiler and Reflection page for tools to inspect generated code and debug kernels.