Maxing out dense FP16 FMA/FP32 accumulation (TFLOPS) on H200 GPU

Accelerated Computing CUDA CUDA Programming and Performance
1

Dear nVidia team and other forum visitors,

On H200 NVL GPU, I am trying to reach TFLOPS as reported on its datasheet:

On page 4, under FP16 Tensor Core (2: with sparsity), there is 1671 TFLOPS.
I assume dense=sparse / 2 (as this was the case in many other more detailed datasheets that I saw).

So I think 835.5 TFLOPS is correct number for dense FP16 FMA with FP32 accumulation.
Please correct me if this is wrong, ideally with a citation to a detailed technical document – I couldn’t find one for the H200 series.

Next, I am doing following simple bechmark.
Make matrix A of size M=K, and matrix B with size K, N with fp16 dtype. I sweep through two variables, M=K and N.
Record the time it takes to make large number of multiplications after a warmup.

The code is roughly this (full code attached):

import torch

def gemm_flops(m: int, n: int, k: int) -> float:
    return 2.0 * m * n * k


a = torch.randn(shape.m, shape.k, device="cuda", dtype=torch.float16)
b = torch.randn(shape.k, shape.n, device="cuda", dtype=torch.float16)
c = torch.empty(shape.m, shape.n, device="cuda", dtype=torch.float16)

flops = gemm_flops(shape.m, shape.n, shape.k)

for _ in range(config.warmup_iters):
    torch.mm(a, b, out=c)
torch.cuda.synchronize()

for _ in range(config.repeats):
    start_event = torch.cuda.Event(enable_timing=True)
    end_event = torch.cuda.Event(enable_timing=True)
    start_event.record()
    for _ in range(config.measure_iters):
        torch.mm(a, b, out=c)
    end_event.record()
    end_event.synchronize()
    
    elapsed_seconds = start_event.elapsed_time(end_event) / 1e3
    seconds_per_iter = elapsed_seconds / config.measure_iters
    tflops = flops / (seconds_per_iter * 1e12)

Maximum I was able to get is 712 TFLOPS, so about 85% of the max GPU TFLOPs.
What am I missing to reach 100% ?

On consumer GPU (RTX 3070) the same benchmark got up to 105%, which I think is due to overclocking.
But I have a hard time to close the 15% gap on the H200.

Some explanations that come to mind:

  1. Torch/other libraries in the stack do not use optimal matmul implementations.
  2. The gemm operations use another algorithm for matmuls than the O(n^3), so gemm_flops is wrong.
  3. There are unaccounted effects, like variable clock speed.
  4. Measurement is somehow wrong.
  5. I have wrong max FLOPs number.
  6. It’s indeed impossible to reach manufacturer’s claims of the theoretical flops.

I report also the power draw during the matmuls.

gemm_efficiency
gemm_efficiency1500×900 257 KB

system.txt (18.5 KB)

gemm_bench.py.txt (11.3 KB)

Thank you very much for any help / pointers!

Sincerly,
Michal Sustr

2

It’s generally not possible to reach a GPU theoretical peak throughput number with a directed test. There are multiple possible contributors to this, but one possible contributor is that when a higher level operation is decomposed to the lower level operation for which a peak throughput is actually specified, for again multiple possible reasons, the code sequence is not purely a sequence of the theoretical op.

You can find other forum posts and internet posts that discuss the potential reasons along with specific cases. Here is one that mentions that for some GPUs, power limits may also play a role.

3

I am not sure if this would change the results, but you could attempt to turn off ECC on the H200, as this is a feature that the consumer GPUs do not have - and it affects memory throughput.

But then it might not help at all if that benchmark is not bandwidth limited

Christian

4

The performance difference in Tensor Core between the H200 and the RTX 3070 is enormous.

You are using a system benchmarking method measuring seconds of a long execution. To realize the H200 Tensor core speed everything at least on the GPU side has to be programmed in an ideal way.

Getting 85% of the theoretical speed of the actual operations out, is a great result. Perhaps more is possible with further optimization, but the higher the performance, the more sensitive the system for anything not running in an ideal way.

5

Thank you very much for the replies. The ECC sounds interesting, it might be hard for me to do though (non-root machine access).

Following up on the theoretical perf. q., I am also curious, is there a link for thorough official tables of technical parameters of the GPUs? Ideally in some format that is friendly for machine processing (CSV/JSON or such).

This would be good for e.g. simple comparisons of evolution of hardware. There are tables on Wikipedia and other sites, but inspecting them closely, they are of low quality: they do not cite where the numbers come from, and often misinterpret categories (for example, sparse/dense, dtypes etc.)

The datasheets, like the one I linked, are sparse on information, and it gets worse over time - V100 was best quality, B-series are the worst :-/

I have painstakingly hand-collected collected some information, but it is only targeted to what I was looking at - FLOPs and power. I had to do this without AI, as it would get blunders as well. And unless I benchmark it, I can’t be sure, but at the same time, practical perf. does not match theoretical one as seen here. So it’s hard to know what numbers are real.

If you can help, thank you very much.

gpu_flops.csv.txt (4.3 KB)

6

Some performance numbers are stored in Nsight Compute for its roof percentages. Not sure, if they can get easily exported.

As I understand your list (with some assumptions by me):

To have comparable numbers, you choose dense, half-precision, FP16 accumulation. The numbers are in FLOP (so multiplications + additions). They only specify the tensor core performance (some marketing material adds scalar and tensor core performance). The numbers are about the whole GPU (not a single SM or tensor core). It is calculated at boost frequency of the Nvidia reference architecture or official products for this graphics chip. The numbers are the theoretical performance, not from benchmarks.

7

Yes you understand correctly.

I don’t know if the roofline numbers are stored, but I suspect they are derived from elementary technical numbers you can get through CUDA API.

It would be nice to have these published by the vendor, rather than having to get hands on every GPU and ask its numbers through the API. Maybe it exists officially somewhere, I just couldn’t find it. If you know about some official link, I would appreciate it a lot!

8

The idea with Nsight Compute was that it could have the data even with the GPU not present.

Normally the numbers are scalable for each compute capability by the number of SMs (and the frequency). So not each GPU is needed separately, just one per compute capability.

(with datagcenter Blackwell Nvidia started to offer TPC-wide MMA instructions, a TPC built as a pair of SMs. Once it goes higher up, e.g. GPC, the simple scaling by SM no longer may be enough)

9

The theoretical peak performance numbers at the lowest level are usually powers of two:

V100: 1024 flops/cycle/SM (F16 mma).
A100: 2048 flops/cycle/SM (dense F16 mma).
H100/200: 4096 flops/cycle/SM (dense FP16 wgmma).
B200/300: 8192 flops/cycle/SM (dense FP16 tcgen05.mma).
Other workstation/consumer Ampere/Ada/Blackwell GPUs: 1024 flops/cycle/SM (dense F16 mma with F16 accumulation. F32 accumulation is half speed on some consumer GPUs).

The claimed overall tensor flops are then derived by multiplying those numbers with the SM count and boost clock (which both may depend on the specific GPU variant).
Most real use cases (including GEMM) have some inefficiencies that mean you’re not getting close to 100%, but you can get very close to this theoretical throughput at the per-cycle level if all you do is issue many mma instructions. You may still see power/thermal throttling, though.