Owen suggested I bring this over from Discord.
Setup: RTX 3090 (GA102, sm_86), Mojo 1.0 / MAX 26.5, clocks pinned at 1695 MHz, median of 30 samples with IQR. Three kernels written twice, once in CUDA and once in Mojo, same algorithm on both sides.
The memory-bound kernels are a dead heat. Reduction at 256M elements: 895 GB/s in Mojo, 889 in CUDA, 895 for cub::DeviceReduce, all 95-96% of the 936 GB/s bus, IQRs overlapping. Softmax on 16384x1024 is the same story.
The fp32 matmul is where the two sides split:
1024³: Mojo 8 497 GFLOP/s vs CUDA 9 279 (92%)
4096³: Mojo 13 780 vs CUDA 17 370 (79%)
Both sides use a 128x128 block tile with BK=8, an 8x8 register tile per thread (256 threads/block), float4-vectorized global loads and a transposed As tile. The compute inner loop is the same shape in both.
The one thing the CUDA kernel does that the Mojo one does not is double buffering: prefetching the next K-tile’s global loads into registers while the current tile’s FMAs run, then writing them to the alternate shared buffer.
That is where I believe the whole gap lives, since it widens with problem size exactly as you would expect from an unhidden global load latency, and it is not memory movement (both sides already vectorize) and not the arithmetic (identical inner loop).
My question is what the idiomatic Mojo expression of that is. In CUDA I write it by hand: two shared buffers, an explicit register staging array, and a barrier placement that lets the loads issue before the FMAs. In Mojo I am not sure whether the answer is to write the same thing manually, or whether there is a construct in the stdlib / layout library intended for this that I have missed.
Repo with methodology, roofline and raw CSVs:
Happy to paste the current Mojo inner loop if that is useful.