inferencekey/deepinfer-sm100-cu132

DeepInfer prebuilt for NVIDIA B200 (SM100), CUDA 13.2, Python 3.12

Prebuilt DeepInfer (DeepSeek's inference engine for DeepSeek-V4.1-Flash) so a GPU node can run it without compiling CUDA kernels for hours. Published by InferenceKey and used by the InferenceKey SDK's DeepInfer custom backend.

Tested on 4x NVIDIA B200 (RunPod, driver 595.91, CUDA 13.2) serving deepseek-ai/DeepSeek-V4.1-Flash-1017 with expert parallelism 4 and data parallelism 4, in a container without RDMA devices.

Contents

Path	What
wheels/	DeepSelect, FlashMLA, DeepEP and DeepGEMM compiled for SM100 (sm_100a and sm_103a where the vendor builds both) with CUDA 13.2, cp312, linux x86_64
deepinfer-9bd9cf2-prefixcache-sm100-cu132-cp312.tar.zst	DeepInfer main 9bd9cf2 built in place (deepinfer/_C), with the three patches below applied
patches/	The patches applied to DeepInfer, also usable on a source checkout
jit-cache/	DeepGEMM and TileLang JIT kernel caches from the runs (EP4 on 4x B200, and a trimmed 1-GPU smoke test)
install.sh	Builds a venv: pinned requirements from PyPI and the PyTorch index, the wheels, and DeepInfer
build-deepinfer.sh	The recipe used to build all of this from source on a CPU-only pod
rebuild-deepinfer-C.sh	Applies patches/deepinfer-prefix-cache.patch to an installed tree and rebuilds only deepinfer._C (seconds; no GPU needed)
licenses/	DeepInfer (Apache-2.0) and the vendors (MIT)

torch and the NVIDIA libraries are NOT included: install.sh installs them from their official indexes.

The prefix-cache build of deepinfer._C was compiled on 2026-10-23 against this package (torch 2.13.0+cu132, Python 3.12) and checked by importing it and its new API; it has not run on a GPU yet. Bounded replay after a cache hit is DeepInfer's own mechanism for a missing SWA state, but its effect on output quality has not been measured here.

Install

# on the GPU node (Ubuntu 24.04, CUDA 13 toolkit with nvcc, libdw-dev, zstd)
huggingface-cli download inferencekey/deepinfer-sm100-cu132 --local-dir /workspace/deepinfer-prebuilt
bash /workspace/deepinfer-prebuilt/install.sh /workspace/di-venv /workspace

Optional, to skip most first-run kernel compilation:

export DG_JIT_CACHE_DIR=/workspace/jit-cache/deep_gemm TILELANG_CACHE_DIR=/workspace/jit-cache/tilelang
zstd -dc /workspace/deepinfer-prebuilt/jit-cache/jit-cache-4xb200-ep4.tar.zst | tar -x -C /workspace

Patches

1. deepinfer-engram-local.patch: an RDMA-free path for Engram (DEEPINFER_ENGRAM_LOCAL=1). Upstream, Engram always goes through DeepEP's EngramBuffer, which requires NCCL GIN with GDAKI (a ConnectX NIC with verbs devices in the container). Rented GPU containers usually have none. The local path keeps each rank's Engram shard in torch symmetric memory and reads peers' shards over NVLink (single node). Same hook outputs as DeepEP.

2. deepinfer-nvls-optout.patch: honours DEEPINFER_NCCL_NVLS_ENABLE (upstream always unsets NCCL_NVLS_ENABLE). Containers holding part of an 8-GPU node (for example 4 of 8) cannot create NVLink multicast objects (CUDA error 401); set DEEPINFER_NCCL_NVLS_ENABLE=0.

3. deepinfer-prefix-cache.patch: the hooks a GPU prefix cache needs, which the open release does not have (it frees every KV block when a request finishes, so a coding agent's next turn prefills its whole context again):

   • Scheduler.add_sequence_with_prefix(seq, blocks, num_cached, swa_valid_start) and LLMEngine.enqueue_with_prefix(...): start a new sequence from cached KV blocks (SWA bounded replay: the last cached block is recomputed, with a zeroed state slot). The blocks are validated before pinning; a bad call raises without side effects.
   • Sequence.block_table (read-only) and BlockManager.num_blocks.
   • prefill_caching_rate (GPU x%) in the engine log now counts the reused tokens (upstream always passes 0).
   • A sequence that finishes while still waiting (aborted before admission) now releases its blocks and state slot; upstream leaked them.

   Nothing changes for requests that do not use the new calls. The cache itself (block index, eviction, session routing) lives in the InferenceKey SDK's DeepInfer custom backend.

Running on a node without RDMA

export DEEPINFER_ENGRAM_LOCAL=1 EP_DISABLE_GIN=1 NCCL_IB_DISABLE=1 DEEPINFER_NCCL_NVLS_ENABLE=0

Keep context parallelism at 1 (CP>1 needs RDMA bandwidth detection). Do not set DEEPINFER_ENGRAM_FORCE_GPU_STORAGE.

Measured (4x B200, EP4/DP4, DeepSeek-V4.1-Flash, 2026-10-23)

• Load: about 30 min from a network disk (most of it reading 510 GB of weights); use local NVMe.
• KV cache: about 24 GiB per GPU, about 29M tokens per GPU (about 200 contexts of 138K tokens).
• Prefill: about 88K tokens/s per GPU (about 352K for the node), matching DeepInfer's published figure.
• 400 simultaneous requests with 138K-token prompts: all served, no errors.
• Decode at high concurrency was limited by the Python serving layer in front of the engine (about 600 output tokens/s for the node), not by DeepInfer; that layer is being reworked.

Provenance

Built from DeepInfer 9bd9cf2 and the vendor commits it pins (DeepEP f6aa38f, DeepGEMM 1fc9c66, FlashMLA 2e5429f, DeepSelect bfa4507) on runpod/pytorch:1.4.1-cu1320-torch2130-ubuntu2404 with build-deepinfer.sh (MAX_JOBS=1 NVCC_THREADS=1 FLASH_MLA_SKIP_REG_SPILL_CHECK=1; FlashMLA's own register-spill check fails with nvcc 13.2 on four head-64 split-KV decode kernels, a small performance cost, not a correctness issue).

DeepInfer is Apache-2.0 (modified as described in patches/); DeepEP, DeepGEMM, FlashMLA and DeepSelect are MIT. Copyright of the original works remains with their authors.