[GH-ISSUE #1956] v0.6.1 does not compile SM89 (Ada Lovelace) kernels unlike v0.5.3

bsakaguchi commented

2026-10-29 14:49:57 +00:00

Owner

Originally created by @y8nqd4jtfn-cell on GitHub (Apr 26, 2026).
Original GitHub issue: https://github.com/kvcache-ai/ktransformers/issues/1956

Reminder

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

System Info

🐛 Environment & Problem

GPU: Nvidia RTX 40xx (SM89)
ktransformers version: 0.6.1
Comparison version: 0.5.3 (works correctly)
Problem: Version 0.6.1 did not compile kernel files for the SM89 architecture, while version 0.5.3 works normally.

Reproduction

🔬 Steps I Tried

In an attempt to fix this, I set the following environment variables when building from source, but no sm_89-related files were generated:

export CMAKE_ARGS="-DCMAKE_POLICY_VERSION_MINIMUM=3.5 -DENABLE_BELOW_SM90=ON"
export TORCH_CUDA_ARCH_LIST="8.9"

Others

🤔 Expectation & Request

I hope future versions will restore compilation support for the Ada architecture, or provide clear compilation options.

Originally created by @y8nqd4jtfn-cell on GitHub (Apr 26, 2026). Original GitHub issue: https://github.com/kvcache-ai/ktransformers/issues/1956 ### Reminder - [x] I have read the above rules and searched the existing issues. ### System Info ### 🐛 Environment & Problem - **GPU**: Nvidia RTX 40xx (SM89) - **ktransformers version**: 0.6.1 - **Comparison version**: 0.5.3 (works correctly) - **Problem**: Version 0.6.1 did not compile kernel files for the SM89 architecture, while version 0.5.3 works normally. ### Reproduction ### 🔬 Steps I Tried In an attempt to fix this, I set the following environment variables when building from source, but no `sm_89`-related files were generated: - `export CMAKE_ARGS="-DCMAKE_POLICY_VERSION_MINIMUM=3.5 -DENABLE_BELOW_SM90=ON"` - `export TORCH_CUDA_ARCH_LIST="8.9"` ### Others ### 🤔 Expectation & Request I hope future versions will restore compilation support for the Ada architecture, or provide clear compilation options.

bsakaguchi commented

2026-10-29 14:49:58 +00:00

Author

Owner

@cadslab commented on GitHub (Apr 26, 2026):

same issue here #1951

@cadslab commented on GitHub (Apr 26, 2026): same issue here #1951

bsakaguchi commented

2026-10-29 14:49:58 +00:00

Author

Owner

@JimmyPeilinLi commented on GitHub (Apr 30, 2026):

Thanks for reporting this. We confirmed two separate issues here:

SM89 / below-SM100 GPUs should route to the sm90 common_ops path when below-SM90 support is enabled.
The installed sgl_kernel binary still has to match the runtime torch/CUDA ABI.

The loader-side fix and torch baseline cleanup are covered by kvcache-ai/sglang#35 and #1962. The current release baseline is torch==2.9.1; for the PyPI inference path, please use the latest kt-kernel + sglang-kt packages, where sglang-kt==0.6.1.post2 is the fixed package.

If it still fails, please attach the import sgl_kernel traceback. A successful SM89 build may use the sm90 common_ops library; it does not necessarily need a separate sm_89 directory.

@JimmyPeilinLi commented on GitHub (Apr 30, 2026): Thanks for reporting this. We confirmed two separate issues here: 1. SM89 / below-SM100 GPUs should route to the `sm90` common_ops path when below-SM90 support is enabled. 2. The installed `sgl_kernel` binary still has to match the runtime torch/CUDA ABI. The loader-side fix and torch baseline cleanup are covered by kvcache-ai/sglang#35 and #1962. The current release baseline is `torch==2.9.1`; for the PyPI inference path, please use the latest `kt-kernel` + `sglang-kt` packages, where `sglang-kt==0.6.1.post2` is the fixed package. If it still fails, please attach the `import sgl_kernel` traceback. A successful SM89 build may use the `sm90` common_ops library; it does not necessarily need a separate `sm_89` directory.

bsakaguchi commented

2026-10-29 14:49:58 +00:00

Author

Owner

@y8nqd4jtfn-cell commented on GitHub (May 1, 2026):

@JimmyPeilinLi

Thank you so much for the prompt reply and the detailed explanation!

I can confirm that the SM89 routing and ABI matching issues are completely resolved with the latest kt-kernel + sglang-kt==0.6.1.post2. The kernel now loads successfully on my RTX 4090 (SM89) without any ABI errors. Thank you for the fixes covered in kvcache-ai/sglang#35 and #1962!

However, I encountered a new issue when attempting to load the Qwen3.5-MoE model using sglang-kt==0.6.1.post2. Here is the traceback:

File “…/sglang/srt/utils/hf_transformers_utils.py”, line 165, in get_hf_text_config
assert hasattr(config.text_config, “num_attention_heads”)
AssertionError

I did some digging and found the exact root cause of this problem. The get_hf_text_config function is called twice during initialization, and the config object changes between the two calls:

1st Call (Success): config is the original HuggingFace object.
type(config) = <class ‘transformers.models.qwen3_5_moe.configuration_qwen3_5_moe.Qwen3_5MoeConfig’>
type(config.text_config) = <class ‘transformers.models.qwen3_5_moe.configuration_qwen3_5_moe.Qwen3_5MoeTextConfig’>
hasattr(config.text_config, ‘num_attention_heads’) = True ✅

2nd Call (Failure): config has been converted into sglang’s internal config class.
type(config) = <class ‘sglang.srt.configs.qwen3_5.Qwen3_5MoeConfig’>
type(config.text_config) = <class ‘dict’>
hasattr(config.text_config, ‘num_attention_heads’) = False ❌

When sglang-kt converts the HuggingFace config into its own sglang.srt.configs.qwen3_5.Qwen3_5MoeConfig, the text_config field is deserialized from a class object into a plain Python dict. The code uses hasattr() to check for num_attention_heads, which works for objects but fails for dicts (where the key exists, but hasattr(dict, key) is False).

After manually patching the assert to handle dicts, I hit a second error stemming from the exact same root cause:

ValueError: User-specified context_length (131072) is greater than the derived context_length (2048).

Because text_config is a dict, downstream logic that tries to access attributes like config.text_config.max_position_embeddings fails or defaults, resulting in an incorrect derived context length.

Current Workaround: Reverting to kt-kernel==0.5.3 works flawlessly for Qwen3.5-MoE on my RTX 4090, as it bypasses this config conversion pipeline entirely.

It seems the sglang.srt.configs.qwen3_5.Qwen3_5MoeConfig class needs a comprehensive fix to properly handle the text_config dict—either by converting it back into a proper PretrainedConfig object, or by updating the downstream logic to use dictionary key lookups (e.g., config.text_config[“num_attention_heads”]) instead of attribute access.

Please let me know if you need any further details or if you’d like me to test any patches. Thanks again for your hard work and support!

@y8nqd4jtfn-cell commented on GitHub (May 1, 2026): @JimmyPeilinLi Thank you so much for the prompt reply and the detailed explanation! I can confirm that the SM89 routing and ABI matching issues are completely resolved with the latest kt-kernel + sglang-kt==0.6.1.post2. The kernel now loads successfully on my RTX 4090 (SM89) without any ABI errors. Thank you for the fixes covered in kvcache-ai/sglang#35 and #1962! However, I encountered a new issue when attempting to load the Qwen3.5-MoE model using sglang-kt==0.6.1.post2. Here is the traceback: File “…/sglang/srt/utils/hf_transformers_utils.py”, line 165, in get_hf_text_config assert hasattr(config.text_config, “num_attention_heads”) AssertionError I did some digging and found the exact root cause of this problem. The get_hf_text_config function is called twice during initialization, and the config object changes between the two calls: 1st Call (Success): config is the original HuggingFace object. type(config) = <class ‘transformers.models.qwen3_5_moe.configuration_qwen3_5_moe.Qwen3_5MoeConfig’> type(config.text_config) = <class ‘transformers.models.qwen3_5_moe.configuration_qwen3_5_moe.Qwen3_5MoeTextConfig’> hasattr(config.text_config, ‘num_attention_heads’) = True ✅ 2nd Call (Failure): config has been converted into sglang’s internal config class. type(config) = <class ‘sglang.srt.configs.qwen3_5.Qwen3_5MoeConfig’> type(config.text_config) = <class ‘dict’> hasattr(config.text_config, ‘num_attention_heads’) = False ❌ When sglang-kt converts the HuggingFace config into its own sglang.srt.configs.qwen3_5.Qwen3_5MoeConfig, the text_config field is deserialized from a class object into a plain Python dict. The code uses hasattr() to check for num_attention_heads, which works for objects but fails for dicts (where the key exists, but hasattr(dict, key) is False). After manually patching the assert to handle dicts, I hit a second error stemming from the exact same root cause: ValueError: User-specified context_length (131072) is greater than the derived context_length (2048). Because text_config is a dict, downstream logic that tries to access attributes like config.text_config.max_position_embeddings fails or defaults, resulting in an incorrect derived context length. Current Workaround: Reverting to kt-kernel==0.5.3 works flawlessly for Qwen3.5-MoE on my RTX 4090, as it bypasses this config conversion pipeline entirely. It seems the sglang.srt.configs.qwen3_5.Qwen3_5MoeConfig class needs a comprehensive fix to properly handle the text_config dict—either by converting it back into a proper PretrainedConfig object, or by updating the downstream logic to use dictionary key lookups (e.g., config.text_config[“num_attention_heads”]) instead of attribute access. Please let me know if you need any further details or if you’d like me to test any patches. Thanks again for your hard work and support!

bsakaguchi commented

2026-10-29 14:49:58 +00:00

Author

Owner

@JimmyPeilinLi commented on GitHub (May 2, 2026):

Thanks for the detailed trace.

I reproduced the same failure on our side with the Qwen3.5 MoE checkpoint.

Root cause: SGLang's internal Qwen3_5MoeConfig path did not consistently normalize nested Qwen3.5 configs, so text_config could stay as a raw dict. That triggered the num_attention_heads assertion and could also make context length derivation fall back to 2048.

I merged the source fix here:

https://github.com/kvcache-ai/sglang/pull/39

The fix makes Qwen3.5 nested text_config / vision_config consistently load as config objects while preserving the existing initialization behavior.

Validation result after the fix:

get_config.text_config_type <class 'sglang.srt.configs.qwen3_5.Qwen3_5MoeTextConfig'>
get_config.text_has_num_attention_heads True
derived_context_length 262144
ModelConfig.ok 131072

I also checked the non-MoE and MoE Qwen3.5 config constructors, default/dict initialization, from_pretrained(...), and propagation of attn_implementation / experts_implementation into nested text_config.

Until the next sglang-kt wheel includes this change, please use the latest kvcache-ai/sglang source or apply the patch from PR #39 to python/sglang/srt/configs/qwen3_5.py. For source installation, you can run:

cd sglang
pip install -e "python"

This reply was co-authored by an AI agent and a human.

@JimmyPeilinLi commented on GitHub (May 2, 2026): Thanks for the detailed trace. I reproduced the same failure on our side with the Qwen3.5 MoE checkpoint. Root cause: SGLang's internal `Qwen3_5MoeConfig` path did not consistently normalize nested Qwen3.5 configs, so `text_config` could stay as a raw dict. That triggered the `num_attention_heads` assertion and could also make context length derivation fall back to `2048`. I merged the source fix here: https://github.com/kvcache-ai/sglang/pull/39 The fix makes Qwen3.5 nested `text_config` / `vision_config` consistently load as config objects while preserving the existing initialization behavior. Validation result after the fix: ```text get_config.text_config_type <class 'sglang.srt.configs.qwen3_5.Qwen3_5MoeTextConfig'> get_config.text_has_num_attention_heads True derived_context_length 262144 ModelConfig.ok 131072 ``` I also checked the non-MoE and MoE Qwen3.5 config constructors, default/dict initialization, `from_pretrained(...)`, and propagation of `attn_implementation` / `experts_implementation` into nested `text_config`. Until the next `sglang-kt` wheel includes this change, please use the latest `kvcache-ai/sglang` source or apply the patch from PR #39 to `python/sglang/srt/configs/qwen3_5.py`. For source installation, you can run: ```bash cd sglang pip install -e "python" ``` --- *This reply was co-authored by an AI agent and a human.*

bsakaguchi commented

2026-10-29 14:49:58 +00:00

Author

Owner

@y8nqd4jtfn-cell commented on GitHub (May 3, 2026):

@JimmyPeilinLi
Thank you once again for your fix in PR #39. I installed the latest source from the repository and confirmed that the Qwen3.5 MoE config issue is resolved.

However, after that I encountered a new initialization error when starting the server with qwen3.5-122b:
Traceback (most recent call last):
...
File ".../sglang/srt/configs/deepseek_v4.py", line 7, in
@ dataclass
...
TypeError: non-default argument 'quantization_config' follows default argument

To work around this, I set the environment variable:
export SGLANG_APPLY_CONFIG_BACKUP=none

With that, the server started loading the Qwen3.5-122B model. However, because I'm placing only 8 experts on GPU (out of the full model's experts), we then encountered a runtime error during weight loading:
Loading safetensors checkpoint shards: 0% Completed | 0/39 [00:00<?, ?it/s]
Scheduler hit an exception: Traceback (most recent call last):
File ".../sglang/srt/managers/scheduler.py", line 3196, in run_scheduler_process
scheduler = Scheduler(...)
File ".../sglang/srt/managers/scheduler.py", line 367, in init
self.init_model_worker()
...
File ".../sglang/srt/model_loader/loader.py", line 686, in load_weights_and_postprocess
model.load_weights(weights)
File ".../sglang/srt/models/qwen3_5.py", line 1314, in load_weights
weight_loader(...)
File ".../sglang/srt/layers/moe/fused_moe_triton/layer.py", line 671, in weight_loader
self._weight_loader_physical(...)
...
File ".../sglang/srt/layers/moe/fused_moe_triton/layer.py", line 565, in _load_w2
loaded_weight = loaded_weight.narrow(...)
RuntimeError: start (0) + length (32) exceeds dimension size (8).

Received sigquit from a child process. It usually means the child failed.
Loading safetensors checkpoint shards: 0% Completed | 0/39 [00:00<?, ?it/s]

The error shows that in _load_w2, the code tries to narrow a tensor with length 32 (the total number of experts in the checkpoint), but the actual loaded weight only has dimension 8 (the number of experts we placed on GPU). It seems the weight loader does not yet support loading with a reduced number of experts.

@y8nqd4jtfn-cell commented on GitHub (May 3, 2026): @JimmyPeilinLi Thank you once again for your fix in PR #39. I installed the latest source from the repository and confirmed that the Qwen3.5 MoE config issue is resolved. However, after that I encountered a new initialization error when starting the server with qwen3.5-122b: Traceback (most recent call last): ... File ".../sglang/srt/configs/deepseek_v4.py", line 7, in <module> @ dataclass ... TypeError: non-default argument 'quantization_config' follows default argument To work around this, I set the environment variable: export SGLANG_APPLY_CONFIG_BACKUP=none With that, the server started loading the Qwen3.5-122B model. However, because I'm placing only 8 experts on GPU (out of the full model's experts), we then encountered a runtime error during weight loading: Loading safetensors checkpoint shards: 0% Completed | 0/39 [00:00<?, ?it/s] Scheduler hit an exception: Traceback (most recent call last): File ".../sglang/srt/managers/scheduler.py", line 3196, in run_scheduler_process scheduler = Scheduler(...) File ".../sglang/srt/managers/scheduler.py", line 367, in __init__ self.init_model_worker() ... File ".../sglang/srt/model_loader/loader.py", line 686, in load_weights_and_postprocess model.load_weights(weights) File ".../sglang/srt/models/qwen3_5.py", line 1314, in load_weights weight_loader(...) File ".../sglang/srt/layers/moe/fused_moe_triton/layer.py", line 671, in weight_loader self._weight_loader_physical(...) ... File ".../sglang/srt/layers/moe/fused_moe_triton/layer.py", line 565, in _load_w2 loaded_weight = loaded_weight.narrow(...) RuntimeError: start (0) + length (32) exceeds dimension size (8). Received sigquit from a child process. It usually means the child failed. Loading safetensors checkpoint shards: 0% Completed | 0/39 [00:00<?, ?it/s] The error shows that in _load_w2, the code tries to narrow a tensor with length 32 (the total number of experts in the checkpoint), but the actual loaded weight only has dimension 8 (the number of experts we placed on GPU). It seems the weight loader does not yet support loading with a reduced number of experts.

bsakaguchi commented

2026-10-29 14:49:58 +00:00

Author

Owner

@y8nqd4jtfn-cell commented on GitHub (May 4, 2026):

TypeError: non-default argument 'quantization_config' follows default argument when importing server_args under Python 3.12

Environment

Python: 3.12

SGLang version: installed from kvcache-ai/sglang as part of ktransformers

Installation method: fresh conda environment, then git clone ktransformers, git submodule update --init --recursive, ./install.sh, followed by an editable install from third_party/sglang/python

Steps to reproduce (within ktransformers context)

Create a new Python 3.12 environment and install ktransformers including its SGLang submodule.

Run the environment health check: kt doctor

Observe that the SGLang kt-kernel entry reports an error, indicating that the installed SGLang does not support kt‑kernel (missing --kt-gpu-prefill-token-threshold).

Observed behavior – initial kt doctor output

text
SGLang kt-kernel │ ERROR │ SGLang does not support kt-kernel (missing --kt-gpu-prefill-token-threshold argument).
)

Investigation
Trying to import sglang.srt.server_args directly fails with:

python
$ python -c "from sglang.srt.server_args import ServerArgs"
Traceback (most recent call last):
...
File ".../sglang/srt/configs/deepseek_v4.py", line 7, in
@ dataclass
^^^^^^^^^
...
TypeError: non-default argument 'quantization_config' follows default argument
The root cause is that DeepSeekV4Config (in deepseek_v4.py) is a @ dataclass that declares architectures: List[str] without a default, followed by many fields with defaults, including quantization_config. Python 3.12’s dataclass machinery enforces that non‑default fields must appear before any field with a default. This ordering violation causes the entire import chain (server_args → hf_transformers_utils → deepseek_v4) to crash, and because kt doctor (and the SGLang CLI) cannot import ServerArgs, the --kt-gpu-prefill-token-threshold argument is never exposed.

Temporary workaround
To restore functionality we manually commented out the import of DeepSeekV4Config in sglang/srt/utils/hf_transformers_utils.py and removed DeepSeekV4Config from the _CONFIG_REGISTRY list. After this change, the import succeeds and the kt argument becomes visible.

Result after workaround – kt doctor output OK status.

Note: Even after this fix, loading a Qwen3.5-122b model still raises a separate error: RuntimeError: start (0) + length (32) exceeds dimension size (8). This issue has not been resolved by the dataclass correction and may require further investigation. Additionally, setting --kt-num-gpu-experts 0 results in: TypeError: fused_topk() got an unexpected keyword argument 'num_token_non_padded'.

@y8nqd4jtfn-cell commented on GitHub (May 4, 2026): TypeError: non-default argument 'quantization_config' follows default argument when importing server_args under Python 3.12 Environment Python: 3.12 SGLang version: installed from kvcache-ai/sglang as part of ktransformers Installation method: fresh conda environment, then git clone ktransformers, git submodule update --init --recursive, ./install.sh, followed by an editable install from third_party/sglang/python Steps to reproduce (within ktransformers context) Create a new Python 3.12 environment and install ktransformers including its SGLang submodule. Run the environment health check: kt doctor Observe that the SGLang kt-kernel entry reports an error, indicating that the installed SGLang does not support kt‑kernel (missing --kt-gpu-prefill-token-threshold). Observed behavior – initial kt doctor output text SGLang kt-kernel │ ERROR │ SGLang does not support kt-kernel (missing --kt-gpu-prefill-token-threshold argument). ) Investigation Trying to import sglang.srt.server_args directly fails with: python $ python -c "from sglang.srt.server_args import ServerArgs" Traceback (most recent call last): ... File ".../sglang/srt/configs/deepseek_v4.py", line 7, in <module> @ dataclass ^^^^^^^^^ ... TypeError: non-default argument 'quantization_config' follows default argument The root cause is that DeepSeekV4Config (in deepseek_v4.py) is a @ dataclass that declares architectures: List[str] without a default, followed by many fields with defaults, including quantization_config. Python 3.12’s dataclass machinery enforces that non‑default fields must appear before any field with a default. This ordering violation causes the entire import chain (server_args → hf_transformers_utils → deepseek_v4) to crash, and because kt doctor (and the SGLang CLI) cannot import ServerArgs, the --kt-gpu-prefill-token-threshold argument is never exposed. Temporary workaround To restore functionality we manually commented out the import of DeepSeekV4Config in sglang/srt/utils/hf_transformers_utils.py and removed DeepSeekV4Config from the _CONFIG_REGISTRY list. After this change, the import succeeds and the kt argument becomes visible. Result after workaround – kt doctor output OK status. Note: Even after this fix, loading a Qwen3.5-122b model still raises a separate error: RuntimeError: start (0) + length (32) exceeds dimension size (8). This issue has not been resolved by the dataclass correction and may require further investigation. Additionally, setting --kt-num-gpu-experts 0 results in: TypeError: fused_topk() got an unexpected keyword argument 'num_token_non_padded'.

bsakaguchi commented

2026-10-29 14:49:59 +00:00

Author

Owner

@yyj6666667 commented on GitHub (May 6, 2026):

@y8nqd4jtfn-cell
pip install "transformers==4.57.1"
transformers 5.x adds default-valued fields to PretrainedConfig that make DeepSeekV4Config's dataclass declaration raise TypeError: non-default argument 'quantization_config' follows default argument at import time. sglang-kt's pyproject does not pin transformers, so a fresh pip install will pull the latest 5.x and break server startup; pinning explicitly to 4.57.1 is required until the upstream fix lands.

other common error can be found in doc/en/DeepSeek-V4-Flash.md

@yyj6666667 commented on GitHub (May 6, 2026): @y8nqd4jtfn-cell pip install "transformers==4.57.1" transformers 5.x adds default-valued fields to PretrainedConfig that make DeepSeekV4Config's dataclass declaration raise TypeError: non-default argument 'quantization_config' follows default argument at import time. sglang-kt's pyproject does not pin transformers, so a fresh pip install will pull the latest 5.x and break server startup; pinning explicitly to 4.57.1 is required until the upstream fix lands. other common error can be found in doc/en/DeepSeek-V4-Flash.md

bsakaguchi commented

2026-10-29 14:49:59 +00:00

Author

Owner

@y8nqd4jtfn-cell commented on GitHub (May 6, 2026):

@yyj6666667
git clone https://github.com/kvcache-ai/ktransformers.git
cd ktransformers
git submodule update --init --recursive
cd kt-kernel && ./install.sh
SGLang installed (kvcache-ai fork):
./install.sh # from ktransformers root
pip install --upgrade flashinfer-python flashinfer-cubin
pip install "transformers==4.57.1"
pip install tilelang

Launch SGLang Server（0.6.2, kt method: LLAMAFILE or FP8, Model: Qwen3.5-122B-A10B):

Traceback (most recent call last):
File "", line 198, in _run_module_as_main
File "", line 88, in _run_code
File "/miniconda3/envs/<env_name>/lib/python3.12/site-packages/sglang/launch_server.py", line 55, in
server_args = prepare_server_args(sys.argv[1:])
^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
File "/miniconda3/envs/<env_name>/lib/python3.12/site-packages/sglang/srt/server_args.py", line 5823, in prepare_server_args
return ServerArgs.from_cli_args(raw_args)
^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
File "/miniconda3/envs/<env_name>/lib/python3.12/site-packages/sglang/srt/server_args.py", line 5298, in from_cli_args
return cls(**{attr: getattr(args, attr) for attr in attrs})
^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
File "", line 344, in init
File "/miniconda3/envs/<env_name>/lib/python3.12/site-packages/sglang/srt/server_args.py", line 750, in post_init
self._handle_gpu_memory_settings(gpu_mem)
File "/miniconda3/envs/<env_name>/lib/python3.12/site-packages/sglang/srt/server_args.py", line 1015, in _handle_gpu_memory_settings
if not self.use_mla_backend():
^^^^^^^^^^^^^^^^^^^^^^^^^^
File "/miniconda3/envs/<env_name>/lib/python3.12/site-packages/sglang/srt/server_args.py", line 5342, in use_mla_backend
model_config = self.get_model_config()
^^^^^^^^^^^^^^^^^^^^^^^^^^^
File "/miniconda3/envs/<env_name>/lib/python3.12/site-packages/sglang/srt/server_args.py", line 5323, in get_model_config
self.model_config = ModelConfig.from_server_args(self)
^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
File "/miniconda3/envs/<env_name>/lib/python3.12/site-packages/sglang/srt/configs/model_config.py", line 318, in from_server_args
return ModelConfig(
^^^^^^^^^^^^
File "/miniconda3/envs/<env_name>/lib/python3.12/site-packages/sglang/srt/configs/model_config.py", line 192, in init
self.hf_config = get_config(
^^^^^^^^^^^
File "/miniconda3/envs/<env_name>/lib/python3.12/site-packages/sglang/srt/utils/common.py", line 3478, in wrapper
result = func(*args, **kwargs)
^^^^^^^^^^^^^^^^^^^^^
File "/miniconda3/envs/<env_name>/lib/python3.12/site-packages/sglang/srt/utils/hf_transformers_utils.py", line 359, in get_config
config = _load_deepseek_temp_model(
^^^^^^^^^^^^^^^^^^^^^^^^^^
File "/miniconda3/envs/<env_name>/lib/python3.12/site-packages/sglang/srt/utils/hf_transformers_utils.py", line 204, in _load_deepseek_temp_model
raise RuntimeError(
RuntimeError: SGLANG_APPLY_CONFIG_BACKUP=auto could not read a numeric num_hidden_layers from ~/models/Qwen3.5-122b-a10b/gpuFP8/config.json (got None).

@y8nqd4jtfn-cell commented on GitHub (May 6, 2026): @yyj6666667 git clone https://github.com/kvcache-ai/ktransformers.git cd ktransformers git submodule update --init --recursive cd kt-kernel && ./install.sh SGLang installed (kvcache-ai fork): ./install.sh # from ktransformers root pip install --upgrade flashinfer-python flashinfer-cubin pip install "transformers==4.57.1" pip install tilelang Launch SGLang Server（0.6.2, kt method: LLAMAFILE or FP8, Model: Qwen3.5-122B-A10B): Traceback (most recent call last): File "<frozen runpy>", line 198, in _run_module_as_main File "<frozen runpy>", line 88, in _run_code File "~/miniconda3/envs/<env_name>/lib/python3.12/site-packages/sglang/launch_server.py", line 55, in <module> server_args = prepare_server_args(sys.argv[1:]) ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ File "~/miniconda3/envs/<env_name>/lib/python3.12/site-packages/sglang/srt/server_args.py", line 5823, in prepare_server_args return ServerArgs.from_cli_args(raw_args) ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ File "~/miniconda3/envs/<env_name>/lib/python3.12/site-packages/sglang/srt/server_args.py", line 5298, in from_cli_args return cls(**{attr: getattr(args, attr) for attr in attrs}) ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ File "<string>", line 344, in __init__ File "~/miniconda3/envs/<env_name>/lib/python3.12/site-packages/sglang/srt/server_args.py", line 750, in __post_init__ self._handle_gpu_memory_settings(gpu_mem) File "~/miniconda3/envs/<env_name>/lib/python3.12/site-packages/sglang/srt/server_args.py", line 1015, in _handle_gpu_memory_settings if not self.use_mla_backend(): ^^^^^^^^^^^^^^^^^^^^^^^^^^ File "~/miniconda3/envs/<env_name>/lib/python3.12/site-packages/sglang/srt/server_args.py", line 5342, in use_mla_backend model_config = self.get_model_config() ^^^^^^^^^^^^^^^^^^^^^^^^^^^ File "~/miniconda3/envs/<env_name>/lib/python3.12/site-packages/sglang/srt/server_args.py", line 5323, in get_model_config self.model_config = ModelConfig.from_server_args(self) ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ File "~/miniconda3/envs/<env_name>/lib/python3.12/site-packages/sglang/srt/configs/model_config.py", line 318, in from_server_args return ModelConfig( ^^^^^^^^^^^^ File "~/miniconda3/envs/<env_name>/lib/python3.12/site-packages/sglang/srt/configs/model_config.py", line 192, in __init__ self.hf_config = get_config( ^^^^^^^^^^^ File "~/miniconda3/envs/<env_name>/lib/python3.12/site-packages/sglang/srt/utils/common.py", line 3478, in wrapper result = func(*args, **kwargs) ^^^^^^^^^^^^^^^^^^^^^ File "~/miniconda3/envs/<env_name>/lib/python3.12/site-packages/sglang/srt/utils/hf_transformers_utils.py", line 359, in get_config config = _load_deepseek_temp_model( ^^^^^^^^^^^^^^^^^^^^^^^^^^ File "~/miniconda3/envs/<env_name>/lib/python3.12/site-packages/sglang/srt/utils/hf_transformers_utils.py", line 204, in _load_deepseek_temp_model raise RuntimeError( RuntimeError: SGLANG_APPLY_CONFIG_BACKUP=auto could not read a numeric num_hidden_layers from ~/models/Qwen3.5-122b-a10b/gpuFP8/config.json (got None).

bsakaguchi commented

2026-10-29 14:49:59 +00:00

Author

Owner

@y8nqd4jtfn-cell commented on GitHub (May 6, 2026):

Description:
Attempting to deploy Qwen3.5 or Qwen3-Coder-Next MoE models using ktransformers 0.6.2 results in a cascade of 5 sequential crashes.
Below is the complete chain of failures, workarounds, and root cause analysis.

Stage 1: DeepSeek Config Misrouting (Crash on Load)

Error: RuntimeError: SGLANG_APPLY_CONFIG_BACKUP=auto could not read a numeric num_hidden_layers from config.json (got None).
Cause: SGLang detects a large MoE architecture and forcefully routes Qwen3 into the DeepSeek-specific _load_deepseek_temp_model path. The DeepSeek config patching attempts to read num_hidden_layers, which doesn’t exist in Qwen3’s config, causing a crash.
Workaround: export SGLANG_APPLY_CONFIG_BACKUP=none

Stage 2: FP8 MoE Weight Shape Mismatch (4x Discrepancy)

Error: RuntimeError: start (0) + length (32) exceeds dimension size (8) in fused_moe_triton/layer.py during _load_w2.
Cause: With config patching disabled, SGLang proceeds to load FP8 weights. The narrow() operation attempts to slice 32 elements from a dimension of 8. This 4x discrepancy indicates the weight loader calculates slice length based on FP32/BF16 bytesize or un-quantized intermediate_size, but applies it to an FP8 tensor.
Workaround: Replacing GPU-side weights from FP8 to BF16 bypasses this shape calculation error.

Stage 3: MoE TopK API Incompatibility (Crash during CUDA Graph)

Error: TypeError: fused_topk() got an unexpected keyword argument 'num_token_non_padded' (also expert_location_dispatch_info, correction_bias, scoring_func).
Cause: SGLang’s topk.py unconditionally passes DeepSeek-V3/V4 specific arguments to the fused_topk kernel, but the compiled Triton kernel does not recognize them. Forcing Qwen3.5 into the Qwen2MoE execution flow (forward_normal_dual_stream) exposes this incompatible API.
Workaround: Manually comment out these 4 arguments in sglang/srt/layers/moe/topk.py.

Stage 4: DeepSeek-V4 2604B Submode Hijacking (Crash during MoE Forward)

Error: AssertionError: swiglu_limit must be non-None iff submode=2604B (got submode='2604B', swiglu_limit=None)
Cause: In sglang/srt/environ.py, SGLANG_DSV4_2604_SUBMODE defaults to "2604B". This forces ALL MoE models into the DeepSeek-V4 2604B kernel path, which requires swiglu_limit=10.0. Qwen3 models do not have this attribute, triggering the assertion.
Workaround: export SGLANG_DSV4_2604_SUBMODE=2604A

Stage 5: Server Warmup Hang / Mamba State Offload Deadlock (Terminal Failure)

Error: After bypassing the above crashes, CUDA graph capture succeeds and Uvicorn starts, but the server never prints "The server is fired up and ready to roll!". The internal warmup request returns 500 Internal Server Error, and the process hangs.
Cause: The warmup triggers the first real forward pass involving CPU↔GPU state migration. SGLang’s native CPU offload mechanism is fundamentally incompatible with Qwen3’s Hybrid Linear Attention / GDN / Mamba architecture. When Mamba states (conv_states / ssm_states) are swapped, a dtype/layout mismatch occurs inside the Mamba kernel. Killing the process reveals the scheduler stuck inside mamba_radix_cache.sanity_check(). (Consistent with SGLang issue #20774).
Workaround: Disabling CPU offload entirely allows the server to start, but this defeats the purpose of ktransformers.

Root Cause Analysis:

DeepSeek “Virus” in Upstream SGLang: The official SGLang repository has tightly coupled DeepSeek-V3/V4 MoE optimizations (config patching, topk routing, 2604B submode) into the base MoE execution flow. Without the protective layer of sglang-kt, non-DeepSeek models trigger these code paths and crash.
Missing FP8 MoE Adaptation: SGLang’s fused_moe_triton weight loader does not correctly compute byte-to-element ratios for FP8 Qwen3 experts.
Mamba State Offloading Incompatibility: SGLang’s generic CPU offloader does not correctly handle the stateful components of Qwen3’s hybrid Mamba architecture during heterogeneous inference.

@y8nqd4jtfn-cell commented on GitHub (May 6, 2026): Description: Attempting to deploy Qwen3.5 or Qwen3-Coder-Next MoE models using ktransformers 0.6.2 results in a cascade of 5 sequential crashes. Below is the complete chain of failures, workarounds, and root cause analysis. Stage 1: DeepSeek Config Misrouting (Crash on Load) Error: RuntimeError: SGLANG_APPLY_CONFIG_BACKUP=auto could not read a numeric num_hidden_layers from config.json (got None). Cause: SGLang detects a large MoE architecture and forcefully routes Qwen3 into the DeepSeek-specific _load_deepseek_temp_model path. The DeepSeek config patching attempts to read num_hidden_layers, which doesn’t exist in Qwen3’s config, causing a crash. Workaround: export SGLANG_APPLY_CONFIG_BACKUP=none Stage 2: FP8 MoE Weight Shape Mismatch (4x Discrepancy) Error: RuntimeError: start (0) + length (32) exceeds dimension size (8) in fused_moe_triton/layer.py during _load_w2. Cause: With config patching disabled, SGLang proceeds to load FP8 weights. The narrow() operation attempts to slice 32 elements from a dimension of 8. This 4x discrepancy indicates the weight loader calculates slice length based on FP32/BF16 bytesize or un-quantized intermediate_size, but applies it to an FP8 tensor. Workaround: Replacing GPU-side weights from FP8 to BF16 bypasses this shape calculation error. Stage 3: MoE TopK API Incompatibility (Crash during CUDA Graph) Error: TypeError: fused_topk() got an unexpected keyword argument 'num_token_non_padded' (also expert_location_dispatch_info, correction_bias, scoring_func). Cause: SGLang’s topk.py unconditionally passes DeepSeek-V3/V4 specific arguments to the fused_topk kernel, but the compiled Triton kernel does not recognize them. Forcing Qwen3.5 into the Qwen2MoE execution flow (forward_normal_dual_stream) exposes this incompatible API. Workaround: Manually comment out these 4 arguments in sglang/srt/layers/moe/topk.py. Stage 4: DeepSeek-V4 2604B Submode Hijacking (Crash during MoE Forward) Error: AssertionError: swiglu_limit must be non-None iff submode=2604B (got submode='2604B', swiglu_limit=None) Cause: In sglang/srt/environ.py, SGLANG_DSV4_2604_SUBMODE defaults to "2604B". This forces ALL MoE models into the DeepSeek-V4 2604B kernel path, which requires swiglu_limit=10.0. Qwen3 models do not have this attribute, triggering the assertion. Workaround: export SGLANG_DSV4_2604_SUBMODE=2604A Stage 5: Server Warmup Hang / Mamba State Offload Deadlock (Terminal Failure) Error: After bypassing the above crashes, CUDA graph capture succeeds and Uvicorn starts, but the server never prints "The server is fired up and ready to roll!". The internal warmup request returns 500 Internal Server Error, and the process hangs. Cause: The warmup triggers the first real forward pass involving CPU↔GPU state migration. SGLang’s native CPU offload mechanism is fundamentally incompatible with Qwen3’s Hybrid Linear Attention / GDN / Mamba architecture. When Mamba states (conv_states / ssm_states) are swapped, a dtype/layout mismatch occurs inside the Mamba kernel. Killing the process reveals the scheduler stuck inside mamba_radix_cache.sanity_check(). (Consistent with SGLang issue #20774). Workaround: Disabling CPU offload entirely allows the server to start, but this defeats the purpose of ktransformers. Root Cause Analysis: DeepSeek “Virus” in Upstream SGLang: The official SGLang repository has tightly coupled DeepSeek-V3/V4 MoE optimizations (config patching, topk routing, 2604B submode) into the base MoE execution flow. Without the protective layer of sglang-kt, non-DeepSeek models trigger these code paths and crash. Missing FP8 MoE Adaptation: SGLang’s fused_moe_triton weight loader does not correctly compute byte-to-element ratios for FP8 Qwen3 experts. Mamba State Offloading Incompatibility: SGLang’s generic CPU offloader does not correctly handle the stateful components of Qwen3’s hybrid Mamba architecture during heterogeneous inference.

bsakaguchi commented

2026-10-29 14:49:59 +00:00

Author

Owner

@yyj6666667 commented on GitHub (May 7, 2026):

@y8nqd4jtfn-cell
Thank you so much for this incredibly detailed write-up — the step-by-step diagnosis with workarounds at each stage is really appreciated.

I took a look through the codebase to verify each stage, and here's what I found:

Stage 1 (Config Misrouting): Confirmed. SGLANG_APPLY_CONFIG_BACKUP defaulting to "auto" forces non-Mistral models into the DeepSeek path. Your =none workaround is spot on.
Stage 3 (TopK API Mismatch): Confirmed. select_experts() passes num_token_non_padded and expert_location_dispatch_info to fused_topk(), which doesn't accept those kwargs. This triggers under the STANDARD output format path.
Stage 4 (2604B Hijacking): Confirmed. SGLANG_DSV4_2604_SUBMODE defaulting to "2604B" causes assertion failures for models without swiglu_limit.

Thanks again for documenting this so clearly — really helpful for prioritizing fixes!

@yyj6666667 commented on GitHub (May 7, 2026): @y8nqd4jtfn-cell Thank you so much for this incredibly detailed write-up — the step-by-step diagnosis with workarounds at each stage is really appreciated. I took a look through the codebase to verify each stage, and here's what I found: * **Stage 1 (Config Misrouting):** Confirmed. `SGLANG_APPLY_CONFIG_BACKUP` defaulting to `"auto"` forces non-Mistral models into the DeepSeek path. Your `=none` workaround is spot on. * **Stage 3 (TopK API Mismatch):** Confirmed. `select_experts()` passes `num_token_non_padded` and `expert_location_dispatch_info` to `fused_topk()`, which doesn't accept those kwargs. This triggers under the **STANDARD** output format path. * **Stage 4 (2604B Hijacking):** Confirmed. `SGLANG_DSV4_2604_SUBMODE` defaulting to `"2604B"` causes assertion failures for models without `swiglu_limit`. Thanks again for documenting this so clearly — really helpful for prioritizing fixes!

bsakaguchi commented

2026-10-29 14:49:59 +00:00

Author

Owner

@y8nqd4jtfn-cell commented on GitHub (May 7, 2026):

Thank you so much for the quick verification and confirmation! Really appreciate you taking the time to go through the codebase and validate each stage. Looking forward to the fixes. Thanks again for the great work on this project!

@y8nqd4jtfn-cell commented on GitHub (May 7, 2026): Thank you so much for the quick verification and confirmation! Really appreciate you taking the time to go through the codebase and validate each stage. Looking forward to the fixes. Thanks again for the great work on this project!

Rows
Columns

[GH-ISSUE #1956] v0.6.1 does not compile SM89 (Ada Lovelace) kernels unlike v0.5.3 #1211

Reminder

System Info

🐛 Environment & Problem

Reproduction

🔬 Steps I Tried

Others

🤔 Expectation & Request