100 DAYS OF GPU Course preview
Day 73Module 8
in-technical-review

mma.sync and ldmatrix

Many mma.sync guides start with m16n8k16 and say they require Ampere. The PTX ISA gives a lower minimum for two related operations: the f16 mma at m16n8k8 "requires sm_75 or higher", and ldmatrix "Requires sm_75 or higher". Only the k16 shape needs sm_80.

This page's program compiles for -arch=sm_75 with exit code 0 and puts only the k16 probe behind a guard. By the end you can name which of a warp's 32 lanes holds which element of A, B and D, load a fragment with one instruction instead of eight scalar loads, and say exactly what the day 72 API was keeping from you.

What a warp owns while the tensor core runs

mma.sync is one instruction executed by a whole warp on the tensor cores. It computes D = A x B + C, and for the shape this page uses, m16n8k8, A is 16x8, B is 8x8, and C and D are 16x8.

None of those operands is a memory address. They are registers, spread across the 32 lanes, and the ISA says exactly which piece of the matrix each register holds.

The matrix sizes determine each lane's fragment size. A is 128 halves over 32 lanes, so four each, delivered as two .f16x2 registers. B is 64 halves, so two each, one register.

D is 128 floats, so four each, four registers.

The tables give the coordinates with two helper values, and they are clearer as code than as prose (PTX ISA 8.5, "Matrix Fragments for mma.m16n8k8", https://docs.nvidia.com/cuda/parallel-thread-execution/index.html , checked 2026-09-01):

groupID           = %laneid >> 2
threadID_in_group = %laneid % 4

A:   row = groupID for a0 and a1, groupID + 8 for a2 and a3
     col = (threadID_in_group * 2) + (i & 0x1)
B:   row = (threadID_in_group * 2) + i for bi where i = {0, 1}
     col = groupID
C/D: row = groupID for c0 and c1, groupID + 8 for c2 and c3
     col = (threadID_in_group * 2) + (i & 0x1)

Read the A and the C/D blocks twice. They are the same rule. A lane's four accumulator floats sit at the same coordinates as its four A halves, which is why this page has two separate checks.

The next problem is loading the fragments. Lane 5 needs A[1][2] and A[1][3], then A[9][2] and A[9][3].

In shared memory those are two 4-byte pairs far apart, and every lane wants a different pair. As scalar loads that is four instructions per lane per k step for A alone, and two more for B.

ldmatrix moves up to four 8x8 tiles of 16-bit elements from shared memory straight into the fragment layout. The addresses come per lane and are row starts, not element addresses: "The eight addresses required for each matrix are provided by eight threads", .x1 taking them from lanes 0 to 7, .x2 from lanes 0 to 15, .x4 from all 32.

"When reading 8x8 matrices, a group of four consecutive threads loads 16 bytes. The matrix addresses must be naturally aligned accordingly" (PTX ISA 8.5, ldmatrix, same URL, checked 2026-09-01).

That 16-bytes-per-four-lanes shape is why the tile in this program is padded by 8 halves per row rather than the one-element pad day 15 used. The ldmatrix-fragment frame below shows the bank indices for that access.

What WMMA does not specify

Day 72's fragment is opaque on purpose. The programming guide says so in one sentence: "The mapping of matrix elements into fragment internal storage is unspecified and subject to change in future architectures" (CUDA C++ Programming Guide 12.6.2, section 7.24.1, https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html , checked 2026-09-01).

Most readers take that to mean the mapping is unknowable. It means WMMA declines to promise one, while mma.sync publishes a table per shape and per type that your code must follow.

The table enables coordinate-dependent work: an epilogue can apply a per-row bias without a shared-memory round trip, you can choose a swizzle, and you can use shapes that WMMA does not expose. In exchange, all 192 values must be in the right registers, and wrong values do not cause a fault.

Having each lane print where it stored its four accumulator floats looks like a check but is not one, because the C/D rule is the A rule. A correct accumulator map is consistent with an A fragment loaded from the wrong eight rows.

Two probes and one matmul

Full program in mma_sync.cu, under code/day73-mma-sync/. It runs one tile alone, then a 2048x2048 FP16-in, FP32-accumulate matmul built from that tile, then the same single tile at m16n8k16 on hardware that has it.

The two checks cannot cover for each other. The probe fills its tiles with one unique value per cell, row * 8 + col for A and 200 + row * 8 + col for B, unpacks the loaded registers back into floats and asks the host to name the coordinate each lane received. That tests the two ldmatrix calls on their own, with no arithmetic in the way. The matmul then compares every output element against a double-precision reference, which is what tests the accumulator formula.

Nothing here needs a tolerance. Probe values are integers below 2048, exact in __half, and no probe dot product reaches 2^24. Matmul inputs are multiples of 0.25 in [-2, 2], so every product is a multiple of 1/16 and every partial sum stays under 2^24 scaled, exact in an FP32 accumulator whatever order the hardware adds in.

Both gates use exact equality, and both are branches that return EXIT_FAILURE. Day 66's tolerance method is the right tool when rounding occurs; here it would only hide an index bug.

The instruction itself is one line: nine dot-separated qualifiers and four operand lists:

// One tensor-core instruction: D (16x8, f32) += A (16x8, f16) x B (8x8,
// f16). The .row.col layout is the only one the shape offers, which is
// why the B loads above carry .trans.
static __device__ void mmaM16N8K8(float acc[4], unsigned a0, unsigned a1,
                                  unsigned b0) {
    asm volatile(
        "mma.sync.aligned.m16n8k8.row.col.f32.f16.f16.f32 "
        "{%0, %1, %2, %3}, {%4, %5}, {%6}, {%0, %1, %2, %3};"
        : "+f"(acc[0]), "+f"(acc[1]), "+f"(acc[2]), "+f"(acc[3])
        : "r"(a0), "r"(a1), "r"(b0));
}

The B load carries .trans because the shape offers only .row.col, so B has to arrive column major. .trans transposes each 8x8 tile on its own, which matters as soon as you ask for four of them:

// Four row-major 8x8 tiles loaded column major in one instruction: the
// four B fragments a 16x32 warp tile needs per k step.
static __device__ void ldmatrixX4Trans(unsigned& r0, unsigned& r1, unsigned& r2,
                                       unsigned& r3, unsigned addr) {
    asm volatile(
        "ldmatrix.sync.aligned.m8n8.x4.trans.shared.b16 "
        "{%0, %1, %2, %3}, [%4];"
        : "=r"(r0), "=r"(r1), "=r"(r2), "=r"(r3)
        : "r"(addr));
}

A warp's inner loop is then two loads and four multiplies, covering a 16x32 slab of output per k step of 8:

        for (int ks = 0; ks < kBlockK; ks += 8) {
            unsigned a0, a1;
            ldmatrixX2(a0, a1, smemAddr(&tileA[rowA * kLdA + ks]));
            unsigned b0, b1, b2, b3;
            ldmatrixX4Trans(b0, b1, b2, b3,
                            smemAddr(&tileB[(ks + rowB) * kLdB + colB]));
            mmaM16N8K8(acc[0], a0, a1, b0);
            mmaM16N8K8(acc[1], a0, a1, b1);
            mmaM16N8K8(acc[2], a0, a1, b2);
            mmaM16N8K8(acc[3], a0, a1, b3);
        }

And the probe's report, which is the whole reason the probe exists. An .f16x2 register holds the lower-numbered fragment element in its low 16 bits, so decoding it makes a0 and a1 separable:

    // a0 carries elements a0 and a1 of the fragment, a1 carries a2 and a3.
    aFrag[lane * 4 + 0] = halfOf(a0, 0);
    aFrag[lane * 4 + 1] = halfOf(a0, 1);
    aFrag[lane * 4 + 2] = halfOf(a1, 0);
    aFrag[lane * 4 + 3] = halfOf(a1, 1);
    bFrag[lane * 2 + 0] = halfOf(b0, 0);
    bFrag[lane * 2 + 1] = halfOf(b0, 1);

It is not a cuBLAS competitor. Tiles arrive as plain uint4 copies with one __syncthreads() per 32-deep step and no overlap between the copy and the arithmetic. That overlap is day 74's cp.async, and leaving it out keeps this page about the fragment.

Results

Re-verified on the same Tesla T4 with driver 580.173.02 and CUDA 13.0 (V13.0.88). Ownership, exact-reference matches, the sm_80 skip and performance all reproduced: 2.470 ms versus 2.472 ms under CUDA 12.6.

Both run transcripts are listed in front matter. The separate older compile capture still records that -Xptxas -v puts matmulMmaSync at 62 registers, 9728 bytes of shared memory and zero spill stores or loads.

what expected shape
fragment ownership, 192 elements all match the PTX tables
m16n8k8 tile against the double reference exact match
m16n8k16 probe skipped, needs sm_80 or newer
matmul at n = 2048 against the reference exact match
matmul time and GFLOP/s 2.470 ms, 6955.2 GFLOP/s

The run checks five predictions for sm_75.

  1. Held. The ownership gate matched all 192 elements against the PTX tables.
  2. Held. Lane 0 printed A at (0,0), (0,1), (8,0), (8,1) and B at (0,0), (1,0). Lane 5 printed A at (1,2), (1,3), (9,2), (9,3) and B at (2,1), (3,1), exactly as predicted.
  3. Held. The m16n8k8 tile and the full n = 2048 matmul both matched their double references exactly.
  4. Held. The matmul took 2.470 ms and reached 6955.2 GFLOP/s. That is 1.708x faster than day 44's 4.219 ms hand-written FP32 kernel on CUDA cores.
  5. Held. The k16 probe printed its skip line and named sm_80 as the minimum against this sm_75 GPU. The process exited 0.

Every card that supports the shape uses the same ownership map. The timings depend on the card.

Run it yourself

The minimum compute capability is 7.5:

nvcc -std=c++17 -O3 -arch=sm_75 -o mma_sync mma_sync.cu

Build with -arch=sm_80 on Ampere or newer and the k16 probe runs instead of printing its skip line. There is no Compiler Explorer embed: the double-precision reference at n = 2048 is 8.6 billion multiply-adds, which does not fit the 20 second execution cap, so Compiler Explorer is used here to compile and nothing else. That is what the evidence file holds and it says so at the top.

Exercise

Delete ldmatrix from the B path. Build the same B fragment with plain per-lane loads out of the shared tile, keeping every other line of matmulMmaSync as it is, and report what it costs.

Time: 30 to 45 minutes. Submit: the diff, both gate lines, the two times at n = 2048, the two ptxas -v register counts, and one sentence naming where the difference comes from.

Check: use the shipped program's two gates. Your version must still print fragment ownership: all 192 elements match the PTX tables and the exact-match line at n = 2048; on failure it names the lane and element, or the first mismatching row and column with both values, and returns EXIT_FAILURE. Passing the gates and running slower is a completed exercise.

Hint 1

You already know the coordinates: lane l needs B at row (l % 4) * 2 + i and column l / 4, per 8x8 tile. The question is not which elements to fetch, it is how many instructions and how many banks it takes to fetch them.

Hint 2

Count the shared-memory instructions per warp per k step, statically, in both versions. The shipped inner loop issues one ldmatrix.x4 for the four B tiles; yours issues two loads per tile per lane. What is the ratio, and which version needs the bank conflict arithmetic run on it?

Solution

Replace the ldmatrixX4Trans call with two __half reads per tile, tileB[(ks + (lane % 4) * 2 + i) * kLdB + colBase + lane / 4], packed with __halves2half2. The gates still pass because ldmatrix is a load method, not a layout, and the mma shape fixes the layout either way.

The ratio from your two ptxas -v lines and two timings price, eight shared loads per lane per k step against one instruction, each of yours pulling one 2-byte element per lane out of a column of a 72-halves-wide tile. That column is the access that padding and swizzling changes. The one ldmatrix instruction reads rows instead.

A fragment layout is a contract about registers. You can always satisfy it by hand, and you will always pay for that in the memory path rather than in the tensor core.

Pitfalls

ptxas rejects a shape you never meant to run on this card. Inline PTX is assembled for the target whether the code path can execute or not, so an unguarded m16n8k16 breaks the sm_75 build: ptxas /app/example.ptx, line 958; error : Feature '.m16n8k16' requires .target sm_80 or higher, captured 2026-09-01 by building this program with its guard forced open. Guard the PTX with #if __CUDA_ARCH__ >= 800, and note that __CUDA_ARCH__ is undefined in host code, so the host needs its own cudaDeviceProp check before it launches.

Everything fails below the minimum capability. At -arch=sm_70 the same file produces 101 errors of two kinds, ptxas /app/example.ptx, line 122; error : Feature 'ldmatrix' requires .target sm_75 or higher and ptxas /app/example.ptx, line 122; error : Modifier '.m8n8' requires .target sm_75 or higher, ending in ptxas fatal : Ptx assembly aborted due to errors. A V100 has tensor cores and neither of these; its shape is m8n8k4.

The answer has the wrong order and nothing faulted.

On a Turing card every lane must hand ldmatrix a valid address, including the lanes whose rows the instruction ignores: "For .target sm_75 or below, all threads must contain valid addresses. Otherwise, the behavior is undefined. For .num = .x1 and .num = .x2, addresses contained in lower threads can be copied to higher threads to achieve the expected behavior."

That is what the lane % 8 and lane % 16 in this program's address arithmetic are for.

The address is a shared-space address, and it is aligned. Pass a generic pointer and the behavior is undefined; convert with __cvta_generic_to_shared. Then keep every row start 16-byte aligned, because four lanes load 16 bytes together. Day 15's one-element pad is 2 bytes and breaks exactly that.

.trans transposed the tiles, not the group. An .x4 returns four independent 8x8 matrices and .trans transposes each on its own. It is not the transpose of the 16x16 they look like on paper.

Your ownership map looked right, so you shipped. A kernel printing where each lane stored its accumulator floats echoes the formula the store used, and m16n8k8's C/D coordinates are A's exactly. Check the values, from tiles whose cells are all distinct, or check nothing.

Go deeper

Next

Day 74 changes how this kernel loads each tile. Every k step here stops the warp at a barrier while the next tile arrives, and cp.async lets the copy for step n+1 run under the arithmetic for step n.

That overlap may account for part of prediction 4's remaining gap. Day 80 then puts the two together, with day 44's measured 67.7 percent of cuBLAS as the result to beat.