[GH-ISSUE #1793] 运行Glm4.7-FP8失败，提示triton.runtime.errors.OutOfResources: out of resource: shared memory #1127

New issue

Closed

opened 2026-10-29 14:48:54 +00:00 by bsakaguchi · 2 comments

bsakaguchi commented

2026-10-29 14:48:54 +00:00

Owner

Originally created by @CodeZ-Hao on GitHub (Jan 13, 2026).
Original GitHub issue: https://github.com/kvcache-ai/ktransformers/issues/1793

Reminder

I have read the above rules and searched the existing issues.

System Info

# git log
commit 6277da4c2b826eb6354d6a0edb0dbe4af927563c (HEAD -> main, origin/main, origin/HEAD)
Author: Oql <1692110604@qq.com>
Date:   Tue Jan 13 17:36:25 2026 +0800
    support GLM 4.7 (#1791)
    
    support GLM 4.7

# nvidia-smi
Wed Jan 14 00:46:04 2026       
+-----------------------------------------------------------------------------------------+
| NVIDIA-SMI 570.172.08             Driver Version: 570.172.08     CUDA Version: 12.8     |
|-----------------------------------------+------------------------+----------------------+
| GPU  Name                 Persistence-M | Bus-Id          Disp.A | Volatile Uncorr. ECC |
| Fan  Temp   Perf          Pwr:Usage/Cap |           Memory-Usage | GPU-Util  Compute M. |
|                                         |                        |               MIG M. |
|=========================================+========================+======================|
|   0  NVIDIA GeForce RTX 4080        Off |   00000000:16:00.0 Off |                  N/A |
|  0%   45C    P8              7W /  320W |    1222MiB /  32760MiB |      0%      Default |
|                                         |                        |                  N/A |
+-----------------------------------------+------------------------+----------------------+
                                                                                         
+-----------------------------------------------------------------------------------------+
| Processes:                                                                              |
|  GPU   GI   CI              PID   Type   Process name                        GPU Memory |
|        ID   ID                                                               Usage      |
|=========================================================================================|
|    0   N/A  N/A            1581      G   /usr/lib/xorg/Xorg                        4MiB |
|    0   N/A  N/A         3394274      C   python                                  306MiB |
|    0   N/A  N/A         3420974      C   ...ma.cpp/build/bin/llama-server        888MiB |
+-----------------------------------------------------------------------------------------+

# lsmem
RANGE                                  SIZE  STATE REMOVABLE BLOCK
0x0000000000000000-0x000000007fffffff    2G online        是     0
0x0000000100000000-0x000000607fffffff  382G online        是 2-192
Memory block size:         2G
Total online memory:     384G
Total offline memory:      0B

# lscpu
架构：                    x86_64
  CPU 运行模式：          32-bit, 64-bit
  Address sizes:          46 bits physical, 57 bits virtual
  字节序：                Little Endian
CPU:                      48
  在线 CPU 列表：         0-47
厂商 ID：                 GenuineIntel
  型号名称：              Intel(R) Xeon(R) Platinum 8461V
    CPU 系列：            6
    型号：                143
    每个核的线程数：      1
    每个座的核数：        48
    座：                  1
    步进：                6
    CPU 最大 MHz：        3700.0000

Reproduction

启动命令：

python3 -m sglang.launch_server --host 0.0.0.0 --port 6399 --model /models/ZhipuAI/GLM-4.7-FP8/ --kt-weight-path /models/ZhipuAI/GLM-4.7-FP8/ --kt-cpuinfer 40 --kt-threadpool-count 1 --kt-num-gpu-experts 0 --kt-method FP8 --kt-gpu-prefill-token-threshold 100 --attention-backend flashinfer --trust-remote-code --mem-fraction-static 0.98 --chunked-prefill-size 8192 --max-running-requests 1 --max-total-tokens 8000 --served-model-name glm-4.7 --tool-call-parser glm47 --reasoning-parser glm45 --disable-shared-experts-fusion

异常信息：

[NativeMoEWrapper Layer 88] load_experts: 199.1ms, prepare_tensors: 0.0ms, build_ptrs: 0.2ms, create_moe: 28.3ms, cpp_load_weights: 2013.8ms, cleanup: 0.0ms, total: 2241.4ms
TP MOE layer 89, pool: 0x265344c0, expert num: 160, num_experts_per_tok: 8
Created AMX_FP8_MOE_TP 0 at numa 0
[NativeMoEWrapper Layer 89] load_experts: 206.7ms, prepare_tensors: 0.0ms, build_ptrs: 0.2ms, create_moe: 7.3ms, cpp_load_weights: 1888.5ms, cleanup: 0.0ms, total: 2102.8ms
TP MOE layer 90, pool: 0x265344c0, expert num: 160, num_experts_per_tok: 8
Created AMX_FP8_MOE_TP 0 at numa 0
[NativeMoEWrapper Layer 90] load_experts: 207.8ms, prepare_tensors: 0.0ms, build_ptrs: 0.2ms, create_moe: 26.0ms, cpp_load_weights: 1954.1ms, cleanup: 0.0ms, total: 2188.2ms
TP MOE layer 91, pool: 0x265344c0, expert num: 160, num_experts_per_tok: 8
Created AMX_FP8_MOE_TP 0 at numa 0
[NativeMoEWrapper Layer 91] load_experts: 205.5ms, prepare_tensors: 0.0ms, build_ptrs: 0.3ms, create_moe: 27.6ms, cpp_load_weights: 1956.7ms, cleanup: 0.0ms, total: 2190.1ms
[2026-01-14 00:41:13] Load weight end. type=Glm4MoeForCausalLM, dtype=torch.bfloat16, avail mem=12.56 GB, mem usage=17.44 GB.
[2026-01-14 00:41:14] Using KV cache dtype: torch.bfloat16
[2026-01-14 00:41:14] KV Cache is allocated. #tokens: 8000, K size: 1.40 GB, V size: 1.40 GB
[2026-01-14 00:41:14] Memory pool end. avail mem=9.69 GB
[2026-01-14 00:41:14] Capture cuda graph begin. This can take up to several minutes. avail mem=9.17 GB
[2026-01-14 00:41:14] Capture cuda graph bs [1]
Capturing batches (bs=1 avail_mem=9.17 GB):   0%|                                                                                                                                          | 0/1 [00:00<?, ?it/s][2026-01-14 00:41:24] Using default MoE kernel config. Performance might be sub-optimal! Config file not found at /work/ktransformers/sglang/python/sglang/srt/layers/moe/fused_moe_triton/configs/triton_3_5_1/E=0,N=1536,device_name=NVIDIA_GeForce_RTX_4080,dtype=fp8_w8a8,per_channel_quant=True.json, you can create them with https://github.com/sgl-project/sglang/tree/main/benchmark/kernels/fused_moe_triton
[2026-01-14 00:41:24] Using MoE kernel config with down_moe=False. Performance might be sub-optimal! Config file not found at /work/ktransformers/sglang/python/sglang/srt/layers/moe/fused_moe_triton/configs/triton_3_5_1/E=0,N=1536,device_name=NVIDIA_GeForce_RTX_4080,dtype=fp8_w8a8,per_channel_quant=True_down.json, you can create them with https://github.com/sgl-project/sglang/tree/main/benchmark/kernels/fused_moe_triton
Capturing batches (bs=1 avail_mem=9.17 GB):   0%|                                                                                                                                          | 0/1 [00:09<?, ?it/s]
[2026-01-14 00:41:24] Scheduler hit an exception: Traceback (most recent call last):
  File "/work/ktransformers/sglang/python/sglang/srt/managers/scheduler.py", line 2902, in run_scheduler_process
    scheduler = Scheduler(
                ^^^^^^^^^^
  File "/work/ktransformers/sglang/python/sglang/srt/managers/scheduler.py", line 336, in __init__
    self.init_model_worker()
  File "/work/ktransformers/sglang/python/sglang/srt/managers/scheduler.py", line 512, in init_model_worker
    self.init_tp_model_worker()
  File "/work/ktransformers/sglang/python/sglang/srt/managers/scheduler.py", line 474, in init_tp_model_worker
    self.tp_worker = TpModelWorker(
                     ^^^^^^^^^^^^^^
  File "/work/ktransformers/sglang/python/sglang/srt/managers/tp_worker.py", line 240, in __init__
    self._init_model_runner()
  File "/work/ktransformers/sglang/python/sglang/srt/managers/tp_worker.py", line 323, in _init_model_runner
    self._model_runner = ModelRunner(
                         ^^^^^^^^^^^^
  File "/work/ktransformers/sglang/python/sglang/srt/model_executor/model_runner.py", line 381, in __init__
    self.initialize(min_per_gpu_memory)
  File "/work/ktransformers/sglang/python/sglang/srt/model_executor/model_runner.py", line 567, in initialize
    self.init_device_graphs()
  File "/work/ktransformers/sglang/python/sglang/srt/model_executor/model_runner.py", line 1983, in init_device_graphs
    self.graph_runner = graph_runners[self.device](self)
                        ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "/work/ktransformers/sglang/python/sglang/srt/model_executor/cuda_graph_runner.py", line 361, in __init__
    self.capture()
  File "/work/ktransformers/sglang/python/sglang/srt/model_executor/cuda_graph_runner.py", line 517, in capture
    _capture_one_stream()
  File "/work/ktransformers/sglang/python/sglang/srt/model_executor/cuda_graph_runner.py", line 504, in _capture_one_stream
    ) = self.capture_one_batch_size(bs, forward, stream_idx)
        ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "/work/ktransformers/sglang/python/sglang/srt/model_executor/cuda_graph_runner.py", line 723, in capture_one_batch_size
    run_once()
  File "/work/ktransformers/sglang/python/sglang/srt/model_executor/cuda_graph_runner.py", line 710, in run_once
    logits_output_or_pp_proxy_tensors = forward(
                                        ^^^^^^^^
  File "/root/miniconda3/envs/kt/lib/python3.11/site-packages/torch/utils/_contextlib.py", line 120, in decorate_context
    return func(*args, **kwargs)
           ^^^^^^^^^^^^^^^^^^^^^
  File "/work/ktransformers/sglang/python/sglang/srt/models/glm4_moe.py", line 1076, in forward
    hidden_states = self.model(
                    ^^^^^^^^^^^
  File "/root/miniconda3/envs/kt/lib/python3.11/site-packages/torch/nn/modules/module.py", line 1775, in _wrapped_call_impl
    return self._call_impl(*args, **kwargs)
           ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "/root/miniconda3/envs/kt/lib/python3.11/site-packages/torch/nn/modules/module.py", line 1786, in _call_impl
    return forward_call(*args, **kwargs)
           ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "/work/ktransformers/sglang/python/sglang/srt/models/glm4_moe.py", line 967, in forward
    hidden_states, residual = layer(
                              ^^^^^^
  File "/root/miniconda3/envs/kt/lib/python3.11/site-packages/torch/nn/modules/module.py", line 1775, in _wrapped_call_impl
    return self._call_impl(*args, **kwargs)
           ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "/root/miniconda3/envs/kt/lib/python3.11/site-packages/torch/nn/modules/module.py", line 1786, in _call_impl
    return forward_call(*args, **kwargs)
           ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "/work/ktransformers/sglang/python/sglang/srt/models/glm4_moe.py", line 805, in forward
    hidden_states = self.mlp(
                    ^^^^^^^^^
  File "/root/miniconda3/envs/kt/lib/python3.11/site-packages/torch/nn/modules/module.py", line 1775, in _wrapped_call_impl
    return self._call_impl(*args, **kwargs)
           ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "/root/miniconda3/envs/kt/lib/python3.11/site-packages/torch/nn/modules/module.py", line 1786, in _call_impl
    return forward_call(*args, **kwargs)
           ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "/work/ktransformers/sglang/python/sglang/srt/models/glm4_moe.py", line 459, in forward
    return self.forward_normal_dual_stream(
           ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "/work/ktransformers/sglang/python/sglang/srt/models/glm4_moe.py", line 484, in forward_normal_dual_stream
    final_hidden_states = self.experts(hidden_states, topk_output)
                          ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "/root/miniconda3/envs/kt/lib/python3.11/site-packages/torch/nn/modules/module.py", line 1775, in _wrapped_call_impl
    return self._call_impl(*args, **kwargs)
           ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "/root/miniconda3/envs/kt/lib/python3.11/site-packages/torch/nn/modules/module.py", line 1786, in _call_impl
    return forward_call(*args, **kwargs)
           ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "/work/ktransformers/sglang/python/sglang/srt/layers/moe/fused_moe_triton/layer.py", line 920, in forward
    return self.forward_impl(hidden_states, topk_output)
           ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "/work/ktransformers/sglang/python/sglang/srt/layers/moe/fused_moe_triton/layer.py", line 939, in forward_impl
    combine_input = self.run_moe_core(
                    ^^^^^^^^^^^^^^^^^^
  File "/work/ktransformers/sglang/python/sglang/srt/layers/moe/fused_moe_triton/layer.py", line 960, in run_moe_core
    return self.quant_method.apply(
           ^^^^^^^^^^^^^^^^^^^^^^^^
  File "/work/ktransformers/sglang/python/sglang/srt/layers/moe/kt_ep_wrapper.py", line 1321, in apply
    gpu_combine_input = self.gpu_method.apply(layer, masked_dispatch_output)
                        ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "/work/ktransformers/sglang/python/sglang/srt/layers/quantization/compressed_tensors/compressed_tensors_moe.py", line 362, in apply
    return self.runner.run(dispatch_output, quant_info)
           ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "/work/ktransformers/sglang/python/sglang/srt/layers/moe/moe_runner/runner.py", line 78, in run
    return self.fused_func(dispatch_output, quant_info, self.config)
           ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "/work/ktransformers/sglang/python/sglang/srt/layers/moe/moe_runner/triton.py", line 337, in fused_experts_none_to_triton
    output = fused_experts(
             ^^^^^^^^^^^^^^
  File "/work/ktransformers/sglang/python/sglang/srt/layers/moe/fused_moe_triton/fused_moe.py", line 213, in fused_experts
    inplace_fused_experts(
  File "/root/miniconda3/envs/kt/lib/python3.11/site-packages/torch/_ops.py", line 1255, in __call__
    return self._op(*args, **kwargs)
           ^^^^^^^^^^^^^^^^^^^^^^^^^
  File "/work/ktransformers/sglang/python/sglang/srt/layers/moe/fused_moe_triton/fused_moe.py", line 91, in inplace_fused_experts
    fused_experts_impl(
  File "/work/ktransformers/sglang/python/sglang/srt/layers/moe/fused_moe_triton/fused_moe.py", line 445, in fused_experts_impl
    invoke_fused_moe_kernel(
  File "/work/ktransformers/sglang/python/sglang/srt/layers/moe/fused_moe_triton/fused_moe_triton_kernels.py", line 746, in invoke_fused_moe_kernel
    fused_moe_kernel[grid](
  File "/root/miniconda3/envs/kt/lib/python3.11/site-packages/triton/runtime/jit.py", line 419, in <lambda>
    return lambda *args, **kwargs: self.run(grid=grid, warmup=False, *args, **kwargs)
                                   ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "/root/miniconda3/envs/kt/lib/python3.11/site-packages/triton/runtime/jit.py", line 756, in run
    launch_metadata = kernel.launch_metadata(grid, stream, *bound_args.values())
                      ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "/root/miniconda3/envs/kt/lib/python3.11/site-packages/triton/compiler/compiler.py", line 490, in launch_metadata
    self._init_handles()
  File "/root/miniconda3/envs/kt/lib/python3.11/site-packages/triton/compiler/compiler.py", line 464, in _init_handles
    raise_(OutOfResources(self.metadata.shared, max_shared, "shared memory"))
  File "/root/miniconda3/envs/kt/lib/python3.11/site-packages/triton/compiler/compiler.py", line 456, in raise_
    raise err
triton.runtime.errors.OutOfResources: out of resource: shared memory, Required: 147456, Hardware limit: 101376. Reducing block sizes or `num_stages` may help.

Others

看提示信息是硬件支持的shared memory是100k，而这里配置的block size大于这个值，这个问题似乎与 https://zhuanlan.zhihu.com/p/1890325961670624218 这里提到的一致，是否是kvcache-ai/sglang仓库中未合入这个PR？
我的kvcache-ai/sglang已拉取最新更新，且已重新安装，仍然存在这个问题。

Originally created by @CodeZ-Hao on GitHub (Jan 13, 2026). Original GitHub issue: https://github.com/kvcache-ai/ktransformers/issues/1793 ### Reminder - [x] I have read the above rules and searched the existing issues. ### System Info ``` # git log commit 6277da4c2b826eb6354d6a0edb0dbe4af927563c (HEAD -> main, origin/main, origin/HEAD) Author: Oql <1692110604@qq.com> Date: Tue Jan 13 17:36:25 2026 +0800 support GLM 4.7 (#1791) support GLM 4.7 # nvidia-smi Wed Jan 14 00:46:04 2026 +-----------------------------------------------------------------------------------------+ | NVIDIA-SMI 570.172.08 Driver Version: 570.172.08 CUDA Version: 12.8 | |-----------------------------------------+------------------------+----------------------+ | GPU Name Persistence-M | Bus-Id Disp.A | Volatile Uncorr. ECC | | Fan Temp Perf Pwr:Usage/Cap | Memory-Usage | GPU-Util Compute M. | | | | MIG M. | |=========================================+========================+======================| | 0 NVIDIA GeForce RTX 4080 Off | 00000000:16:00.0 Off | N/A | | 0% 45C P8 7W / 320W | 1222MiB / 32760MiB | 0% Default | | | | N/A | +-----------------------------------------+------------------------+----------------------+ +-----------------------------------------------------------------------------------------+ | Processes: | | GPU GI CI PID Type Process name GPU Memory | | ID ID Usage | |=========================================================================================| | 0 N/A N/A 1581 G /usr/lib/xorg/Xorg 4MiB | | 0 N/A N/A 3394274 C python 306MiB | | 0 N/A N/A 3420974 C ...ma.cpp/build/bin/llama-server 888MiB | +-----------------------------------------------------------------------------------------+ # lsmem RANGE SIZE STATE REMOVABLE BLOCK 0x0000000000000000-0x000000007fffffff 2G online 是 0 0x0000000100000000-0x000000607fffffff 382G online 是 2-192 Memory block size: 2G Total online memory: 384G Total offline memory: 0B # lscpu 架构： x86_64 CPU 运行模式： 32-bit, 64-bit Address sizes: 46 bits physical, 57 bits virtual 字节序： Little Endian CPU: 48 在线 CPU 列表： 0-47 厂商 ID： GenuineIntel 型号名称： Intel(R) Xeon(R) Platinum 8461V CPU 系列： 6 型号： 143 每个核的线程数： 1 每个座的核数： 48 座： 1 步进： 6 CPU 最大 MHz： 3700.0000 ``` ### Reproduction - 启动命令： ```text python3 -m sglang.launch_server --host 0.0.0.0 --port 6399 --model /models/ZhipuAI/GLM-4.7-FP8/ --kt-weight-path /models/ZhipuAI/GLM-4.7-FP8/ --kt-cpuinfer 40 --kt-threadpool-count 1 --kt-num-gpu-experts 0 --kt-method FP8 --kt-gpu-prefill-token-threshold 100 --attention-backend flashinfer --trust-remote-code --mem-fraction-static 0.98 --chunked-prefill-size 8192 --max-running-requests 1 --max-total-tokens 8000 --served-model-name glm-4.7 --tool-call-parser glm47 --reasoning-parser glm45 --disable-shared-experts-fusion ``` - 异常信息： ``` [NativeMoEWrapper Layer 88] load_experts: 199.1ms, prepare_tensors: 0.0ms, build_ptrs: 0.2ms, create_moe: 28.3ms, cpp_load_weights: 2013.8ms, cleanup: 0.0ms, total: 2241.4ms TP MOE layer 89, pool: 0x265344c0, expert num: 160, num_experts_per_tok: 8 Created AMX_FP8_MOE_TP 0 at numa 0 [NativeMoEWrapper Layer 89] load_experts: 206.7ms, prepare_tensors: 0.0ms, build_ptrs: 0.2ms, create_moe: 7.3ms, cpp_load_weights: 1888.5ms, cleanup: 0.0ms, total: 2102.8ms TP MOE layer 90, pool: 0x265344c0, expert num: 160, num_experts_per_tok: 8 Created AMX_FP8_MOE_TP 0 at numa 0 [NativeMoEWrapper Layer 90] load_experts: 207.8ms, prepare_tensors: 0.0ms, build_ptrs: 0.2ms, create_moe: 26.0ms, cpp_load_weights: 1954.1ms, cleanup: 0.0ms, total: 2188.2ms TP MOE layer 91, pool: 0x265344c0, expert num: 160, num_experts_per_tok: 8 Created AMX_FP8_MOE_TP 0 at numa 0 [NativeMoEWrapper Layer 91] load_experts: 205.5ms, prepare_tensors: 0.0ms, build_ptrs: 0.3ms, create_moe: 27.6ms, cpp_load_weights: 1956.7ms, cleanup: 0.0ms, total: 2190.1ms [2026-01-14 00:41:13] Load weight end. type=Glm4MoeForCausalLM, dtype=torch.bfloat16, avail mem=12.56 GB, mem usage=17.44 GB. [2026-01-14 00:41:14] Using KV cache dtype: torch.bfloat16 [2026-01-14 00:41:14] KV Cache is allocated. #tokens: 8000, K size: 1.40 GB, V size: 1.40 GB [2026-01-14 00:41:14] Memory pool end. avail mem=9.69 GB [2026-01-14 00:41:14] Capture cuda graph begin. This can take up to several minutes. avail mem=9.17 GB [2026-01-14 00:41:14] Capture cuda graph bs [1] Capturing batches (bs=1 avail_mem=9.17 GB): 0%| | 0/1 [00:00<?, ?it/s][2026-01-14 00:41:24] Using default MoE kernel config. Performance might be sub-optimal! Config file not found at /work/ktransformers/sglang/python/sglang/srt/layers/moe/fused_moe_triton/configs/triton_3_5_1/E=0,N=1536,device_name=NVIDIA_GeForce_RTX_4080,dtype=fp8_w8a8,per_channel_quant=True.json, you can create them with https://github.com/sgl-project/sglang/tree/main/benchmark/kernels/fused_moe_triton [2026-01-14 00:41:24] Using MoE kernel config with down_moe=False. Performance might be sub-optimal! Config file not found at /work/ktransformers/sglang/python/sglang/srt/layers/moe/fused_moe_triton/configs/triton_3_5_1/E=0,N=1536,device_name=NVIDIA_GeForce_RTX_4080,dtype=fp8_w8a8,per_channel_quant=True_down.json, you can create them with https://github.com/sgl-project/sglang/tree/main/benchmark/kernels/fused_moe_triton Capturing batches (bs=1 avail_mem=9.17 GB): 0%| | 0/1 [00:09<?, ?it/s] [2026-01-14 00:41:24] Scheduler hit an exception: Traceback (most recent call last): File "/work/ktransformers/sglang/python/sglang/srt/managers/scheduler.py", line 2902, in run_scheduler_process scheduler = Scheduler( ^^^^^^^^^^ File "/work/ktransformers/sglang/python/sglang/srt/managers/scheduler.py", line 336, in __init__ self.init_model_worker() File "/work/ktransformers/sglang/python/sglang/srt/managers/scheduler.py", line 512, in init_model_worker self.init_tp_model_worker() File "/work/ktransformers/sglang/python/sglang/srt/managers/scheduler.py", line 474, in init_tp_model_worker self.tp_worker = TpModelWorker( ^^^^^^^^^^^^^^ File "/work/ktransformers/sglang/python/sglang/srt/managers/tp_worker.py", line 240, in __init__ self._init_model_runner() File "/work/ktransformers/sglang/python/sglang/srt/managers/tp_worker.py", line 323, in _init_model_runner self._model_runner = ModelRunner( ^^^^^^^^^^^^ File "/work/ktransformers/sglang/python/sglang/srt/model_executor/model_runner.py", line 381, in __init__ self.initialize(min_per_gpu_memory) File "/work/ktransformers/sglang/python/sglang/srt/model_executor/model_runner.py", line 567, in initialize self.init_device_graphs() File "/work/ktransformers/sglang/python/sglang/srt/model_executor/model_runner.py", line 1983, in init_device_graphs self.graph_runner = graph_runners[self.device](self) ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ File "/work/ktransformers/sglang/python/sglang/srt/model_executor/cuda_graph_runner.py", line 361, in __init__ self.capture() File "/work/ktransformers/sglang/python/sglang/srt/model_executor/cuda_graph_runner.py", line 517, in capture _capture_one_stream() File "/work/ktransformers/sglang/python/sglang/srt/model_executor/cuda_graph_runner.py", line 504, in _capture_one_stream ) = self.capture_one_batch_size(bs, forward, stream_idx) ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ File "/work/ktransformers/sglang/python/sglang/srt/model_executor/cuda_graph_runner.py", line 723, in capture_one_batch_size run_once() File "/work/ktransformers/sglang/python/sglang/srt/model_executor/cuda_graph_runner.py", line 710, in run_once logits_output_or_pp_proxy_tensors = forward( ^^^^^^^^ File "/root/miniconda3/envs/kt/lib/python3.11/site-packages/torch/utils/_contextlib.py", line 120, in decorate_context return func(*args, **kwargs) ^^^^^^^^^^^^^^^^^^^^^ File "/work/ktransformers/sglang/python/sglang/srt/models/glm4_moe.py", line 1076, in forward hidden_states = self.model( ^^^^^^^^^^^ File "/root/miniconda3/envs/kt/lib/python3.11/site-packages/torch/nn/modules/module.py", line 1775, in _wrapped_call_impl return self._call_impl(*args, **kwargs) ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ File "/root/miniconda3/envs/kt/lib/python3.11/site-packages/torch/nn/modules/module.py", line 1786, in _call_impl return forward_call(*args, **kwargs) ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ File "/work/ktransformers/sglang/python/sglang/srt/models/glm4_moe.py", line 967, in forward hidden_states, residual = layer( ^^^^^^ File "/root/miniconda3/envs/kt/lib/python3.11/site-packages/torch/nn/modules/module.py", line 1775, in _wrapped_call_impl return self._call_impl(*args, **kwargs) ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ File "/root/miniconda3/envs/kt/lib/python3.11/site-packages/torch/nn/modules/module.py", line 1786, in _call_impl return forward_call(*args, **kwargs) ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ File "/work/ktransformers/sglang/python/sglang/srt/models/glm4_moe.py", line 805, in forward hidden_states = self.mlp( ^^^^^^^^^ File "/root/miniconda3/envs/kt/lib/python3.11/site-packages/torch/nn/modules/module.py", line 1775, in _wrapped_call_impl return self._call_impl(*args, **kwargs) ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ File "/root/miniconda3/envs/kt/lib/python3.11/site-packages/torch/nn/modules/module.py", line 1786, in _call_impl return forward_call(*args, **kwargs) ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ File "/work/ktransformers/sglang/python/sglang/srt/models/glm4_moe.py", line 459, in forward return self.forward_normal_dual_stream( ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ File "/work/ktransformers/sglang/python/sglang/srt/models/glm4_moe.py", line 484, in forward_normal_dual_stream final_hidden_states = self.experts(hidden_states, topk_output) ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ File "/root/miniconda3/envs/kt/lib/python3.11/site-packages/torch/nn/modules/module.py", line 1775, in _wrapped_call_impl return self._call_impl(*args, **kwargs) ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ File "/root/miniconda3/envs/kt/lib/python3.11/site-packages/torch/nn/modules/module.py", line 1786, in _call_impl return forward_call(*args, **kwargs) ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ File "/work/ktransformers/sglang/python/sglang/srt/layers/moe/fused_moe_triton/layer.py", line 920, in forward return self.forward_impl(hidden_states, topk_output) ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ File "/work/ktransformers/sglang/python/sglang/srt/layers/moe/fused_moe_triton/layer.py", line 939, in forward_impl combine_input = self.run_moe_core( ^^^^^^^^^^^^^^^^^^ File "/work/ktransformers/sglang/python/sglang/srt/layers/moe/fused_moe_triton/layer.py", line 960, in run_moe_core return self.quant_method.apply( ^^^^^^^^^^^^^^^^^^^^^^^^ File "/work/ktransformers/sglang/python/sglang/srt/layers/moe/kt_ep_wrapper.py", line 1321, in apply gpu_combine_input = self.gpu_method.apply(layer, masked_dispatch_output) ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ File "/work/ktransformers/sglang/python/sglang/srt/layers/quantization/compressed_tensors/compressed_tensors_moe.py", line 362, in apply return self.runner.run(dispatch_output, quant_info) ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ File "/work/ktransformers/sglang/python/sglang/srt/layers/moe/moe_runner/runner.py", line 78, in run return self.fused_func(dispatch_output, quant_info, self.config) ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ File "/work/ktransformers/sglang/python/sglang/srt/layers/moe/moe_runner/triton.py", line 337, in fused_experts_none_to_triton output = fused_experts( ^^^^^^^^^^^^^^ File "/work/ktransformers/sglang/python/sglang/srt/layers/moe/fused_moe_triton/fused_moe.py", line 213, in fused_experts inplace_fused_experts( File "/root/miniconda3/envs/kt/lib/python3.11/site-packages/torch/_ops.py", line 1255, in __call__ return self._op(*args, **kwargs) ^^^^^^^^^^^^^^^^^^^^^^^^^ File "/work/ktransformers/sglang/python/sglang/srt/layers/moe/fused_moe_triton/fused_moe.py", line 91, in inplace_fused_experts fused_experts_impl( File "/work/ktransformers/sglang/python/sglang/srt/layers/moe/fused_moe_triton/fused_moe.py", line 445, in fused_experts_impl invoke_fused_moe_kernel( File "/work/ktransformers/sglang/python/sglang/srt/layers/moe/fused_moe_triton/fused_moe_triton_kernels.py", line 746, in invoke_fused_moe_kernel fused_moe_kernel[grid]( File "/root/miniconda3/envs/kt/lib/python3.11/site-packages/triton/runtime/jit.py", line 419, in <lambda> return lambda *args, **kwargs: self.run(grid=grid, warmup=False, *args, **kwargs) ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ File "/root/miniconda3/envs/kt/lib/python3.11/site-packages/triton/runtime/jit.py", line 756, in run launch_metadata = kernel.launch_metadata(grid, stream, *bound_args.values()) ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ File "/root/miniconda3/envs/kt/lib/python3.11/site-packages/triton/compiler/compiler.py", line 490, in launch_metadata self._init_handles() File "/root/miniconda3/envs/kt/lib/python3.11/site-packages/triton/compiler/compiler.py", line 464, in _init_handles raise_(OutOfResources(self.metadata.shared, max_shared, "shared memory")) File "/root/miniconda3/envs/kt/lib/python3.11/site-packages/triton/compiler/compiler.py", line 456, in raise_ raise err triton.runtime.errors.OutOfResources: out of resource: shared memory, Required: 147456, Hardware limit: 101376. Reducing block sizes or `num_stages` may help. ``` ### Others 看提示信息是硬件支持的shared memory是100k，而这里配置的block size大于这个值，这个问题似乎与 https://zhuanlan.zhihu.com/p/1890325961670624218 这里提到的一致，是否是kvcache-ai/sglang仓库中未合入这个PR？我的kvcache-ai/sglang已拉取最新更新，且已重新安装，仍然存在这个问题。

bsakaguchi closed this issue

2026-10-29 14:48:54 +00:00

bsakaguchi commented

2026-10-29 14:48:55 +00:00

Author

Owner

@CodeZ-Hao commented on GitHub (Jan 15, 2026):

做了如下修改跑通了：

diff --git a/python/sglang/srt/layers/moe/fused_moe_triton/fused_moe_triton_config.py b/python/sglang/srt/layers/moe/fused_moe_triton/fused_moe_triton_config.py
index 90a6a31c7..faf37421e 100644
--- a/python/sglang/srt/layers/moe/fused_moe_triton/fused_moe_triton_config.py
+++ b/python/sglang/srt/layers/moe/fused_moe_triton/fused_moe_triton_config.py
@@ -10,7 +10,7 @@ import torch
 import triton
 
 from sglang.srt.server_args import get_global_server_args
-from sglang.srt.utils import get_device_name, is_hip
+from sglang.srt.utils import get_device_capability, get_device_name, is_hip
 
 logger = logging.getLogger(__name__)
 _is_hip = is_hip()
@@ -153,35 +153,71 @@ def get_default_config(
             "GROUP_SIZE_M": 8,
         }
         return config
+    
+    # Check for SM89 architecture (Ada Lovelace) which has limited shared memory
+    device_capability = get_device_capability()
+    is_sm89 = device_capability == (8, 9)
+    
     if dtype == "fp8_w8a8":
         if block_shape is None:
-            config = {
-                "BLOCK_SIZE_M": 128,
-                "BLOCK_SIZE_N": 256,
-                "BLOCK_SIZE_K": 128,
-                "GROUP_SIZE_M": 32,
-                "num_warps": 8,
-                "num_stages": 2 if _is_hip else 4,
-            }
-            if M <= E:
+            # SM89 has limited shared memory (101376 bytes), use smaller block sizes
+            if is_sm89:
                 config = {
                     "BLOCK_SIZE_M": 64,
                     "BLOCK_SIZE_N": 128,
-                    "BLOCK_SIZE_K": 128,
-                    "GROUP_SIZE_M": 1,
+                    "BLOCK_SIZE_K": 64,
+                    "GROUP_SIZE_M": 32,
                     "num_warps": 4,
+                    "num_stages": 2 if _is_hip else 3,
+                }
+                if M <= E:
+                    config = {
+                        "BLOCK_SIZE_M": 32,
+                        "BLOCK_SIZE_N": 64,
+                        "BLOCK_SIZE_K": 64,
+                        "GROUP_SIZE_M": 1,
+                        "num_warps": 2,
+                        "num_stages": 2 if _is_hip else 3,
+                    }
+            else:
+                config = {
+                    "BLOCK_SIZE_M": 128,
+                    "BLOCK_SIZE_N": 256,
+                    "BLOCK_SIZE_K": 128,
+                    "GROUP_SIZE_M": 32,
+                    "num_warps": 8,
                     "num_stages": 2 if _is_hip else 4,
                 }
+                if M <= E:
+                    config = {
+                        "BLOCK_SIZE_M": 64,
+                        "BLOCK_SIZE_N": 128,
+                        "BLOCK_SIZE_K": 128,
+                        "GROUP_SIZE_M": 1,
+                        "num_warps": 4,
+                        "num_stages": 2 if _is_hip else 4,
+                    }
         else:
             # Block-wise quant: BLOCK_SIZE_K must be divisible by block_shape[1]
-            config = {
-                "BLOCK_SIZE_M": 64,
-                "BLOCK_SIZE_N": block_shape[0],
-                "BLOCK_SIZE_K": block_shape[1],
-                "GROUP_SIZE_M": 32,
-                "num_warps": 4,
-                "num_stages": 2 if _is_hip else 3,
-            }
+            # SM89 has limited shared memory, use smaller block sizes
+            if is_sm89:
+                config = {
+                    "BLOCK_SIZE_M": 32,
+                    "BLOCK_SIZE_N": block_shape[0],
+                    "BLOCK_SIZE_K": block_shape[1],
+                    "GROUP_SIZE_M": 16,
+                    "num_warps": 2,
+                    "num_stages": 2 if _is_hip else 2,
+                }
+            else:
+                config = {
+                    "BLOCK_SIZE_M": 64,
+                    "BLOCK_SIZE_N": block_shape[0],
+                    "BLOCK_SIZE_K": block_shape[1],
+                    "GROUP_SIZE_M": 32,
+                    "num_warps": 4,
+                    "num_stages": 2 if _is_hip else 3,
+                }
     else:
         config = {
             "BLOCK_SIZE_M": 64,

@CodeZ-Hao commented on GitHub (Jan 15, 2026): 做了如下修改跑通了： ``` diff --git a/python/sglang/srt/layers/moe/fused_moe_triton/fused_moe_triton_config.py b/python/sglang/srt/layers/moe/fused_moe_triton/fused_moe_triton_config.py index 90a6a31c7..faf37421e 100644 --- a/python/sglang/srt/layers/moe/fused_moe_triton/fused_moe_triton_config.py +++ b/python/sglang/srt/layers/moe/fused_moe_triton/fused_moe_triton_config.py @@ -10,7 +10,7 @@ import torch import triton from sglang.srt.server_args import get_global_server_args -from sglang.srt.utils import get_device_name, is_hip +from sglang.srt.utils import get_device_capability, get_device_name, is_hip logger = logging.getLogger(__name__) _is_hip = is_hip() @@ -153,35 +153,71 @@ def get_default_config( "GROUP_SIZE_M": 8, } return config + + # Check for SM89 architecture (Ada Lovelace) which has limited shared memory + device_capability = get_device_capability() + is_sm89 = device_capability == (8, 9) + if dtype == "fp8_w8a8": if block_shape is None: - config = { - "BLOCK_SIZE_M": 128, - "BLOCK_SIZE_N": 256, - "BLOCK_SIZE_K": 128, - "GROUP_SIZE_M": 32, - "num_warps": 8, - "num_stages": 2 if _is_hip else 4, - } - if M <= E: + # SM89 has limited shared memory (101376 bytes), use smaller block sizes + if is_sm89: config = { "BLOCK_SIZE_M": 64, "BLOCK_SIZE_N": 128, - "BLOCK_SIZE_K": 128, - "GROUP_SIZE_M": 1, + "BLOCK_SIZE_K": 64, + "GROUP_SIZE_M": 32, "num_warps": 4, + "num_stages": 2 if _is_hip else 3, + } + if M <= E: + config = { + "BLOCK_SIZE_M": 32, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 64, + "GROUP_SIZE_M": 1, + "num_warps": 2, + "num_stages": 2 if _is_hip else 3, + } + else: + config = { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 256, + "BLOCK_SIZE_K": 128, + "GROUP_SIZE_M": 32, + "num_warps": 8, "num_stages": 2 if _is_hip else 4, } + if M <= E: + config = { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "GROUP_SIZE_M": 1, + "num_warps": 4, + "num_stages": 2 if _is_hip else 4, + } else: # Block-wise quant: BLOCK_SIZE_K must be divisible by block_shape[1] - config = { - "BLOCK_SIZE_M": 64, - "BLOCK_SIZE_N": block_shape[0], - "BLOCK_SIZE_K": block_shape[1], - "GROUP_SIZE_M": 32, - "num_warps": 4, - "num_stages": 2 if _is_hip else 3, - } + # SM89 has limited shared memory, use smaller block sizes + if is_sm89: + config = { + "BLOCK_SIZE_M": 32, + "BLOCK_SIZE_N": block_shape[0], + "BLOCK_SIZE_K": block_shape[1], + "GROUP_SIZE_M": 16, + "num_warps": 2, + "num_stages": 2 if _is_hip else 2, + } + else: + config = { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": block_shape[0], + "BLOCK_SIZE_K": block_shape[1], + "GROUP_SIZE_M": 32, + "num_warps": 4, + "num_stages": 2 if _is_hip else 3, + } else: config = { "BLOCK_SIZE_M": 64, ```

bsakaguchi commented

2026-10-29 14:48:55 +00:00

Author

Owner

@CodeZ-Hao commented on GitHub (Jan 17, 2026):

验证已在新版本修复

@CodeZ-Hao commented on GitHub (Jan 17, 2026): 验证已在新版本修复