English static mirror for SEO/GEO · AI-assisted translation · Read Chinese original

DeepGEMM: Why DeepSeek Built Its Own BLAS Kernels

Forum topic · ✨步子哥 · 2026-10-06

Summary

DeepGEMM is DeepSeek's open-source GEMM kernel library that achieves up to 1550 TFLOPS on FP8 GEMM on the H800 GPU — roughly 95% of theoretical peak — something general-purpose libraries like cuBLAS cannot match. This post analyzes why DeepSeek wrote its own kernels: modern LLMs now require non-standard primitives (fused MoE, MQA scoring, HyperConnection layers, sparse indexing) that vendor libraries don't support. The library uses DeepJIT, a runtime compilation approach that generates raw CUDA code (leveraging TMA, wgmma, tcgen05 concepts) without requiring a CUDA toolkit at install time. It deliberately trades generality for performance and simplicity, containing only a handful of clean, maintainable core kernels. A September 2026 update added Ascend (Huawei) support, signaling serious preparation for NVIDIA alternatives. The post frames DeepGEMM as a key piece of DeepSeek's vertical integration strategy — owning model architecture, inference framework, and kernels — which enables the extreme inference cost efficiency behind their low API pricing.

> Project: deepseek-ai/DeepGEMM · 363 stars/day · CUDA

A Counterintuitive Choice

NVIDIA provides cuBLAS — a BLAS library refined over more than a decade, supporting nearly all matrix shapes and data types. CUTLASS goes further with composable GEMM templates. The open-source ecosystem has Triton, FlashAttention, and xFormers. So why did DeepSeek write its own GEMM kernels?

The answer hides in one specific number: 1550 TFLOPS on the H800. That is DeepGEMM's peak performance on FP8 GEMM, approaching 95% of the H800's theoretical FP8 compute ceiling.

cuBLAS cannot reach that number. Not because cuBLAS isn't good, but because it's a general-purpose library — it must support every shape, every data type, every GPU architecture. DeepSeek only needs the shapes and data types its own models use. The cost of generality is performance; DeepGEMM trades away generality for extreme performance.

But the story is more complicated than "in-house is more efficient."

Not Just GEMM Anymore

Despite the name, DeepGEMM is no longer just GEMM. The September 2026 update turned it into a unified library covering the core compute primitives of modern LLMs:

| Primitive | Purpose | Why in-house | |------|------|---------------| | FP8 / FP4 / BF16 GEMM | Core matrix multiplication | Mixed-precision training and inference | | Fused MoE (Mega MoE) | Mixture-of-Experts models | Communication/compute overlap | | MQA Scoring | Lightning Indexer | DeepSeek V3.2's compressed attention | | HyperConnection (HC) | Inter-layer connections in DeepSeek V4 | Non-standard architecture, no existing implementation | | Sparse Indexer | Sparse attention indexing | Deep customization |

Note the last row — HyperConnection. This is an inter-layer connection mechanism specific to the DeepSeek V4 architecture. It's not standard attention, not a standard MLP — it's DeepSeek's own design. No existing library supports it, because it's an entirely new compute primitive.

This exposes a structural problem: when model architecture innovation outpaces library updates, writing your own kernels isn't an optimization choice — it's the only choice. cuBLAS won't write a dedicated kernel for your new architecture unless you're one of NVIDIA's most important customers. CUTLASS offers templates, but template composition is complex and can't necessarily cover non-standard operations.

DeepSeek's choice: write it yourself. Then open-source it.

DeepJIT: No Compile at Install, JIT at Runtime

DeepGEMM's most interesting engineering decision is DeepJIT — runtime JIT compilation.

Traditional CUDA libraries: write CUDA code → compile with NVCC into .so or .cubin → users get prebuilt binaries at install → load at runtime. The problem: different GPU architectures (SM80, SM90, SM100) need different build artifacts, CUDA versions have compatibility quirks, and users frequently hit "compile failed" or "architecture mismatch" errors.

DeepGEMM's approach: compile no CUDA code at install time. All kernels are dynamically compiled at runtime via DeepJIT. Users pip install deep_gemm without needing a CUDA toolkit, without NVCC version conflicts, without architecture mismatch worries — the JIT detects the GPU architecture at runtime, generates the corresponding PTX, and compiles/loads it.

This isn't a new concept — Triton does something similar. But DeepGEMM's JIT differs in one key way: it generates raw CUDA code, not Triton IR. This means it can leverage CUTLASS and CuTe concepts (TMA, wgmma, tcgen05) without depending on their template systems.

An analogy: Triton is like Python — you write high-level code and the compiler optimizes it. DeepJIT is like C++ templates — you write close-to-hardware code and the compiler instantiates it at runtime. The former is easier to use; the latter is more flexible. DeepSeek chose the latter because they need to squeeze out the last drop of performance.

Simplicity as a Design Goal

The DeepGEMM README contains a sentence that's easy to overlook:

> The library is designed for simplicity, with only a limited number of core kernel functions, making it a clean and accessible resource for learning NVIDIA GPU kernel optimization techniques.

"Designed for simplicity" doesn't sound like something a library chasing peak performance would say. But on reflection, it makes sense.

DeepGEMM has only a handful of core kernel functions, each mapping to a well-defined compute primitive (dense GEMM, grouped GEMM, masked GEMM, MoE GEMM). No deep template hierarchies, no code structures requiring ten layers of abstraction to modify.

This contrasts with CUTLASS, whose template depth intimidates many people — modifying one kernel requires understanding concepts like TiledMma, ThrMma, SmemCopy, R2S, and S2R. DeepGEMM simplifies these concepts down to the essentials.

Simplicity isn't just about ease of use — it's about maintainability. DeepSeek needs to iterate fast: when a new model architecture (like V4's HyperConnection) needs a new kernel, they must be able to implement and deploy it quickly. If the library itself were overly complex, with every new primitive requiring edits across ten template layers, iteration speed would become unacceptable.

There's a trade-off here: generality vs. performance vs. simplicity — you can pick at most two. cuBLAS picked generality + performance (sacrificing simplicity, since it's closed-source). CUTLASS picked generality + simplicity (sacrificing some performance due to template overhead). DeepGEMM picked performance + simplicity (sacrificing generality, supporting only specific shapes and architectures).

DeepSeek can make this choice because they control both the model architecture and the inference framework. The architecture determines which primitives are needed; the framework determines how they're invoked. With both in-house, generality stops being a requirement — you don't need to support every shape, only the shapes your models use.

Ascend Support: Beyond NVIDIA

The September 30, 2026 update: DeepGEMM-Ascend is released.

This means DeepGEMM is no longer NVIDIA-only — it now supports Huawei's Ascend chips. This is not a simple port but an architectural extension: Ascend's Da Vinci architecture and NVIDIA's Tensor Cores are fundamentally different programming models.

Why does this matter? It shows DeepSeek is seriously considering NVIDIA alternatives — not just in rhetoric, but with preparation at the lowest level of the compute stack.

If DeepSeek V5 were to run inference on Ascend, it wouldn't be a matter of moving model weights over — the entire compute stack would need adaptation. From GEMM to MoE to attention, every primitive needs a high-performance Ascend implementation. DeepGEMM-Ascend is the first step in that adaptation.

It also vindicates the DeepJIT design. If DeepGEMM shipped prebuilt NVIDIA binaries, extending to Ascend would mean rewriting the entire compilation pipeline. With JIT, you only add a new code generation backend — the flexibility of runtime compilation becomes a real advantage when scaling across architectures.

What This Means

On the surface, DeepGEMM is a GEMM library. In reality, it's a foundational piece of a vertical integration strategy.

DeepSeek controls the whole stack: model architecture (V3/V4 design) → training framework → inference framework → compute kernels (DeepGEMM) → hardware adaptation (from NVIDIA to Ascend). Each layer is optimized for the others, with no dependency on "waiting for upstream library support."

The cost of vertical integration is engineering investment — maintaining your own GEMM library requires a steady supply of GPU kernel experts. The payoff: nobody can choke you off. NVIDIA doesn't update cuBLAS? Doesn't matter. CUTLASS doesn't support your new architecture? Doesn't matter. Ascend needs support? Add it yourself.

While most AI companies still rely on third-party libraries for inference, DeepSeek has dug down to the GPU kernel layer. This isn't over-engineering — when your monthly inference bill runs into millions of dollars, a 5% performance difference is real money. DeepSeek's inference efficiency is a core reason they can offer APIs at one-tenth the price.

DeepGEMM's open-source strategy is also worth noting. Open-sourcing a library you deeply depend on lets the community find bugs, contribute optimizations, and extend it to more hardware — while core model architectures and training data remain closed. Open-source infrastructure, closed-source upper-layer value — DeepSeek isn't the first to use this strategy, but the depth of execution at the kernel layer is rare.

What's worth watching next isn't whether DeepGEMM can get another 5% faster, but whether DeepSeek will run inference on Ascend clusters once Ascend support matures. If that day comes, NVIDIA's pricing power in the AI inference market will face a genuine challenge — not from policy sanctions, but from technical substitution.

Tags

#deepseek#deepgemm#cuda#fp8#gpu-kernels#jit-compilation#mixture-of-experts#huawei-ascend

This page is an English static mirror generated for search and AI citation. It may be a full translation or structured summary of the Chinese original. Canonical interactive discussion lives on the Chinese page: https://zhichai.net/topic/178635667