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

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.

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.

Basic Kernel Example

using oneAPI

function vadd_kernel!(a, b, c)
    i = get_global_id()
    if i <= length(a)
        @inbounds c[i] = a[i] + b[i]
    end
    return
end

N = 1024
a = 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 operations
N = length(array)
items = 256  # Typical workgroup size
groups = 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 kernel
kernel = @oneapi launch=false vadd_kernel!(a, b, c)

# Launch later with different configurations
kernel(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 index
i = get_global_id(0)     # X dimension
j = get_global_id(1)     # Y dimension
k = get_global_id(2)     # Z dimension

# Local ID (within workgroup)
local_i = get_local_id()   # 1D linear index
local_i = get_local_id(0)  # X dimension
local_j = get_local_id(1)  # Y dimension
local_k = get_local_id(2)  # Z dimension

# Workgroup ID
group_i = get_group_id(0)  # X dimension
group_j = get_group_id(1)  # Y dimension
group_k = get_group_id(2)  # Z dimension

# Workgroup size
local_size = get_local_size()   # Total work-items in workgroup
local_size_x = get_local_size(0)
local_size_y = get_local_size(1)

# Global size
global_size = get_global_size()   # Total work-items
global_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
    return
end

M, N, K = 256, 256, 256
A = oneArray(rand(Float32, M, K))
B = oneArray(rand(Float32, K, N))
C = oneArray{Float32}(undef, M, N)

# Launch with 2D configuration
items = (16, 16)  # 16x16 work-items per workgroup
groups = (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 workgroup
barrier()

# Memory fences (ensure memory operations are visible)
mem_fence()     # Both local and global memory
local_mem_fence()   # Local memory only
global_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
    return
end

Atomic Operations

For thread-safe operations on shared data:

# Atomic add
oneAPI.atomic_add!(ptr, value)

# Atomic exchange
old_value = oneAPI.atomic_xchg!(ptr, new_value)

# Atomic compare-and-swap
old_value = oneAPI.atomic_cas!(ptr, compare, new_value)

# Atomic min/max
oneAPI.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
    return
end

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 value
function bad_kernel(a)
    return a[1] + 1
end

# ✅ Good: Returns nothing, uses output parameter
function good_kernel!(result, a)
    @inbounds result[1] = a[1] + 1
    return
end

KernelAbstractions.jl

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

using KernelAbstractions
using oneAPI

@kernel function generic_kernel!(a, b)
    i = @index(Global)
    @inbounds a[i] = a[i] + b[i]
end

a = 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.