100 DAYS OF GPU Course preview
Day 76Module 8
draft

The Tensor Memory Accelerator (TMA)

Every tile loaded since day 16 has required each of 256 threads to compute a global address before moving data. Day 44's fastest hand-written kernel still uses instructions for that address arithmetic.

The Tensor Memory Accelerator (TMA) moves this work to a dedicated engine. One thread supplies a descriptor and coordinates, and a barrier reports completion in bytes. This lesson builds the descriptor on the host, issues the copy in a kernel, and checks an exact tile round trip.

TMA requires compute capability 9.0.

The address arithmetic moves into a descriptor

TMA splits a tile copy into a description and a request. The description is a CUtensorMap, the 128 opaque bytes cuda.h declares as cuuint64_t opaque[16], built on the host by cuTensorMapEncodeTiled: element type, tensor rank, base pointer, the tensor's dimensions, its row stride in bytes, and the box (tile) shape to cut from it. The encode call is strict, and the rules are documented, not folklore: each box dimension is capped at 256, global strides must be multiples of 16 bytes, and the base address must be 16-byte aligned (https://docs.nvidia.com/cuda/archive/12.6.2/cuda-driver-api/group__CUDA__TENSOR__MEMORY.html , checked 2026-09-01).

The kernel receives the map by value as a __grid_constant__ parameter and treats it as opaque.

The request is one PTX instruction, cp.async.bulk.tensor, issued by a single thread. It names the map, a coordinate pair for the tile's corner in element units, a destination in shared memory (128-byte aligned), and a barrier. The engine does the rest: address generation, bounds handling, the whole 16384-byte transfer.

No register holds tile data in flight, and only one thread issues an instruction for the load (https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#data-movement-and-conversion-instructions-cp-async-bulk-tensor , checked 2026-08-29).

Completion uses day 75's async barrier with a transaction count. A thread-count barrier cannot know when a hardware engine is done, so the issuing thread registers an expected byte count with barrier_arrive_tx, and the engine itself ticks the barrier as bytes land. When the count hits 16384, the wait releases and the tile is visible to every thread in the block.

Diagram: three ways to fill a 64x64 float tile. Three horizontal bands, each showing global memory on the left, shared memory on the right, and what sits between them. Band 1, day 16's loads: 256 thread arrows, each through a register. Caption "4096 per-thread addresses computed, data through registers." Band 2, day 74's cp.async: 256 thread arrows, straight into shared memory. Caption "same 4096 addresses, registers out of the path." Band 3, TMA: one thread arrow carrying a map and an (x, y); the engine owns the transfer. Caption "one instruction, zero per-thread addresses, 16384 bytes counted at the barrier." Alt text: "Three ways to fill a sixty-four by sixty-four tile: every thread computing addresses through registers, every thread copying straight to shared memory, and one thread handing hardware a map that moves all sixteen thousand bytes."

The copy does not use lane-level loads

The instinct this course has trained since day 11 says a good global load is 32 lanes reading 32 consecutive elements. Day 74's cp.async kept that shape and removed the registers; every thread still computed where its elements lived, so the coalescing rules still graded the kernel.

TMA is not a coalesced warp load. The warp that issues it continues, and the transfer belongs to an engine with its own path to memory. Lane participation does not describe this transfer.

cudaMemcpyAsync also hides the CPU thread that submitted the copy from the device-side transfer.

The work moves to setup. Saving address arithmetic, load-slot pressure, and registers requires an encode call on the host, a barrier with a byte count, and a fence, because the engine reads and writes through the async proxy while ordinary loads and stores use the generic proxy. Without the fence, those accesses can race.

Single-CTA TMA is not Hopper-only. The feature table marks the TMA unit Yes for compute capability 9.0, 10.x, 11.0 and 12.x, so an RTX 5090 runs everything on this page (https://docs.nvidia.com/cuda/cuda-programming-guide/05-appendices/compute-capabilities.html , Table 29, checked 2026-08-29). What a GeForce card lacks is TMA multicast, one copy feeding several blocks of a cluster: the PTX manual warns .multicast::cluster "may have substantially reduced performance" off the listed datacenter targets, and CUTLASS pins GeForce cluster shapes to 1x1x1 outright (https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#data-movement-and-conversion-instructions-cp-async-bulk-tensor and https://github.com/NVIDIA/cutlass/blob/main/media/docs/cpp/blackwell_functionality.md#cluster-size , both checked 2026-08-29).

An exact round trip

Full program in code/day76-tma/tma_tile.cu. A 256x256 float matrix, sixteen tiles, one TMA load per block, +1.0f per element in shared memory, plain per-thread stores back out, and a bit-exact host comparison. The check tolerates nothing on purpose: the copy and the single float add are both deterministic, so the first mismatching index is always a bug, and its row and column say whether a whole tile landed wrong (bad coordinates) or scattered elements did (a race with the engine).

Everything device-side comes from libcu++ as shipped in CUDA 12.6; nothing here needs a 13.x toolkit or a hand-written PTX block.

The host builds the map once. The kernel needs no per-tile pointer math, leading-dimension parameter, or device-side offset bookkeeping:

    // The tensor map is the whole address story, built once on the host:
    // element type, tensor rank, base pointer, tensor shape, row stride in
    // bytes (one stride for rank 2: the innermost dimension is implicitly
    // contiguous), tile shape, and no interleave, swizzle, L2 promotion or
    // out-of-bounds fill. The kernel receives the map by value and never
    // computes a global address.
    CUtensorMap tileMap{};
    const cuuint64_t gmemDim[2] = {kGmemWidth, kGmemHeight};
    const cuuint64_t gmemStrideBytes[1] = {kGmemWidth * sizeof(float)};
    const cuuint32_t boxDim[2] = {kTileWidth, kTileHeight};
    const cuuint32_t elemStride[2] = {1, 1};
    const CUresult mapResult = encodeTiled(
        &tileMap, CU_TENSOR_MAP_DATA_TYPE_FLOAT32, 2, d_in, gmemDim,
        gmemStrideBytes, boxDim, elemStride, CU_TENSOR_MAP_INTERLEAVE_NONE,
        CU_TENSOR_MAP_SWIZZLE_NONE, CU_TENSOR_MAP_L2_PROMOTION_NONE,
        CU_TENSOR_MAP_FLOAT_OOB_FILL_NONE);

encodeTiled is cuTensorMapEncodeTiled fetched through cudaGetDriverEntryPointByVersion, which exposes the driver API through the runtime, so the build line needs no -lcuda. In the kernel, one thread issues the copy and everyone meets at the barrier:

    // One thread asks for the whole tile. The coordinates are element
    // offsets into the global tensor, x (the contiguous dimension) first.
    // Everyone else just arrives; the engine, not any thread, completes
    // the barrier by counting 16384 copied bytes against the expectation
    // that barrier_arrive_tx registered.
    BlockBarrier::arrival_token token;
    if (threadIdx.x == 0) {
        const int32_t coords[2] = {
            static_cast<int32_t>(blockIdx.x) * kTileWidth,
            static_cast<int32_t>(blockIdx.y) * kTileHeight};
        cuda::ptx::cp_async_bulk_tensor(
            cuda::ptx::space_cluster, cuda::ptx::space_global, &tile, &tileMap,
            coords, cuda::device::barrier_native_handle(bar));
        token = cuda::device::barrier_arrive_tx(bar, 1, sizeof(tile));
    } else {
        token = bar.arrive();
    }
    bar.wait(cuda::std::move(token));

The space_cluster destination is correct even though the program uses no cluster: the PTX state space for a TMA destination is named shared::cluster because a plain block is a cluster of one, a naming choice used in day 77. After the wait, the program uses plain stores:

    const size_t colBase = blockIdx.x * static_cast<size_t>(kTileWidth);
    const size_t rowBase = blockIdx.y * static_cast<size_t>(kTileHeight);
    for (int k = static_cast<int>(threadIdx.x); k < kTileWidth * kTileHeight;
         k += kThreadsPerBlock) {
        const int r = k / kTileWidth;
        const int c = k % kTileWidth;
        out[(rowBase + r) * kGmemWidth + colBase + c] = tile[r][c] + 1.0f;
    }

Keeping the store plain isolates the TMA load. Turning this store into TMA's write-back path is the exercise. This page makes no timing claim; a 256x256 copy checks the mechanism, not its speed.

Results

Not yet run. TMA needs compute capability 9.0, and the program's gate rejects older GPUs. Run it on a compatible GPU and save its transcript, SASS and PTX dumps, sanitizer output, and gate transcript in code/day76-tma/evidence/, per the capture list in the README. Until then the front matter stays draft and every hardware field stays null.

The hardware availability and price snapshot used during planning is in FACT-SHEET.md.

Step On a CC 9.0+ card On the T4
capability gate ? ?
tensor map encode ? ?
round trip check ? ?

Checks for the first run

  1. The round trip passes bit-exactly, first try. The only synchronization on the load path is the barrier wait; the lesson claims that wait alone makes the engine's writes visible. Any mismatch rejects that claim, and the printed row and column say how: a whole wrong tile means bad coordinates, scattered elements mean the visibility story is wrong.
  2. The PTX dump shows exactly one cp.async.bulk.tensor.2d in loadTileTma, and no ld.global touches tile data. Checkable with the toolkit alone, without a GPU.
  3. On the T4, the program prints the capability message and exits nonzero before allocating anything. The gate is a branch, not an assert, so the Release build keeps it.
  4. All four compute-sanitizer tools come back clean, racecheck included: the engine's shared-memory writes are ordered by the barrier, and the tool should agree that ordering is airtight.

A TMA copy uses a descriptor built once, a request from one thread, and completion counted in bytes. Tile size, matrix size, and GPU model do not change that structure.

Run it yourself

Use a GPU with compute capability 9.0 or newer. An RTX 50 runs the program as built because -arch=sm_90 embeds PTX that JIT-compiles forward. An sm_90a build is an architecture-specific target and loads on Hopper alone.

Build lines, inspection commands, and the artifact list are in the README. The setup guide lists remote GPU options.

Exercise

Make the write-back symmetric: encode a second tensor map over d_out, and replace the plain store loop with TMA's shared-to-global direction, so the kernel's only global traffic in either direction is two bulk tensor instructions. The bit-exact check must still pass.

Time: 30 to 45 minutes. Submit: the diff, the transcript with the passing check, and the one PTX line your change added.

Check: harness-style, and already in the program: the round trip comparison exits nonzero on the first mismatching element and prints its index, row and column. A stale output buffer (all values off by exactly 1.0f from zero) means your store never ran; scattered garbage means it ran before the data was ready. Without a CC 9.0 card, do the reading version against the shipped PTX dump: write down the instruction sequence you would emit and compare against evidence/tma_tile.ptx once the verification run lands.

Hint 1

The load path needed a barrier because the engine wrote shared memory and threads read it. Your change inverts producer and consumer: now threads write shared memory and the engine reads it. What has to be true about those writes before the engine may start, and which side of the proxy line did they happen on?

Hint 2

Four steps, in order: every thread finishes its shared-memory writes (__syncthreads()), one thread publishes them across the proxy line (cuda::ptx::fence_proxy_async), the same thread issues the space_global, space_shared overload of cp_async_bulk_tensor (map and coordinates first, source last), then commits and waits on the bulk async group before the kernel may end. The completion mechanism is not the mbarrier this time; look at cp_async_bulk_commit_group and cp_async_bulk_wait_group_read.

Solution

Compute into tile, then __syncthreads(), then from thread 0 only: cuda::ptx::fence_proxy_async(cuda::ptx::space_shared), then cuda::ptx::cp_async_bulk_tensor(cuda::ptx::space_global, cuda::ptx::space_shared, &outMap, coords, &tile), then cuda::ptx::cp_async_bulk_commit_group() and cuda::ptx::cp_async_bulk_wait_group_read(cuda::ptx::n32_t<0>()). The host encodes outMap over d_out with the same shape arguments, and the + 1.0f moves into a shared-memory update between load and store.

The two directions complete differently on purpose, and picking the wrong mechanism is not a style error but a race: an mbarrier counts bytes arriving in shared memory, a bulk async group tracks reads leaving it. Every TMA kernel you read from here on, CUTLASS included, is built on that asymmetry.

Pitfalls

You built with -arch=sm_90a and an RTX 50 refuses to load it. Architecture-specific targets do not JIT forward, and nothing on this page needs one. Build sm_90 and let the embedded PTX carry the binary to 10.x and 12.x. Day 78 maps which feature actually requires which target.

Every thread issues the copy. The data can even come out right, which is the trap: you moved sixteen tiles' worth of bytes per tile, and the barrier's byte count no longer matches what barrier_arrive_tx promised, so the wait's release timing is no longer coupled to your data. One thread issues; everyone else only arrives.

No fence between the barrier's init and the copy. init(&bar, ...) is a generic-proxy write, and the engine reads the barrier through the async proxy. Skip the fence_proxy_async and the program can pass for weeks on one card and corrupt on another, with no error string to search for. The fence is one line; the bug it prevents has none.

cuTensorMapEncodeTiled returns CUDA_ERROR_INVALID_VALUE and tells you nothing else. The encode call enforces every rule at once (box dimensions at most 256, strides multiples of 16 bytes, base address 16-byte aligned, rank at most 5), and one failed rule fails the whole call. Check the rules against the driver API page rather than guessing; this program's static_asserts pin the compile-time-checkable ones.

Coordinates in the wrong order. The coordinate array is innermost dimension first: x, then y. Swap them and off-diagonal blocks load each other's transposed positions with no error raised; the round-trip check catches it as a whole tile of wrong values, first mismatch inside the first off-diagonal tile.

You wrote .multicast::cluster because your card "is Blackwell". On a GeForce card it compiles and crawls: the qualifier is optimized for the datacenter targets and CUTLASS fixes GeForce clusters at 1x1x1 precisely because the hardware feature is absent. Multicast is day 78's story; nothing on this page uses it.

Go deeper

Next

Day 77 uses the shared::cluster state space when blocks form a cluster and read each other's shared memory. Day 78 maps these mechanisms across Blackwell GPU families.