A learning-oriented comparison of the programming model, data path, synchronization, and kernel design implications for GEMM and attention.
cp.async is an explicit asynchronous copy into shared memory. Intel prefetch is a non-blocking cache-prefetch request before a later load. They are related latency-hiding tools, but they are not direct equivalents.
cp.async creates a program-visible copy destination in shared memory. The kernel later waits for completion and consumes the staged tile.
Prefetch asks the memory system to bring data closer. A later block or global load is still required to move the tile into GRF registers.
| Property | NVIDIA classic cp.async |
Intel Xe prefetch |
|---|---|---|
| Primitive type | Asynchronous memory copy | Asynchronous cache-prefetch request |
| Typical source | Global memory | Global memory |
| Program-visible destination | Shared memory | No explicit user-addressable destination buffer |
| Does it place data in registers? | No | No |
| Later load required? | Yes, shared memory → registers | Yes, cache/global memory → GRF |
| Explicit completion/wait? | Yes, via async-group or barrier mechanisms | No comparable “prefetch complete” wait in the normal programming model |
| Deterministic software staging | Strong fit | Weaker; cache residency is not a program-visible staging contract |
| Typical kernel use | Double-buffered shared-memory pipeline | Prefetch distance + block load + GRF reuse |
// iteration k
cp.async global_tile[k + 1] -> smem[next];
compute(smem[current]);
cp.async.commit_group();
cp.async.wait_group();
swap(current, next);The destination buffer is explicit and addressable by the kernel.
// iteration k
prefetch(global_tile[k + distance]);
compute(tile_in_grf[k]);
// later
tile = block_load(global_tile[k + distance]);
// tile now arrives in GRFPrefetch and load are separate operations.
SLM is Intel’s work-group-local scratchpad memory and is conceptually similar to CUDA shared memory. However, Intel prefetch does not automatically fill SLM.
// conceptual Intel SLM staging
tile = block_load(global_ptr); // global/cache -> GRF
slm_store(slm_ptr, tile); // GRF -> SLM
work_group_barrier();
tile2 = slm_load(slm_ptr); // SLM -> GRF
compute(tile2);| NVIDIA concept | Closest Intel-side concept |
|---|---|
| Shared memory | SLM |
| Registers | GRF |
cp.async latency hiding | Prefetch + later block/global load |
cp.async explicit staging | Explicit load + SLM store + barrier |
| MMA | DPAS / matrix instructions |
cp.async plus shared memory.
cp.async bypasses L1.cp.async, cuda::memcpy_async, and newer TMA / cp.async.bulk are related but distinct mechanisms.local_mem_size is not L2
In SYCL, local_mem_size means the size of the local memory arena.
For GPU programming, this maps conceptually to CUDA shared memory or Intel SLM
rather than L2 cache.
dev.get_info<sycl::info::device::local_mem_size>()Means: how much local/shared memory is available for a work-group.
dev.get_info<sycl::info::device::global_mem_cache_size>()Means: the runtime-reported global-memory cache size. Depending on backend and device, this may correspond to an L2/global cache, but portable SYCL does not guarantee the exact cache-level name.
local_mem_size == L2 cache size // wrongLocal memory is software-managed scratchpad-style memory. L2 is hardware-managed cache. They are different concepts and normally different physical resources.
local_mem_size ≈ shared memory / SLM capability.
global_mem_cache_size ≈ reported global-memory cache capability.
Usually, L2 does not share the same physical memory as shared memory / SLM. The resource that may share or trade off with shared memory is usually the local L1/data-cache side, not the lower-level L2 cache.
On NVIDIA, shared memory and L1 can be carved out from a shared local on-chip resource on architectures that support configurable carveout. L2 is separate from that local pool.
Intel exposes SLM as work-group local memory. Some Intel APIs expose configuration choices around SLM-vs-data-cache behavior, but this is not the same as making SLM equal to L2.
This matters for Intel prefetch. Prefetch does not populate local_mem_size / SLM.
It requests that future global-memory data be brought into the cache hierarchy before a later load.
prefetch(global_tile[k + distance]);
// later
tile = block_load(global_tile[k + distance]); // cache/HBM -> GRFtile = block_load(global_ptr); // global/cache -> GRF
slm_store(slm_ptr, tile); // GRF -> SLM
work_group_barrier();
tile2 = slm_load(slm_ptr); // SLM -> GRF| Term / Query | Meaning | Managed by | Equivalent mental model |
|---|---|---|---|
local_mem_size |
Local memory arena available to a work-group | Software / kernel | CUDA shared memory, Intel SLM |
global_mem_cache_size |
Reported cache size for global-memory accesses | Hardware cache system | Global cache, often L2-like depending on backend |
| CUDA shared memory | Explicit per-block scratchpad | Software / kernel | Intel SLM |
| Intel SLM | Explicit per-work-group scratchpad | Software / kernel | CUDA shared memory |
| L1 / data cache | Local cache level near compute units | Hardware, sometimes carveout/config influenced | May share local resources with shared memory / SLM on some architectures |
| L2 cache | Lower-level/global cache | Hardware, sometimes policy influenced | Not a software-managed scratchpad |