Case study. Kernel correctness on Kaggle GPUs

A softmax wrapper fault on Tesla T4, fixed to pass 48 of 48.

The Triton fused softmax tutorial crashed on narrow rows before the kernel ever ran. We traced it to the host wrapper, corrected only that, and measured the result against PyTorch on the same GPU.

Book a demo See the engine

Softmax cases, run 02 and run 04Measured

Run 02, original wrapper36 passedRun 04, corrected wrapper48 of 48

Same shapes, input patterns, dtypes and seed in both runs. Tolerances unchanged.

The workload

Five runs, 80 distinct configurations.

Vector add and fused softmax from the Triton tutorials, each against a PyTorch control, in float32 and float16. Every notebook had two Tesla T4 GPUs, compute capability 7.5, with the work on cuda:0.

  • Coverage. 80 distinct configurations of operation, shape, dtype, input pattern and implementation, 128 passing case records and 384 repeat observations.
  • Shapes. Add lengths from 1 to 1,048,577. Softmax from 1×1 to 256×1025 with random, ×100 and zero inputs.
  • Stack. PyTorch 2.11.0+cu128, CUDA 12.8, Triton 3.6.0, driver 580.178.04, pinned container image.

Run sequenceMeasured

01-addVector add, lengths 1 to 65,53720 of 20 passed02-softmaxSoftmax 1×1 to 31×102536 passed03-benchTiming, add and softmax16 of 16 passed04-correctedRun 02 again, fixed wrapper48 of 48 passed05-fairEqual preallocated outputs8 of 8 passed

Run 02 errors were preserved, then rerun on the corrected wrapper in run 04.

What changed

One division by zero, fixed on the host.

The tutorial wrapper estimates occupancy by dividing by the shared memory each program needs. Softmax widths 1 and 7 compiled to zero shared memory bytes in both dtypes, so the wrapper raised ZeroDivisionError before dispatch. Width 129 used 32 bytes and ran.

  • Change. When the requirement is zero, skip only the shared memory term. Keep the register limit and cap residency by the thread limit the device reports.
  • Kernel. GPU arithmetic, precision and the kernel algorithm stay the same. The masked row loop still visits each row once.
  • Upstream. The current official tutorial at revision 9fbab754 still carries the same division.

Occupancy pathRecorded

Original wrapper, occupancy estimateoccupancy = shared memory / size_smemwidths 1 and 7, size_smem = 0 bytes, ZeroDivisionErrorCorrected wrappersize_smem = 0, skip the shared memory termregister limit keptcapped by device thread limit

Six static wrapper checks passed before the corrected GPU run.

Before and after

Where the time goes, host or GPU.

The fair rerun gave both sides the same inputs and preallocated outputs, alternated the order, and timed two ways. End to end includes Python dispatch and synchronization. GPU graph timing captures only the GPU work.

End to end p50, µsMeasured

add 1,048,577 float32PyTorch 71.08 µsTriton 98.94 µsadd 1,048,577 float16PyTorch 42.61 µsTriton 90.65 µssoftmax 256×1025 float32PyTorch 36.73 µsTriton 82.57 µssoftmax 256×1025 float16PyTorch 37.59 µsTriton 85.33 µs
PyTorchTriton

Two Tesla T4 allocated, run on cuda:0. 30 synchronized host samples per side, includes Python dispatch and sync.

GPU graph p50, µsMeasured

add 1,048,577 float32PyTorch 53.87 µsTriton 65.06 µsadd 1,048,577 float16PyTorch 25.54 µsTriton 54.60 µssoftmax 256×1025 float32PyTorch 17.28 µsTriton 13.18 µssoftmax 256×1025 float16PyTorch 19.20 µsTriton 14.19 µs
PyTorchTriton

20 GPU graph averages per side from triton.testing.do_bench_cudagraph. Excludes Python dispatch and allocation.

The checks

Every case against a float64 reference.

Inputs used a fixed seed. A CPU reference ran on the same quantized inputs in NumPy float64. Each case ran three times and had to match tolerance, stay finite and repeat bit for bit.

  • Tolerances. float32 atol 1e−6, rtol 1e−5. float16 atol 1e−3, rtol 2e−3. Set before the first run and never loosened.
  • Largest add error. 2.3841858e−7 in float32 and 0.001953125 in float16, both inside tolerance.
  • Provenance. Original and corrected modules each carry their own SHA-256, with all 88 evidence files hash checked.

CorrectnessMeasured

48 of 48 cases passed after the correction24 Triton softmax cases and 24 PyTorch controls
8 of 8 fair comparison checks passed4 Triton and 4 PyTorch, equal preallocated outputs

Three repeats per case, exact repeat equality required.

FAQ

Questions engineers ask.

What exactly was measured?

Correctness of Triton vector add and fused softmax against PyTorch on Tesla T4 GPUs, and p50 time per call for add at 1,048,577 elements and softmax at 256×1025, end to end and in GPU graph replay.

Did the fix change the kernel?

No. Only the host wrapper changed. GPU arithmetic, precision and the kernel algorithm are the same, and the tolerances stayed fixed.

Why is PyTorch faster end to end?

End to end timing includes Python dispatch and synchronization. In GPU graph replay, which removes that host cost, Triton softmax is about 1.31× and 1.35× faster.

How can an engineer check this?

Each module and notebook carries a SHA-256 hash, and every run is tied to a pinned container image and a fixed seed.

Next step

Bring the kernel that fails or stalls.

The engine runs it on real hardware, finds where the time and the faults sit, and corrects them layer by layer. See how the engine works or book a demo.

Book a demo See the engine