DeepGEMM: clean and efficient FP8 GEMM kernels with fine-grained scaling
50 comments
·February 26, 2025Bimos
> FFMA SASS interleaving
> 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.
shaklee3
scott grey figured out this exact thing and more back in 2015 for maxwell, and it's been written about many times since by other people.
blackeyeblitzar
From what I read elsewhere, this is the type of typical performance optimization for matrix math you would see when performance is critical. It’s just not been applied yet to this specific problem by other AI players since it wasn’t a necessity for other companies. But eventually everyone would probably end up here regardless.
mitthrowaway2
How many people does it take to implement this? A 10% gain in performance could pay for a lot of people's salaries when your company is spending hundreds of millions on GPU clusters.
fulafel
If you think how many people who looked and failed to realize this optimization in the preceding performance efforts of the community, you could argue for quite a big number.
rfoo
Uh, three? I worked at $CORP where we had a three people sub-team, they reverse engineered most of Volta's SASS instruction encoding, built a working SASS assembler (before the open source one of course), with the ultimate goal of making GEMM / Conv faster. And they did it. Though it wasn't applied to a high-profile enough big picture so we never heard about it :>
If you don't believe me, previous open source SASS assemblers were mostly from university, they surely didn't have that many people.
saagarjha
I mean it’s not a significant change so one? But that isn’t to say anyone could do it.
Bimos
I think most AI players rely on high performance GEMM. But most people would be satisfied with cutlass or cublas, and the others implement gemm themselves, but not necessarily use undocumented features?
creato
Using undocumented features is not rare. People reverse engineered Apple's undocumented AMX instructions on their CPU, and I know people use undocumented/private extensions for several different kinds of GPUs.
Zacharias030
I‘ve only seen it done by hedge funds so far. What were you referring to?
fracon
talk and commenting is quite cheap, isnt it
ETH_start
[flagged]
dang
Literally literally means not literally.
I love it when words turn into their opposites!
tough
I think he might mean hyperbolically figuratively so
Bimos
I edited it.
kneegerman
orthogonally
jmward01
I'm not sure the lower and lower precision optimization is a good idea long term. It indicates that models are really sparse and that may be true right now but I think that is likely just because we have some bad ideas about how to train them and not because they really should be that sparse.
rfoo
Well, let's enjoy free "sparsity" until it doesn't. Being able to train a really good model but in higher precision only is a research problem. Low precision training and inference is an engineering one.
We've been doing this since CNN days (9 years ago if not more), and I believe we have a good few years left.
sudosysgen
The activation function throws away much of the dynamic range of floating point numbers, it's relatively clear that having a lot of range where the activation is already saturated is unlikely to be useful.
jmward01
That depends on the activation function. I personally think Layernorm is destroying density (and have some solid evidence for this) but it is in use because there is a lot of missing supporting structure to really help pump data into the weights so it helps so long as we are using simple linear combinations.
nbonaparte
This might be rendered moot by native microscaling support in Blackwell (MXFP). They've manually done a coarser-grained version of that for Hopper, but with full FP32 scaling factors.
WiSaGaN
I think these kind of open-source is really showing their objective of achieving efficiency in the industry. The reason is this kind of software benefits a lot to the big guys serving the model (competitors to Deekseek themselves if they are interested in being a provider) rather than to the general open-source community that wants to learn and tinker or serve model in consumer hardware.
fspeech
Efficiency could lead to cheaper hardware for everyone, themselves included.
null
fulafel
It seems mostly Python, which is nice: https://github.com/deepseek-ai/DeepGEMM/blob/main/deep_gemm/...
azurezyq
The secret is here https://github.com/deepseek-ai/DeepGEMM/blob/main/deep_gemm%...
fulafel
And the yield bit modification: https://github.com/deepseek-ai/DeepGEMM/blob/main/deep_gemm/...
dr_kretyn
Honestly, this is beyond my usage and understanding. But I really appreciate such sharing findings and improvements so that everyone can benefit from them. It's a refreshment.
greenavocado
FFMA (Fused Floating-point Multiply-Add) is a fundamental GPU instruction that performs D = A*B + C in a single operation. This instruction is critical for matrix multiplication and deep learning workloads.
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.
jarbus
This is crazy insightful, thanks! I’d really love to learn how to get to this level of understanding, but can’t seem to figure out what curriculum I’d follow where I’d end up with this level of technical competence.
randomNumber7
You need to understand how the gpu architecture works on a abstract level. Try to understand the SIMT (Single Instruction Multiple Threads) principle. Doing some shader programming or writing a cuda kernel could be a nice exercise. In a nutshell, if you want to add two vectors with hundred elements, instead of looping from 0 to 99 you would call a function called "kernel" (or "shader" in graphics programming) 100 times and pass it different indices.
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.
apples_oranges
Perhaps I know as much as you, but to begin, I would dive into CUDA and running code on GPUs.
fotta
They should call the warp that is yielded to the weft.
dekhn
No, that doesn't make sense; both the yielder and yieldee are warps, the PC is the weft (approximately).
loginx
I'm gonna ask my chatgpt to write like you ;)
I went from understanding none of it, to everything making sense. Thanks!
Zacharias030
HN at its best! What do you do for a living, Sir,
rramadass
Very Nice!
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?
m3kw9
The 20$ question, what can I do with this?
devit
Multiply FP8 matrices with FP32 scaling factors giving a bfloat16 matrix result on an nVidia Hopper or newer GPU.
Maxious
Just tested and it doesn't work out of the box on the consumer 50 series ie. 5080:
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
devit
Perhaps your card has less per-SM shared memory than the GPUs DeepSeek uses.
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".
wenc
It says:
> DeepGEMM exclusively supports NVIDIA Hopper tensor cores
null
dbfudyfg
[dead]
cde-v
Interesting timing with NVDA releasing results tomorrow.
wenc
Lol. I don't think the people that buy/sell NVDA even know what this is about.
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.
porridgeraisin
Careful, they're going to tell you it's "priced in" next.
This kind of stuff is an intersting demonstration of how far compilers are from extracting high performance from hardware based on high level code.
What would it take for traditional compiler tech or AI assisted optimization agents to come up with something like it?