perf: route Qwen3.5 FSGDR model path to TileLang.

[ShareableXue]

Description

This PR depends on #1873.

#1873 adds the TileLang implementation of the FSGDR operator. This PR routes the Qwen3.5 FSGDR model path from the existing Triton backend to the TileLang backend.

Main changes:

• Dispatch NPU fused_sigmoid_gating_delta_rule_update to the TileLang FSGDR wrapper.
• Reshape Qwen3.5 Gated Delta Net tensors from model layout into the TileLang wrapper layout.
• Keep this PR scoped to model-path integration; the TileLang FSGDR operator itself is introduced in feat: add TileLang FSGDR kernel for NPU. #1873.

Validation Summary

Validated with Qwen3.5-27B TP4.

Accuracy

Benchmark	Setting	Score
GSM8K	0-shot CoT chat prompt	95.91
C-Eval	0-shot CoT chat prompt, naive average	90.94
C-Eval	0-shot CoT chat prompt, weighted average	90.34

Performance

Synthetic streaming benchmark, 2048 input tokens / 1024 output tokens, 128 requests, max concurrency 16.
The summary below uses aggregate benchmark metrics. The raw per-request latency and throughput distributions are listed in the appendix.

Metric	Triton FSGDR	TileLang FSGDR	Change
Average E2EL	27342.1372 ms	26674.0222 ms	-2.44%
Average TTFT	463.7918 ms	481.3325 ms	+3.78%
Average TPOT	26.2740 ms	25.6038 ms	-2.55%
Average ITL	26.2129 ms	25.5400 ms	-2.57%
Request throughput	0.5851 req/s	0.5998 req/s	+2.51%
Aggregate output token throughput	599.1590 token/s	614.1678 token/s	+2.50%
Total token throughput	1797.4769 token/s	1842.5035 token/s	+2.50%

The TileLang path improves end-to-end latency, TPOT/ITL, and aggregate output throughput in this setup. Average TTFT increases in this run.

Related Issues

Depends on #1873.

Change Type

• Bug fix
• New feature
• Performance improvement
• Refactor
• Documentation
• Test
• Build or CI

Pull Request Checklist

Thank you for contributing to xLLM. Before requesting review, please make sure the following items are complete.

PR Title and Commit Messages

• The PR title and each commit message follow the xLLM commit format: <type>: <subject>.

Pre-commit Checks

• I have installed pre-commit by running pip install pre-commit or an equivalent command.
• I have installed the hooks with pre-commit install.
• I have run pre-commit run --all-files and fixed any reported issues.

Self Review

• I have self-reviewed the code according to .agents/skills/code-review/references/custom-code-style.md, especially code written or assisted by AI.
• I have rebased this PR onto the latest main branch.

Build and Test Coverage

• Tests have been added or updated as needed.
• CUDA: python setup.py build test has passed on a CUDA machine.
• NPU: python setup.py build test has passed on an NPU machine.
• MLU: python setup.py build test has passed on an MLU machine.

Reviewer Notes

Please focus review on:

• The model-side tensor reshaping into the TileLang FSGDR wrapper layout.
• State-cache semantics through initial_state_source, initial_state_indices, and cu_seqlens.
• The NPU dispatch change from the Triton FSGDR backend to the TileLang FSGDR backend.

Appendix: Validation Details

xLLM Server Configuration

Launcher key settings:

• NNODES=4
• MAX_SEQS_PER_BATCH=16
• MAX_CONCURRENT_REQUESTS=256
• MAX_TOKENS_PER_BATCH=8192
• MAX_MEMORY_UTILIZATION=0.7
• BLOCK_SIZE=128
• ENABLE_GRAPH=true
• ENABLE_GRAPH_MODE_DECODE_NO_PADDING=true
• ENABLE_SCHEDULE_OVERLAP=true
• ENABLE_CHUNKED_PREFILL=true
• ENABLE_PREFIX_CACHE=true
• NPU_KERNEL_BACKEND=TORCH
• COMMUNICATION_BACKEND=hccl

Performance Benchmark Configuration

Benchmark model config:

• Model: Qwen3.5-27B
• API model name: xllm-qwen35-27b-general-stream
• Dataset: synthetic
• Batch size / max concurrency: 16
• Request rate: 0
• Max output length: 1024
• Generation kwargs: ignore_eos=true, temperature=0.0, top_p=1.0
• Requests: 128
• Input tokens: 2048 per request
• Output tokens: 1024 per request
• Failed requests: 0 for both runs

Raw Performance Metrics

TileLang FSGDR

Performance Parameters	Stage	Average	Min	Max	Median	P75	P90	P99	N
E2EL	total	26674.0222 ms	26536.5461 ms	26738.4477 ms	26690.7902 ms	26703.6746 ms	26715.5259 ms	26733.4346 ms	128
TTFT	total	481.3325 ms	130.8807 ms	556.1582 ms	501.8586 ms	513.6867 ms	536.5274 ms	554.8065 ms	128
TPOT	total	25.6038 ms	25.5556 ms	25.837 ms	25.5918 ms	25.605 ms	25.6208 ms	25.8297 ms	128
ITL	total	25.54 ms	0.011 ms	278.7728 ms	25.5568 ms	25.7586 ms	25.9315 ms	26.4636 ms	128
InputTokens	total	2048.0	2048.0	2048.0	2048.0	2048.0	2048.0	2048.0	128
OutputTokens	total	1024.0	1024.0	1024.0	1024.0	1024.0	1024.0	1024.0	128
OutputTokenThroughput	total	38.3895 token/s	38.2969 token/s	38.5883 token/s	38.3653 token/s	38.4238 token/s	38.5041 token/s	38.5384 token/s	128

Triton FSGDR (baseline)

Performance Parameters	Stage	Average	Min	Max	Median	P75	P90	P99	N
E2EL	total	27342.1372 ms	27283.9227 ms	27409.943 ms	27336.5717 ms	27357.0447 ms	27394.6988 ms	27406.6229 ms	128
TTFT	total	463.7918 ms	111.1144 ms	528.5816 ms	512.5106 ms	517.6789 ms	520.599 ms	526.8069 ms	128
TPOT	total	26.274 ms	26.1886 ms	26.5728 ms	26.2351 ms	26.2774 ms	26.4404 ms	26.5526 ms	128
ITL	total	26.2129 ms	0.0121 ms	245.1767 ms	26.2104 ms	26.3655 ms	26.5297 ms	27.0245 ms	128
InputTokens	total	2048.0	2048.0	2048.0	2048.0	2048.0	2048.0	2048.0	128
OutputTokens	total	1024.0	1024.0	1024.0	1024.0	1024.0	1024.0	1024.0	128
OutputTokenThroughput	total	37.4514 token/s	37.3587 token/s	37.5313 token/s	37.459 token/s	37.4797 token/s	37.4939 token/s	37.5037 token/s	128

Raw Accuracy Metrics

GSM8K:

dataset	version	metric	mode	xllm-qwen35-27b-tilelang
gsm8k	7cd45e	accuracy	gen	95.91

C-Eval:

dataset	version	metric	mode	xllm-qwen35-27b-tilelang
ceval-computer_network	6fdaa9	accuracy	gen	89.47
ceval-operating_system	1f61dd	accuracy	gen	100.00
ceval-computer_architecture	96481d	accuracy	gen	100.00
ceval-college_programming	268efc	accuracy	gen	91.89
ceval-college_physics	264c12	accuracy	gen	94.74
ceval-college_chemistry	e2aec8	accuracy	gen	95.83
ceval-advanced_mathematics	910069	accuracy	gen	73.68
ceval-probability_and_statistics	9376fc	accuracy	gen	83.33
ceval-discrete_mathematics	52c40a	accuracy	gen	62.50
ceval-electrical_engineer	a554da	accuracy	gen	72.97
ceval-metrology_engineer	6a3e90	accuracy	gen	79.17
ceval-high_school_mathematics	0a2955	accuracy	gen	83.33
ceval-high_school_physics	d654f1	accuracy	gen	100.00
ceval-high_school_chemistry	8e4f48	accuracy	gen	100.00
ceval-high_school_biology	b72e30	accuracy	gen	100.00
ceval-middle_school_mathematics	6536db	accuracy	gen	100.00
ceval-middle_school_biology	78de25	accuracy	gen	85.71
ceval-middle_school_physics	71fcea	accuracy	gen	100.00
ceval-middle_school_chemistry	09a74b	accuracy	gen	100.00
ceval-veterinary_medicine	1dbb02	accuracy	gen	95.65
ceval-college_economics	31c49d	accuracy	gen	87.27
ceval-business_administration	96b9ce	accuracy	gen	84.85
ceval-marxism	e797d7	accuracy	gen	94.74
ceval-mao_zedong_thought	4d15db	accuracy	gen	95.83
ceval-education_science	1637c2	accuracy	gen	96.55
ceval-teacher_qualification	995db1	accuracy	gen	93.18
ceval-high_school_politics	89bf67	accuracy	gen	94.74
ceval-high_school_geography	ada66e	accuracy	gen	100.00
ceval-middle_school_politics	f4c957	accuracy	gen	100.00
ceval-middle_school_geography	a8e72c	accuracy	gen	91.67
ceval-modern_chinese_history	a917ca	accuracy	gen	91.30
ceval-ideological_and_moral_cultivation	681251	accuracy	gen	100.00
ceval-logic	937358	accuracy	gen	95.45
ceval-law	e26a5f	accuracy	gen	91.67
ceval-chinese_language_and_literature	fb9d94	accuracy	gen	91.30
ceval-art_studies	009e1a	accuracy	gen	90.91
ceval-professional_tour_guide	2dda62	accuracy	gen	93.10
ceval-legal_professional	c15753	accuracy	gen	82.61
ceval-high_school_chinese	c94ef3	accuracy	gen	78.95
ceval-high_school_history	f5aa51	accuracy	gen	95.00
ceval-middle_school_history	ac2319	accuracy	gen	100.00
ceval-civil_servant	083900	accuracy	gen	91.49
ceval-sports_science	e8eecf	accuracy	gen	94.74
ceval-plant_protection	eb7ad1	accuracy	gen	100.00
ceval-basic_medicine	b0fb42	accuracy	gen	94.74
ceval-clinical_medicine	6f192e	accuracy	gen	81.82
ceval-urban_and_rural_planner	c3e73c	accuracy	gen	82.61
ceval-accountant	d3491c	accuracy	gen	91.84
ceval-fire_engineer	bd1824	accuracy	gen	70.97
ceval-environmental_impact_assessment_engineer	729e25	accuracy	gen	83.87
ceval-tax_accountant	99b77f	accuracy	gen	89.80
ceval-physician	00a46d	accuracy	gen	89.80
ceval-stem	-	naive_average	gen	90.41
ceval-social-science	-	naive_average	gen	93.88
ceval-humanities	-	naive_average	gen	91.85
ceval-other	-	naive_average	gen	88.33
ceval-hard	-	naive_average	gen	86.68
ceval	-	naive_average	gen	90.94
ceval-weighted	-	weighted_average	gen	90.34

[gemini-code-assist]

Code Review

This pull request updates the NPU kernel implementation for fused_sigmoid_gating_delta_rule_update to use the TileLang wrapper fused_sigmoid_gating_delta_rule and adjusts the input tensor shapes in Qwen3GatedDeltaNetBaseImpl::forward by replacing .contiguous() with appropriate .reshape(...) calls. A critical issue was identified where attn_metadata.q_cu_seq_lens is not explicitly cast to torch::kInt32 before being passed to the kernel, which could lead to a runtime crash due to type assertions.

Important

The consumer version of Gemini Code Assist on GitHub is being sunset. Starting June 18, 2026, new organization installations will be blocked, and all code review activity will officially cease on July 17, 2026.
For more details on the timeline and next steps, please review the Help Documentation.

[gemini-code-assist]

The attn_metadata.q_cu_seq_lens tensor (assigned to params.cu_seqlens on line 863) is not guaranteed to be of type torch::kInt32 and can often be torch::kInt64/torch::kLong. However, the TileLang wrapper fused_sigmoid_gating_delta_rule has an explicit assertion CHECK_EQ(cu_seqlens.scalar_type(), torch::kInt32) which will fail and cause a runtime crash if a 64-bit integer tensor is passed. Please explicitly cast attn_metadata.q_cu_seq_lens to torch::kInt32 before calling .contiguous(), similar to how it is handled in other kernels (e.g., causal_conv1d_update or chunk_gated_delta_rule).