API Reference

Vectorized Memory Access

Basic Operations

KernelIntrinsics.vload — Function

vload(A::AbstractArray{T}, idx, ::Val{Nitem}, ::Val{Rebase}=Val(true)) -> NTuple{Nitem,T}

Load Nitem elements from array A as a tuple, using vectorized memory operations on GPU.

Arguments

A: Source array
idx: Starting index
Nitem: Number of elements to load (must be a power of 2)
Rebase: Indexing mode (default: Val(true))

Indexing Modes

Val(true) (rebased): Index is multiplied by Nitem, so idx=2 loads elements [5,6,7,8] for Nitem=4. This mode generates optimal aligned vector loads (ld.global.v4).
Val(false) (direct): Loads starting directly at idx, so idx=2 loads elements [2,3,4,5]. Handles misaligned access by decomposing into smaller aligned loads.

Example

a = CuArray{Int32}(1:16)

# Rebased indexing (default): idx=2 → loads elements 5,6,7,8
values = vload(a, 2, Val(4))  # returns (5, 6, 7, 8)

# Direct indexing: idx=2 → loads elements 2,3,4,5
values = vload(a, 2, Val(4), Val(false))  # returns (2, 3, 4, 5)

Dynamic Alignment Operations

These functions handle arbitrary starting indices by computing alignment at runtime and dispatching to the appropriate vectorized instruction pattern.

KernelIntrinsics.vload_multi — Function

vload_multi(A::AbstractArray{T}, i, ::Val{Nitem}) -> NTuple{Nitem,T}

Load Nitem elements from array A starting at index i, automatically handling arbitrary alignment.

Computes alignment at runtime (mod = (pointer_offset + i - 1) % Nitem + 1) and dispatches to a statically-compiled load pattern via a switch table. This generates a mix of ld.global.v4, ld.global.v2, and scalar loads to maximize throughput.

Example

src = cu(Int32.(1:100))

# Works for any starting index — alignment handled automatically
values = vload_multi(src, 7, Val(8))  # loads elements 7:14

Memory Ordering

KernelIntrinsics.@access — Macro

@access [Scope] [Ordering] expr

Perform a memory load or store with specified scope and ordering semantics.

This macro provides fine-grained control over memory ordering for lock-free synchronization patterns on GPU. It generates appropriate ld.acquire or st.release PTX instructions.

Arguments

Scope (optional): Visibility scope, one of Device (default), Workgroup, or System
Ordering (optional): Memory ordering (see below)
expr: Either a load (var = array[idx]) or store (array[idx] = value) expression

Orderings

For loads (default: Acquire):

Acquire: Subsequent reads see all writes before the corresponding release
Relaxed: No ordering guarantees
Volatile: Volatile load (scope-less)
Weak: Weak load (scope-less)

For stores (default: Release):

Release: Prior writes are visible before this store
Relaxed: No ordering guarantees
Volatile: Volatile store (scope-less)
Weak: Weak store (scope-less)

AcqRel and SeqCst are not valid for individual loads/stores (use @fence instead). Volatile and Weak cannot have an explicit scope.

Syntax Forms

@access array[idx] = value          # Release store (default)
@access var = array[idx]            # Acquire load (default)
@access array[idx]                  # Acquire load, returns value directly

@access Release array[idx] = value  # Explicit ordering
@access Acquire var = array[idx]    # Explicit ordering
@access Device Release array[idx] = value  # Explicit scope and ordering

Example

@kernel function producer_consumer(X, Flag)
    if @index(Global, Linear) == 1
        X[1] = 42
        @access Flag[1] = 1  # Release store: X[1]=42 visible before Flag[1]=1
    end

    # Other threads wait
    while (@access Acquire Flag[1]) != 1
    end
    # Now X[1] is guaranteed to be 42
end

Warp Operations

KernelIntrinsics.@shfl — Macro

@shfl(direction, val, src, [warpsize=32], [mask=0xffffffff])

Perform a warp shuffle operation, exchanging values between lanes within a warp.

Arguments

direction: Shuffle direction (Up, Down, Xor, or Idx)
val: Value to shuffle (supports primitives, structs, and NTuples)
src: Offset (for Up/Down), XOR mask (for Xor), or source lane (for Idx)
warpsize: Warp size (default: 32)
mask: Lane participation mask (default: 0xffffffff for all lanes)

Example

@kernel function shfl_kernel(dst, src)
    I = @index(Global, Linear)
    val = src[I]
    
    shuffled = @shfl(Up, val, 1)      # Get value from lane below
    shuffled = @shfl(Down, val, 1)    # Get value from lane above
    shuffled = @shfl(Xor, val, 1)     # Swap with adjacent lane
    shuffled = @shfl(Idx, val, 0)     # Broadcast lane 0 to all
    
    dst[I] = shuffled
end