feat: implement single-process mode on npu device. by XuZhang99 · Pull Request #1894 · xLLM-AI/xllm

XuZhang99 · 2026-07-06T09:11:31Z

Description

Add single-node single-process serving for NPU: one OS process drives multiple NPUs (one Worker thread per device) instead of one process per device wired together over the network. Selected by leaving --master_node_addr unset; DistManager::setup_single_node_workers builds the in-process worker fleet and their process groups.

Tensor-parallel collectives run through a new HcclProcessGroup, an HcclCommInitAll-backed process group that issues raw HCCL collectives on each device's current NPU stream. It deliberately bypasses c10d_npu::ProcessGroupHCCL because two ProcessGroupHCCL instances in one process can dispatch collectives in mismatched orders and deadlock the local_tp_group under load; a single HcclCommInitAll bootstrap sidesteps that. HcclCommInitAll needs a device bound on the calling thread, set via c10_npu::SetDevice (not raw aclrtSetDevice, which desyncs torch_npu's thread-local device cache and leaves worker streams in the wrong context).

Single-process multi-device requires the TORCH kernel backend; the ATB backend routes tensor-parallel comms through a process-global comm manager that cannot rendezvous two ranks in one process. setup_single_node_workers CHECKs for --npu_kernel_backend=TORCH and fails fast with guidance to use the multi-process launcher for ATB. Bump the torch_npu_ops submodule to the device-aware ATB context/operation caches the TORCH path's custom ops need.

Verified on NPU (910B, 2 cards, Qwen3-8B, TP=2): single-process and multi-process produce byte-identical greedy output, and throughput matches across the two topologies (single 5253 tok/s vs multi 5242 tok/s total, at 96 prompts / concurrency 48 / 1024-in / 512-out).

Performance (96 prompts, concurrency 48, 1024 in / 512 out):

  ┌─────────────────────────┬────────────────┬───────────────┬───────┐
  │         Metric          │ Single-process │ Multi-process │   Δ   │
  ├─────────────────────────┼────────────────┼───────────────┼───────┤
  │ Total token throughput  │ 5252.94 tok/s  │ 5242.25 tok/s │ +0.2% │
  ├─────────────────────────┼────────────────┼───────────────┼───────┤
  │ Output token throughput │ 1750.98 tok/s  │ 1747.42 tok/s │ +0.2% │
  ├─────────────────────────┼────────────────┼───────────────┼───────┤
  │ Benchmark duration      │ 28.07 s        │ 28.13 s       │ −0.2% │
  ├─────────────────────────┼────────────────┼───────────────┼───────┤
  │ Successful requests     │ 96/96          │ 96/96         │ =     │
  ├─────────────────────────┼────────────────┼───────────────┼───────┤
  │ Mean TPOT               │ 21.83 ms       │ 22.32 ms      │ −2.2% │
  ├─────────────────────────┼────────────────┼───────────────┼───────┤
  │ Mean TPOT               │ 21.83 ms       │ 22.32 ms      │ −2.2% │
  ├─────────────────────────┼────────────────┼───────────────┼───────┤
  │ Median TPOT             │ 21.14 ms       │ 21.13 ms      │ =     │
  ├─────────────────────────┼────────────────┼───────────────┼───────┤
  │ Mean TTFT               │ 2612 ms        │ 2388 ms       │ +9%   │
  └─────────────────────────┴────────────────┴───────────────┴───────┘

Related Issues

Change Type

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>.

Allowed types: feat, bugfix, docs, test, refactor, chore, style, revert, perf, model, build, release.
The subject should use clear English, start with a verb, include at least 4 words, and end with ..

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.

If you are unsure how to set up pre-commit, see the pre-commit documentation.

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

weizhehuang0827 · 2026-07-07T08:34:07Z

Does the service currently support multi-machine startup? How to start it?

XuZhang99 · 2026-07-07T08:42:00Z

Does the service currently support multi-machine startup? How to start it?

one machine one process.

I think the only difference between multi-process and single-process is that we don't need for loop when starting xllm server.

weizhehuang0827 · 2026-07-07T08:46:27Z

Does the service currently support multi-machine startup? How to start it?

one machine one process.

I think the only difference between multi-process and single-process is that we don't need for loop when starting xllm server.

Previously, the master node's address was required for synchronization; you can check that again.

XuZhang99 · 2026-07-08T06:25:22Z

here is the torch-npu-ops pr: xLLM-AI/torch-npu-ops#1

Add single-node single-process serving for NPU: one OS process drives multiple NPUs (one Worker thread per device) instead of one process per device wired together over the network. Selected by leaving --master_node_addr unset; DistManager::setup_single_node_workers builds the in-process worker fleet and their process groups. Tensor-parallel collectives run through a new HcclProcessGroup, an HcclCommInitAll-backed process group that issues raw HCCL collectives on each device's current NPU stream. It deliberately bypasses c10d_npu::ProcessGroupHCCL because two ProcessGroupHCCL instances in one process can dispatch collectives in mismatched orders and deadlock the local_tp_group under load; a single HcclCommInitAll bootstrap sidesteps that. HcclCommInitAll needs a device bound on the calling thread, set via c10_npu::SetDevice (not raw aclrtSetDevice, which desyncs torch_npu's thread-local device cache and leaves worker streams in the wrong context). Single-process multi-device requires the TORCH kernel backend; the ATB backend routes tensor-parallel comms through a process-global comm manager that cannot rendezvous two ranks in one process. setup_single_node_workers CHECKs for --npu_kernel_backend=TORCH and fails fast with guidance to use the multi-process launcher for ATB. Bump the torch_npu_ops submodule to the device-aware ATB context/operation caches the TORCH path's custom ops need. Verified on NPU (910B, 2 cards, Qwen3-8B, TP=2): single-process and multi-process produce byte-identical greedy output, and throughput matches across the two topologies (single 5253 tok/s vs multi 5242 tok/s total, at 96 prompts / concurrency 48 / 1024-in / 512-out). Co-Authored-By: Claude <noreply@anthropic.com>

…ingle_process Single-process serving (one OS process driving multiple local NPUs through a shared HcclCommInitAll world) was selected implicitly by leaving master_node_addr empty. That conflated two orthogonal axes: the in-process parallelism model (single-process multi-device vs one process per device) and whether the deployment spans nodes (which requires master_node_addr for rendezvous). The conflation made "multi-node, one process per machine driving its local cards" inexpressible: setting master_node_addr always forced the one-process-per-device path. Introduce an independent boolean --enable_single_process (DistributedConfig, defaulting false) and drive the single-process decision off it instead of inferring from master_node_addr. Thread it end-to-end: both Options structs (common + runtime), create_options, and all five runtime::Options builders in master.cpp. DistManager forks on options.enable_single_process(); with neither that flag nor master_node_addr set it now fails fast with actionable guidance. xllm.cpp derives is_local from the new flag and keys master-vs-assistant purely on node_rank (single-process always runs at node_rank 0). Scope: single-node single-process only. Multi-node single-process (nnodes > 1) is explicitly rejected in setup_single_node_workers with a "not yet supported" message, leaving the seam for a future intra-node HcclCommInitAll + inter-node CollectiveService two-level group build. BEHAVIOR CHANGE: single-node single-process now requires --enable_single_process=true (previously triggered by omitting --master_node_addr). Start scripts updated; the master_node_addr flag help text no longer claims empty implies single-process. Verified on NPU (910B, 2 cards, Qwen3-8B, TP=2): single-process (new flag) and multi-process both serve and produce byte-identical greedy output; both misconfig guards (neither flag; enable_single_process with nnodes>1) fail fast with clear messages. Co-Authored-By: Claude <noreply@anthropic.com>

…env var Replace the --enable_single_process gflag with the XLLM_ENABLE_SINGLE_PROCESS environment variable so single-process serving can be toggled from the launch environment without threading a CLI flag through every entrypoint. The DistributedConfig field is now populated from the env var (via util::get_bool_env) in from_flags(); the gflag DEFINE, its JSON binding, and its option-category entry are removed. All downstream consumers are unchanged — they still read options.enable_single_process(). Fail-fast and help messages updated to reference XLLM_ENABLE_SINGLE_PROCESS instead of --enable_single_process. Start scripts export the env var. Verified on NPU (910B, 2 cards, Qwen3-8B, TP=2): XLLM_ENABLE_SINGLE_PROCESS=1 serves single-process (Single-node serving world_size=2, coherent output); unset with no master_node_addr fails fast pointing to the env var; the removed --enable_single_process flag is now rejected as unknown; multi-process (master_node_addr set) regression unchanged with byte-identical output. Co-Authored-By: Claude <noreply@anthropic.com>

pjgao · 2026-07-11T01:23:06Z

在910C上 TP=2 上补测了这个 PR：Qwen3-8B 的单进程模式能跑通，但 Qwen3.5-27B 会在第一条生成请求上崩掉。

测试结果：

Qwen3-8B：从 ModelScope 下载官方 Qwen/Qwen3-8B 权重，冷启动一次，首请求成功；随后又跑了 16 个请求，并发为 8，全部成功，服务进程一直存活。
Qwen3.5-27B：模型加载、HCCL 初始化、KV cache 分配和 HTTP ready 都正常。第一条真实请求发出后，两个 device worker 同时注册相同的 Triton kernel，服务最终收到 signal 11。
同一个 PR binary、同一个模型、同样的两张卡和启动参数，改回多进程 TP2 后可以正常完成请求。因此这次失败不像是模型加载、OOM 或 HCCL 初始化问题，差异点就是两个 rank 是否共享同一个进程地址空间。

环境： Atlas 800T A3 / Ascend910_93，物理卡 2、3，CANN 9.0.0，driver 26.0.rc1，ATB 9.0.0，torch 2.9.0，torch_npu 2.9.0.post2，需注意我这里是910C PR描述里是 910B。

测试代码是 merge preview 7c167166，torch_npu_ops 使用公开 PR xLLM-AI/torch-npu-ops#1 的 head b65c7918。

子模块问题

当前 PR 把 third_party/torch_npu_ops 固定在 84a9596ccd436858084997ee7fa98ee3dd2f1715，但公开 remote 里取不到这个 object。 CI 也停在 checkout 阶段，错误是：

upload-pack: not our ref 84a9596ccd436858084997ee7fa98ee3dd2f1715

复现时我把 torch_npu_ops 切到了 PR #1 的 b65c7918。

Qwen3.5-27B 复现方法

权重可以从 ModelScope 下载：

env -u HTTP_PROXY -u HTTPS_PROXY -u ALL_PROXY \
    -u http_proxy -u https_proxy -u all_proxy \
  modelscope download \
    --model Qwen/Qwen3.5-27B \
    --local_dir /path/to/Qwen3.5-27B \
    --max-workers 8

切换torch_npu_ops:

git -C third_party/torch_npu_ops fetch \
  https://github.com/xLLM-AI/torch-npu-ops.git \
  refs/pull/1/head
git -C third_party/torch_npu_ops checkout \
  b65c7918d70e3f1d46eac11bed751d23eb95ba46

MAX_JOBS=32 CTEST_PARALLEL=16 SKIP_TEST=1 SKIP_EXPORT=1 \
  python3 setup.py build --device npu --tilelang-jobs 16

单进程 TP2 启动命令：

source /usr/local/Ascend/ascend-toolkit/set_env.sh
source /usr/local/Ascend/nnal/atb/set_env.sh

export ASCEND_RT_VISIBLE_DEVICES=2,3
export HCCL_IF_BASE_PORT=43432
export XLLM_ENABLE_SINGLE_PROCESS=1

./build/xllm/core/server/xllm \
  --model=/path/to/Qwen3.5-27B \
  --port=29000 \
  --devices=npu:0,npu:1 \
  --nnodes=1 \
  --node_rank=0 \
  --npu_kernel_backend=TORCH \
  --block_size=32 \
  --max_tokens_per_batch=81920 \
  --max_seqs_per_batch=2048 \
  --max_memory_utilization=0.8 \
  --enable_prefix_cache=false \
  --enable_chunked_prefill=false \
  --enable_schedule_overlap=true \
  --enable_graph=false \
  --enable_prefill_piecewise_graph=false \
  --use_contiguous_input_buffer=false \
  > /tmp/xllm-qwen35-single.log 2>&1 &

等 /v1/models ready 后发第一条请求：

MODEL_ID=$(curl -fsS http://127.0.0.1:29000/v1/models | jq -r '.data[0].id')

jq -n --arg model "$MODEL_ID" '{
  model: $model,
  messages: [{role:"user", content:"请只回答：北京是中国的首都。"}],
  temperature: 0,
  max_tokens: 32,
  stream: false,
  chat_template_kwargs: {enable_thinking:false}
}' | curl -fsS --max-time 180 \
  -H 'Content-Type: application/json' \
  -H 'Authorization: Bearer EMPTY' \
  --data-binary @- \
  http://127.0.0.1:29000/v1/chat/completions

我这里 curl 收到 empty reply，服务随后退出。最早的一组错误是两个线程同时注册 LayerNorm kernel：

Successfully registered kernel
'layer_norm_fwd_kernel_fast_rms_bf16_z_nobias'

rtFunctionRegister failed for kernel
'layer_norm_fwd_kernel_fast_rms_bf16_z_nobias': error=507028

整个失败请求里一共出现：

507028：49 次
Kernel stub is null：6 次
rtKernelLaunch failed：54 次
最终结果：signal 11

直接原因

CANN 头文件里，507028 对应 ACL_ERROR_RT_KERNEL_DUPLICATE。

当前 torch_npu_ops 的 KernelRegistry 是进程级单例，内部 unordered_map 没有同步。两个 worker 线程会走出下面这条竞态：

两个线程同时通过 kernel_infos_.find()。
两边都执行 emplace()。代码没有处理 inserted == false，两个线程继续操作同一个 entry。
两边对相同函数调用 rtFunctionRegister()。一边成功，另一边收到 507028。
失败分支执行 kernel_infos_.erase(it)，把成功线程依赖的共享 entry 也删了，后面开始出现 null stub 和重复注册。
stackcore 的 raw PC 0x34a1874 用本次实际 binary 反汇编后，落在 KernelRegistry::register_kernel() 展开的 unordered_map::erase 路径，对应 kernel_registry.cpp:177。这就是本次 signal 11 的直接崩点。

建议的修复和回归门槛

这个问题不能靠忽略 507028 处理。registry 的查询、首次注册、结果发布、读取和失败清理需要放在同一套同步及 device ownership 设计里，失败线程也不能删除其他线程已经发布的 entry。

我顺手看了 OperationFactory、AclNNGlobalCache 和 ExecutorManager，这些地方也有进程级可变状态。本次 stackcore 没有证明它们已经触发，但单进程改造合入前最好一起审计。

建议至少补下面几项：

先把 torch_npu_ops gitlink 换成公开可获取的 commit。
加一个带 barrier 的双线程、双 NPU 首次注册测试，固定制造同时注册。
Qwen3.5-27B TP2 单进程连续冷启动 3 次，要求 507028=0、null stub=0、launch failure=0。
smoke 稳定后，再做 single/multi greedy token-by-token 对齐和性能测试。

… fix Serializes the process-global triton KernelRegistry with a shared_mutex so single-process multi-device (e.g. Qwen3.5) no longer crashes on concurrent kernel registration (ACL_ERROR_RT_KERNEL_DUPLICATE 507028). Co-Authored-By: Claude <noreply@anthropic.com>

XuZhang99 · 2026-07-13T04:55:46Z

在910C上 TP=2 上补测了这个 PR：Qwen3-8B 的单进程模式能跑通，但 Qwen3.5-27B 会在第一条生成请求上崩掉。

测试结果：

Qwen3-8B：从 ModelScope 下载官方 Qwen/Qwen3-8B 权重，冷启动一次，首请求成功；随后又跑了 16 个请求，并发为 8，全部成功，服务进程一直存活。

Qwen3.5-27B：模型加载、HCCL 初始化、KV cache 分配和 HTTP ready 都正常。第一条真实请求发出后，两个 device worker 同时注册相同的 Triton kernel，服务最终收到 signal 11。

同一个 PR binary、同一个模型、同样的两张卡和启动参数，改回_多_进程 TP2 后可以正常完成请求。因此这次失败不像是模型加载、OOM 或 HCCL 初始化问题，差异点就是两个 rank 是否共享同一个进程地址空间。

环境： Atlas 800T A3 / Ascend910_93，物理卡 2、3，CANN 9.0.0，driver 26.0.rc1，ATB 9.0.0，torch 2.9.0，torch_npu 2.9.0.post2，需注意我这里是910C PR描述里是 910B。

测试代码是 merge preview 7c167166，torch_npu_ops 使用公开 PR xLLM-AI/torch-npu-ops#1 的 head b65c7918。

子模块问题

当前 PR 把 third_party/torch_npu_ops 固定在 84a9596ccd436858084997ee7fa98ee3dd2f1715，但公开 remote 里取不到这个 object。 CI 也停在 checkout 阶段，错误是：
upload-pack: not our ref 84a9596ccd436858084997ee7fa98ee3dd2f1715
复现时我把 torch_npu_ops 切到了 PR #1 的 b65c7918。

Qwen3.5-27B 复现方法

权重可以从 ModelScope 下载：
env -u HTTP_PROXY -u HTTPS_PROXY -u ALL_PROXY \
    -u http_proxy -u https_proxy -u all_proxy \
  modelscope download \
    --model Qwen/Qwen3.5-27B \
    --local_dir /path/to/Qwen3.5-27B \
    --max-workers 8
切换torch_npu_ops:
git -C third_party/torch_npu_ops fetch \
  https://github.com/xLLM-AI/torch-npu-ops.git \
  refs/pull/1/head
git -C third_party/torch_npu_ops checkout \
  b65c7918d70e3f1d46eac11bed751d23eb95ba46

MAX_JOBS=32 CTEST_PARALLEL=16 SKIP_TEST=1 SKIP_EXPORT=1 \
  python3 setup.py build --device npu --tilelang-jobs 16
单进程 TP2 启动命令：
source /usr/local/Ascend/ascend-toolkit/set_env.sh
source /usr/local/Ascend/nnal/atb/set_env.sh

export ASCEND_RT_VISIBLE_DEVICES=2,3
export HCCL_IF_BASE_PORT=43432
export XLLM_ENABLE_SINGLE_PROCESS=1

./build/xllm/core/server/xllm \
  --model=/path/to/Qwen3.5-27B \
  --port=29000 \
  --devices=npu:0,npu:1 \
  --nnodes=1 \
  --node_rank=0 \
  --npu_kernel_backend=TORCH \
  --block_size=32 \
  --max_tokens_per_batch=81920 \
  --max_seqs_per_batch=2048 \
  --max_memory_utilization=0.8 \
  --enable_prefix_cache=false \
  --enable_chunked_prefill=false \
  --enable_schedule_overlap=true \
  --enable_graph=false \
  --enable_prefill_piecewise_graph=false \
  --use_contiguous_input_buffer=false \
  > /tmp/xllm-qwen35-single.log 2>&1 &
等 /v1/models ready 后发第一条请求：
MODEL_ID=$(curl -fsS http://127.0.0.1:29000/v1/models | jq -r '.data[0].id')

jq -n --arg model "$MODEL_ID" '{
  model: $model,
  messages: [{role:"user", content:"请只回答：北京是中国的首都。"}],
  temperature: 0,
  max_tokens: 32,
  stream: false,
  chat_template_kwargs: {enable_thinking:false}
}' | curl -fsS --max-time 180 \
  -H 'Content-Type: application/json' \
  -H 'Authorization: Bearer EMPTY' \
  --data-binary @- \
  http://127.0.0.1:29000/v1/chat/completions
我这里 curl 收到 empty reply，服务随后退出。最早的一组错误是两个线程同时注册 LayerNorm kernel：
Successfully registered kernel
'layer_norm_fwd_kernel_fast_rms_bf16_z_nobias'

rtFunctionRegister failed for kernel
'layer_norm_fwd_kernel_fast_rms_bf16_z_nobias': error=507028
整个失败请求里一共出现：

507028：49 次

Kernel stub is null：6 次

rtKernelLaunch failed：54 次

最终结果：signal 11

直接原因

CANN 头文件里，507028 对应 ACL_ERROR_RT_KERNEL_DUPLICATE。

当前 torch_npu_ops 的 KernelRegistry 是进程级单例，内部 unordered_map 没有同步。两个 worker 线程会走出下面这条竞态：

两个线程同时通过 kernel_infos_.find()。

两边都执行 emplace()。代码没有处理 inserted == false，两个线程继续操作同一个 entry。

两边对相同函数调用 rtFunctionRegister()。一边成功，另一边收到 507028。

失败分支执行 kernel_infos_.erase(it)，把成功线程依赖的共享 entry 也删了，后面开始出现 null stub 和重复注册。

stackcore 的 raw PC 0x34a1874 用本次实际 binary 反汇编后，落在 KernelRegistry::register_kernel() 展开的 unordered_map::erase 路径，对应 kernel_registry.cpp:177。这就是本次 signal 11 的直接崩点。

相关代码在：

triton_npu/kernel_registry.cpp:112-178

triton_npu/kernel_registry.h:116

triton_npu/operation_base.h:218-232

Qwen3-8B 能跑和这个结论并不冲突。Qwen3-8B 的成功运行里没有任何 Triton kernel registration 日志，它没有进入 Qwen3.5 这条懒注册路径。我只验证了 Qwen3-8B 的功能可用，没有复测 PR 描述中的单/多进程吞吐数字。

建议的修复和回归门槛

这个问题不能靠忽略 507028 处理。registry 的查询、首次注册、结果发布、读取和失败清理需要放在同一套同步及 device ownership 设计里，失败线程也不能删除其他线程已经发布的 entry。

我顺手看了 OperationFactory、AclNNGlobalCache 和 ExecutorManager，这些地方也有进程级可变状态。本次 stackcore 没有证明它们已经触发，但单进程改造合入前最好一起审计。

建议至少补下面几项：

先把 torch_npu_ops gitlink 换成公开可获取的 commit。

加一个带 barrier 的双线程、双 NPU 首次注册测试，固定制造同时注册。

Qwen3.5-27B TP2 单进程连续冷启动 3 次，要求 507028=0、null stub=0、launch failure=0。

smoke 稳定后，再做 single/multi greedy token-by-token 对齐和性能测试。

fixed in xLLM-AI/torch-npu-ops#1

pjgao · 2026-07-14T09:47:43Z

补充验证了 torch-npu-ops#1 最新的 bebc9fe1623e89742576905c80c4593f1604c635。
结论：Qwen3.5-27B TP2 单进程现在可以稳定跑通，之前的并发注册崩溃已经消失；不过性能和多进程基本持平，目前看不到明确的性能优化。

测试环境仍是上一条评论中的机器：Atlas 800T A3 / Ascend910_93（机器侧标称 910C），物理卡 2、3，CANN 9.0.0、ATB 9.0.0、torch_npu 2.9.0.post2。为了隔离验证这次修复，我保留了之前的 xLLM merge preview 7c167166，只把 third_party/torch_npu_ops 从 b65c7918 更新到 bebc9fe。

测试和构建方式如下：

git -C third_party/torch_npu_ops fetch \
  https://github.com/xLLM-AI/torch-npu-ops.git \
  refs/pull/1/head
git -C third_party/torch_npu_ops checkout \
  bebc9fe1623e89742576905c80c4593f1604c635

# setup.py 会检查父仓库 gitlink；测试 lane 需要先把 gitlink 对齐到 bebc9fe
git add third_party/torch_npu_ops

MAX_JOBS=32 CTEST_PARALLEL=16 SKIP_TEST=1 SKIP_EXPORT=1 \
  python3 setup.py build --device npu --tilelang-jobs 16

启动参数与上一条评论完全相同，关键配置是：

export ASCEND_RT_VISIBLE_DEVICES=2,3
export XLLM_ENABLE_SINGLE_PROCESS=1

./build/xllm/core/server/xllm \
  --model=/home/data/weights/Qwen35-27B \
  --port=29200 \
  --devices=npu:0,npu:1 \
  --nnodes=1 \
  --node_rank=0 \
  --npu_kernel_backend=TORCH \
  --enable_graph=false \
  --enable_chunked_prefill=false \
  --enable_schedule_overlap=true

稳定性方面，做了 3 次冷启动。每次 ready 后先发 1 个请求，再发 16 个请求、并发 8：三轮共 51/51 成功，服务在请求完成后仍然存活。每轮日志都是：

507028=0
Kernel stub is null=0
rtKernelLaunch failed=0
signal 11=0

日志里可以看到一个 worker 完成首次注册，另一个 worker 命中 Kernel ... is already registered，符合这次 shared_mutex 修复的预期。

精度使用与上次相同的 C-Eval 57 题、5-shot、temperature=0：

模式	computer_network	operating_system	marxism	总分
PR single + `bebc9fe`	16/19	16/19	18/19	50/57（87.72%）
上次 PR multi	16/19	16/19	18/19	50/57（87.72%）
上次 main multi	16/19	16/19	18/19	50/57（87.72%）

single 与两组 multi 的 57 条模型输出逐条一致。

性能 workload 也与上次一致：parallel=48，warmup=48，measured=96，随机输入目标 1024 tokens，输出固定 512 tokens，temperature=0，ignore_eos=true。single 三轮全部 96/96 成功：

模式	Output tok/s	TTFT(ms)	TPOT(ms)	ITL(ms)
PR single + `bebc9fe`，median	673.35	7008.63	57.63	57.42
上次 PR multi，median	669.83	7005.42	58.01	57.79
上次 main multi，median	680.00	6696.56	57.17	56.96

single 相比 PR multi 的输出吞吐高 0.53%，相比 main multi 低 0.98%。这个差异在 1% 左右，而且 multi 数据来自 7 月 11 日，不是当天 ABBA 配对，因此只能判断“没有明显性能损失”，不能认为单进程带来了稳定性能收益。

所以目前的判断是：

bebc9fe 修复了 Qwen3.5 TP2 单进程的硬崩溃，功能和精度验证通过。
单进程和多进程性能基本持平，没有观察到明确优化。如果这个 PR 的目标包含性能提升，当前数据还不支持；如果目标是增加单进程多设备运行模式，那么目标路径已经能正常使用。
xLLM PR 当前的 gitlink 仍然是公开 remote 取不到的 84a9596...。合入前请把父仓库 third_party/torch_npu_ops 更新到公开可 fetch 的提交，否则 CI checkout 和独立复现仍会失败。

…ocess mode Single-node single-process mode has no transport between worker and engine, so forward_output_to_raw no longer routes through the forward_output_to_proto -> proto_to_forward_output serialize/deserialize round-trip. It now decodes the host-side tensors straight into the raw vectors, reusing build_token (the primitive the brpc path calls internally) so output remains byte-identical to the multi-node path. Co-Authored-By: Claude <noreply@anthropic.com>

XuZhang99 · 2026-07-15T03:04:51Z

In single-process mode, LLMEngine::step prepares the ForwardInput on the engine thread and hands it straight to the worker as a C++ object:

llm_engine.cpp:932 → worker_clients_[rank]->step_remote_async(forward_inputs[dp_rank])
WorkerClient::step_remote_async (worker_client.cpp:215-218) captures input into a lambda and calls worker_->step_async(input) directly — worker_ is a local in-process Worker* (dist_manager.cpp:255-258).

There is no forward_input_to_packed_proto, no input_write, no shm on this path. The ForwardInput is just copied as a struct (shallow — the tensors are refcounted handles, so it's a pointer/refcount bump, not a data copy).

Compare multi-process, where exactly the transport you describe happens:

RemoteWorker::step_remote_async → channel_->execute_model_async →
brpc: forward_input_to_packed_proto (comm_channel.cpp:568) → serialize → network → packed_proto_to_forward_input (worker_service.cpp:734)
shm: input_write (shm_channel.cpp:52) → shared memory → input_read(worker_service.cpp:311)

XuZhang99 changed the title ~~feat: implement single-node single-process mode on npu device.~~ feat: implement single-process mode on npu device. Jul 6, 2026

XuZhang99 force-pushed the feat/single_node_npu branch from a416ffa to a6384a5 Compare July 7, 2026 02:23

XuZhang99 marked this pull request as ready for review July 7, 2026 02:42

XuZhang99 requested review from Clement-Wang26, DongheJin, DragonFive, JimHsiung, Kang-Meng, RobbieLeung, liujinguang0125, liutongxuan, walsonyang, xiao-yu-chen, yingxudeng, yq33victor and zhang-minchao as code owners July 7, 2026 02:42

weizhehuang0827 previously approved these changes Jul 8, 2026

View reviewed changes

XuZhang99 and others added 3 commits July 10, 2026 10:58

XuZhang99 force-pushed the feat/single_node_npu branch from 11128c7 to 9d00af5 Compare July 10, 2026 02:58

XuZhang99 dismissed weizhehuang0827’s stale review via b5604fa July 14, 2026 12:43

Provide feedback

Saved searches

Use saved searches to filter your results more quickly

Uh oh!

feat: implement single-process mode on npu device.#1894

feat: implement single-process mode on npu device.#1894
XuZhang99 wants to merge 5 commits into
xLLM-AI:mainfrom
XuZhang99:feat/single_node_npu

XuZhang99 commented Jul 6, 2026 •

edited

Loading

Uh oh!

weizhehuang0827 commented Jul 7, 2026

Uh oh!

XuZhang99 commented Jul 7, 2026 •

edited

Loading

Uh oh!

weizhehuang0827 commented Jul 7, 2026

Uh oh!

XuZhang99 commented Jul 8, 2026

Uh oh!

pjgao commented Jul 11, 2026 •

edited

Loading

Uh oh!

XuZhang99 commented Jul 13, 2026

子模块问题

Qwen3.5-27B 复现方法

直接原因

建议的修复和回归门槛

Uh oh!

pjgao commented Jul 14, 2026 •

edited

Loading

Uh oh!

XuZhang99 commented Jul 15, 2026

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

3 participants

Uh oh!

Conversation

XuZhang99 commented Jul 6, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Description

Related Issues

Change Type

Pull Request Checklist

PR Title and Commit Messages

Pre-commit Checks

Self Review

Build and Test Coverage

Reviewer Notes

Uh oh!

weizhehuang0827 commented Jul 7, 2026

Uh oh!

XuZhang99 commented Jul 7, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

weizhehuang0827 commented Jul 7, 2026

Uh oh!

XuZhang99 commented Jul 8, 2026

Uh oh!

pjgao commented Jul 11, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

子模块问题

Qwen3.5-27B 复现方法

直接原因

建议的修复和回归门槛

Uh oh!

XuZhang99 commented Jul 13, 2026

子模块问题

Qwen3.5-27B 复现方法

直接原因

建议的修复和回归门槛

Uh oh!

pjgao commented Jul 14, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

XuZhang99 commented Jul 15, 2026

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

3 participants

XuZhang99 commented Jul 6, 2026 •

edited

Loading

XuZhang99 commented Jul 7, 2026 •

edited

Loading

pjgao commented Jul 11, 2026 •

edited

Loading

pjgao commented Jul 14, 2026 •

edited

Loading