Our Benchmark Read an H200 at 48% of Spec. The Benchmark Was Wrong.

An NVIDIA H200 joined the network on 11 October. Our own verification measured its memory at 2,319 GB/s, 48% of the 4,800 on the spec sheet, and the page a renter sees said so, with a warning. The card was fine. Our benchmark was not.

This is the measurement that showed it, the eleven lines of CUDA that were wrong, and what we run now.

What the page said

Before a machine can be rented here, a verification job runs on it: it allocates the card's memory, measures compute with cuBLAS, and measures memory bandwidth. The results go on the machine's page, next to the figure from the manufacturer's spec sheet, and a card that measures under 60% of its spec gets a visible warning.

For the H200 that page read:

Memory bandwidth, measured by Paralon
2,319 GB/s
48% of spec
Spec sheet 4,800 GB/s per GPU

Everything else about the card looked right: 143 GB of memory reported, 128 GB of it allocatable, 647 TFLOPS in FP16. Only the bandwidth was off, and by half.

The pattern we should have seen

The same test, on the cards we see most:

The faster the memory, the further behind the test falls. A defective card does not produce that shape. A test with a ceiling of its own does.

That is a pattern, not a proof. An H200 in someone else's machine can also be power-capped, throttled by a hot chassis, or sitting behind a hypervisor that costs it something. So we rented the card, the way a customer would, and looked.

Ruling out the card

From a terminal inside the rental:

Nothing capped. Then the measurement itself, done another way. PyTorch was already in the rental image, and a copy between two tensors on the same device moves bytes through the card's memory and nothing else:

src = torch.empty(n, dtype=torch.uint8, device="cuda").fill_(1)
dst = torch.empty(n, dtype=torch.uint8, device="cuda")
dst.copy_(src)   # reads n bytes, writes n bytes

4,110 GB/s is 86% of the spec sheet, in line with what the consumer cards reach. The card does what it says.

The test that was wrong

This is what our verification ran, on a 256 MB buffer:

__global__ void memory_bandwidth_kernel(float *data, int N) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < N) {
        data[idx] = data[idx] * 2.0f + 1.0f;
    }
}

One thread per float, in blocks of 256: 64 million threads, each loading four bytes and storing four bytes. It counts two memory operations per element and divides by the time.

Our first guess was the buffer: 256 MB is small for a card with 141 GB. The table above rules that out. PyTorch reaches 4,013 GB/s on a quarter of a gigabyte.

The second check was the arithmetic. The same operation, x * 2 + 1 in place on float32, through PyTorch:

Same card, same buffer size, same arithmetic, and it moves memory at 4,000 GB/s. So it is neither the size nor the operation. It is the shape of the kernel: one thread per element. On a card whose memory is slower than the rate at which it can start and retire those threads, the memory is the limit and the test reads the memory. On an H200 the memory is no longer the slowest thing, and the test reads something else.

We did not profile where the time goes inside that kernel. What we can say is what the measurements show: the kernel tops out near 2,380 GB/s on this card while the card moves 4,100 GB/s.

The fix, and the two side by side

The test is now a copy inside the card, done by the driver:

cudaMemcpyAsync(dst, src, bytes, cudaMemcpyDeviceToDevice, 0);

Two buffers of 1 GB, read plus write, twenty times. On a card without 2 GB free it falls back to two of 256 MB and then two of 64 MB, so it never needs more memory than the kernel it replaced: a failed bandwidth test fails the whole verification, and we did not want a stricter test by accident.

Before shipping it we built both versions into one small binary, copied it to the rented H200, and ran them back to back:

The old kernel reproduces the figure that was on the page, so the binary measures what production measured. The new one lands on the PyTorch number.

After the new image went out, the card's next verification stored 4,062 GB/s, 85% of spec, and the warning is gone.

What we have not measured

• Consumer cards with the new test. As this is written, the H200 is the only card verified with it. An RTX 4090 that read 91% under the old kernel will read something a little different under the copy, and we do not know yet in which direction. Each machine keeps its old figure until its next verification.
• Other fast-memory cards. One H200 is one data point. We expect an H100 to have been under-read by the old test too, by less.
• The mechanism. We know the kernel was the ceiling. We have not shown which part of it.

What this changes for a renter

Nothing about how a rental works, and one thing about what to trust. The figure on a machine's page is a measurement, and a measurement has a method. When a number on this site looks wrong for a card, it may be the card, the host, or us. This time it was us.

The verification never refused the machine over this: the pass criterion is how much of the memory can actually be allocated. But until we checked, the page told renters that one of the most expensive cards on the network was running at half speed, which is the kind of error that costs a provider rentals they earned.