> We observe a performance improvement in the CUTLASS FP8 kernel between NVCC 12.2 and 12.3. By comparing the compiled SASS, we discover that one bit in a series of FADD instructions is flipped in an interleaving pattern. After referencing some open-source CUDA assembler implementations, we identified that this bit controls yield, which may enhance warp-level parallelism (just a guess, yielding the current warp and let other warps work).

> To leverage this, we develop a similar script to modify the FFMA instructions in the compiled binary. Besides simply modifying the yield bit, we also flip the reuse bit (registers cannot be reused if the warp is yielded). This adjustment improves performance (10%+ in some cases) for fine-grained scaling FP8 GEMMs by creating more opportunities to overlap MMA instructions with promotion FFMA instructions.

I would say it is really mind-blowing.

If you don't believe me, previous open source SASS assemblers were mostly from university, they surely didn't have that many people.

I guess since the majority here are blown away by the very low-level code involved, it tells me that they're likely not ready to use it or have been stuck on very high level tools that abstract this away.

All cutlass results I have seen so far for Gemm are within ~10% of cuBLAS. If the 2x-2.5x speedup they report holds up that would be extremely impressive.

First of all, cuBLAS needs the cuBLASLt extension API for mixed-precision workloads to handle FP8. Second, some adequate type combinations, like E5M2 x E5M2 for A x B, are not supported, while others, like E5M2 x E4M3, are! Moreover, matrix A must always come in a transposed layout for Ampere, Hopper, and Blackwell... and the list of constraints goes on.

I've integrated FP8 cuBLASLt benchmarks into my "Less Slow C++" repository <https://github.com/ashvardanian/less_slow.cpp>, adding to the list of existing cuBLAS and hand-rolled CUDA and PTX benchmarks. I'm running them on H200 GPUs, which should have the same performance as H100. For square inputs, the throughput peaks around 1.35 Peta-ops.

  --------------------------------------------------------------------------------------------------
  Benchmark                                        Time             CPU   Iterations UserCounters...
  --------------------------------------------------------------------------------------------------
  cublaslt_tops/256         12496 ns        12496 ns        56284 TOP=2.67999T/s
  cublaslt_tops/512         13089 ns        13089 ns        53100 TOP=20.4883T/s
  cublaslt_tops/1024        14882 ns        14882 ns        46918 TOP=144.23T/s
  cublaslt_tops/2048        25802 ns        25802 ns        26869 TOP=665.679T/s
  cublaslt_tops/4096       109316 ns       109313 ns         6021 TOP=1.25715P/s
  cublaslt_tops/8192       821080 ns       821050 ns          629 TOP=1.33907P/s
  cublaslt_tops/16384     7135472 ns      7135461 ns           93 TOP=1.23269P/s
  cublaslt_tops_BigO         0.00 N^3        0.00 N^3  
  cublaslt_tops_RMS             2 %             2 % 

What would it take for traditional compiler tech or AI assisted optimization agents to come up with something like it?

We've been doing this since CNN days (9 years ago if not more), and I believe we have a good few years left.

You make a good point about LayerNorm, it's probably even worse.

These are very good and high profile public demonstrations of where $NVDA's moat is: that GPGPU is very flexible and you can program to do a lot of stuff that makes perfect sense but wasn't in the mind of hardware vendors.

Now, if you predict the future to eventually converge on more and more dedicated hardware support, to the point that there's no more software optimizations like these, then the so-called "CUDA moat" breaks.

To stay in this game, NVIDIA is breaking down their own moat :p

Wouldn’t it make sense to provide these to the user? Even if they might not be perfectly reliable.

This stuff must be documented internally, why not just release it?

Security by obscurity does not work: Your competitor reverse engineer everything you do anyways.

Probably no. They are likely only documented in architectural design doc / spec etc which you surely do not want to share.

In NVIDIA's SASS (Streaming Assembly), FFMA instructions are encoded as 64-bit or 128-bit instructions with various control bits that determine their exact behavior.

When the yield bit is set the bit tells the warp scheduler that the current warp can yield execution after this instruction. The hardware can then schedule a different warp to execute, potentially hiding latency.

GPUs achieve high throughput through massive parallelism. When one warp stalls (e.g., waiting for memory), others can proceed. The yield bit creates explicit opportunities for the scheduler to switch warps.

This bit indicates whether the source registers can be reused immediately in subsequent operations. When the yield bit is set, the reuse bit must be cleared. If a warp yields, it might not be the next one to execute. Another warp might modify the register file state. The hardware cannot guarantee register values will remain unchanged across yields.

By setting the yield bit in an alternating pattern across FFMA instructions, the compiler creates explicit scheduling points where other warps can make progress. When modifying the yield bit, they also had to clear the reuse bit for affected instructions to maintain correctness. This modification specifically helps overlap two types of operations: MMA (Matrix Multiply-Accumulate) instructions: Heavy compute operations that form the core of matrix multiplication, and Promotion FFMA instructions: Operations that convert between precision formats (likely FP8 to higher precision for accumulation)

FP8 (8-bit floating point) GEMM operations have specific characteristics that make this optimization particularly effective. FP8 calculations typically require conversion to higher precision for accumulation and back, creating additional FFMA operations. FP8 reduces memory bandwidth requirements but creates complex computation patterns with promotion/demotion operations. The mention of "fine-grained scaling" suggests these are operations where precision is carefully managed at multiple points in the calculation.

The yield bit manipulation creates a more optimal interleaving of compute operations and format conversions, allowing the GPU to utilize its execution units more efficiently. Without this optimization, the warp scheduler might not find natural opportunities to switch between warps, leading to underutilization of compute resources.

Then research how it is realized on the hardware with "warp"s or "wavefront"s (on AMD i think). How the cache works is also very important here. Sadly the information on the internet is relatively sparse here.

Can you recommend some good resources/books on GPU/TPU/ML Accelerators/etc. architecture/ISA where i can read the above details? Also on Computer Math where one can study how FP8/etc. works?

Amazon search brought up the following two interesting books, perhaps somebody who has browsed/read them can chime in;

1) Advanced GPU Assembly Programming: A Technical Reference for NVIDIA and AMD Architectures by Gareth Thomas.

2) Numerical Computations with GPUs edited by Volodymyr Kindratenko.

    Testing GEMM:
    Assertion failed: 
    deep_gemm/jit/../include/deep_gemm/fp8_gemm.cuh:369, condition: cudaFuncSetAttribute(kernel, 
 cudaFuncAttributeMaxDynamicSharedMemorySize, smem_size) == cudaSuccess
    terminate called after throwing an instance of 
    'AssertionException'
      what():  Assertion failed: cudaFuncSetAttribute(kernel, cudaFuncAttributeMaxDynamicSharedMemorySize, smem_size) == cudaSuccess

Try to lower the sm90_capacity value in gemm.py: I think 128KB is the correct value for RTX 5080 compared to 256KB for the H100/H800.

And probably add ", 3, 2, 1" after "6, 5, 4".

This is a highly specialized linear algebra library to do general matrix-matrix multiplications for low-precision floats (FP8, vs FP32 (float), FP64 (double)) while maintaining accuracy.