SYCL*TLA · CuTe · Intel Xe · Audited against current repository docs

From Layout Algebra to Intel Xe XMX GEMM

A practical tutorial for engineers who know GPU programming but are new to the Intel Xe path in SYCL*TLA. The core mental model is Layout → Tensor → Tile → 2D block copy → register fragment → DPAS/XMX → store.

16-work-item subgroupsXE_LOAD_2D*XE_DPAS_TTTiledMMAHelperGlobal memory → GRF
0 · Start here

The one picture to keep in your head

CuTe is a C++ template library for describing and operating on hierarchical multidimensional layouts of data and execution. Its central abstractions are Layout and Tensor; Intel's Xe path adds Xe-specific copy and MMA atoms that plug into the same CuTe layout machinery.

Scope note: statements such as “16-work-item subgroup” and the copy/MMA constraints below refer to the Intel Xe path documented by the current SYCL*TLA repository. They should not be read as universal statements about every Intel GPU API or every possible kernel implementation.
Global tensorsA(M,K), B(N,K), C(M,N)
→
local_tile()select a work-group tile
→
2D block loadglobal/L1 → GRF
→
reorder() if neededcopy layout → MMA layout
→
gemm()DPAS on XMX
→
2D storeGRF partition → global
Beginner rule: when a CuTe type looks intimidating, first ask: what is its shape? what is its layout/stride? and which work-items or subgroups participate in it?
1 · Translation layer

CUDA/CUTLASS words vs Intel/SYCL words

The SYCL*TLA documentation inherits substantial terminology from upstream CUTLASS. The table below is a programming analogy, not a claim that NVIDIA and Intel hardware structures are physically identical.

CUDA / CUTLASS termIntel / SYCL*TLA termCorrect mental model for this tutorial
threadwork-itemOne SYCL execution item.
warpsub_groupThe lockstep execution group used by this Xe path; the documented kernels require 16 work-items per subgroup.
CTA / threadblockwork-groupA group of work-items that can synchronize and, when used, share local memory.
shared memory (SMEM)Shared Local Memory (SLM)On-chip work-group-shared scratchpad memory.
register file / register fragmentsGRFGeneral Register File storage holding copy fragments, MMA fragments, and accumulators.
Tensor CoreXMXXe Matrix Extensions, the hardware matrix engine used by DPAS.
mma.sync / HMMADPAS / XE_DPAS_TTThe matrix multiply-accumulate instruction/atom used by this CuTe Xe path.
cp.async / TMAXe 2D block operationsAn analogy for optimized data movement only; Xe 2D loads/prefetches have their own semantics and constraints.
2 · CuTe core

Layout = Shape + Stride = coordinate → index mapping

A CuTe Layout is a pair of Shape and Stride. Semantically, it maps a logical coordinate inside the shape to an index. Shapes and strides can contain static compile-time integers, dynamic runtime integers, or hierarchical tuples of them.

Example

// 4 × 4; mode 0 has stride 1
using L = Layout<
  Shape<_4,_4>,
  Stride<_1,_4>
>;

For coordinate (m,n), this layout maps to m*1 + n*4. Swapping the two strides changes the mapping without changing the logical coordinate interface.

Interactive layout explorer

Tap a cell to inspect its coordinate → index mapping.

Why layout algebra matters

composition

Functional composition is central to CuTe. For scalar coordinates, if R = A ∘ B, then R(c) = A(B(c)).

tile / divide

Layout algebra can factor a larger logical domain into tiles and residual coordinates. local_tile() is a convenient tensor-level operation for selecting a tile from a larger tensor.

partition

TiledMMA and TiledCopy partition functions map logical tensor coordinates onto the work-items/subgroups participating in an operation.

Key insight: CuTe moves much of the address, tiling, and ownership logic into composable layouts instead of handwritten index arithmetic.
3 · Put data behind the mapping

Tensor = Engine + Layout

A CuTe Tensor combines a layout with an engine. The engine is a wrapper around an iterator or array; the layout maps logical coordinates to indices used to access that engine. Data may live in global memory, shared/local memory, register memory, or be transformed/generated on the fly.

coordinate(m,k)
→
Layoutshape + stride
→
indexlogical → linear
→
Engineiterator / storage
→
valueA(m,k)

CuTe GEMM mode convention

For GEMM, CuTe uses the logical tensor modes A(M,K), B(N,K), and C(M,N). This differs from the conventional BLAS view of B as (K,N). Keeping K as the second mode of both inputs lets generic implementations reduce over the same logical mode for A and B. The physical memory layout is still determined by each tensor's strides.

4 · Intel-specific pieces

What the Intel Xe path adds to upstream CuTe

Work-groupcomputes an output tile
16-work-item subgroupcollective execution unit used here
2D copy atomrectangular global↔register operation
SubgroupTensorsubgroup-distributed register abstraction
XE_DPAS_TTXMX MMA atom

XE_LOAD_2D*, XE_PREFETCH_2D, XE_STORE_2D

Parameterized Xe copy atoms wrap hardware 2D block operations. Loads can be normal, VNNI-transformed, or transposed where supported; prefetch has no destination fragment and hints data into cache.

TiledMMAHelper

Constructs a TiledMMA from an MMA atom, a work-group tile layout, and a subgroup layout, including the permutation used to map subgroup work onto the tile.

SubgroupTensor

An Intel-specific tensor abstraction that models storage distributed across subgroup lanes. It appears in Intel Xe partitioning, copy, and MMA machinery.

Standard direct-to-GRF mainloop

The documented standard Xe GEMM mainloop streams data from global memory (often after prefetch to L1/L2) into GRF with 2D block loads. SLM is skipped in that path; it may still be useful for designs that need cross-subgroup sharing, tiles beyond 2D-load limits, or explicit multi-buffer staging.

5 · Compute atom

DPAS: the matrix operation underneath gemm()

XE_DPAS_TT<M,...> is the current CuTe Xe atom wrapping a DPAS matrix multiply-accumulate. In the documented API, M is 1–8, N is fixed at 16, and K = 256 / max(bit-width(A), bit-width(B)).

DPAS K explorer

K = 16
K = 256 / max(input element bit widths)

One DPAS atom computes

D[M × 16] += A[M × K] × B[K × 16]

Examples from the documented formula: BF16/FP16 → K=16, TF32 → K=8, INT8 → K=32, INT4 → K=64. With BF16 and M=8, the atom shape is therefore 8×16×16.

From one atom to a tiled MMA

// Documented BF16 starting configuration
using WGTile   = Shape<_256,_256,_32>;
using SGLayout = Layout<Shape<_8,_4,_1>,
                        Stride<_4,_1,_0>>;

using Tiled = typename TiledMMAHelper<
    MMA_Atom<XE_DPAS_TT<8, float, bfloat16_t>>,
    Layout<WGTile>,
    SGLayout
>::TiledMMA;

This BF16 starting configuration uses a 256×256 output tile and 32 subgroups arranged as an 8×4 grid over M×N. In the pure xe_gemm.cpp tutorial, however, the K tile is selected programmatically as either op.K or 2*op.K depending on operand type/layout conditions; _32 is therefore a useful BF16 baseline, not a universal K tile.

6 · End-to-end kernel

Walk through the Intel Xe CuTe GEMM dataflow

1
Receive CuTe tensors A, B, and C.
The tutorial kernel sees A(M,K), B(N,K), and C(M,N), then creates identity-coordinate proxy tensors for tiling.
2
Choose the MMA atom and tiled MMA.
choose_mma_op() selects a supported XE_DPAS_TT (or a BF16/FP16 fallback), and choose_tiled_mma() chooses the K tile and subgroup layout.
3
Use local_tile() to select this work-group's coordinates.
The proxy tensors become gA(BLK_M,BLK_K,k), gB(BLK_N,BLK_K,k), and gC(BLK_M,BLK_N).
4
Create 2D TiledCopy objects.
Factory functions select copy atoms according to element width, tensor layout, and MMA requirements.
5
Slice MMA and copy operations with local_id.
get_slice(local_id) maps each work-item into its role in the collective MMA/copy operation.
6
Allocate separate copy and MMA register fragments.
Copy fragments match the layout produced by the 2D load; MMA fragments match the layout required by DPAS.
7
Warm the prefetch pipeline.
XE_PREFETCH_2D issues cache hints for future K tiles and has no register destination.
8
Run the K loop.
Split-barrier arrive → 2D loads → future prefetch → optional register reorder → DPAS accumulation → split-barrier wait.
9
Store the accumulator partitions.
The output TiledCopy maps each register-resident accumulator partition to the appropriate 2D store operation(s); this should not be interpreted as one instruction storing the entire work-group tile.

The K-loop as a pipeline

1
barrier arrive
2
2D load
3
prefetch ahead
4
reorder if needed
5
DPAS
6
barrier wait
// Guarded teaching version of the documented mainloop
int prefetch_k = 0;

for (; prefetch_k < min(prefetch_dist, k_tile_count); ++prefetch_k) {
  prefetch(prefetch_a, A_prefetch_tile(prefetch_k));
  prefetch(prefetch_b, B_prefetch_tile(prefetch_k));
}

for (int k_tile = 0; k_tile < k_tile_count; ++k_tile, ++prefetch_k) {
  barrier_arrive(barrier_scope);

  copy(copy_a, A_tile(k_tile), a_copy_regs);
  copy(copy_b, B_tile(k_tile), b_copy_regs);

  if (prefetch_k < k_tile_count) {
    prefetch(prefetch_a, A_prefetch_tile(prefetch_k));
    prefetch(prefetch_b, B_prefetch_tile(prefetch_k));
  }

  reorder(a_copy_regs, a_mma_regs);
  reorder(b_copy_regs, b_mma_regs);
  gemm(mma, a_mma_regs, b_mma_regs, accum);

  barrier_wait(barrier_scope);
}

copy(copy_c, accum, C_partition);
Why two fragment layouts? The 2D copy atom and the MMA atom can require different register distributions. If the layouts match, the compiler can elide reorder(); otherwise it becomes a subgroup-scope register-to-register shuffle. The Intel performance guide describes this shuffle as replacing an SLM store-and-reload alternative.
7 · Data movement

Understand Xe 2D block copies before tuning GEMM

Xe 2D block I/O is subgroup-cooperative: 16 work-items cooperate on a rectangular memory region. CuTe exposes these operations through parameterized copy atoms and maps their results into subgroup-distributed register fragments.

API-version note: xe_2d_copy.md still explains the legacy XE_2D_* names. The current API uses parameterized atoms such as XE_LOAD_2D, XE_LOAD_2D_VNNI, XE_LOAD_2D_TRANSPOSE, XE_PREFETCH_2D, and XE_STORE_2D. The data-distribution concepts in the legacy document still apply.
Current CuTe atomPurposeImportant current constraint / role
XE_LOAD_2DNormal rectangular 2D loadGeneral row-major-style block load; supports block counts 1, 2, or 4.
XE_LOAD_2D_VNNI2D load with VNNI transformUsed for B-style VNNI packing; Bits must be 8 or 16.
XE_LOAD_2D_TRANSPOSE2D load with transposeCurrent parameterized atom supports 32- or 64-bit elements and has additional width/shape limits.
XE_PREFETCH_2D2D cache prefetchNo destination fragment; hints future data into L1/L2.
XE_STORE_2DRectangular 2D storeStore height is limited to 8 rows.

Legacy 32×16 U16 distribution example

For the documented unpacked XE_2D_U16x32x16_LD_N example, a 16-work-item subgroup loads 32 rows × 16 columns, and each work-item receives one column containing 32 elements.

work-item 0 → column 0
work-item 1 → column 1
…
work-item 15 → column 15

What VNNI means here

VNNI combines elements from multiple rows of one column and packs them into 32-bit values. It changes the register/storage format; it does not numerically transform the element values.

Current 2D-operation limits highlighted by the performance guide

ConstraintValueNotes
Base pointer alignment64 bytesAll 2D operations.
Pitch/stride alignment16 bytesThe guide notes that 4 bytes may work on some PVC configurations.
Width alignment4 bytesAll 2D operations.
Maximum load/prefetch height32 rowsNormal, VNNI, transpose loads, and prefetch.
Maximum store height8 rowsXE_STORE_2D.
Maximum total width64 bytesBits × Width ≤ 512.
Supported element widths8 / 16 / 32 / 64 bitsBits describes element width for the memory operation, not necessarily a unique C++ numeric type.
8 · First performance model

Tune only after you can explain the dataflow

KnobWhy increase/change it?What to watch
M / N tileMore compute per work-group can improve XMX utilization.Larger accumulator footprint, possible GRF spill, and fewer concurrent work-groups.
K tileFewer 2D-load issues across the full K reduction can amortize load overhead.Larger copy fragments and GRF pressure; K must be a multiple of the DPAS K dimension.
Subgroup count/layoutChanges how the work-group tile is distributed and can improve utilization/locality.Smaller per-subgroup tiles and diminishing returns; the documented baseline uses 32 subgroups.
Prefetch depthPrefetch future K blocks while current data computes on XMX.Too few stages expose memory latency; deeper staged pipelines can increase live state/GRF pressure. XE_PREFETCH_2D itself has no register destination.
Copy ↔ MMA layout matchA matching layout lets the compiler remove reorder().A mismatch becomes a register-to-register subgroup shuffle.
SLM stagingUseful for designs needing sharing, explicit buffering, or data beyond direct 2D-load limits.Extra copies and synchronization can negate the benefit.
Documented BF16/BMG starting point: Shape<_256,_256,_32>, an 8×4 subgroup layout (32 subgroups), and PipelineStages = 2 in the referenced BMG GEMM example. The pure xe_gemm.cpp tutorial separately uses a prefetch distance of 3. These are starting points, not universal optima.

Correctness-first tuning sequence

  1. Verify the required 16-work-item subgroup configuration.
  2. Verify 2D-copy base-pointer, pitch, width, and shape constraints.
  3. Profile whether the kernel is bandwidth-bound or compute-bound.
  4. Adjust M/N/K tile dimensions one change at a time and watch GRF spill.
  5. Adjust prefetch depth and measure rather than assuming more stages are better.
  6. Inspect whether reorder() is compiled away or remains a shuffle.
9 · Source-reading guide

How to read examples/cute/tutorial/xe_gemm.cpp

For a first pass, follow the dataflow instead of reading every template declaration in file order:

A
gemm_cute() — inspect launch geometry and the explicit kernel properties sub_group_size<16> and grf_size<256>.
B
choose_mma_op() — see how A/B/C element types select a supported XE_DPAS_TT<8,...>, with BF16 or FP16 fallback paths.
C
choose_tiled_mma() — see how the source chooses _K = op.K or 2*op.K and chooses an 8×4 or 4×8 subgroup layout.
D
gemm_device() setup — identify the identity tensors, local_tile(), copy factories, slices, partitions, and register fragments.
E
The K loop — follow split barrier → 2D load → prefetch → reorder → gemm() → split-barrier wait.
F
Output copy — inspect tCrC, tCgC, and copy(copy_c, tCrC, tCgC).

Three source-faithful lines that unlock the file

auto wg_tile = mma.tile_mnk();          // work-group M/N/K tile
auto thr_mma = mma.get_slice(local_id); // this work-item's MMA slice
gemm(mma, tCrA, tCrB, tCrC);           // DPAS-backed accumulation
10 · Check yourself

Five questions before you write your own kernel

CuTe keeps K in the second mode of both A and B, so generic GEMM code reduces over the same logical mode. Physical row/column organization is still determined by strides.
Its subgroups can load their required tiles directly from global memory/L1 into GRF with 2D block loads. SLM remains useful for other designs, such as cross-subgroup sharing or tiles beyond direct 2D-load limits.
The copy atom's register layout can differ from the MMA atom's required layout. Matching layouts let the compiler remove the reorder; otherwise it becomes a subgroup register shuffle.
BF16/FP16: 16; TF32: 8; INT8: 32; INT4: 64.
Its logical shape, layout/stride, memory/storage role, and which work-items/subgroups own or participate in each fragment.
References · audited 2026-08-17

Source trail and audit scope

Technical statements in this page were rechecked against the current main branch of Intel SYCL*TLA on August 17, 2026, plus the upstream CuTe GEMM convention where needed. The code blocks labeled as teaching versions are intentionally simplified; source-faithful snippets are labeled explicitly.