Skip to content

Kimi k3 sglang blackwell - #6497

Open
ezhong0211 wants to merge 7 commits into
mainfrom
kimi-k3-sglang-blackwell
Open

Kimi k3 sglang blackwell#6497
ezhong0211 wants to merge 7 commits into
mainfrom
kimi-k3-sglang-blackwell

Conversation

@ezhong0211

@ezhong0211 ezhong0211 commented Aug 6, 2026

Copy link
Copy Markdown
Contributor

Purpose

Enables Kimi-K3 on the SGLang AL2023 images and fixes two Blackwell defects found while
verifying it on 8x B300.

Commit 1, Blackwell correctness. Three changes, each turning a load-time failure on
B200/B300 into either a correct build or a failed one. No effect on pre-Blackwell GPUs.

  • torch_cuda_arch_list 10.3 to 10.3f. MXFP4 kernels emit cvt.e2m1x2, which ptxas
    accepts only under an a- or f-suffixed target. A bare 10.3 compiles and links without
    error but emits no usable FP4 device code for sm_103, so the gap appears only when a
    model with MXFP4 weights loads on the hardware. This diverges from upstream's own
    Dockerfile, which still uses a bare 10.3. Deliberate, not a catch-up.
  • Assert cutlass._mlir imports. nvidia-cutlass-dsl ships pure Python and its MLIR
    extensions come from a separate -libs-cu13 wheel. Both are declared, yet a prior image
    shipped without _mlir, so import cutlass failed. That is a hard startup failure on
    Blackwell only, since bf16_gemm_backend='auto' selects the CuTe DSL path on SM100/SM103
    and cuBLAS elsewhere. Root cause is unresolved; this assertion does not fix it, it stops
    the same silent gap from reaching a released image again.
  • Record reachable MXFP4 MoE backends to /etc/dlc-mxfp4-moe-backends.txt. With
    moe_runner_backend='auto' SGLang picks trtllm-gen on SM100/SM103 when its cubins are
    present and silently drops to marlin when they are not, so an image can ship on a slower
    expert path unnoticed. Warns rather than fails, because trtllm-gen's MoE cubins are
    unmerged upstream and a version bump does not help.

Also folded in, all found while building this image and all inert on a green build: a
single-tree pip bootstrap via get-pip.py, libz3 taken from the z3-solver wheel, and a
conditional tilelang cudart.cc patch.

Commit 2, the ref bump. sglang_ref moves to abddb1c7, with sglang_kernel_version
0.4.5 and flashinfer_version 0.6.15.post1. Kimi-K3 needs srt/models/kimi_k3.py, which
is in no released tag, so this is a bare commit hash. The two version pins are ABI-coupled
to the ref and are what upstream's own Dockerfile pins at that exact commit, read from the
tree rather than assumed. Both amzn2023 variants move together. The ubuntu variants pin no
sglang_ref (they build the 0.5.16 tag) and are untouched.

The two commits must land together and must not be split. This ref without the 10.3f
change produces an image that advertises K3 support, builds green, then fails at load time
on B300 because MXFP4 has no working device code.

Open items for the reviewer:

  • framework_version stays 0.5.14+dlc1 and dlc_minor_version stays 2. Both feed
    image tags and are arguably stale. The only required_image_pattern in
    sglang-model-tests.yml keys on os_version, so neither breaks test matching. Flagged as
    a release-versioning decision rather than silently picked.
  • Whether the trtllm-gen fallback costs accuracy or only throughput is unmeasured.

Test Plan

CI cannot validate Kimi-K3 at any point. It needs roughly 1.51 TiB of GPU memory against
640 GB on the largest runner, about 2.4x short, and B300 is not in the fleet. No model-test
entry is added, deliberately. So the plan splits into what CI can prove and what only
hardware can.

CI, this PR:

  1. Image builds at abddb1c7 with sglang-kernel 0.4.5 and flashinfer 0.6.15.post1.
  2. The cutlass._mlir assertion passes, proving the CuTe DSL native libs landed.
  3. The MXFP4 probe runs and records a backend list.
  4. Existing sanity, security, telemetry and model tests pass unchanged, confirming no
    regression on non-Blackwell hardware.

Local, already done:

  1. resolve_build_args.py emits all four keys as upper-case build args on both variants,
    confirming the config overrides actually bind rather than silently falling back to the
    Dockerfile defaults.
  2. Both configs parse as YAML and hold identical values.

Hardware, out of band on 8x B300, not gated on this PR:

  1. Serve Kimi-K3 end to end on the built image.
  2. cuobjdump --list-elf on the sgl_kernel shared objects, bare 10.3 against 10.3f,
    to replace the ptxas argument with direct evidence of emitted FP4 code.
  3. Compare deep_gemm against marlin for throughput and output equivalence.

Test Result

Pending. Will update with the CI run once the PR is open.

Local checks pass:

  • resolve_build_args.py emits SGLANG_REF=abddb1c7e9d61ddddeaf016d885c2f20aab426e8,
    SGLANG_KERNEL_VERSION=0.4.5, FLASHINFER_VERSION=0.6.15.post1,
    TORCH_CUDA_ARCH_LIST=9.0;10.0;10.3f on both ec2 and sagemaker.
  • Both config files parse and agree.

Prior hardware evidence, for context rather than as validation of this exact tree: Kimi-K3
served on 8x B300 (p6-b300.48xlarge, compute capability 10.3) on 2026-08-04, roughly
229 GiB per rank under SGLang. That run used 40feea27, a dev-branch commit of the same
work, so the hardware evidence is one commit removed from the abddb1c7 shipped here. A
main commit was chosen anyway because it is immutable, where a dev branch can be rebased or
deleted. Item 7 above closes that gap.


Toggle if you are merging into main Branch

PR Checklist

  • I ran pre-commit run --all-files locally before creating this PR. (Read DEVELOPMENT.md for details).

Eric Zhong added 2 commits August 6, 2026 12:53
…aging gaps

Three Blackwell-specific correctness fixes to the amzn2023 sglang image. All
three convert failures that surface at load time on B200/B300 into either a
correct build or a failed one. No behavior change on pre-Blackwell GPUs.

1. torch_cuda_arch_list 10.3 -> 10.3f

MXFP4 kernels emit cvt.e2m1x2, which ptxas accepts only under an a- or
f-suffixed target. A bare 10.3 compiles and links without error but emits no
working FP4 device code for sm_103, so the gap appears only when a model with
MXFP4 weights is loaded on the hardware. The f (family) suffix is preferred
over a for forward compatibility across the sm_10x family. This diverges from
upstream's own Dockerfile, which still uses a bare 10.3 -- deliberate, not a
catch-up.

2. Assert cutlass._mlir imports

nvidia-cutlass-dsl ships pure Python; its MLIR extension modules come from a
separate -libs-cu13 wheel. Both are declared, yet an image built 2026-08-03
shipped the Python tree with no _mlir package, so `import cutlass` failed. That
is a hard startup failure on Blackwell and only there: bf16_gemm_backend='auto'
selects the CuTe DSL path on SM100/SM103 and cuBLAS elsewhere, so the server
died in init_moe_gemm_config on B200/B300 while every other GPU was fine. Root
cause is not established; this assertion does not fix it, it just prevents the
same silent gap from reaching a released image again.

3. Record reachable MXFP4 MoE backends at build time

With moe_runner_backend='auto' sglang picks trtllm-gen on SM100/SM103 when its
cubins are present and drops to marlin when they are not, logging nothing a
build or release check would notice -- so an image can ship on a slower expert
path unnoticed. This probe writes the reachable set to
/etc/dlc-mxfp4-moe-backends.txt. It warns rather than fails, because a missing
trtllm-gen path is not fixable here: its MoE cubins are unmerged upstream and a
version bump provably does not help. A positive is necessary-not-sufficient --
the symbol resolving does not prove cubins exist for a given shape.

Also folded in, all discovered while building this image and all inert on a
green build: single-tree pip bootstrap via get-pip.py (the runtime pip and the
builder site-packages COPY spliced two pip trees together, breaking as
ImportError: get_runnable_pip), libz3 from the z3-solver wheel (AL2023 packages
only 4.8, tilelang needs 4.15), and a conditional tilelang cudart.cc patch
(>=0.1.10 ships a prebuilt stub and no C++ sources).

CI cannot validate any of this end to end: Kimi-K3 needs ~1.51 TiB of GPU
memory against 640 GB on the largest runner, and B300 is not in the fleet. The
signal here is that the image builds and the two guards pass.
Kimi-K3 needs python/sglang/srt/models/kimi_k3.py, absent from the previous pin
bc8b3ab1 and from every released tag (v0.5.14/v0.5.15/v0.5.16), so this is a
bare commit hash rather than a version: the merge to main, #32541, 2026-08-04.
Move to the first released tag containing kimi_k3.py once one exists.

sglang_kernel_version and flashinfer_version are ABI/JIT-coupled to sglang_ref
and must move in lockstep. 0.4.5 and 0.6.15.post1 are what upstream's own
docker/Dockerfile pins at this exact commit, read from the tree at abddb1c7
rather than assumed, and both are confirmed present in the built image
(sglang-kernel==0.4.5+cu130, flashinfer-python==0.6.15.post1).

Both amzn2023 variants move together. Letting ec2 and sagemaker sit on
different upstream trees is a known divergence trap. The ubuntu variants pin no
sglang_ref at all -- they build the 0.5.16 tag -- so they are untouched and
gain no K3 support.

Depends on the preceding 10.3f commit: this ref without that arch list yields an
image that advertises K3 support, builds green, and then fails at load time on
B300 because MXFP4 has no working device code. Do not land these out of order or
separately.

Caveats a reviewer should weigh:

- The B300 verification on 2026-08-04 ran 40feea27, a dev-branch commit of the
  same work, so the hardware evidence is one commit removed from what ships
  here. A main commit was chosen anyway because it is immutable, where a dev
  branch can be rebased or deleted.
- CI cannot validate K3 at any point: it needs ~1.51 TiB of GPU memory against
  640 GB on the largest runner, ~2.4x short, and B300 is not in the fleet. No
  model-test entry is added, deliberately. The available signal is that the
  image builds at this ref and existing models still pass.
- framework_version stays 0.5.14+dlc1 and dlc_minor_version stays 2. Both feed
  image tags and are arguably stale now; the only required_image_pattern in
  sglang-model-tests.yml keys on os_version, so neither breaks test matching.
  Flagged as a release-versioning decision, not silently picked.
@ezhong0211
ezhong0211 force-pushed the kimi-k3-sglang-blackwell branch from cd675c9 to d54a8d6 Compare August 6, 2026 19:57
Eric Zhong added 5 commits August 6, 2026 14:57
The probe was written as a RUN with a shell brace group wrapping a heredoc. The
RUN line had no trailing backslash, so BuildKit ended the instruction at the
heredoc terminator and then parsed the closing brace as a new instruction:

  Dockerfile.amzn2023:371
  ERROR: failed to solve: dockerfile parse error on line 371: unknown instruction: }

Move the probe to scripts/docker/sglang/probe_mxfp4_moe_backends.py and COPY it
in, matching how every other script in this Dockerfile is handled. This also
avoids introducing the repo's only Dockerfile heredoc in a file that declares no
`# syntax=` frontend, so a release build does not depend on the daemon's default
BuildKit accepting one.

Behaviour is unchanged: still writes /etc/dlc-mxfp4-moe-backends.txt and still
exits 0 on every path, since a missing trtllm-gen is expected today.
… CUDA comment

An unset FRAMEWORK_VERSION expands to a leading-dot version string that
packaging rejects, but only after the sglang source build starts -- the
failure surfaces as a 60-line setuptools_scm traceback that never names
the missing build-arg. Check it directly instead.

The comment above the install also still described a cu129 index and a
CUDA 12.9 toolkit; the config has been on CUDA 13 for some time.
TORCH_CUDA_ARCH_LIST is consumed only by FlashMLA and DeepEP, both of which
build through torch's cpp_extension. Its _get_cuda_arch_flags validates each
entry against a fixed list that carries a and +PTX variants but no f variants,
so 10.3f fails the DeepEP build with "Unknown CUDA arch (10.3f) or GPU not
supported". 10.3a is in that list and yields
-gencode=arch=compute_103a,code=sm_103a, which is what ptxas needs to accept
cvt.e2m1x2 for Kimi-K3's MXFP4 routed experts.

f remains correct for CMake-based builds -- vLLM reaches sm_103 that way via
10.0f family compat -- but is not usable here. Also drops the claim that this
list governs MXFP4 kernel emission in sgl_kernel: sgl_kernel is installed as
a prebuilt wheel, not compiled in this Dockerfile.
The runtime stage reinstalled nvidia-cutlass-dsl and -libs-cu13 with a >=4.5.2
floor under --force-reinstall, so it floated onto 4.7.0 and away from the
==4.6.0 that sglang pins at SGLANG_REF. 4.7.0 relocated the package to
nvidia_cutlass_dsl/dsl_packages/cutlass/ and its -libs wheel no longer lands
_mlir there, so `import cutlass._mlir` fails.

Caught by the existing DSL assertion, which had been added against an earlier
image that shipped without _mlir for what is now clearly the same reason. That
comment claimed the root cause was unresolved; it was this floating floor.

Only bf16_gemm_backend=auto on SM100/SM103 reaches the CuTe DSL path, so this
is invisible outside Blackwell and CI could not have caught it.

Pin exactly via a new CUTLASS_DSL_VERSION arg, plumbed through both image
configs so the pin is visible next to the other version pins.
nvidia-cutlass-dsl delivers one logical package via two wheels that disagree
on their wheel tag:

  nvidia-cutlass-dsl-libs-core  Root-Is-Purelib: true   -> lib/site-packages
  nvidia-cutlass-dsl-libs-base  Root-Is-Purelib: false  -> lib64/site-packages

Both write into nvidia_cutlass_dsl/dsl_packages/cutlass/, so on AL2023 pip puts
cutlass/__init__.py under lib and cutlass/_mlir/ under lib64. nvidia-cutlass-dsl
itself is only a .pth that appends a single dsl_packages path, so one half is
importable and `import cutlass._mlir` fails.

Upstream bug, present in 4.6.0 and 4.7.0 alike, and invisible on distros where
purelib == platlib (Debian, Ubuntu, most manylinux images). It reproduces only
on the lib/lib64-split distros, AL2023 among them.

This is the actual cause of the missing _mlir that the DSL assertion was
originally added for. My previous commit blamed a 4.7.0 relocation; that was
wrong, though the exact pin it added is still correct on its own terms since
sglang requires ==4.6.0.

Symlink the halves together after install. Verifies co-location structurally
rather than by link count so a rerun on a cached layer is not a failure.
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Projects

None yet

Development

Successfully merging this pull request may close these issues.

1 participant