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.
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 term | Intel / SYCL*TLA term | Correct mental model for this tutorial |
|---|---|---|
| thread | work-item | One SYCL execution item. |
| warp | sub_group | The lockstep execution group used by this Xe path; the documented kernels require 16 work-items per subgroup. |
| CTA / threadblock | work-group | A 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 fragments | GRF | General Register File storage holding copy fragments, MMA fragments, and accumulators. |
| Tensor Core | XMX | Xe Matrix Extensions, the hardware matrix engine used by DPAS. |
mma.sync / HMMA | DPAS / XE_DPAS_TT | The matrix multiply-accumulate instruction/atom used by this CuTe Xe path. |
cp.async / TMA | Xe 2D block operations | An analogy for optimized data movement only; Xe 2D loads/prefetches have their own semantics and constraints. |
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
Why layout algebra matters
Functional composition is central to CuTe. For scalar coordinates, if R = A ∘ B, then R(c) = A(B(c)).
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.
TiledMMA and TiledCopy partition functions map logical tensor coordinates onto the work-items/subgroups participating in an operation.
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.
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.
What the Intel Xe path adds to upstream CuTe
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.
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
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.
Walk through the Intel Xe CuTe GEMM dataflow
The tutorial kernel sees
A(M,K), B(N,K), and C(M,N), then creates identity-coordinate proxy tensors for tiling.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.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).TiledCopy objects.Factory functions select copy atoms according to element width, tensor layout, and MMA requirements.
local_id.get_slice(local_id) maps each work-item into its role in the collective MMA/copy operation.Copy fragments match the layout produced by the 2D load; MMA fragments match the layout required by DPAS.
XE_PREFETCH_2D issues cache hints for future K tiles and has no register destination.Split-barrier arrive → 2D loads → future prefetch → optional register reorder → DPAS accumulation → split-barrier wait.
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
barrier arrive
2D load
prefetch ahead
reorder if needed
DPAS
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);
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.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.
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 atom | Purpose | Important current constraint / role |
|---|---|---|
XE_LOAD_2D | Normal rectangular 2D load | General row-major-style block load; supports block counts 1, 2, or 4. |
XE_LOAD_2D_VNNI | 2D load with VNNI transform | Used for B-style VNNI packing; Bits must be 8 or 16. |
XE_LOAD_2D_TRANSPOSE | 2D load with transpose | Current parameterized atom supports 32- or 64-bit elements and has additional width/shape limits. |
XE_PREFETCH_2D | 2D cache prefetch | No destination fragment; hints future data into L1/L2. |
XE_STORE_2D | Rectangular 2D store | Store 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 15What 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
| Constraint | Value | Notes |
|---|---|---|
| Base pointer alignment | 64 bytes | All 2D operations. |
| Pitch/stride alignment | 16 bytes | The guide notes that 4 bytes may work on some PVC configurations. |
| Width alignment | 4 bytes | All 2D operations. |
| Maximum load/prefetch height | 32 rows | Normal, VNNI, transpose loads, and prefetch. |
| Maximum store height | 8 rows | XE_STORE_2D. |
| Maximum total width | 64 bytes | Bits × Width ≤ 512. |
| Supported element widths | 8 / 16 / 32 / 64 bits | Bits describes element width for the memory operation, not necessarily a unique C++ numeric type. |
Tune only after you can explain the dataflow
| Knob | Why increase/change it? | What to watch |
|---|---|---|
| M / N tile | More compute per work-group can improve XMX utilization. | Larger accumulator footprint, possible GRF spill, and fewer concurrent work-groups. |
| K tile | Fewer 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/layout | Changes 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 depth | Prefetch 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 match | A matching layout lets the compiler remove reorder(). | A mismatch becomes a register-to-register subgroup shuffle. |
| SLM staging | Useful for designs needing sharing, explicit buffering, or data beyond direct 2D-load limits. | Extra copies and synchronization can negate the benefit. |
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
- Verify the required 16-work-item subgroup configuration.
- Verify 2D-copy base-pointer, pitch, width, and shape constraints.
- Profile whether the kernel is bandwidth-bound or compute-bound.
- Adjust M/N/K tile dimensions one change at a time and watch GRF spill.
- Adjust prefetch depth and measure rather than assuming more stages are better.
- Inspect whether
reorder()is compiled away or remains a shuffle.
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:
gemm_cute() — inspect launch geometry and the explicit kernel properties sub_group_size<16> and grf_size<256>.choose_mma_op() — see how A/B/C element types select a supported XE_DPAS_TT<8,...>, with BF16 or FP16 fallback paths.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.gemm_device() setup — identify the identity tensors, local_tile(), copy factories, slices, partitions, and register fragments.gemm() → split-barrier wait.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
Five questions before you write your own kernel
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.
- Intel — CuTe in SYCL*TLA: Intel Overview
- Intel — Intel Xe GPU GEMM Companion
- Intel — Performance Tuning Guide for CuTe on Xe
- CuTe Layout tutorial
- CuTe Layout Algebra tutorial
- CuTe Tensor tutorial
- Intel Xe 2D Copy Operations (legacy naming, still useful for data distribution)
- Reference source — examples/cute/tutorial/xe_gemm.cpp
- Current Xe MMA atom definitions — mma_xe.hpp
- Current parameterized Xe 2D copy atoms — copy_xe_2d.hpp
- Xe 2D copy selection/traits — copy_traits_xe_2d.hpp
- Xe register reorder implementation — reorder_xe.hpp
- Upstream CuTe dense GEMM tutorial — A(M,K), B(N,K), C(M,N) convention
- Khronos — SPV_INTEL_2d_block_io