Please help me correct my understanding of H100 Tensor Core and the WGMMA instruction

Accelerated Computing CUDA CUDA Programming and Performance
1

I have recently started working on optimizing tensor core–related algorithms on the H100, but I keep encountering various performance issues. Here I summarize my current understanding of the H100 architecture (mainly the Tensor Core–related aspects).

  1. Each SM is capable of sustaining up to 128 register reads and 128 register writes per cycle.

  2. For the instruction wgmma.mma_async.sync.aligned.shape{.satfinite}.s32.s8.s8, the theoretical peak throughput on Tensor Cores is 4096 MAC operations per cycle per SM.

  3. The WGMMA instruction does not necessarily write the updated C values back to registers after every execution, as this would impose a significant bandwidth pressure on the register file. For example, for the ss mode instruction m64n128k32.s32.s8.s8, if each execution required reading from and writing back the C matrix in the register file, it would on average require 128 register accesses per cycle (since the C matrix contains 64×128 elements; the average throughput is 64×128×32 / 4 / 1024 = 64 cycles).

  4. In some special cases, redundant reads and writes of the C registers can be avoided. For example, when two consecutive WGMMA instructions use the same C input and output, the first WGMMA does not need to write back the C matrix, and the second WGMMA does not need to read it again.

    As for the timing of writing the C matrix back to registers, I believe there are at least two cases:

    1. When a wait_group instruction is explicitly executed, the C data must be written back to ensure visibility to subsequent instructions.

    2. When subsequent WGMMA operations work on different matrices, and the register cache cannot hold multiple matrices simultaneously, the old matrix must be evicted (written back).

    3. The submitted WGMMA instructions are reordered within the group by the async proxy or other mechanisms, while still ensuring correctness, so that the output of a previous WGMMA can be used as the input to the next one as much as possible, thereby reducing register file accesses.

      Additional question: is reordering allowed across groups

  5. Tensor Cores execute the matrix multiplication described by WGMMA at a finer granularity. It takes some time for the internal Tensor Core pipeline to reach full utilization (based on discussions with engineers working on other AI accelerator designs). Pipeline disruptions can significantly impact performance. For instance, issuing a WGMMA followed immediately by a wait_group 0, and repeating this pattern, is highly inefficient. Therefore, WGMMA instructions should be batched, and synchronization should also be performed in batches.

Final question: are there any best practices for developing Tensor Core programs on the H100? So far, I have only found PTX documentation, but I have not come across any materials similar to best-practice guides. If such resources exist, could you please share them with me?

2

The register file is split into even and odd banks. Each SMSP is capable of reading two warp-wide registers per cycle, one from the even and one from the odd bank, for a total of 4(SMSPs)x2(banks)x32(warpsize)=256 register reads (32b) per cycle. Then there’s the register reuse caches on top of that.

There is a possibility that TC instructions can write to both banks in the same cycle. However, as far as I know all instructions producing wide results (64+b) take more than one cycle so this is hard to measure.

Is this actually the case? If so, I would expect it to be done by ptxas at compile time so the order is static and observable in SASS code.
Reordering across groups makes little sense IMO.