Global-to-LDS builtins for AMD GPUs

Global-to-LDS builtins for AMD GPUs#

AMD GPUs provide a family of builtins that load data from global or buffer memory directly into Local Data Share (LDS) without staging the values in Vector General-Purpose Registers (VGPRs). Because the loaded data never occupies a VGPR, these builtins reduce register pressure and can enable higher occupancy in VGPR-constrained kernels.

All builtins in this family are wavefront-wide operations: each lane provides its own source address, and the hardware writes lane k’s value to the LDS base address plus an implicit per-lane stride.

Architecture availability#

The following table summarizes architecture support for each builtin.

Builtin	CDNA	CDNA2	CDNA3	CDNA4	RDNA2	RDNA3	RDNA3.5	RDNA4
`__builtin_amdgcn_global_load_lds`	No	No	Yes	Yes	No	No	No	No
`__builtin_amdgcn_load_to_lds`	No	No	Yes	Yes	No	No	No	No
`__builtin_amdgcn_raw_ptr_buffer_load_lds`	No	No	Yes	Yes	Yes	No	No	No
`__builtin_amdgcn_struct_ptr_buffer_load_lds`	No	No	Yes	Yes	Yes	No	No	No

Flat-addressed builtins#

These builtins take a pointer directly as the source address.

`__builtin_amdgcn_global_load_lds`#

Signature and parameters for this builtin.

void __builtin_amdgcn_global_load_lds(const void* src,
                                       __shared__ void* dst_base,
                                       std::int32_t size,
                                       std::int32_t offset,
                                       std::int32_t aux);

Wavefront-wide gather from global memory to LDS, bypassing VGPRs. Each lane loads size bytes from its own global address; the hardware writes lane k’s value to dst_base + offset + k * stride where stride = 4 for size <= 4 and (on CDNA4 GPUs) 16 for size > 4.

Parameter	Description
`src`	Pointer to global address space. Each lane can hold a different address.
`dst_base`	Pointer to LDS address space base. Must be wave-uniform. The pointer is implicitly offset by `lane_id * stride`. `stride` is `4` for `size <= 4` and (on CDNA4 GPUs) `16` for `size > 4`.
`size`	Transfer size per lane in bytes. Must be a compile-time constant. CDNA3 supports 1, 2, and 4. CDNA4 additionally supports 12 and 16.
`offset`	Signed byte offset applied to both `src` and `dst_base`. Must be a compile-time constant. Encoded as a 13-bit signed immediate (valid range: `-4096` to `4095`).
`aux`	Cache policy bits. Must be a compile-time constant. See Cache policy (aux parameter).

`__builtin_amdgcn_load_to_lds`#

Signature and parameters for this builtin.

void __builtin_amdgcn_load_to_lds(const void* src,
                                   __shared__ void* dst_base,
                                   std::int32_t size,
                                   std::int32_t offset,
                                   std::int32_t aux);

Wavefront-wide gather from generic, global, or buffer memory to LDS, bypassing VGPRs. Each lane loads size bytes from its own address; the hardware writes lane k’s value to dst_base + offset + k * stride where stride = 4 for size <= 4 and (on CDNA4 GPUs) 16 for size > 4.

Parameter	Description
`src`	Generic, global, or buffer address pointer. Each lane can hold a different address.
`dst_base`	Pointer to LDS address space base. Must be wave-uniform. The pointer is implicitly offset by `lane_id * stride`. `stride` is `4` for `size <= 4` and (on CDNA4 GPUs) `16` for `size > 4`.
`size`	Transfer size per lane in bytes. Must be a compile-time constant. CDNA3 supports 1, 2, and 4. CDNA4 additionally supports 12 and 16.
`offset`	Signed byte offset applied to both `src` and `dst_base`. Must be a compile-time constant. Encoded as a 13-bit signed immediate (valid range: `-4096` to `4095`).
`aux`	Cache policy bits. Must be a compile-time constant. See Cache policy (aux parameter).

Buffer-addressed builtins#

These builtins take a buffer resource descriptor as the source, created with __builtin_amdgcn_make_buffer_rsrc. The three offset parameters (per-lane, wave-uniform, and compile-time) are combined to form the final source address.

Note

The buffer-addressed builtins are available on CDNA, CDNA2, CDNA3, CDNA4, and RDNA2. They’re not supported on RDNA3, RDNA3.5, or RDNA4.

`__builtin_amdgcn_raw_ptr_buffer_load_lds`#

Signature and parameters for this builtin.

void __builtin_amdgcn_raw_ptr_buffer_load_lds(
    __amdgpu_buffer_rsrc_t src,
    __shared__ void*       dst_base,
    std::int32_t           size,
    std::int32_t           voffset,
    std::int32_t           soffset,
    std::int32_t           offset,
    std::int32_t           aux);

For each lane in a wavefront, the load address (global source) is:

src_addr = src_base + soffset + offset + voffset + stride * lane_id

src_base and stride come from the buffer resource descriptor. The lane_id term is active only if the corresponding flag was set when creating the resource descriptor. Otherwise, it evaluates to zero.

For each lane in a wavefront, the store address (LDS destination) is:

dst_addr = dst_base + offset + lane_id * 4    // size <= 4
dst_addr = dst_base + offset + lane_id * 16   // size > 4 (CDNA4 only)

The lane_id term in the LDS address is always active.

Parameter	Description
`src`	Buffer resource descriptor (`__amdgpu_buffer_rsrc_t`). Created with `__builtin_amdgcn_make_buffer_rsrc`. The descriptor wraps a base pointer, stride, extent, and flags.
`dst_base`	Pointer to LDS address space base. Must be wave-uniform. The hardware adds `lane_id * stride` implicitly.
`size`	Transfer size per lane in bytes. Must be a compile-time constant. CDNA, CDNA2, CDNA3, and RDNA2 support 1, 2, and 4. CDNA4 additionally supports 12 and 16.
`voffset`	Per-lane byte offset into the buffer. Each lane can provide a different value, enabling gather-style reads from the buffer.
`soffset`	Wave-uniform byte offset into the buffer.
`offset`	Unsigned compile-time byte offset. Encoded as a 12-bit unsigned immediate (valid range: `0` to `4095`).
`aux`	Cache policy bits (compile-time constant). Bit 3 enables buffer swizzling on all supported architectures. The remaining bits differ by architecture. See Cache policy (aux parameter).

`__builtin_amdgcn_struct_ptr_buffer_load_lds`#

Signature and parameters for this builtin.

void __builtin_amdgcn_struct_ptr_buffer_load_lds(
    __amdgpu_buffer_rsrc_t src,
    __shared__ void*       dst_base,
    std::int32_t           size,
    std::int32_t           vindex,
    std::int32_t           voffset,
    std::int32_t           soffset,
    std::int32_t           offset,
    std::int32_t           aux);

Structured-buffer direct-to-LDS load. The source is described by a buffer resource descriptor created with __builtin_amdgcn_make_buffer_rsrc. The vindex parameter provides a per-lane index that is combined with lane_id in the source address calculation, enabling structured access patterns over the buffer.

For each lane in a wavefront, the load address (global source) is:

src_addr = src_base + soffset + offset + voffset + stride * (vindex + lane_id)

src_base and stride come from the buffer resource descriptor. The lane_id term is active only if the corresponding flag was set when creating the resource descriptor. Otherwise, it evaluates to zero.

For each lane in a wavefront, the store address (LDS destination) is:

dst_addr = dst_base + offset + lane_id * 4    // size <= 4
dst_addr = dst_base + offset + lane_id * 16   // size > 4 (CDNA4 only)

The lane_id term in the LDS address is always active.

Parameter	Description
`src`	Buffer resource descriptor (`__amdgpu_buffer_rsrc_t`). Created with `__builtin_amdgcn_make_buffer_rsrc`. The descriptor wraps a base pointer, stride, extent, and flags.
`dst_base`	Pointer to LDS address space base. Must be wave-uniform. The hardware adds `lane_id * stride` implicitly.
`size`	Transfer size per lane in bytes. Must be a compile-time constant. CDNA, CDNA2, CDNA3, and RDNA2 support 1, 2, and 4. CDNA4 additionally supports 12 and 16.
`vindex`	Per-lane index value.
`voffset`	Per-lane byte offset into the buffer. Each lane can provide a different value, enabling gather-style reads from the buffer.
`soffset`	Wave-uniform byte offset into the buffer.
`offset`	Unsigned compile-time byte offset. Encoded as a 12-bit unsigned immediate (valid range: `0` to `4095`).
`aux`	Cache policy bits (compile-time constant). Bit 3 enables buffer swizzling on all supported architectures. The remaining bits differ by architecture. See Cache policy (aux parameter).

Cache policy (`aux` parameter)#

All direct-to-LDS builtins accept an aux parameter whose individual bits control cache scope and temporal reuse hints. The bit layout differs by architecture generation. For the buffer-addressed builtins, bit 3 enables buffer swizzling on all supported architectures.

For background on AMD GPU cache hierarchy and coherence scopes, see the CDNA3 Instruction Set Architecture white paper and the Instruction Set Architecture (ISA) reference guides published on GPUOpen.

In a single-GPU kernel, wave scope with temporal reuse (aux = 0) is the typical choice.

CDNA3 and CDNA4#

Three control bits are relevant:

SC0 (Scope Control 0, bit 0) and SC1 (Scope Control 1, bit 4): together define the coherence scope.
NT (Non-Temporal, bit 1): 0 = expect temporal reuse, 1 = do not expect temporal reuse.

SC1	SC0	NT	L1 cache behavior	L2 cache behavior	Last-level cache behavior
0	0	0	Hit LRU	Hit LRU	Hit LRU
0	0	1	Miss Evict	Hit Stream	Hit Evict
0	1	0	Hit LRU / Miss LRU [1]	Hit LRU	Hit Evict / Hit LRU [1]
0	1	1	Miss Evict	Hit Stream	Hit Evict
1	0	0	Miss Evict	Hit LRU (1 L2) or coherent cache bypass (>1 L2)	Hit LRU
1	0	1	Miss Evict	Hit Stream (1 L2) or coherent cache bypass (>1 L2)	Hit Evict
1	1	0	Miss Evict	Coherent cache bypass	Hit LRU
1	1	1	Miss Evict	Coherent cache bypass	Hit Evict

The following terms describe cache line behavior in the table above:

Hit LRU – cache the data and replace the least recently used line when the cache is full. This is the normal temporal caching behavior.
Miss Evict – bypass the cache level; fetch data directly from the next level.
Hit Stream – on a cache hit, serve the data and evict the line; on a cache miss, fetch the data into the cache, use it once, then discard it.
Hit Evict – on a cache hit, immediately discard the line after serving the request.
Coherent cache bypass – when reading a cache line: if the line is modified, write it back to memory and reissue the read; if not modified, discard it and fetch a new line.

CDNA and CDNA2#

One control bit is relevant:

GLC (Global Level Coherent, bit 0): controls the first-level cache (L1).

GLC	L1 cache behavior	Notes
0	The load can read data from the L1 cache.
1	The load intentionally misses the L1 cache and reads from L2. If a matching line exists in L1, it is invalidated and L2 is re-read.	Depending on alignment, L2 might not be re-read for every lane in the same wavefront. If the address is aligned correctly, the first lane brings in the line from L2 (or beyond), and all other lanes in the wavefront read from the same L1 cache line.

RDNA2#

Three control bits are relevant:

GLC (Global Level Coherent, bit 0): controls the first-level cache (L0).
SLC (System Level Coherent, bit 1) and DLC (Device Level Coherent, bit 2): together control the behavior of L1 and L2.

GLC behavior:

GLC	L0 cache behavior	Notes
0	The load can read data from the L0 cache.
1	The load intentionally misses the L0 cache and reads from L2. If a matching line exists in L0, it is invalidated and L2 is re-read.	Depending on alignment, L2 might not be re-read for every lane in the same wavefront. If the address is aligned correctly, the first lane brings in the line from L2 (or beyond), and all other lanes in the wavefront read from the same L0 cache line.

SLC and DLC behavior:

SLC	DLC	L2 cache	L1 cache
0	0	Hit LRU	Hit LRU
0	1	Hit LRU	Miss Evict
1	0	Hit Stream	Hit LRU
1	1	Hit No Allocate	Miss Evict

The following terms describe cache line behavior in the table above:

Hit LRU – cache the data and replace the least recently used line when the cache is full. This is the normal temporal caching behavior.
Miss Evict – bypass the cache level; fetch data directly from the next level.
Hit Stream – on a cache hit, serve the data and evict the line; on a cache miss, fetch the data into the cache, use it once, then discard it.
Hit No Allocate – on a cache hit, serve the data and evict the line; on a cache miss, fetch the data directly from the next memory level without allocating a cache line.

Tip

When choosing between Hit Stream and Hit No Allocate, consider the spatial locality of the access pattern. For contiguous or regularly-strided reads, prefer Hit Stream: even though each cache line is discarded after one use, a single memory fetch serves multiple lanes that share the same line, so the fetch cost is amortized. For scattered or random reads, prefer Hit No Allocate: each fetch is unlikely to benefit more than one lane, so skipping cache allocation avoids polluting L2 with lines that will never be reused.

Global-to-LDS builtins for AMD GPUs

Contents

Global-to-LDS builtins for AMD GPUs#

Architecture availability#

Flat-addressed builtins#

__builtin_amdgcn_global_load_lds#

__builtin_amdgcn_load_to_lds#

Buffer-addressed builtins#

__builtin_amdgcn_raw_ptr_buffer_load_lds#

__builtin_amdgcn_struct_ptr_buffer_load_lds#

Cache policy (aux parameter)#

CDNA3 and CDNA4#

CDNA and CDNA2#

RDNA2#

`__builtin_amdgcn_global_load_lds`#

`__builtin_amdgcn_load_to_lds`#

`__builtin_amdgcn_raw_ptr_buffer_load_lds`#

`__builtin_amdgcn_struct_ptr_buffer_load_lds`#

Cache policy (`aux` parameter)#