Can RTX 6000 Match H100?

Lesson

Hypothesis: exposing independent FP32 multiply-add chains can change the balance between instruction throughput, dependency latency and register pressure. We expect the RTX 6000 Ada implementation may reach at least 80% of a pinned H100 implementation. Parity and a higher score are stretch goals, not promised outcomes.

Evaluate the same degree-512 polynomial at 67,108,864 points with Horner's method, using explicit FP32 fmaf in the prescribed order. The deterministic inputs, coefficients and precision stay fixed. The harness also checks boundary sizes and alternate coefficients, then verifies every large-case output outside timing. You may change how work is assigned to threads and reuse coefficients; do not change the mathematical workload or specialize to the fixture values.

Start with one Horner chain per thread. Try 2, 4 or 8 independent chains, block sizes 128/256/512, and different grid sizes. Each experiment should change one variable. More chains can hide dependency latency, but use more registers and may spill. Measure native CUDA-event latency first; use the profile to investigate the cause.

Score = 100 × pinned H100 median time ÷ your RTX 6000 Ada median time. Correct, stable runs at 80% or more pass. 100% is parity. 80% throughput means your latency is 1.25× H100 latency. The reference and threshold are frozen for each submitted run. Until an operator publishes measured H100 calibration, results remain ungraded.

A challenge score compares your implementation with the published reference implementation. To investigate hardware alone, run exactly the same code and launch configuration on both devices. To demonstrate performance portability, measure both variants on both devices: A/H100, A/RTX, B/H100, B/RTX. Repeat each cell. Only claim a ranking reversal if the winner changes consistently beyond run-to-run noise. H100 winning every cell is also a valid experimental result.

This is a hypothesis about one pinned polynomial workload, not a general GPU ranking. Keep the calibration version and IQR with each comparison and review the run cost before capture.

What you will learn

• Separate algorithm choice, implementation tuning, and GPU effects.
• Use native timing for a performance claim and profiling to explain it.
• Measure a two-variant, two-GPU matrix before claiming that the best implementation depends on the device.

Starter code

    const int stride = blockDim.x * gridDim.x;
    for (int i = blockIdx.x * blockDim.x + threadIdx.x; i < N; i += stride) {
        const float x = X[i];
        float y = coefficients[degree];
        for (int k = degree - 1; k >= 0; --k)
            y = fmaf(y, x, coefficients[k]);
        Y[i] = y;
    }