Day 79Module 8

draft

Tile programming with cuTile

This lesson makes no cuTile performance claim. The public benchmarks checked for this course were either refuted or published by NVIDIA (FACT-SHEET.md section 8).

The lesson covers the cuTile programming model. You will write and check a tile matmul and tile softmax in Python, then read the C++ equivalents. You will also apply the tile-shape, memory-model, and compute-capability constraints.

A kernel written for a tile block

Every earlier kernel was a program for one thread. Day 44 reached 67.7 percent of cuBLAS on a T4 by managing 128 registers, float4 loads, and the tiling scheme from day 16. Day 73 used mma.sync with values assigned to individual lanes.

cuTile removes thread and lane control from the source. A tile kernel is a program for one tile block.

The core object is the tile: a fixed-size multidimensional array whose shape is known at compile time. The Python docs state the constraint plainly: "Tile dimensions must be compile-time constants that are powers of two" (https://docs.nvidia.com/cuda/cutile-python/index.html , checked 2026-09-01). You load a tile from a global array by tile-space index, not by address.

For example, select tile (3, 2) of a matrix partitioned into 32 by 16 tiles. The launch grid is that tile space. For the 256 by 128 output in this lesson's program, tiles of 32 by 32 give a grid of 8 by 4, thirty-two tile blocks, and the programmer does not pick a thread count: in C++ the second chevron argument "must be 1" because "the compiler determines the thread count internally" (https://docs.nvidia.com/cuda/cuda-programming-guide/02-basics/writing-tile-kernels.html , checked 2026-09-01).

The compiler controls how tile operations map to instructions. ct.mma becomes tensor core instructions where the card has the right ones; loads through a partition view can become TMA transfers on Hopper and Blackwell; the same kernel source targets everything from compute capability 8.0 up. This model handles the fragment layouts that day 72's WMMA API hid and day 73's mma.sync required in source.

Diagram: the grid is a tile space. Left: a 256 by 128 matrix C drawn as an 8 by 4 grid of 32 by 32 tiles, one tile shaded, labelled with its tile-space index (2, 1). Right: the launch, one box per tile block, 32 boxes, the shaded tile's box highlighted; no thread rows anywhere. Band 1: the matmul: the shaded block's K loop walks 8 tiles of A and 8 of B. Caption "8 K-steps of (32 x 16) @ (16 x 32) into one FP32 tile." Band 2: the softmax: one block loads 8 full rows as one (8, 128) tile. Caption "the reduction never leaves the tile." Alt text: "A 256 by 128 matrix partitioned into thirty-two 32 by 32 tiles, with one launch block per tile. Blocks index tiles, not addresses, and no thread count appears anywhere."

Constraints of the tile model

A tile kernel has no threadIdx, so code cannot use register blocking, per-lane loads, or warp shuffles from days 43 and 44. Use the day 73 model when a kernel needs per-lane control.

cuTile also has a different memory model from SIMT CUDA. Its contract states: "cuTile's memory model permits the compiler and hardware to reorder operations for performance. Without explicit synchronization, the ordering of memory accesses across blocks is not guaranteed" (https://docs.nvidia.com/cuda/cutile-python/memory_model.html , checked 2026-09-01).

Unlike a SIMT store followed by __syncthreads(), this contract does not define cross-block visibility without explicit synchronization.

Coordinate blocks with atomics that have an explicit memory order and memory scope. Both kernels avoid cross-block communication, so no block reads data written by another block.

Python kernels and C++ equivalents

Both programs live in code/day79-cutile/, Python in tile_matmul_softmax.py and the C++ mirror in tile_matmul_softmax.cu. Each gates on compute capability 8.0 with a branch, checks both kernels against a float64 reference with the mixed tolerance discipline from day 66, and exits nonzero at the first mismatch. The programs do not measure run time.

The Python matmul is the whole algorithm from day 16 in fourteen lines, with an FP32 accumulator type held as a tile:

@ct.kernel
def matmul(A, B, C, tm: ct.Constant[int], tn: ct.Constant[int],
           tk: ct.Constant[int]):
    # One tile block owns the (bm, bn) tile of C. No thread index exists;
    # the block is the unit and the compiler picks the lanes.
    bm = ct.bid(0)
    bn = ct.bid(1)
    acc = ct.full((tm, tn), 0, dtype=ct.float32)
    zero_pad = ct.PaddingMode.ZERO
    for k in range(ct.num_tiles(A, axis=1, shape=(tm, tk))):
        a = ct.load(A, index=(bm, k), shape=(tm, tk), padding_mode=zero_pad)
        b = ct.load(B, index=(k, bn), shape=(tk, tn), padding_mode=zero_pad)
        acc = ct.mma(a, b, acc)
    # C is float32 and so is acc, so the store needs no conversion. Give C a
    # narrower dtype and this line becomes ct.astype(acc, ct.float16).
    ct.store(C, index=(bm, bn), tile=acc)

The softmax gives each block eight whole rows, so the max and sum are tile reductions. It needs no cross-block ordering:

@ct.kernel
def softmax(X, Y, tb: ct.Constant[int], n: ct.Constant[int]):
    # One tile block owns tb full rows, so the max and the sum never leave
    # the tile and no cross-block ordering question exists. Python drops
    # the reduced axis unless keepdims=True; keeping it is what lets the
    # (tb, 1) results broadcast back over the (tb, n) rows.
    rows = ct.load(X, index=(ct.bid(0), 0), shape=(tb, n))
    num = ct.exp(rows - ct.max(rows, axis=1, keepdims=True))
    den = ct.sum(num, axis=1, keepdims=True)
    ct.store(Y, index=(ct.bid(0), 0), tile=num / den)

The C++ mirror says the same things in the 13.3 toolkit's dialect. __tile_global__ is a new execution space specifier beside __global__, views are built from tensor_span plus partition_view, and the one behavioural difference worth memorising is in the comment: C++ reductions keep the reduced axis, Python drops it unless you ask. The C++ reference says the result shape "matches S except at dimension d where its length is 1", with no keepdims to turn that off (https://docs.nvidia.com/cuda/cuda-tile-cpp-api-reference/reductions_and_scans.html , checked 2026-09-01), while cuda.tile.sum is sum(x, /, axis=None, *, keepdims=False, ...) and defaults to dropping it.

__tile_global__ void softmaxRows(const float* __restrict__ in,
                                 float* __restrict__ out, std::size_t rows) {
    namespace ct = cuda::tiles;
    using namespace ct::literals;

    in = ct::assume_aligned(in, 16_ic);
    out = ct::assume_aligned(out, 16_ic);

    constexpr auto tr = 8_ic;    // keep in step with kRowsPerTile
    constexpr auto tc = 128_ic;  // keep in step with kN

    auto inView = ct::partition_view{ct::tensor_span{in, ct::extents{rows, tc}},
                                     ct::shape{tr, tc}};
    auto outView = ct::partition_view{
        ct::tensor_span{out, ct::extents{rows, tc}}, ct::shape{tr, tc}};

    int bx = ct::bid().x;
    auto rowsTile = inView.load_masked(bx, 0);

    auto mx = ct::reduce_max(rowsTile, 1_ic);  // shape (8, 1), axis kept
    auto num = ct::exp(rowsTile - mx);         // (8, 1) broadcasts to (8, 128)
    auto den = ct::sum(num, 1_ic);             // shape (8, 1)
    outView.store_masked(num / den, bx, 0);
}

The masked loads and stores are what replace every bounds check you have written since day 5: load_masked zero-pads a short edge tile and store_masked drops the out-of-bounds lanes. The matmul checks against float64 over the identical FP16-rounded inputs, so its tolerance covers accumulation order alone; the softmax reference consumes the GPU's own matmul output, so each check isolates one kernel. Neither program measures speed.

Results

Not yet run. The C++ file compiles: NVCC 13.3.0 targeting sm_80 accepts it at exit code 0, captured 2026-09-01 with execution off and kept verbatim in code/day79-cutile/evidence/compile-2026-09-01.txt. That is a compile and nothing more. No GPU has executed either program, so every row below is a question mark and stays one.

The quickstart requires two things: "a GPU with compute capability 8.x, 9.x, 10.x, 11.x or 12.x" and "NVIDIA Driver r580 or later" (https://docs.nvidia.com/cuda/cutile-python/quickstart.html , checked 2026-09-01).

The project's Tesla T4 node clears the second, at driver 595.84, and fails the first at 7.5. Its CUDA 12.6 cannot build the C++ half, which ends at fatal error: cuda_tile.h: No such file or directory in the same capture.

Run both programs on a compatible GPU and save the transcripts in code/day79-cutile/evidence/. Until then the front matter stays draft with every hardware field null.

The environment and availability snapshots are recorded in FACT-SHEET.md section 4 and FACT-SHEET.md section 8.

Program	Check	Result
`tile_matmul_softmax.py`, matmul	float64, rtol 1e-4, atol 1e-3	?
`tile_matmul_softmax.py`, softmax	float64, rtol 1e-5, atol 1e-6	?
`tile_matmul_softmax.cu`, both	same bounds	?
either program on the T4	capability gate, exit 1	?

Checks for the first sm_80 run

Both matmuls pass their mixed bound. FP16 products are exact in FP32, so the only divergence from the float64 reference is accumulation order, and the absolute bound derived in the source comment holds with a margin. If the error scales past it, the derivation is wrong and the tolerance was a guess, which is a day 66 failure and must be reported.
Both softmaxes pass the harness-default bound. Rows sum to 1 within float32. A failure here with a passing matmul points at the reduction or the broadcast, not the data.
The T4 gate is captured. The Python program prints its capability line and exits 1 before any cuTile call; that refusal transcript is itself an artifact, evidence/t4-gate-YYYY-MM-DD.txt.
API compatibility. cuTile Python is at 1.5.0 and moving. If the pinned install resolves an API that no longer accepts these kernels (ct.max keepdims, ct.num_tiles, PaddingMode), the run records the actual failure and the lesson must match the released API.

No throughput number will be added after the run. One measured correctness pass says nothing about speed, and a single unprofiled timing on one GPU would not support a performance claim.

Run it yourself

Use a GPU with compute capability 8.0 or newer and NVIDIA Driver r580 or later.

Compiler Explorer publishes its CUDA runner metadata at (https://godbolt.org/api/compilers/cuda?fields=id,name,supportsExecute). The setup guide lists remote GPU options.

Install the Python package with pip install --upgrade cuda-tile[tileiras]; to pin a tileiras line the quickstart gives pip install cuda-toolkit[tileiras,nvvm,nvcc]>=13.3, with no upper bound as of 2026-09-01, and requires the three cuda packages it pulls to "match up to the same major.minor version". An earlier project note said the extra was pinned below 13.4; the current quickstart no longer states that limit.

The C++ half needs CUDA Toolkit 13.3 or newer and builds with nvcc -std=c++20 -O3 -arch=sm_80 -enable-tile. The pitfalls show the errors caused by omitting the required flags.

Exercise

Fuse the two kernels: one cuTile Python kernel that computes softmax(A @ B) row-wise, so each tile block accumulates whole rows of C and applies the softmax before a single store. This is the kernel fusion argument from day 48 restated in tile terms.

Time: 30 to 45 minutes. Submit: fused.py and one sentence naming which cuTile limit you ran into first.

Check: the harness compares your output against softmax(A @ B) in numpy float64 at rtol 1e-4, atol 1e-3, and prints the first failing (row, column) with both values. On a card below compute capability 8.0 it prints the capability line and exits 1 without judging your kernel. A wrong answer concentrated in tile-sized bands is a tile-index bug; a wrong answer everywhere by a small margin is a tolerance or dtype bug.

Hint 1

The softmax needs a whole row before it can start. Which choice of the matmul's output tile shape hands a block a whole row of C, and what does that do to the grid?

Hint 2

Set tn equal to N. 128 is a power of two, so (32, 128) is a legal accumulator tile. After the K loop, acc is eight-times-four whole rows sitting in the block, and the three softmax lines from this page apply to it unchanged, before the store instead of after a round trip.

Solution

Launch a grid of (ct.cdiv(M, tm), 1, 1). Keep the K loop exactly as in matmul but load B tiles with index=(k, 0) and shape=(tk, N), so acc has shape (tm, N). After the loop, apply the softmax lines to acc (max, exp, sum along axis 1 with keepdims) and store once at index=(bm, 0).

The intermediate C never touches global memory, and the program needs no second kernel or synchronization point. In a tile kernel, fusion composes tile expressions in one body. The compiler chooses whether and how to stage the data through shared memory, so profile the fused kernel before making a speed claim.

Pitfalls

Your tile dimension is 48, or comes from a runtime variable. The model's first hard rule: "Tile dimensions must be compile-time constants that are powers of two". In Python that means a ct.Constant[int] argument or a literal, and 48 is out at any position. Sizes of the arrays are free; sizes of the tiles are not.

nvcc rejected the C++ file from inside a NVIDIA header. Two flags that section 2.4 does not name decide whether this file builds at all, and both were found by compiling it rather than by reading.

Under -std=c++17, the build stops at the header's own #error "This file needs C++20 features. Please compile with c++20 or later dialect". Fix that and drop -enable-tile and you get warning #20359-D, "parsing for Tile constructs is not enabled. Tile annotations (including __tile__ and __tile_global__) will be ignored", followed by many "calling a __tile__ function from a __host__ function is not allowed" errors whose first line numbers sit inside crt/cuda_tile.h.

The errors do not name the missing flag: with tile parsing off, __tile_global__ is dropped and your kernel is a host function, so every tile call is illegal.

The working line is nvcc -std=c++20 -O3 -arch=sm_80 -enable-tile, and all four transcripts are in the day's evidence file.

You launched the C++ kernel with a thread count. Writing gemm<<<grid, 256>>> treats a tile kernel as a SIMT kernel. The second chevron argument "must be 1"; the compiler owns the thread count, and a tile kernel gives you no threadIdx to spend those threads on. The Python ct.launch grid tuple counts tile blocks, not threads, for the same reason.

The package installed, then the launch failed. pip install cannot check the GPU's compute capability, so an unsupported device fails at run time. Both of this lesson's programs gate first and say which card they found; run them before blaming your kernel.

One block reads what another block wrote, without an atomic. The memory model "permits the compiler and hardware to reorder operations for performance", and cross-block ordering is "not guaranteed" without explicit synchronization. Coordinate through the atomics and their memory order and scope attributes, or restructure so blocks share nothing, as both kernels here do.

pip resolved mismatched cuda packages. The quickstart requires nvidia-cuda-tileiras, nvidia-cuda-nvcc and nvidia-nvvm to "match up to the same major.minor version". A half-upgraded environment breaks in ways that look like kernel bugs. Record pip freeze next to every run; the README shows the line.

You read a cuTile-versus-Triton benchmark and believed it. As of this writing no independent post-13.3 comparison exists; the widely shared numbers predate Hopper support or come from the vendor. Day 89 is this course's tool-choice day, and it holds the same line.

Go deeper

CUDA Programming Guide 2.4, "Writing Tile Kernels", the C++ and Python pairing this lesson mirrors: https://docs.nvidia.com/cuda/cuda-programming-guide/02-basics/writing-tile-kernels.html (checked 2026-09-01)
cuTile Python docs: quickstart, memory model and operations: https://docs.nvidia.com/cuda/cutile-python/quickstart.html and https://docs.nvidia.com/cuda/cutile-python/memory_model.html (both checked 2026-09-01)
NVIDIA's cuTile samples, including the swizzled and persistent matmul this lesson deliberately simplified: https://github.com/NVIDIA/cutile-python/tree/main/samples (checked 2026-09-01)
CUDA Tile C++ API reference, reductions and scans: https://docs.nvidia.com/cuda/cuda-tile-cpp-api-reference/reductions_and_scans.html (checked 2026-09-01)

Day 80 builds a tensor core GEMM by hand with WMMA, targeting the same layer that cuTile abstracts. It must beat day 44's measured result of 67.7 percent of cuBLAS. Day 86 writes this softmax in Triton, and day 89 compares the two programming models.