feat(ascend): add Huawei Ascend NPU support

[ningweikang]

feat(ascend): add Huawei Ascend NPU support

Overview

Add Huawei Ascend NPU inference support to RTP-LLM, based on torch_npu + CANN ACLNN operators. This is the first upstream merge of the Ascend adaptation, covering build infrastructure, device management, operator adaptation, attention / KV cache, sampling, and module factories.

Guiding principles:

1. Follow the NVIDIA and AMD GPU integration pattern — add Ascend code side by side
2. All changes are behind --config=ascend / --define=using_ascend=true — DO NOT break existing CUDA/ROCm code logic
3. Explicit fail-fast for capabilities not yet supported on Ascend — no silent degradation to wrong results

Changes

1. Build Infrastructure

• Bazel Ascend toolchain autoconfiguration (3rdparty/gpus/ascend_configure.bzl, analogous to cuda_configure): probes CANN headers/libs (acl/acl.h, libascendcl.so, libhccl.so) via TF_NEED_ASCEND / ASCEND_TOOLKIT_PATH / TF_ASCEND_VERSION, and generates the @local_config_ascend repository; falls back to a dummy repo when Ascend is not configured
• .bazelrc --config=ascend preset with -DUSING_ASCEND=1, --define=using_ascend=true, -lascendcl -lhccl linking and CANN LD_LIBRARY_PATH; dedicated test:ascend env (ASCEND_CUSTOM_OPP_PATH, TEST_USING_DEVICE=ASCEND, torch/torch_npu lib paths)
• Ascend pip requirements (deps/requirements_ascend_x86.txt: torch 2.9.0+cpu, torch_npu 2.9.0, numpy 1.26.4) and lockfile (requirements_lock_ascend.txt); pip_ascend_torch repo is gated behind TF_NEED_ASCEND so non-Ascend builds never resolve torch_npu wheels
• AscendC custom-op build (3rdparty/aclnn_custom_ops): compiles the apply_top_k_top_p_custom top-k/top-p sampling kernel from an external source repo via build_for_bazel.sh (soc version ascend950), exporting aclnn runtime libs + OPP packages
• BUILD.torch_npu to wire torch_npu headers/libs into Bazel targets
• -Wl,-Bsymbolic-functions link option to resolve the CXX-ABI mismatch across CANN libraries (fixes a std::regex runtime memory error)
• arch_select.bzl Ascend branch for wheel dependency resolution (drops CUDA-only deps such as flash_attn, decord, av)

2. C++ Device Generalization

• USING_ASCEND branches in ExecOps: runtimeSyncAndCheck via aclrtSynchronizeDevice, events on torch::kPrivateUse1 bound to NPU streams, no-op DeviceGuard
• CUDA graph shims (cuda_graph_device_shims.h): GRAPH_DEVICE_TYPE = torch::kPrivateUse1 and graph-stream stubs — CUDA-graph capture is not yet available on Ascend; the shim keeps the graph runner compiling
• Device-agnostic routing across BlockPool, MemoryEvaluationHelper (HBM stats via aclrtGetMemInfo(ACL_HBM_MEM)), MemoryLayoutStrategy, KVCacheAllocator, KV-cache connectors, Sampler, ModelTypes, PyWrappedModel, NormalEngine / NormalExecutor, input gatherers / output dispatcher, and the MTP (speculative) executor
• ascend_host_utils.cc: ACL error checking (ASCEND_CHECK), device count / current device / memory queries, graph-capture flag helpers
• Simplified fusedCopy for Ascend host/device copies; an ops stub keeps CUDA kernels out of Ascend builds
• GenerateStream: NPU RNG generator support for deterministic sampling
• Shared-code portability fixes: CRC-instruction compile error and std::regex runtime memory error

3. Python Device & Attention

• DeviceType.Ascend detection via torch_npu (device_type.py, device_impl.py) with automatic device routing
• Attention ascend_impl registered in the attention factory:
  • Prefill and decode via npu_fused_infer_attention_score_v2 (FIA) with TND layout + block_table over the 3D paged KV cache — the same default path as vLLM-Ascend;
  • Pure PyTorch RoPE (apply_rope_pos_ids_nhd), paged KV-cache write op with slot mapping, write_cache_store integration

4. Sampling (Ascend)

• NPU-accelerated greedy and top-k/top-p sampling: temperature via torch_npu-mapped div_(), top-k/top-p filtering via the AscendC custom op applyTopKTopP, final sampling via torch::multinomial (same fallback strategy as ROCm)
• NPU generator for deterministic seeded sampling
• Not yet available on Ascend: repetition penalty (AscendC kernel TBD), rejection sampling

5. Module Factories & Server Integration

• Linear factory: AscendF16Linear for F16/BF16 non-quantized linear via F.linear; the NV cutlass implementation is moved into the cuda/ directory to keep the factory device-clean
• Fused MoE: AscendBf16FallbackStrategy — pure PyTorch batched-experts fallback (no quantization), plus Ascend select_topk / moe_gating fallbacks
• Base modules under modules/base/ascend/: RMSNorm / LayerNorm / QKRMSNorm (pure PyTorch), Embedding with TP all_gather, activation fallbacks
• Explicit fail-fast stubs: GroupTopK, FakeBalanceExpert, IndexerOp
• Server startup and configuration adaptations (server_config_setup.py, start_backend_server.py); Ascend-aware weight loading and memory cleanup (model_loader/loader.py, weight_manager.py)

Test Environment

• NPU: Ascend 950
• Software: PyTorch 2.9.0+cpu, torch_npu 2.9.0, CANN toolkit, Python 3.10

How to Build

source /usr/local/Ascend/cann/set_env.sh
export TF_NEED_ASCEND=1
export ASCEND_TOOLKIT_PATH=/usr/local/Ascend/ascend-toolkit/latest

bazel build //rtp_llm:rtp_llm --config=ascend --verbose_failures

Known Limitations

Capability	Status
CUDA graph capture	Stubbed, not yet available
MLA attention	Not supported (MHA only)
Repetition penalty / rejection sampling	Not yet implemented (AscendC kernels TBD)
Weight-only quantized linear / quantized MoE	Not yet supported (BF16/F16 only)
Performance	Current linear / MoE / norm paths are functional PyTorch fallbacks; optimized AscendC kernels will land in follow-up PRs

[CLAassistant]

Thank you for your submission! We really appreciate it. Like many open source projects, we ask that you all sign our Contributor License Agreement before we can accept your contribution.
1 out of 2 committers have signed the CLA.

✅ QQzhouQ
❌ nwk_bince

---

nwk_bince seems not to be a GitHub user. You need a GitHub account to be able to sign the CLA. If you have already a GitHub account, please add the email address used for this commit to your account.
_{You have signed the CLA already but the status is still pending? Let us recheck it.}

[LLLLKKKK]

AI Code Review - PR #1349

Status: NEEDS REBASE

该 PR 当前无法 rebase（与目标分支存在冲突 / dirty），已跳过本次自动代码审查。请 rebase 到最新目标分支并解决冲突后重新推送，审查会在新 head 上自动重新触发。

[LLLLKKKK]

AI Code Review - PR #1349

Status: BLOCKING

Summary: P0/2 · P1/6 · P2/17 · P3/2

Reviewed: commit e9f7c1ed25e0 · 2026-08-28 23:25 UTC+8

Blocking Issues

P0

• ROCm 下 graphCaptureBegin 不再调用 capture_begin，既有平台的图捕获被静默破坏 @ rtp_llm/cpp/cuda_graph/cuda_graph_device_shims.cc:212
  • 建议：改为三分支：#if USING_CUDA 保留 graph.capture_begin(pool)，#elif USING_ROCM 恢复 graph.capture_begin()，#else 才空实现，避免 Ascend 适配顺带改变 ROCm 行为。注意 GraphPoolHandle 在 shims.h:49-53 也被收窄为 #if USING_CUDA，ROCm 需要能接受空 pool 的重载路径。合入前请在 MI308X 上跑一次开启 CUDA Graph 的 ROCm 冒烟用例确认 capture/replay 正常，并补一条 ROCm 图捕获测试，避免该路径回归后长期无人发现。
• Ascend pip lock 文件未随 PR 提交且存在三种不一致命名，顶层无条件 load 阻塞所有平台 @ deps/pip.bzl:89
  • 建议：补交由 bazel run //deps:requirements_ascend_x86.update 生成的 deps/requirements_lock_ascend_x86.txt，三处命名统一到该名字（与 deps/requirements_ascend_x86.txt 对齐），并在 deps/BUILD 的 exports_files 中显式导出；删除无任何消费者的根级悬空软链，以免引入第四种名字。若短期无法提供 lock，请把 arch_select.bzl:8 的 load 与 WORKSPACE:80-81 的 pip_ascend_torch_install_deps() 一并延后，不要把新平台的未完成状态扩散成全平台阻塞。合入前请至少用非 ascend 配置跑一次 bazel query //rtp_llm/... 验证 loading 通过。顺带修正 deps/pip.bzl:90 混入的 TAB 缩进。

P1

• base_model.py 为 ROCm 返回 hip:N 设备串，与全栈 torch.cuda 语义冲突 @ rtp_llm/models/base_model.py:175
  • 建议：删除 ROCm 分支，仅保留 Ascend 的 npu:N 与默认 cuda:N——本 PR 目标是 Ascend，不应顺带改变 ROCm 的设备串约定。更进一步，把设备串生成下沉到已有的 rtp_llm/device 抽象统一提供，并补一个 mock get_device_type 的单测覆盖 Cuda/ROCm/Ascend 三种取值。若确认需要区分，必须先在 ROCm 环境实测 torch.device("hip:N") 可用于张量分配与 torch.cuda.Stream 构造，再合入。
• Ascend 下 MtpExecutor 的 AsyncRunner::sync 被整段编译掉，丢失 host join 与异常回抛 @ rtp_llm/cpp/normal_engine/speculative/MtpExecutor.cc:1404
  • 建议：采用本 PR 已在 NormalExecutor.cc 建立的范式：Ascend 仍调用 sync()，只替换 stream 实参；或为 AsyncRunner 增加只做 host join + 异常回抛的 sync() 无参重载（Ascend 上 graphGetCurrentStream() 返回 nullptr 无法作实参，这正是当前被编译掉的根因）。host join 与异常回抛在任何平台都必须保留。注意 NormalExecutor.cc:206 中 c10::Device(c10::DeviceType::PrivateUse1, 0) 硬编码 index 0，多卡 TP 下会指向错误设备，请一并改为 local_rank。若短期无法支持，请在 MtpExecutor 构造处对 Ascend 显式 fail-fast 并说明原因。
• Ascend 上事件 record 退化为空操作，transfer_done_event 的 D2H 完成不变量被静默破坏 @ rtp_llm/cpp/cuda_graph/cuda_graph_device_shims.h:140
  • 建议：不要让 record/block 在新平台静默降级：Ascend 分支应基于 aclrtEvent / c10_npu 事件实现真实 record/block；若事件语义暂时无法保证，请在该 shim 内改为显式 stream 同步（而非空操作），使 synchronize() 的调用方契约在所有平台一致。若本阶段不支持 Ascend 投机解码，请在引擎初始化处对 propose_params_/sp_type 增加 fail-fast 校验。同时把注释改为"Ascend 依赖单流执行顺序，跨流事件依赖被丢弃"，避免后续开启 useStreamAsync() 时静默产生数据竞争。另 makeGraphEvent 的 Ascend 分支与 #else 分支（GRAPH_DEVICE_TYPE 在 :10 即为同一值）完全等价，属冗余可删。
• start_backend_server 的 Ascend 设备数分支不可达，多卡 NPU 启动静默退化为单 rank @ rtp_llm/start_backend_server.py:438
  • 建议：把 434 行的早退条件从"CUDA 不可用"改为"CUDA 与 NPU 均不可用"，建议复用本 PR 已新增的 rtp_llm/device/device_type.py:48 的 is_ascend()。注意解除早退后 126/148 行会真正求值 torch.npu.*，而本文件全程没有 import torch_npu，需先补守卫；同时 _create_rank_processes:178 只写 CUDA_VISIBLE_DEVICES，Ascend 需改写对应的 NPU 可见设备环境变量，否则即使进入多 rank 路径也无法正确分卡。请补一条多卡 Ascend 启动用例。若本期不支持 Ascend 多卡，请回退 126/148/438 三处 NPU 三元表达式，并在早退处显式 raise/warning 说明限制。
• Ascend 依赖链全程无完整性校验：空 sha256、关闭 hash、第三方个人仓库缩写 commit @ deps/http.bzl:84
  • 建议：为两个 http_archive 填入真实 sha256（bazel fetch 报错信息会给出实际值），并为 torch_npu_ascend 补内部镜像 URL 作首选源、公网直链作回退；requirements 两条直链加上 #sha256= 片段并移除 deps/BUILD:112 的 generate_hashes = False，与 rocm/cuda 保持同等校验强度。aclnn 源码请镜像到受控存储并改用带 sha256 的 http_archive；若必须保留 git 依赖，至少改用 40 位完整 SHA 并把声明收敛到 deps/git.bzl 以复用统一镜像通道。该依赖的 genrule 会执行外部仓库脚本并把产物 .so 带入运行期，请在 PR 描述中记录其许可协议、来源归属与在 CI 沙箱内执行的安全边界。
• Ascend 依赖清单缺少 torchvision，而 //rtp_llm:torchvision 是无条件依赖 @ deps/requirements_ascend_x86.txt:4
  • 建议：二选一并保持自洽：若 ascend 运行时需要 torchvision（多模态预处理路径会 import），请在 deps/requirements_ascend_x86.txt 补入与 torch 2.9.0 匹配的版本；若 ascend 阶段性不支持多模态，请把 "torchvision" 加入 _ascend_excluded，并确认 rtp_llm/BUILD:266 的消费方在 torchvision 缺失时不会在 import 期崩溃。合入前建议跑一次 bazel query 'deps(//rtp_llm:models)' --config=ascend，把这类 select 分支缺包问题在分析期暴露。

Non-blocking Suggestions

P2

• Ascend 专属链接选项 -Bsymbolic-functions 被无条件施加到全平台首个加载的 .so @ BUILD:152
  • 建议：按平台 gate：linkopts = ["-Wl,-rpath='$$ORIGIN'"] + if_ascend(["-Wl,-Bsymbolic-functions"])（def.bzl:21 已导出 if_ascend）。在注释中说明为何仅 config 库需要而 rtp_compute_ops/th_transformer 不需要，便于后续回滚判断；若确认 CUDA/ROCm 也需要，请在 PR 描述中给出验证依据（例如对比两个 .so 的动态符号绑定）。同时修正 Fox → For。
• ExecOps.h 被 cpp 核心层新增引用但两个 Bazel 目标未声明该依赖，且属反向跨层依赖 @ rtp_llm/cpp/cuda_graph/cuda_graph_runner.h:3
  • 建议：分别在 cuda_graph_hdrs_lib 与 cuda_copy_util 的 deps 中补上 //rtp_llm/models_py/bindings/core:exec_ops_hdr（该 header-only 目标已存在并在 pybind/BUILD:139,144 被同样引用），让头文件目标自带完整依赖闭包，否则一旦切到沙箱/远端执行或启用 layering_check 就会编译失败。更彻底的做法是把 getTorchCudaDevice() 下沉到 rtp_llm/cpp/utils 之类的中立层——本 PR 已让该函数成为全仓设备入口，放在 models_py/bindings/core 会让层次边界持续劣化。
• BlockInfo::is_cuda 语义被扩为"设备内存"，但字段名、注释与部分消费者未同步 @ rtp_llm/cpp/cache/MemoryLayoutStrategy.cc:307
  • 建议：把字段重命名为 is_device，或改为直接携带 MemoryType（与本 PR 新增的 MEMORY_NPU、torchDeviceToMemoryType 对齐），并全仓搜索消费者统一改造、同步更新 BlockInfo.h 的注释；至少修掉 KVCacheManager.cc:765，并明确 DecodeRpcServer.cc 两处在 NPU 下传给 cache store 的期望语义。建议补一条 NPU/CPU 混合 BlockInfo 的单测锁定新语义。
• NPU 保 ACL format 的 from_blob 包装无法复用，KV 拷贝路径仍直接用 torch::from_blob @ rtp_llm/cpp/cache/MemoryLayoutStrategy.cc:17
  • 建议：把 makeCacheTypedBlob 提到共享头（如 models_py/bindings/core/torch_utils/TypeConvert.h），让所有以 NPU 裸指针构造 tensor 的位置统一走它。若确认 execNoBlockCopy 的纯 uint8 字节视图不依赖 ACL format tag，请在 wrapRawPtr 与 KVCacheMemoryConnector 两处补一行注释说明豁免理由与边界条件，避免后续读者误用或误删。
• KVCacheAllocator 新增未使用的 include，且未复用同 PR 新增的 allocationTypeToMemoryType @ rtp_llm/cpp/cache/KVCacheAllocator.cc:294
  • 建议：将 :294 替换为 allocationTypeToMemoryType(allocation_type_)（这样新增的 include 也有了实际用途，且 Ascend 下内存类型标注正确）；否则应删除该未使用的 include，避免留下"看起来已改造"的假象。
• PyWrappedModel 同一函数内三种设备判定口径并存，NPU 张量会走到非法 pin_memory 路径 @ rtp_llm/cpp/models/PyWrappedModel.cc:170
  • 建议：抽取一个 isDeviceResident(const torch::Tensor&) 内联函数（is_cuda() || is_privateuseone()，或直接用 !t.device().is_cpu()）供本文件所有判定点复用，统一提升块、normalize_i32、to_device_i32 与 :211/:254/:316 的判定，避免同函数内多种口径并存导致 Ascend 在运行期随机位置失败。Torch_ext.h 的 CHECK_CUDA 已采用同一口径，可作命名参考。
• MtpExecutor 设备判定仅部分转换，Ascend 下投机路径既不支持也不显式拒绝 @ rtp_llm/cpp/normal_engine/speculative/MtpExecutor.cc:276
  • 建议：二选一：要么完成该文件内设备判定的一致转换（统一引入 isDeviceTensor(t) 配合 getTorchCudaDevice()，并把断言文案改为设备无关描述），要么在 MtpExecutor 构造处对 Ascend 显式 fail-fast 并说明原因。避免留下半转换状态——它既让读者无法判断支持边界，也把失败点推迟到运行期随机位置。
• Ascend 不链接 cuda_graph_impl，shims.cc 全部 Ascend 分支为未经验证的死代码 @ rtp_llm/cpp/cuda_graph/cuda_graph_device_shims.cc:143
  • 建议：明确二选一并落实：若 Ascend 计划支持 graph 捕获，请在 pybind/BUILD:142 与 normal_engine/BUILD 的 ascend 分支补上 cuda_graph_impl、为 cuda_graph/BUILD 补 acl 依赖，并修正实现（graphMemcpyAsync 改用 aclrtMemcpyAsync，参照 FusedCopyOp.cc:30-53 的 memcpyKindFor；aclrtSynchronizeDevice 接入 ASCEND_CHECK；graphCheck 真正校验错误码），再补一条 Ascend 捕获/回放用例。若暂不支持，请删除这些 Ascend 分支与 shims.h:175 的无实现重载声明，用 RTP_LLM_CHECK_WITH_INFO(false, ...) 明确表达 fail-fast（与 PyWrappedModel.h 既有范式一致），并在 BUILD 的 ascend 空分支处加注释说明"Ascend 不链接 cuda_graph_impl，graph 路径关闭"。顺带清理已无引用的 cuda_graph_srcs filegroup。
• getTorchCudaDevice 返回不带 index 的设备，叠加 Ascend 侧 setDevice/流绑定空实现，多卡落卡与流排序均无保证 @ rtp_llm/models_py/bindings/core/ExecOps.cc:599
  • 建议：在 getTorchCudaDevice() 中显式带上当前设备号（复用 ascend::getDevice() 构造 torch::Device(torch::kPrivateUse1, index)），或提供带 index 的重载供需要绑定的调用点使用；async runner 的流 device index 显式绑定为 local_rank，NormalExecutor.cc:206 的硬编码 0 一并改掉。Ascend 分支的 setDevice 改用 aclrtSetDevice（ascend_host_utils 已有对应能力）、stream guard 改用 c10_npu 实现；若暂不实现，至少在首次进入时打印一次 WARNING 或直接 fail-fast，不要静默产生错卡执行。请补一个 Ascend 多卡（至少 2 卡）P2P 传输与 MTP 用例，或在 PR 描述中说明该路径当前不在 Ascend 支持范围内。
• cuda_host_utils.h 新增的 Ascend 分支只有注释，头文件在 Ascend 下依然无法编译 @ rtp_llm/models_py/bindings/cuda/cuda_host_utils.h:29
  • 建议：要么删除这个空分支（保持"此头文件仅限 CUDA/ROCm"的事实），要么把 timing_function 等依赖 cudaStream_t 的声明也用 #if USING_CUDA || USING_ROCM 包裹，使头文件在 Ascend 下真正可安全 include；两种做法都比留一段注释更清晰。
• Sampler.cc 把原 ROCm 分支嵌入 #if USING_CUDA，形成不可达死代码并丢失关键说明 @ rtp_llm/cpp/models/Sampler.cc:126
  • 建议：删除已不可达的内层 #if USING_ROCM 分支，把其中关于 hipMemcpyAsync 与 record_stream() 在 at::hip 流上 abort 的说明移到 #else 分支上方（ROCm 现在正走该分支），使注释与实际编译路径对齐；或把结构改为互斥三分支 #if USING_CUDA && !USING_ROCM / #elif USING_ROCM / #else，恢复 ROCm 分支可达性。不要留下永不编译的分支。
• DEVICE_TYPE 全局宏在两个新增头文件中语义冲突，-Werror 下会编译失败 @ rtp_llm/cpp/normal_engine/test/MockEngine.h:27
  • 建议：删除测试侧这两处宏，直接调用 getTorchCudaDevice()（本身已是平台无关入口，无需再包一层），或在 rtp_llm/cpp/testing/TestBase.h 提供 inline 辅助（如 testDevice()）供各测试复用；若确需保留宏，请改用带前缀的名字（如 RTP_TEST_DEVICE_TYPE）。同时建议把 op_api_common.h 的 DEVICE_TYPE 重命名为 ASCEND_DEVICE_TYPE，不要在公共头文件里占用如此通用的全局宏名。
• torch_npu 探测逻辑多处内联重复，未复用同 PR 新增的 is_ascend()，且设备绑定失败被静默吞掉 @ rtp_llm/config/server_config_setup.py:268
  • 建议：在两个文件顶部 from rtp_llm.device.device_type import is_ascend，把内联块统一替换为 if is_ascend():；设备数获取同理，建议在 device_type.py 补一个 device_count() 供两处共用。若担心重复走 get_device_type()，可加 functools.lru_cache(maxsize=1)。set_device 分支在已识别为 Ascend 但设置失败时至少 logging.warning 输出 local_rank 与原因，并把捕获范围扩到 (ImportError, OSError, AttributeError)（缺底层 .so 抛 OSError、torch.npu 未注册抛 AttributeError，都会穿透当前的 except ImportError）。setup_cuda_device_and_accl_env 的函数名与 docstring 仍为 CUDA-only 语义却已承担 NPU 职责，建议一并更新表述。
• _get_cuda_device_list 用 dir() 成员检测做控制流，条件恒为 False @ rtp_llm/start_backend_server.py:148
  • 建议：删除 _dev_count if "_dev_count" in dir() else 这层判断，改为把设备计数作为显式参数传入（_get_cuda_device_list(dev_count)），或调用统一的设备计数辅助函数（与前一条 finding 的修复共用）并在其中用 is_ascend() 守卫 torch.npu。不要用 dir()/hasattr 之类的内省做控制流分支。函数命名与 CUDA_VISIBLE_DEVICES 语义也建议随之泛化。
• Ascend 侧新增分支基本零测试覆盖，测试改造未接入承载 target @ rtp_llm/cpp/normal_engine/speculative/test/MtpExecutorTest.cc:474
  • 建议：参照本 PR 已对 rtp_llm/cpp/utils/test/BUILD:30-33 采用的做法，把两个 test BUILD 的 env/deps/exec_properties 改为 select({"@//:using_ascend": ..., "//conditions:default": ...})，让相关用例在 Ascend 下可构建可执行；若本期不打算在 Ascend 跑，建议暂缓这 4 个文件的设备无关化改动，避免留下无覆盖代码。Python 侧在 server_config_setup_test.py 补用例：unittest.mock.patch 让 torch.cuda.is_available() 返回 False 并向 sys.modules 注入伪 torch_npu 与 torch.npu，参数化断言（a）seq_size_per_block 被置 128 且用户显式配置时不被覆盖，（b）local_world_size == min(npu_count, world_size)，（c）torch_npu 缺失时回落原有行为不抛异常（注意 mock target 是使用处）。并在 PR 描述中说明 Ascend 当前如何验证（机型、CANN 版本、bazel 命令与结果）与接入 CI 的计划。
• _ascend_excluded 排除名单不完整，ascend 配置下通配构建会解析到不存在的 pip 包 @ arch_config/arch_select.bzl:26
  • 建议：以 ascend lock 的实际包集合为准补全排除名单，或改为"白名单 + 默认解析为空"，避免每加一个平台都要重复维护一份 denylist。同时把一条 ascend 配置下的 analysis 级校验纳入 CI（例如 bazel query 'deps(//rtp_llm:models)' --config=ascend），否则这类缺口只有在真正跑 ascend 通配构建时才会暴露。
• aclnn 自定义算子目标结构不规范、SoC 硬编码、运行期路径依赖相对 bazel-bin 与私有 conda 环境 @ 3rdparty/aclnn_custom_ops/BUILD:30
  • 建议：把 :build_aclnn_ops 从 deps 移除（产物已在 data，genrule 会被正常触发；如需聚合请用 filegroup），消除重复的输出路径列表，并给 cc_library 补上与 genrule 相同的 target_compatible_with；把 ascend950 提取为可配置项（--define 或 select）并在非目标 SoC 上 fail-fast；自定义算子路径改由 target 的 data + runfiles 解析（$(location) 展开或 runfiles API），不要在配置文件里写相对 bazel-bin 路径；PYTHONHOME 与 torch/torch_npu 库路径统一到仓库标准解释器位置（/opt/conda310）或改为 action_env 注入，避免绑定单机私有环境。

P3

• 多处新增的 using_ascend select 空分支与 default 等价，且 cuda_graph_impl 收敛顺带改变了 cpu/arm 链接图 @ rtp_llm/cpp/utils/BUILD:216
  • 建议：删除与 //conditions:default 等价的空分支；若刻意保留是为"显式声明该平台已评估过"，请统一加一行简短注释说明意图。models/BUILD:191 若意图仅排除 Ascend，请写成 select({"@//:using_ascend": [], "//conditions:default": ["//rtp_llm/cpp/cuda_graph:cuda_graph_impl"]}) 以保持 cpu/arm 行为不变；若确实要一并收敛，请在 PR 描述中说明并至少跑一次 bazel build //rtp_llm/cpp/models:models --config=cpu 验证链接通过。
• 条件编译缩进错乱、测试行超长与注释拼写错误 @ rtp_llm/cpp/normal_engine/NormalExecutor.cc:320
  • 建议：统一 #if/#endif 顶格书写（与文件既有风格一致）；把重复的 host 拷贝判断抽成测试内小工具函数（例如 toHost(const torch::Tensor&)，或直接用 !t.device().is_cpu()），既消除重复也自然满足 120 列限制，随后对这两个测试文件执行一次 clang-format，避免后续提交混入无关格式化 diff；修正 Fox → For。

Checklist Findings (18 fail / 58 total)

General Principles Checklist

• [6.1] Architecture — 依赖方向：无循环依赖/跨层惊喜 → issue ExecOps.h 被 cpp 核心层新增引用但两个 Bazel 目标未声明该依赖，且属反向跨层依赖
  两处新增 #include "rtp_llm/models_py/bindings/core/ExecOps.h" 的目标都缺声明：（1）cuda_graph_runner.h:3 新增该 include（置于标准库之前）并在 inline 构造函数中调用 getTorchCudaDevice()，但导出该头的 cuda_graph_hdrs_lib（cuda_graph/BUILD:45-57）deps 只有 //rtp_llm/models_py/bindings:op_defs，没有 exec_ops_hdr（对比同文件 cuda_graph_impl 在 :78 已显式声明）。（2）CudaCopyUtil.cc:5 新增该 include，但 cuda_copy_util（.../tcp/BUILD:5-15，本 PR 未修改）deps 仅 //:rtp_compute_ops（linkshared=1 cc_binary，不传播编译上下文）、no_block_copy_hdr、core_utils。当前能编译仅因
• [6.1] Architecture — 兼容性：外部 HTTP/RPC API、持久数据、配置、环境迁移安全 → issue aclnn 自定义算子目标结构不规范、SoC 硬编码、运行期路径依赖相对 bazel-bin 与私有 conda 环境
  :30-33 把 genrule :build_aclnn_ops 放进 cc_library.deps，而 genrule 不提供 CcInfo，存在 analysis 期报错风险；即使当前 Bazel 容忍，它与 data:24-29 已列出的同一批 .so 完全重复。target_compatible_with 只加在 genrule 上（:15-18），cc_library 自身无平台约束，误引用时报错会指向平台不兼容而非依赖写错。cmd:12-13 把 SoC 目标硬编码为 ascend950，在 910B 等芯片上会产出不匹配的算子包且无构建期校验。产物仅通过相对路径暴露给运行期：.bazelrc:417 的 LD_LIBRARY_PATH 与 :421 的 ASCEND_CUSTOM_OPP_PATH 都写 bazel-bin/3rdparty/aclnn_custom_ops/...，但测试进程 cwd 为 runfiles 目录而非 workspace 根；:417-418 还把 torch/torch_npu 库路径与
• [6.1] Architecture — 分层边界：新概念在正确层级，不泄漏内部 → issue cuda_host_utils.h 新增的 Ascend 分支只有注释，头文件在 Ascend 下依然无法编译
  新增的 #elif USING_ASCEND（:29-32）分支体内只有两行注释，不引入任何声明。但该头文件 :49-51 的 timing_function 无条件使用 cudaStream_t（两处），Ascend 下没有任何 include 提供该类型，因此头文件在 USING_ASCEND 下仍编译失败。这个空分支只会让读者误以为该头已完成 Ascend 适配，掩盖了"Ascend 必须改用 ascend_host_utils.h"这一真实约束（cuda_graph_device_shims.cc:2-10 的 include 分支正是这样做的）。
• [6.1] Architecture — 可观测性：日志/指标/超时可操作、非噪声 → issue torch_npu 探测逻辑多处内联重复，未复用同 PR 新增的 is_ascend()，且设备绑定失败被静默吞掉
server_config_setup.py 在 :265-276（set_parallelism_config）、:435-441（setup_default_args）、:540-548（setup_cuda_device_and_accl_env）三处重复了完全相同的 try: import torch_npu / if torch.npu.is_available() / except ImportError: pass；start_backend_server.py :126/:148/:438 又各自内联了同类三元（且全程无 import torch_npu）。而同一 PR 已在 rtp_llm/device/device_type.py:27-33 实现同样探测并对外暴露 is_ascend()（:48-49），与既有 is_cuda()/is_hip()/is_ppu() 并列，且仅依赖 os/enum/torch 是叶子模块。其中 :543-548 的 torch.npu.set_device(local_rank)
• [6.1] Architecture — 回滚路径：风险行为存在运维回滚手段 → issue Ascend 专属链接选项 -Bsymbolic-functions 被无条件施加到全平台首个加载的 .so
BUILD:150-153 给 cc_binary(name="th_transformer_config") 无条件加上 -Wl,-Bsymbolic-functions，注释自述动机是"import torch_npu 引入 CANN 库时 CXX11 ABI 混用导致 std::regex 内存问题"，属纯 Ascend 场景（.bazelrc:404 也只为 ascend 设 _GLIBCXX_USE_CXX11_ABI=1），但该选项对 cuda12/cuda12_9/cuda13/rocm/arm/cpu 全部生效（相邻 th_grammar_tokenizer_info:162-166 仍只有 rpath）。libth_transformer_config.so 是所有平台最先加载的 .so，而本仓库明确依赖跨 .so 的 load-time 符号解析（pybind/BUILD:134-136）。我未能证实"双份函数级 static 单例"会实际发生（该 .so 以 Python 扩展模块 dlopen，默认 RTLD_LOCA
• [6.1] Architecture — 状态不变量：创建/更新/失败/重试/回滚路径有效 → issue getTorchCudaDevice 返回不带 index 的设备，叠加 Ascend 侧 setDevice/流绑定空实现，多卡落卡与流排序均无保证
ExecOps.cc:597-603 在 Ascend 下返回 torch::Device(torch::kPrivateUse1)，不带 device index（index=-1）。CUDA 侧 index=-1 由 torch 解析为"当前设备"，与改动前等价、无回归；但 PrivateUse1 后端由 torch_npu 注册，其对 index=-1 的解析未在本 PR 中验证。放大风险的是 shims.h 中一批活代码空实现：GraphStream 退化为 void*（:56），GraphStreamGuard 丢弃入参（:57-59），setDevice（:79-81）、graphSetCurrentStream（:120-122）为空，toGraphStream/graphGetCurrentStream/graphGetStreamFromPool 一律返回 nullptr。MtpExecutor.cc:552-557 用它构造 5 个 runner 流，AsyncRunner.cc:12 以 `stream_
• [6.1] Architecture — 错误语义：fail-fast/retry/fallback/silent 行为显式 → issue aclnn 自定义算子目标结构不规范、SoC 硬编码、运行期路径依赖相对 bazel-bin 与私有 conda 环境
  :30-33 把 genrule :build_aclnn_ops 放进 cc_library.deps，而 genrule 不提供 CcInfo，存在 analysis 期报错风险；即使当前 Bazel 容忍，它与 data:24-29 已列出的同一批 .so 完全重复。target_compatible_with 只加在 genrule 上（:15-18），cc_library 自身无平台约束，误引用时报错会指向平台不兼容而非依赖写错。cmd:12-13 把 SoC 目标硬编码为 ascend950，在 910B 等芯片上会产出不匹配的算子包且无构建期校验。产物仅通过相对路径暴露给运行期：.bazelrc:417 的 LD_LIBRARY_PATH 与 :421 的 ASCEND_CUSTOM_OPP_PATH 都写 bazel-bin/3rdparty/aclnn_custom_ops/...，但测试进程 cwd 为 runfiles 目录而非 workspace 根；:417-418 还把 torch/torch_npu 库路径与
• [6.1] Quality — Commit 原子、message 与行为匹配 → issue Ascend pip lock 文件未随 PR 提交且存在三种不一致命名，顶层无条件 load 阻塞所有平台
  _deps/pip.bzl:87-89 的 pip_parse(name="pip_ascend_torch") 引用 @rtp_deps//:requirements_lock_ascend_x86.txt，deps/BUILD:111 同名引用；但 deps/ 下经枚举只有 8 个 lock（cuda12_arm/cuda13_arm/rocm/torch_arm/torch_cpu/torch_gpu_cuda12{,9}/cuda13），该文件不存在。PR 另新增根级 requirements_lock_ascend.txt，读取直接报"文件不存在"（悬空软链），构成第三个名字。关键在于 WORKSPACE:80-81 与 arch_config/arch_select.bzl:8 都是顶层无条件 load，与 --config 无关，且 arch_select.bzl 被几乎所有包加载，故 cuda12/cuda12_9/rocm/arm/cpu 任一配置的 build/test/query 都会在 loading 阶段中断。
• [6.1] Quality — 逻辑变更未混入无关格式化 → issue 条件编译缩进错乱、测试行超长与注释拼写错误
  新增的条件编译指令缩进不一致：NormalExecutor.cc:320 为 #if !USING_ASCEND 而其 #endif（:324）在第 0 列；MtpExecutor.cc:1305/1387/1404/1497/1699 同为缩进 #if 配 0 列 #endif，而 :1558/1708 又改成缩进 #endif，同一文件内三种写法并存。MtpExecutorTest.cc:479 与 MtpBatchStreamProcessorTest.cc:148 改动后单行约 150 字符，超过 .clang-format 的 120 列上限，而相邻代码均为已格式化状态，pre-commit 会再次改写这两行；同一 host 拷贝判断 is_cuda() || is_privateuseone() ? t.cpu() : t 在两文件中重复 4 次。另 BUILD:152 注释 Fox 应为 For。
• [6.1] Software Engineering — DRY：重复非平凡逻辑被抽取或显式复用 → issue 条件编译缩进错乱、测试行超长与注释拼写错误
  新增的条件编译指令缩进不一致：NormalExecutor.cc:320 为 #if !USING_ASCEND 而其 #endif（:324）在第 0 列；MtpExecutor.cc:1305/1387/1404/1497/1699 同为缩进 #if 配 0 列 #endif，而 :1558/1708 又改成缩进 #endif，同一文件内三种写法并存。MtpExecutorTest.cc:479 与 MtpBatchStreamProcessorTest.cc:148 改动后单行约 150 字符，超过 .clang-format 的 120 列上限，而相邻代码均为已格式化状态，pre-commit 会再次改写这两行；同一 host 拷贝判断 is_cuda() || is_privateuseone() ? t.cpu() : t 在两文件中重复 4 次。另 BUILD:152 注释 Fox 应为 For。
• [6.1] Software Engineering — KISS/YAGNI：无投机性抽象 → issue 多处新增的 using_ascend select 空分支与 default 等价，且 cuda_graph_impl 收敛顺带改变了 cpu/arm 链接图
  五处新增的 "@//:using_ascend": [] 与同一 select 已有的 "//conditions:default": [] 语义完全相同、不改变任何构建结果：rtp_llm/cpp/utils/BUILD:216、embedding_engine/BUILD:28、models/BUILD:95、models/BUILD:191、normal_engine/BUILD:41；空分支会让后续读者误以为 Ascend 有特殊处理。同时 diff 显示 models/BUILD 删除了原本无条件的 cuda_graph_impl（diff 第 2077 行 -），改为只在 using_cuda/using_rocm 下提供，纯 cpu/arm 配置也一并落入 default 丢掉该依赖，超出 PR 声明的 Ascend 范围且无注释说明（核查确认当前不产生未解析符号，但 CI 不构建 cpu/arm，该变化无任何验证）。
• [6.1] Tests — 分布式/跨平台变更有对应覆盖 → issue 多处新增的 using_ascend select 空分支与 default 等价，且 cuda_graph_impl 收敛顺带改变了 cpu/arm 链接图
  五处新增的 "@//:using_ascend": [] 与同一 select 已有的 "//conditions:default": [] 语义完全相同、不改变任何构建结果：rtp_llm/cpp/utils/BUILD:216、embedding_engine/BUILD:28、models/BUILD:95、models/BUILD:191、normal_engine/BUILD:41；空分支会让后续读者误以为 Ascend 有特殊处理。同时 diff 显示 models/BUILD 删除了原本无条件的 cuda_graph_impl（diff 第 2077 行 -），改为只在 using_cuda/using_rocm 下提供，纯 cpu/arm 配置也一并落入 default 丢掉该依赖，超出 PR 声明的 Ascend 范围且无注释说明（核查确认当前不产生未解析符号，但 CI 不构建 cpu/arm，该变化无任何验证）。
• [6.1] Tests — 新逻辑有聚焦单测 + 相关集成/smoke 测试 → issue _ascend_excluded 排除名单不完整，ascend 配置下通配构建会解析到不存在的 pip 包
arch_select.bzl:52 对不在 _ascend_excluded（:26-36，仅 9 项）中的每个名字生成 requirement_ascend(name)。但 requirement() 的调用点 rtp_llm/BUILD:81-145 还包含 tilelang/flash-mla/fast-hadamard-transform(:37-41)、deep_gemm/deep_ep(:43-46)、rtp-kernel(:48-50)、flash_attn/flash-attn-3(:52-55) 以及 flashinfer_with_cache 各项，已 grep 确认这些既不在 deps/requirements_base.txt 也不在 deps/requirements_ascend_x86.txt 中。虽然消费者 select 让 ascend 目标不依赖它们，但 requirement() 仍会为每个名字生成 py_library 并在其 ascend 分支引用缺失包，`bazel b
• [6.1] Tests — 边界 case 覆盖（空、单元素、最大值） → issue Ascend 侧新增分支基本零测试覆盖，测试改造未接入承载 target
  C++ 侧：本 PR 改 4 个测试文件以支持 PrivateUse1（MtpExecutorTest.cc:474,479、MtpBatchStreamProcessorTest.cc:144,148 增加 is_privateuseone()；MockEngine.h:27、NormalBatchStreamProcessorTest.cc:8 引入 DEVICE_TYPE），但承载它们的 normal_engine/test/BUILD 与 speculative/test/BUILD 均不在 diff_paths（diff 中相关测试 BUILD 仅 rtp_llm/cpp/utils/test/BUILD）：后者经核查仍无条件依赖 @local_config_cuda//cuda:cuda_headers/cudart（:63-64,89-90）、env 硬编码 TEST_USING_DEVICE: "CUDA"（:40,67,93）、exec_properties 指定 H20，故改造在 Ascend 上无法构建、在

RTP-LLM Checklist

• [I] 代码质量 — 删除或重命名内部 file、registry entry、model name、metric enum、op binding、plugin symbol 时，必须全仓搜索消费者，并提供替代实现、迁移说明或 smoke 覆盖；只有暴露到 HTTP/RPC/config/persisted format 时才按外部兼容性处理 → issue BlockInfo::is_cuda 语义被扩为"设备内存"，但字段名、注释与部分消费者未同步
MemoryLayoutStrategy.cc:307（dev.is_cuda() || dev.is_privateuseone()）与 LayerBlockConverterImpl.h:35-36（同时纳入 BufferKey 去重键）把 BlockInfo::is_cuda 扩展为"设备内存"语义。但 BlockInfo.h:14-16 未随本 PR 修改，注释仍写"Torch device of the backing storage (CPU/CUDA)"，字段名与含义自此不一致。消费者亦未全部同步：KVCacheManager.cc:765 仍 dst_block.is_cuda ? torch::kCUDA : torch::kCPU，在 Ascend 上会用 NPU 指针构造 CUDA 张量（该函数为 writeKVBlockForTest，见 :756 报错文案，影响面限于测试）；model_rpc/DecodeRpcServer.cc:1007,1153 仍把该布尔作为 GPU 标志传入 `load_layer_cache-
• [I] 代码质量 — 同一功能用统一工具函数 → issue _get_cuda_device_list 用 dir() 成员检测做控制流，条件恒为 False
  :148 写作 [str(i) for i in range(_dev_count if "_dev_count" in dir() else (torch.npu.device_count() if not torch.cuda.is_available() else torch.cuda.device_count()))]。无参 dir() 返回当前作用域的名字，而 _dev_count 在 _get_cuda_device_list 内从未赋值（列表推导另有独立作用域），故条件永远为 False、第一支是死分支；反之若条件成立则直接 NameError。该写法看似想读取 _get_local_world_size:126 或 start_backend_server:438 的同名局部变量，但 Python 不存在跨函数作用域穿透，且一旦后续有人在本函数内绑定同名局部变量，分支会静默翻转。本文件全程无 import torch_npu，torch.npu 属未守卫访问。

Python Static-First Checklist

• [P.A] 静态结构与类型纪律 — 禁止 hasattr 做控制流分支 → issue _get_cuda_device_list 用 dir() 成员检测做控制流，条件恒为 False
  :148 写作 [str(i) for i in range(_dev_count if "_dev_count" in dir() else (torch.npu.device_count() if not torch.cuda.is_available() else torch.cuda.device_count()))]。无参 dir() 返回当前作用域的名字，而 _dev_count 在 _get_cuda_device_list 内从未赋值（列表推导另有独立作用域），故条件永远为 False、第一支是死分支；反之若条件成立则直接 NameError。该写法看似想读取 _get_local_world_size:126 或 start_backend_server:438 的同名局部变量，但 Python 不存在跨函数作用域穿透，且一旦后续有人在本函数内绑定同名局部变量，分支会静默翻转。本文件全程无 import torch_npu，torch.npu 属未守卫访问。
• [P.B] 错误处理 — 禁止 bare except 或静默吞异常 → issue torch_npu 探测逻辑多处内联重复，未复用同 PR 新增的 is_ascend()，且设备绑定失败被静默吞掉
server_config_setup.py 在 :265-276（set_parallelism_config）、:435-441（setup_default_args）、:540-548（setup_cuda_device_and_accl_env）三处重复了完全相同的 try: import torch_npu / if torch.npu.is_available() / except ImportError: pass；start_backend_server.py :126/:148/:438 又各自内联了同类三元（且全程无 import torch_npu）。而同一 PR 已在 rtp_llm/device/device_type.py:27-33 实现同样探测并对外暴露 is_ascend()（:48-49），与既有 is_cuda()/is_hip()/is_ppu() 并列，且仅依赖 os/enum/torch 是叶子模块。其中 :543-548 的 torch.npu.set_device(local_rank)

Strengths

• 设备句柄收敛到 getTorchCudaDevice()（ExecOps.cc:597-603）单一入口，而非在调用点散落 #if；已核实 CUDA/ROCm 返回值与原 torch::kCUDA 隐式转换所得设备完全等价，属零行为变更重构，新增平台只需改一处。
• MemoryLayoutStrategy.cc:3-8 明确注释了 torch::from_blob 会丢失 NPU tensor 的 ACL format 并导致下游 FIA/scatter attention kernel 读错，是必要且到位的 WHY 注释——正是这条注释让评审能反向定位到其他未做同样处理的调用点。
• TypeConvert.h:81-104 抽出 allocationTypeToMemoryType/torchDeviceToMemoryType/memoryTypeToTorchDevice 并新增 MEMORY_NPU，没有把 NPU 硬塞进 MEMORY_GPU，让 BatchCopyParams 的 D2D/D2H/H2D 分类在新平台自然成立。
• 平台接入严格复用既有扩展点：requirement()、torch_deps()、platform_deps()、select_py_bindings()、no_block_copy_link_deps() 均是在已有 select() 中追加 using_ascend 分支并保留 //conditions:default，未改写中心分发逻辑，符合 OCP。
• NormalExecutor.cc:203-207 为 Ascend 保留了 sync() 的 host 侧语义、只替换 stream 实参，是本 PR 内处理该问题的正确范式。
• fuse_copy_kernel.h:15-19 用 #if USING_CUDA || USING_ROCM 精确收敛依赖 cudaStream_t 的声明；Ascend 侧在 FusedCopyOp.cc:30-53 以 aclrtPointerGetAttributes 判定方向 + aclrtMemcpyAsync + ASCEND_CHECK 实现，分层清晰且是 fail-fast 正确范式。
• 3rdparty/aclnn_custom_ops/BUILD:15-18 对 genrule 使用 target_compatible_with，非 Ascend 配置自动跳过而非编译失败；FusedCopyOp.cc:16-19 对 torch_npu 头文件用 #pragma GCC diagnostic 精确收敛告警，未全局关闭。
• rtp_llm/cpp/pybind/BUILD:142-145 的 Ascend 装配正确且有注释说明跨 .so 符号解析约定：exec_ctx_ops 提供实现、exec_ops_hdr 提供头，不产生重复符号。
• seq_size_per_block = 128 带明确理由（FIA v2 paged attention 推荐 block_size=128），沿用 MI308X/alixpu 已有的"仅当 == 0 才覆盖"守卫保留用户覆盖能力，且插入位置在 alixpu 的 256 之后、通用兜底 64 之前，优先级正确。
• PrefixToCandidateTokens.h:504-505 把 unique_ptr<FILE, decltype(&fclose)> 换成无捕获 lambda deleter，规避不同 libc 下 fclose 声明差异，RAII 语义完全不变。

[LLLLKKKK]

📍 实际位置 rtp_llm/cpp/cuda_graph/cuda_graph_device_shims.cc:212（不在 diff 展示范围内，就近挂载）

[P0] ROCm 下 graphCaptureBegin 不再调用 capture_begin，既有平台的图捕获被静默破坏

diff 中该函数 #else 分支为 - graph.capture_begin(); / + (void)graph;，即本 PR 明确删掉了非 CUDA 平台的捕获起始调用。.bazelrc:351-352 对 rocm 设置 USING_CUDA=0/USING_ROCM=1，故 ROCm 走 #else，捕获从不开始；而 cuda_graph_runner.cc:1186 无平台守卫地调用 graphCaptureBegin，紧接 1197 行执行 graph.capture_end()，形成无 begin 的 end 配对。该路径线上有效：models/BUILD:190、normal_engine/BUILD:39、pybind/BUILD:140、embedding_engine/BUILD:27 均为 ROCm 链接 cuda_graph_impl；全仓 grep 确认 capture_begin 仅此一处调用点。

建议: 改为三分支：#if USING_CUDA 保留 graph.capture_begin(pool)，#elif USING_ROCM 恢复 graph.capture_begin()，#else 才空实现，避免 Ascend 适配顺带改变 ROCm 行为。注意 GraphPoolHandle 在 shims.h:49-53 也被收窄为 #if USING_CUDA，ROCm 需要能接受空 pool 的重载路径。合入前请在 MI308X 上跑一次开启 CUDA Graph 的 ROCm 冒烟用例确认 capture/replay 正常，并补一条 ROCm 图捕获测试，避免该路径回归后长期无人发现。

[LLLLKKKK]

[P0] Ascend pip lock 文件未随 PR 提交且存在三种不一致命名，顶层无条件 load 阻塞所有平台

deps/pip.bzl:87-89 的 pip_parse(name="pip_ascend_torch") 引用 @rtp_deps//:requirements_lock_ascend_x86.txt，deps/BUILD:111 同名引用；但 deps/ 下经枚举只有 8 个 lock（cuda12_arm/cuda13_arm/rocm/torch_arm/torch_cpu/torch_gpu_cuda12{,_9}/cuda13），该文件不存在。PR 另新增根级 requirements_lock_ascend.txt，读取直接报"文件不存在"（悬空软链），构成第三个名字。关键在于 WORKSPACE:80-81 与 arch_config/arch_select.bzl:8 都是顶层无条件 load，与 --config 无关，且 arch_select.bzl 被几乎所有包加载，故 cuda12/cuda12_9/rocm/arm/cpu 任一配置的 build/test/query 都会在 loading 阶段中断。

建议: 补交由 bazel run //deps:requirements_ascend_x86.update 生成的 deps/requirements_lock_ascend_x86.txt，三处命名统一到该名字（与 deps/requirements_ascend_x86.txt 对齐），并在 deps/BUILD 的 exports_files 中显式导出；删除无任何消费者的根级悬空软链，以免引入第四种名字。若短期无法提供 lock，请把 arch_select.bzl:8 的 load 与 WORKSPACE:80-81 的 pip_ascend_torch_install_deps() 一并延后，不要把新平台的未完成状态扩散成全平台阻塞。合入前请至少用非 ascend 配置跑一次 bazel query //rtp_llm/... 验证 loading 通过。顺带修正 deps/pip.bzl:90 混入的 TAB 缩进。

Checklist: [6.1] Commit 原子、message 与行为匹配

[LLLLKKKK]

[P1] Ascend 上事件 record 退化为空操作，transfer_done_event 的 D2H 完成不变量被静默破坏

graphRecordEvent/graphBlockEvent（:140-156）在 USING_ASCEND 下仅 (void)event; (void)stream;——既不 record 也不插入任何同步，真实语义是"丢弃跨流依赖"，而 :138-139 注释称 "degrade to stream-synchronous semantics"，描述与行为不符。这两个是 inline 头文件实现，models/BUILD:186-187 无条件链接 cuda_graph_base/cuda_graph_hdrs_lib，故在 Ascend 上是活代码。调用点未按平台隔离：SpeculativeSampler.cc:255-257 先做两次 to(kCPU, true) 非阻塞 D2H 再 record；消费侧 MtpBatchStreamProcessor.cc:1278 注释 "wait for the transfer to complete" 后无条件 synchronize()，:1279-1280 立即读 `accep...

建议: 不要让 record/block 在新平台静默降级：Ascend 分支应基于 aclrtEvent / c10_npu 事件实现真实 record/block；若事件语义暂时无法保证，请在该 shim 内改为显式 stream 同步（而非空操作），使 synchronize() 的调用方契约在所有平台一致。若本阶段不支持 Ascend 投机解码，请在引擎初始化处对 propose_params_/sp_type 增加 fail-fast 校验。同时把注释改为"Ascend 依赖单流执行顺序，跨流事件依赖被丢弃"，避免后续开启 useStreamAsync() 时静默产生数据竞争。另 makeGraphEvent 的 Ascend 分支与 #else 分支（GRAPH_DEVICE_TYPE 在 :10 即为同一值）完全等价，属冗余可删。

[LLLLKKKK]

[P1] Ascend 依赖清单缺少 torchvision，而 //rtp_llm:torchvision 是无条件依赖

rtp_llm/BUILD:101 把 torchvision 传入 requirement([...])，:266 又在 multimodal py_library 的 deps 中以 ":torchvision" 无条件（不在任何 select 内，select 从 :269 才开始）依赖它，而 :303 使 //rtp_llm:models 依赖 :multimodal。经 arch_select.bzl:52,61 处理后，--config=ascend 下该目标会解析到 requirement_ascend("torchvision")。但 deps/requirements_ascend_x86.txt 只列出 torch/torch_npu/numpy 三项，已 grep 确认 deps/requirements_base.txt 无 torchvision，且 torchvision 也不在 _ascend_excluded（arch_select.bzl:26-36）中。对比 rocm：...

建议: 二选一并保持自洽：若 ascend 运行时需要 torchvision（多模态预处理路径会 import），请在 deps/requirements_ascend_x86.txt 补入与 torch 2.9.0 匹配的版本；若 ascend 阶段性不支持多模态，请把 "torchvision" 加入 _ascend_excluded，并确认 rtp_llm/BUILD:266 的消费方在 torchvision 缺失时不会在 import 期崩溃。合入前建议跑一次 bazel query 'deps(//rtp_llm:models)' --config=ascend，把这类 select 分支缺包问题在分析期暴露。

[LLLLKKKK]

AI Code Review - PR #1349 (non-blocking suggestions)

19 条 P2/P3 建议，不阻塞合并。阻塞判定与完整摘要见上一条 review。

[LLLLKKKK]

[P2] ExecOps.h 被 cpp 核心层新增引用但两个 Bazel 目标未声明该依赖，且属反向跨层依赖

两处新增 #include "rtp_llm/models_py/bindings/core/ExecOps.h" 的目标都缺声明：（1）cuda_graph_runner.h:3 新增该 include（置于标准库之前）并在 inline 构造函数中调用 getTorchCudaDevice()，但导出该头的 cuda_graph_hdrs_lib（cuda_graph/BUILD:45-57）deps 只有 //rtp_llm/models_py/bindings:op_defs，没有 exec_ops_hdr（对比同文件 cuda_graph_impl 在 :78 已显式声明）。（2）CudaCopyUtil.cc:5 新增该 include，但 cuda_copy_util（.../tcp/BUILD:5-15，本 PR 未修改）deps 仅 //:rtp_compute_ops（linkshared=1 cc_binary，不传播编译上下文）、no_block_copy_hdr、core_utils。当前能编译...

建议: 分别在 cuda_graph_hdrs_lib 与 cuda_copy_util 的 deps 中补上 //rtp_llm/models_py/bindings/core:exec_ops_hdr（该 header-only 目标已存在并在 pybind/BUILD:139,144 被同样引用），让头文件目标自带完整依赖闭包，否则一旦切到沙箱/远端执行或启用 layering_check 就会编译失败。更彻底的做法是把 getTorchCudaDevice() 下沉到 rtp_llm/cpp/utils 之类的中立层——本 PR 已让该函数成为全仓设备入口，放在 models_py/bindings/core 会让层次边界持续劣化。

Checklist: [6.1] 依赖方向：无循环依赖/跨层惊喜

[LLLLKKKK]

[P2] BlockInfo::is_cuda 语义被扩为"设备内存"，但字段名、注释与部分消费者未同步

MemoryLayoutStrategy.cc:307（dev.is_cuda() || dev.is_privateuseone()）与 LayerBlockConverterImpl.h:35-36（同时纳入 BufferKey 去重键）把 BlockInfo::is_cuda 扩展为"设备内存"语义。但 BlockInfo.h:14-16 未随本 PR 修改，注释仍写"Torch device of the backing storage (CPU/CUDA)"，字段名与含义自此不一致。消费者亦未全部同步：KVCacheManager.cc:765 仍 dst_block.is_cuda ? torch::kCUDA : torch::kCPU，在 Ascend 上会用 NPU 指针构造 CUDA 张量（该函数为 writeKVBlockForTest，见 :756 报错文案，影响面限于测试）；model_rpc/DecodeRpcServer.cc:1007,1153 仍把该布尔作为 GPU 标志传入 `load_layer_cac...

建议: 把字段重命名为 is_device，或改为直接携带 MemoryType（与本 PR 新增的 MEMORY_NPU、torchDeviceToMemoryType 对齐），并全仓搜索消费者统一改造、同步更新 BlockInfo.h 的注释；至少修掉 KVCacheManager.cc:765，并明确 DecodeRpcServer.cc 两处在 NPU 下传给 cache store 的期望语义。建议补一条 NPU/CPU 混合 BlockInfo 的单测锁定新语义。

Checklist: [I] 删除或重命名内部 file、registry entry、model name、metric enum、op binding、plugin symbol 时，必须全仓搜索消费者，并提供替代实现、迁移说明或 smoke 覆盖；只有暴露到 HTTP/RPC/config/persisted format 时才按外部兼容性处理

[LLLLKKKK]

[P2] NPU 保 ACL format 的 from_blob 包装无法复用，KV 拷贝路径仍直接用 torch::from_blob

MemoryLayoutStrategy.cc:3-8 明确记录"torch::from_blob 会丢掉 NPU tensor 的 ACL format，导致下游 FIA/scatter attention kernel 读取出错"，并在 :17-23 用 at_npu::native::from_blob 做了 makeCacheTypedBlob 包装。但该函数位于 .cc 的匿名 namespace（:15-24），天然无法被外部复用；同一批改动中其它以 NPU 裸指针建 tensor 的位置仍直接用 torch::from_blob：CudaCopyUtil.cc:11-14 的 wrapRawPtr（调用方 :31/:56 已传入 getTorchCudaDevice()，Ascend 下即 kPrivateUse1）、KVCacheMemoryConnector.cc:2757 附近。同一陷阱只在一处被防住，且未说明为何纯字节视图可豁免该约束。

建议: 把 makeCacheTypedBlob 提到共享头（如 models_py/bindings/core/torch_utils/TypeConvert.h），让所有以 NPU 裸指针构造 tensor 的位置统一走它。若确认 execNoBlockCopy 的纯 uint8 字节视图不依赖 ACL format tag，请在 wrapRawPtr 与 KVCacheMemoryConnector 两处补一行注释说明豁免理由与边界条件，避免后续读者误用或误删。

[LLLLKKKK]

📍 实际位置 rtp_llm/cpp/cache/KVCacheAllocator.cc:294（不在 diff 展示范围内，就近挂载）

[P2] KVCacheAllocator 新增未使用的 include，且未复用同 PR 新增的 allocationTypeToMemoryType

本 PR 在 KVCacheAllocator.cc:13 新增 #include ".../torch_utils/TypeConvert.h"，但 grep 确认全文件对其导出符号（allocationTypeToMemoryType/torchDeviceToMemoryType/memoryTypeToTorchDevice/dataTypeToTorchType）零引用。同时 :294 仍硬编码 allocation_type_ == AllocationType::DEVICE ? rtp_llm::MEMORY_GPU : rtp_llm::MEMORY_CPU——正是该头文件 TypeConvert.h:81-90 中 allocationTypeToMemoryType() 要解决的映射（其 Ascend 分支返回 MEMORY_NPU）。Ascend 下设备池因此被标记为 MEMORY_GPU，与同 PR 新引入的 MEMORY_NPU 建模自相矛盾，目前仅因 isDeviceMemory() 同时接受两者才未出错。

建议: 将 :294 替换为 allocationTypeToMemoryType(allocation_type_)（这样新增的 include 也有了实际用途，且 Ascend 下内存类型标注正确）；否则应删除该未使用的 include，避免留下"看起来已改造"的假象。

[LLLLKKKK]

[P2] Ascend 侧新增分支基本零测试覆盖，测试改造未接入承载 target

C++ 侧：本 PR 改 4 个测试文件以支持 PrivateUse1（MtpExecutorTest.cc:474,479、MtpBatchStreamProcessorTest.cc:144,148 增加 is_privateuseone()；MockEngine.h:27、NormalBatchStreamProcessorTest.cc:8 引入 DEVICE_TYPE），但承载它们的 normal_engine/test/BUILD 与 speculative/test/BUILD 均不在 diff_paths（diff 中相关测试 BUILD 仅 rtp_llm/cpp/utils/test/BUILD）：后者经核查仍无条件依赖 @local_config_cuda//cuda:cuda_headers/cudart（:63-64,89-90）、env 硬编码 TEST_USING_DEVICE: "CUDA"（:40,67,93）、exec_properties 指定 H20，故改造在 Ascend 上无法构建...

建议: 参照本 PR 已对 rtp_llm/cpp/utils/test/BUILD:30-33 采用的做法，把两个 test BUILD 的 env/deps/exec_properties 改为 select({"@//:using_ascend": ..., "//conditions:default": ...})，让相关用例在 Ascend 下可构建可执行；若本期不打算在 Ascend 跑，建议暂缓这 4 个文件的设备无关化改动，避免留下无覆盖代码。Python 侧在 server_config_setup_test.py 补用例：unittest.mock.patch 让 torch.cuda.is_available() 返回 False 并向 sys.modules 注入伪 torch_npu 与 torch.npu，参数化断言（a）seq_size_per_block 被置 128 且用户显式配置时不被覆盖，（b）local_world_size == min(npu_count, world_size)，（c）torch_npu 缺失时回落原有行为不抛异常（注意 mock target 是使用处）。并在 PR 描述中说明 Ascend 当前如何验证（机型、CANN 版本、bazel 命令与结果）与接入 CI 的计划。

Checklist: [6.1] 边界 case 覆盖（空、单元素、最大值）

[LLLLKKKK]

[P2] _ascend_excluded 排除名单不完整，ascend 配置下通配构建会解析到不存在的 pip 包

arch_select.bzl:52 对不在 _ascend_excluded（:26-36，仅 9 项）中的每个名字生成 requirement_ascend(name)。但 requirement() 的调用点 rtp_llm/BUILD:81-145 还包含 tilelang/flash-mla/fast-hadamard-transform(:37-41)、deep_gemm/deep_ep(:43-46)、rtp-kernel(:48-50)、flash_attn/flash-attn-3(:52-55) 以及 flashinfer_with_cache 各项，已 grep 确认这些既不在 deps/requirements_base.txt 也不在 deps/requirements_ascend_x86.txt 中。虽然消费者 select 让 ascend 目标不依赖它们，但 requirement() 仍会为每个名字生成 py_library 并在其 ascend 分支引用缺失包，`baze...

建议: 以 ascend lock 的实际包集合为准补全排除名单，或改为"白名单 + 默认解析为空"，避免每加一个平台都要重复维护一份 denylist。同时把一条 ascend 配置下的 analysis 级校验纳入 CI（例如 bazel query 'deps(//rtp_llm:models)' --config=ascend），否则这类缺口只有在真正跑 ascend 通配构建时才会暴露。

Checklist: [6.1] 新逻辑有聚焦单测 + 相关集成/smoke 测试

[LLLLKKKK]

[P2] aclnn 自定义算子目标结构不规范、SoC 硬编码、运行期路径依赖相对 bazel-bin 与私有 conda 环境

:30-33 把 genrule :build_aclnn_ops 放进 cc_library.deps，而 genrule 不提供 CcInfo，存在 analysis 期报错风险；即使当前 Bazel 容忍，它与 data:24-29 已列出的同一批 .so 完全重复。target_compatible_with 只加在 genrule 上（:15-18），cc_library 自身无平台约束，误引用时报错会指向平台不兼容而非依赖写错。cmd:12-13 把 SoC 目标硬编码为 ascend950，在 910B 等芯片上会产出不匹配的算子包且无构建期校验。产物仅通过相对路径暴露给运行期：.bazelrc:417 的 LD_LIBRARY_PATH 与 :421 的 ASCEND_CUSTOM_OPP_PATH 都写 bazel-bin/3rdparty/aclnn_custom_ops/...，但测试进程 cwd 为 runfiles 目录而非 workspace 根；:417-418 还把 torch/torch_npu 库...

建议: 把 :build_aclnn_ops 从 deps 移除（产物已在 data，genrule 会被正常触发；如需聚合请用 filegroup），消除重复的输出路径列表，并给 cc_library 补上与 genrule 相同的 target_compatible_with；把 ascend950 提取为可配置项（--define 或 select）并在非目标 SoC 上 fail-fast；自定义算子路径改由 target 的 data + runfiles 解析（$(location) 展开或 runfiles API），不要在配置文件里写相对 bazel-bin 路径；PYTHONHOME 与 torch/torch_npu 库路径统一到仓库标准解释器位置（/opt/conda310）或改为 action_env 注入，避免绑定单机私有环境。

Checklist: [6.1] 兼容性：外部 HTTP/RPC API、持久数据、配置、环境迁移安全；[6.1] 错误语义：fail-fast/retry/fallback/silent 行为显式

[LLLLKKKK]

[P3] 多处新增的 using_ascend select 空分支与 default 等价，且 cuda_graph_impl 收敛顺带改变了 cpu/arm 链接图

五处新增的 "@//:using_ascend": [] 与同一 select 已有的 "//conditions:default": [] 语义完全相同、不改变任何构建结果：rtp_llm/cpp/utils/BUILD:216、embedding_engine/BUILD:28、models/BUILD:95、models/BUILD:191、normal_engine/BUILD:41；空分支会让后续读者误以为 Ascend 有特殊处理。同时 diff 显示 models/BUILD 删除了原本无条件的 cuda_graph_impl（diff 第 2077 行 -），改为只在 using_cuda/using_rocm 下提供，纯 cpu/arm 配置也一并落入 default 丢掉该依赖，超出 PR 声明的 Ascend 范围且无注释说明（核查确认当前不产生未解析符号，但 CI 不构建 cpu/arm，该变化无任何验证）。

建议: 删除与 //conditions:default 等价的空分支；若刻意保留是为"显式声明该平台已评估过"，请统一加一行简短注释说明意图。models/BUILD:191 若意图仅排除 Ascend，请写成 select({"@//:using_ascend": [], "//conditions:default": ["//rtp_llm/cpp/cuda_graph:cuda_graph_impl"]}) 以保持 cpu/arm 行为不变；若确实要一并收敛，请在 PR 描述中说明并至少跑一次 bazel build //rtp_llm/cpp/models:models --config=cpu 验证链接通过。

Checklist: [6.1] KISS/YAGNI：无投机性抽象；[6.1] 分布式/跨平台变更有对应覆盖

[LLLLKKKK]

AI Code Review - PR #1349

Status: BLOCKING

Summary: P0/2 · P1/5 · P2/16 · P3/3

Reviewed: commit 547522b35454 · 2026-08-29 00:35 UTC+8

Blocking Issues

P0

• ROCm 下 graphCaptureBegin 不再调用 capture_begin，既有平台的图捕获被静默破坏 @ rtp_llm/cpp/cuda_graph/cuda_graph_device_shims.cc:212
  • 建议：改为三分支：#if USING_CUDA 调 graph.capture_begin(pool)；新增 #elif USING_ROCM 恢复 ROCm 原有的 capture_begin() 调用；仅在真正不支持的分支（Ascend/CPU）做空操作或 fail-fast。同时补一条 ROCm 侧图捕获+replay 的最小回归用例，避免同类 #else 合并再次吞掉既有平台行为；仓库已有 rocm 配置，可在 CI 中实际验证。
• Ascend pip lock 文件未随 PR 提交且存在三种不一致命名，顶层无条件 load 阻塞所有平台 @ deps/pip.bzl:89
  • 建议：提交由 bazel run //deps:requirements_ascend_x86.update 生成的 deps/requirements_lock_ascend_x86.txt 并入库，统一 deps/pip.bzl:89、deps/BUILD:111 与根 symlink 三处命名（建议与 src 对齐用 _x86 后缀），并像 requirements_lock_rocm.txt 一样加入 deps/BUILD 的 exports_files。合并前请在不叠加 --override_repository @rtp_deps 的树上跑一次 bazel query //rtp_llm/...（任一非 Ascend config）确认 WORKSPACE 可正常求值——内部 override 会掩盖此问题使开源侧单向破损。若短期无法提供 lock 文件，应把 pip_ascend_torch 与两处 load 一并延后引入。另请修正 deps/pip.bzl:90 的制表符缩进。

P1

• base_model.py 为 ROCm 返回 hip:N 设备串，与全栈 torch.cuda 语义冲突 @ rtp_llm/models/base_model.py:175
  • 建议：删除 DeviceType.ROCm 分支，让 ROCm 继续返回 cuda:{local_rank}，只保留 Ascend 的 npu:{local_rank}。若确认 ROCm 需改用 hip:，请拆成独立 PR，并先在 ROCm 环境验证 torch.empty(device="hip:0") 与 torch.cuda.Stream(device=torch.device("hip:0")) 均可用。同时补一条覆盖 _get_device_str() 三平台返回值的单测，并考虑把设备串推导下沉到 rtp_llm/device/ 公共函数，避免与 weight_manager.py:122 的 Stream 类型分派逻辑各自演化。
• Ascend 上事件 record 退化为空操作，transfer_done_event 的 D2H 完成不变量被静默破坏 @ rtp_llm/cpp/cuda_graph/cuda_graph_device_shims.h:140
  • 建议：Ascend 分支不要留空：graphRecordEvent 应调用 event.record(<当前 NPU stream>)（FusedCopyOp.cc:49 已在用 c10_npu::getCurrentNPUStream()），或在无法建立事件时退化为该 stream 的 synchronize()，使消费侧 synchronize() 的契约成立；同时把 .h:138-139 注释改成与实现一致的描述（「Ascend 当前只使用默认流，暂不提供跨流定序」）。若短期无法实现事件语义，请把 accept_*_cpu 的 D2H 改为 non_blocking=false 并在 shim 头部注明该退化，并补一条 Ascend MTP 解码的 smoke 用例锁定 accept 序列。
• Ascend 下 MtpExecutor 的 AsyncRunner::sync 被整段编译掉，丢失 host join 与异常回抛 @ rtp_llm/cpp/normal_engine/speculative/MtpExecutor.cc:1404
  • 建议：不要在调用点屏蔽 sync()。对齐 NormalExecutor 保留 sync() 并替换 wait stream，或把 AsyncRunner::sync() 内的 event_.block(wait_stream) 改为 cuda_graph::graphBlockEvent(event_, wait_stream)（Ascend 下即 no-op），保留 cv_done_.wait 与异常重抛，然后移除全部 8 处 #if !USING_ASCEND。若 Ascend 确实暂不支持这些异步路径，应在入口处显式关闭该特性（让 useStreamAsync()/useAsyncPrepare() 在 Ascend 返回 false 并打印一次告警），使 launch() 与 sync() 成对不执行。另注意 NormalExecutor.cc:81/206/219/323 把 PrivateUse1 的 device index 写死为 0，多卡 TP 应取 local_rank，不要照抄。
• start_backend_server 的 Ascend 设备数分支不可达，多卡 NPU 启动静默退化为单 rank @ rtp_llm/start_backend_server.py:438
  • 建议：把早退条件改为「既无 CUDA 也无 NPU」（建议复用 rtp_llm.device.device_type.is_ascend()），仅当设备数为 0 时才走 CPU 单进程；把设备计数抽成模块级 _device_count() helper 供 :126/:148/:438 共用；_create_rank_processes 按设备类型写入对应可见设备环境变量。建议补一条以 device_type 为 patch 点、mock「CUDA 不可用 + NPU 可用」的单测断言进入 multi_rank_start。若本期只支持 Ascend 单卡，则应删除 :438/440/444/456 的 NPU 分支，并在早退处显式记录「Ascend 暂不支持多卡」。
• Ascend 依赖链全程无完整性校验：空 sha256、关闭 hash、第三方个人仓库缩写 commit @ deps/http.bzl:84
  • 建议：为两个 http_archive 填入真实 sha256；在 requirements_ascend_x86.txt 两条直链后补 #sha256=<digest>；复核 deps/BUILD:112 的 generate_hashes = False 是否必需，若受上游 metadata 限制无法开启请就地加注释说明。按 ROCm 做法把 torch_npu wheel 转存到团队自有对象存储并在 urls 保留「自有镜像 + 上游」双源。aclnn 算子源码请镜像到组织内可控仓库、改用带 sha256 的 http_archive（或至少用 40 位完整 SHA 并补 shallow_since），并迁到 deps/git.bzl 的 git_deps() 使内部构建可通过 --override_repository 切换来源；同时评估能否去掉 tags = ["local"] 恢复沙箱执行。

Non-blocking Suggestions

P2

• Ascend 专属链接选项 -Bsymbolic-functions 被无条件施加到全平台首个加载的 .so @ BUILD:152
  • 建议：用 select({"@//:using_ascend": ["-Wl,-Bsymbolic-functions"], "//conditions:default": []})（或 if_ascend(...)）门控该 linkopt，使非 Ascend 平台链接语义保持不变；若确认 CUDA/ROCm 也需要，请在 PR description 给出这三个 .so 之间符号共享不受影响的验证结论。同时补注释说明为何只需给 th_transformer_config 加，并修正拼写 Fox → Fix。
• ExecOps.h 被 cpp 核心层新增引用但两个 Bazel 目标未声明该依赖，且属反向跨层依赖 @ rtp_llm/cpp/cuda_graph/cuda_graph_runner.h:3
  • 建议：在 cuda_graph_hdrs_lib.deps 与 cuda_copy_util.deps 中显式加入 //rtp_llm/models_py/bindings/core:exec_ops_hdr（与 th_transformer_lib 的做法一致），并在 Ascend 与 CUDA 两个 config 下各验证一次这两个 target 的构建。更好的做法是把 getTorchCudaDevice() 这类纯设备查询下沉到 rtp_llm/cpp/utils 的小头文件，避免引擎层/cache 层头文件反向依赖 models_py/bindings 层。
• NPU 保 ACL format 的 from_blob 包装无法复用，KV 拷贝路径仍直接用 torch::from_blob @ rtp_llm/cpp/cache/MemoryLayoutStrategy.cc:17
  • 建议：把 makeCacheTypedBlob() 提升为共享工具（例如放到 rtp_llm/cpp/cache/ 公共头或 models_py/bindings/core/torch_utils/TypeConvert.h），在所有包装设备侧裸指针的位置统一使用：CudaCopyUtil.cc:13、KVCacheMemoryConnector.cc:2757/2760。若这两条 connector 路径在 Ascend 上确实不启用，请显式 fail-fast 或加注释说明未支持，而不是把设备改成 NPU 却保留 CUDA 时代的包装方式；若确认纯 memcpy 语义与 format tag 无关，请在 makeCacheTypedBlob 注释里写清适用边界。（BlockPool.cc:272 的 from_blob 位于 #if USING_CUDA 内，不受影响。）
• BlockInfo::is_cuda 语义被扩为"设备内存"，但字段名、注释与部分消费者未同步 @ rtp_llm/cpp/cache/connector/p2p/LayerBlockConverterImpl.h:35
  • 建议：抽一个共享 helper（如 isDeviceTensor(const torch::Tensor&)）放入 rtp_llm/cpp/utils 供全部调用点使用，并把 :36 的重复求值合并为一个局部 const bool on_device；把 BlockInfo::is_cuda 更名为 is_device/is_accelerator 并同步 BlockInfo.h:14 注释与 KVCacheMemoryConnector.cc:2751 的形参命名。同时确认 RDMA regUserMr 在收到 NPU 指针且 gpu=true 时行为可接受，并为以 is_cuda 为条件直调 CUDA API 的测试辅助函数补平台断言。
• BlockPool::where() 的 Ascend 分支丢失 PINNED 语义，NPU 设备池被按 host 内存注册 MR @ rtp_llm/cpp/cache/BlockPool.cc:751
  • 建议：Ascend 分支保留 pinned 判定，例如统一为 if (is_privateuseone() || is_cuda()) return <device type>; return is_pinned() ? MEMORY_CPU_PINNED : MEMORY_CPU;（可直接复用同 PR 新增的 TypeConvert.h:92-96 torchDeviceToMemoryType()）。把 :580/:614 的判定改为「是否设备内存」语义——将 OpData.cc:80-82 的 isDeviceMemory()（MEMORY_GPU || MEMORY_NPU）从文件内 static 提升为公共 helper 并在此复用，避免 MEMORY_NPU 在各消费点被逐个漏判；KVCacheAllocator.cc:294 硬编码的 MEMORY_GPU 也应改用同 PR 新增的 allocationTypeToMemoryType(allocation_type_)。若 Ascend 暂不支持 cache_store，请在该路径显式 fail-fast。
• PyWrappedModel 同一函数内多种设备判定口径并存，NPU 张量会走到非法 pin_memory 路径 @ rtp_llm/cpp/models/PyWrappedModel.cc:170
  • 建议：抽出统一的 isDeviceResident(const torch::Tensor&) 谓词（与上一条的 isDeviceTensor 合并为同一 helper）并在该函数内全部替换，包括 normalize_i32、to_device_i32、:195-205 回落循环、:211-212、:254；同时修正 :196-199 的注释使之与实际匹配的谓词一致。若 Ascend 上确实需要 host 回落，请确认 pin_memory() 对 NPU 后端可用，否则在该分支前显式排除 PrivateUse1；若 Ascend 暂不走 MTP device-state 快路径，请回退 :298 的改动或加注释说明该分支目前不可达。
• MtpExecutor 设备判定仅部分转换，Ascend 下投机路径既不支持也不显式拒绝 @ rtp_llm/cpp/normal_engine/speculative/MtpExecutor.cc:280
  • 建议：将所有设备构造统一收敛到 getTorchCudaDevice()（含 ensureModelInputsOnCuda 三处、checkModelInputsOnCuda 的 is_cuda() 判定、TP≠0 分支，并删除 :2218-2219 的局部 #if USING_ASCEND 改用同一 helper），并全局 grep 确认无残留。若本 PR 不打算让 MTP 在 Ascend 上工作，请在 PR 描述与代码中明确声明该限制，并在 MtpExecutor 构造/入口处 fail-fast 报出可操作错误，避免留下「能编译但运行必崩」的半移植状态。
• Ascend 不链接 cuda_graph_impl，shims.cc 全部 Ascend 分支为未经验证的死代码 @ rtp_llm/cpp/cuda_graph/cuda_graph_device_shims.cc:143
  • 建议：二选一：要么删除 cuda_graph_device_shims.cc 中的 USING_ASCEND 分支与 .h:175 的 void* 重载声明（仅保留 header 中确实被 models/normal_engine 使用的 inline shim）；要么在本 PR 内真正打通 Ascend 的 cuda_graph_impl（补 BUILD select 的 ascend 依赖、修正流获取、统一 graphCaptureBegin 签名），并加入 Ascend 构建目标以获得编译验证。graphMemcpyAsync 的 Ascend 实现应改用 aclrtMemcpyAsync 并按 kind 映射 ACL_MEMCPY_*（方向判定可复用 FusedCopyOp.cc 已有逻辑并提取为共享 helper）；graphCheck 应比较 ACL_SUCCESS 并用 RTP_LLM_CHECK_WITH_INFO 输出 call_expr 与 aclGetRecentErrMsg()，而不是用 host memcpy 与静默返回伪装设备语义。
• getTorchCudaDevice 返回不带 index 的设备，叠加 Ascend setDevice 空实现，多卡落卡与流排序均无保证 @ rtp_llm/models_py/bindings/core/ExecOps.cc:599
  • 建议：给 getTorchCudaDevice() 增加可选 device_index 参数（默认沿用当前设备），并在需要确定落卡的位置（AsyncRunner 构造、KV cache 分配、MoE stats 分配）显式传入 local_rank；为 cuda_graph::setDevice 的 Ascend 分支实现 c10_npu::set_device(rank)（或 aclrtSetDevice），graphGetStreamFromPool/graphGetCurrentStream 改为基于 c10_npu::getCurrentNPUStream() 实现，暂不实现时至少加一次性 WARNING 说明多卡/流隔离不受支持。建议补一条 TP>1 的 Ascend 端到端用例确认各 rank 张量落在预期设备上。
• DEVICE_TYPE 全局宏在测试公共头文件中定义并与 CANN/第三方头文件语义冲突 @ rtp_llm/cpp/normal_engine/test/MockEngine.h:27
  • 建议：删除该宏，改为具名的 inline 常量或函数（如 static const torch::Device kTestDevice = getTorchCudaDevice();）以规避与 CANN 及第三方头文件的冲突，并去掉 .cc 中的重复定义；同时建议把 getTorchCudaDevice() 改名为设备中立的名字（如 getTorchAcceleratorDevice()），因其在 Ascend 上返回 PrivateUse1，函数名与契约不符。另建议 normal_engine/test 的 mock_engine 显式依赖 //rtp_llm/models_py/bindings/core:exec_ops_hdr 而非依赖传递头文件。
• torch_npu 探测逻辑多处内联重复，未复用同 PR 新增的 is_ascend()，且降级路径无日志 @ rtp_llm/config/server_config_setup.py:268
  • 建议：统一 from rtp_llm.device.device_type import is_ascend（直接从叶子模块导入，避免 rtp_llm/device/__init__.py 连带加载 .so），把三处内联探测替换为 elif is_ascend():，并在 device_type.py 内补 device_count()/set_device(local_rank) 两个平台无关封装，把 import torch_npu 的副作用集中在一处；同步整理 loader.py:595。异常语义显式化：仅 import torch_npu 处于 try 内并捕获 (ImportError, OSError)，降级时留一条 logging.warning 记录原因与实际生效值；:540 的 set_device 属必须成功的启动步骤，判定为 Ascend 后不应静默跳过。另建议把 setup_cuda_device_and_accl_env 更名为设备无关的名字。
• _get_cuda_device_list 用 dir() 成员检测做控制流，条件恒为 False @ rtp_llm/start_backend_server.py:148
  • 建议：删除 _dev_count if "_dev_count" in dir() else 前缀，改为调用与 :126/:438 共享的模块级 _device_count() helper（内部基于 device_type.is_ascend() 判定，避免裸触碰 torch.npu）；需要跨函数共享数值时用参数传递，不要依赖作用域内省。补充覆盖「CUDA 不可用 + NPU 可用」与「两者都不可用」两条分支的单测。
• Ascend 侧新增分支基本零测试覆盖，测试改造未接入承载 target @ rtp_llm/cpp/normal_engine/speculative/test/MtpExecutorTest.cc:474
  • 建议：把这两个 test 包的 env/deps/exec_properties 迁到 device_test_envs()/device_impl_target()（而不是手写 select，utils/test/BUILD:30-33 目前也是手写），并把文件内剩余设备写法一次性统一到 getTorchCudaDevice() 与统一的 isDeviceTensor()，给出 Ascend 实跑结果。若本期无法在 CI 验证，请在 PR 描述中说明是为后续 Phase 预留，并给出计划在哪个 stage 打开 --config=ascend，避免留下「看起来已支持、实际未支持」的半成品状态。
• Ascend 默认 seq_size_per_block=128 缺少单测，且探测方式与相邻平台默认值不一致 @ rtp_llm/config/server_config_setup.py:437
  • 建议：补一组参数化单测，mock 收敛后的 is_ascend() 与设备计数，断言四种决议结果（kfd→16、alixpu→256、ascend→128、无匹配→64）以及「用户已显式设置非 0 值时不被覆盖」这一边界，另加一条 Ascend 下 local_world_size 推导用例，使默认值不再取决于测试机的 site-packages 状态。同时确认在 Ascend 上于 spawn 前调用 torch.npu.is_available() 不影响子进程设备初始化；若存在风险，改用与相邻分支一致的无副作用探测，把运行期初始化推迟到 backend 进程内。
• aclnn 自定义算子目标结构不规范、SoC 硬编码、测试环境 Python 前缀与仓库统一运行时冲突 @ 3rdparty/aclnn_custom_ops/BUILD:3
  • 建议：1) 若 opp 是目录产物，改用 declare_directory 风格的自定义 rule，或把 opp 包打成单个 tar 作为唯一 out；2) 给 genrule 追加 no-remote-cache（或把 CANN 版本纳入输入）；3) 把 ascend950 提到 --define 或 config_setting；4) 从 cc_library.deps 移除 :build_aclnn_ops；5) ASCEND_CUSTOM_OPP_PATH 改为 runfiles 相对路径并验证 libcust_opapi.so 在测试中确实可被加载；6) ascend 统一使用与 --python_top 一致的 Python 运行时（或说明为何必须独立环境并在镜像中固化）；7) 逐项为 -Wno-error=*/-Wno-* 加注释说明触发文件与清理计划。
• ascend 分支把 cache_store 与 stack_tracer 引入 compute .so，依赖分层不对称且无说明 @ rtp_llm/cpp/pybind/BUILD:94
  • 建议：定位 ascend 下实际缺失的符号来源，把依赖补在真正引用它的目标上（例如 ascend_host_utils 或 exec_ctx_ops），而不是在 pybind 汇聚层补洞；若确实无法避免，请像 :67-70 那样注明是哪个符号、来自哪个 TU。同时把 exec_ctx_rocm 泛化并改名（如 exec_ctx_combined），让 Ascend/CUDA/ROCm 统一走同一 deps 形状，由该目标自己维护 ExecOps.cc 的依赖清单；kv_cache_utils 若只是 ascend 需要也应放进 select 分支。

P3

• 条件编译结构冗余与死代码，ROCm 的关键说明被埋入不可达分支 @ rtp_llm/cpp/models/Sampler.cc:126
  • 建议：Sampler.cc 改为平铺的 #if USING_CUDA ... #elif USING_ROCM ... #else 三分支，把 ROCm 说明放回可达分支并把 #else 注释改为准确描述 Ascend/CPU 的阻塞拷贝；删除 ComputeInit.cc:11-12 的空块；makeGraphEvent() 统一 return torch::Event(GRAPH_DEVICE_TYPE); 并给 graphGetStreamFromPool 的 Ascend 分支补 (void)is_high_priority;；SpecLogitsVerifyRunner.cc 二选一（去掉外层 #if USING_CUDA 或恢复直接 record/block）；cuda_host_utils.h 的 Ascend 注释改为明确「该头不支持 Ascend，请使用 ascend_host_utils.h」；删除与 default 等价的 using_ascend 空分支。
• _ascend_excluded 与 whl_deps 的 Ascend 分支不符合既有约定 @ arch_config/arch_select.bzl:26
  • 建议：改名为 _ASCEND_EXCLUDED 并补注释说明排除依据（哪些包在 Ascend lock 中不存在、为何不存在）；剔除 decord/xfastertransformer_devel/xfastertransformer_devel_icx 三项冗余条目，若排除确有意图则在注释中写明；:52 改为单空格。whl_deps() 的 Ascend 分支与 cuda13 保持一致改为直链形式，使产出 wheel 在无额外索引配置的环境也可安装；若团队约定依赖内部 index 解析 local version，请在该行补注释指明所需 index，并顺带记录 :108 using_cuda12 的相同遗留以便后续统一。lock 文件入库后请用 bazel query 在 ascend 配置下复核清单完整性。
• 新增 Python/C++ 代码未过仓库既有 black / isort / flake8 / clang-format 约定 @ rtp_llm/config/server_config_setup.py:439
  • 建议：提交前对改动文件跑 pre-commit run --files <changed>（black + isort + clang-format）并跑一遍 flake8；超长日志字符串拆成隐式拼接的多行；条件编译指令统一顶格或统一跟随所在块缩进。若按前述建议把探测收敛到 is_ascend()，import torch_npu 的 F401/E501 问题会一并消失；否则需补 # noqa: F401 并注明其用途是注册 torch.npu。

Checklist Findings (19 fail / 58 total)

General Principles Checklist

• [6.1] Architecture — 依赖方向：无循环依赖/跨层惊喜 → issue ascend 分支把 cache_store 与 stack_tracer 引入 compute .so，依赖分层不对称且无说明
pybind/BUILD:94-99 为 th_compute_lib（librtp_compute_ops.so）的 ascend 分支加入 //rtp_llm/cpp/disaggregate/cache_store:cache_store_interface 与 @havenask//aios/autil:stack_tracer，而 using_cuda12（:88-92）与 using_rocm（:93）分支都不需要。compute .so 的职责是计算算子注册，cache_store 属传输/存储层，这种只在单平台出现的反向依赖通常是为消除未定义符号临时补上的，且缺少像 :67-70 那样的注释说明。同时「ExecOps.cc 只进 compute .so」这一不变量现在有三种写法：CUDA 用 srcs = exec_ops_srcs；ROCm 用组合目标 exec_ctx_rocm；Ascend（:70 + :94-99）用第三种 exec_ops_srcs + exec_ctx_ops，因而必
• [6.1] Architecture — 兼容性：外部 HTTP/RPC API、持久数据、配置、环境迁移安全 → issue _ascend_excluded 与 whl_deps 的 Ascend 分支不符合既有约定
:26-36 的 _ascend_excluded 与紧邻的 _CUDA13_DEFERRED（:17-19）、_DSV4_PLATFORM_ONLY（:21-24）三点不一致：命名用小写而既有同类常量用 UPPER_SNAKE；完全没有注释说明每项为何排除，而既有两个常量都写清了原因；条目与实际解析结果不符——decord、xfastertransformer_devel、xfastertransformer_devel_icx 在 deps/requirements_base.txt:8/56/57 均带 platform_machine == "x86_64" 标记而 Ascend 目标是 x86_64，且 rtp_llm/BUILD:57-65 的 xft_dep 全部注释掉、:70/77 对 ascend 返回 []，这三个 target 并不在 Ascend 依赖图内。另 :52 ascend_deps = ... 用多空格对齐（buildifier 应为单空格）；:119 为 Ascend 输出
• [6.1] Architecture — 分层边界：新概念在正确层级，不泄漏内部 → issue ascend 分支把 cache_store 与 stack_tracer 引入 compute .so，依赖分层不对称且无说明
pybind/BUILD:94-99 为 th_compute_lib（librtp_compute_ops.so）的 ascend 分支加入 //rtp_llm/cpp/disaggregate/cache_store:cache_store_interface 与 @havenask//aios/autil:stack_tracer，而 using_cuda12（:88-92）与 using_rocm（:93）分支都不需要。compute .so 的职责是计算算子注册，cache_store 属传输/存储层，这种只在单平台出现的反向依赖通常是为消除未定义符号临时补上的，且缺少像 :67-70 那样的注释说明。同时「ExecOps.cc 只进 compute .so」这一不变量现在有三种写法：CUDA 用 srcs = exec_ops_srcs；ROCm 用组合目标 exec_ctx_rocm；Ascend（:70 + :94-99）用第三种 exec_ops_srcs + exec_ctx_ops，因而必
• [6.1] Architecture — 可观测性：日志/指标/超时可操作、非噪声 → issue torch_npu 探测逻辑多处内联重复，未复用同 PR 新增的 is_ascend()，且降级路径无日志
  同一份 try: import torch_npu / if torch.npu.is_available() 探测在本文件内联三次（:265-275 推导 local_world_size、:435-441 设 seq_size_per_block、:540-548 调 set_device），写法互不一致（一处 else 兜底、两处 except ImportError: pass）；model_loader/loader.py:594-600 是第四份，start_backend_server.py:126/148/438 又各自裸写。而本 PR 已在 device/device_type.py:17-49 提供规范实现 get_device_type()/is_ascend()（只依赖 os/enum/torch，叶子模块，无循环依赖风险）。异常语义也不安全：except ImportError 过窄（CANN 无法 dlopen 抛 OSError、版本错配抛 AttributeError 都会逃逸）；而
• [6.1] Architecture — 回滚路径：风险行为存在运维回滚手段 → issue Ascend 专属链接选项 -Bsymbolic-functions 被无条件施加到全平台首个加载的 .so
BUILD:150-153 给 cc_binary(name = "th_transformer_config")（即最先被 Python 加载的 libth_transformer_config.so）的 linkopts 无条件追加 -Wl,-Bsymbolic-functions，注释自述动机是「import torch_npu 会链接 CXXabi=1/0 混用的 CANN 库，导致 std::regex 内存问题」——纯 Ascend workaround，对 cuda12/cuda12_9/cuda13/rocm/arm/cpu 无收益。该选项使库内函数引用在链接期绑定自身副本，不再可被主程序或先加载的 .so 统一；而 th_transformer_config_lib 静态链接的 core_utils/config_modules 同时也被 librtp_compute_ops.so 链接，pybind/BUILD:134-136 明确说明项目依赖跨 .so 的 load-time 符号解析。同文件 `rtp_comput
• [6.1] Architecture — 状态不变量：创建/更新/失败/重试/回滚路径有效 → issue getTorchCudaDevice 返回不带 index 的设备，叠加 Ascend setDevice 空实现，多卡落卡与流排序均无保证
ExecOps.cc:597-603 的 getTorchCudaDevice() 在 Ascend 下返回不带 index 的 torch::Device(torch::kPrivateUse1)，依赖「线程当前设备」隐式确定落卡。但 cuda_graph_device_shims.h:79-90 的 setDevice(int rank) 在 Ascend 分支是 (void)rank; 空实现，而 AsyncRunner.cc:11-14 的 worker 线程正是靠 cuda_graph::setDevice(stream_.device_index()) 绑定设备——CANN 的 aclrtSetDevice 是线程局部的，worker 不绑定即无法确定落卡。同文件 toGraphStream/graphGetStreamFromPool/graphGetCurrentStream（:68-118）在 Ascend 一律返回 nullptr，GraphStreamGuard（:57-59）为空构造，流隔离也不存在；`
• [6.1] Architecture — 错误语义：fail-fast/retry/fallback/silent 行为显式 → issue aclnn 自定义算子目标结构不规范、SoC 硬编码、测试环境 Python 前缀与仓库统一运行时冲突
  genrule（:3-20）只声明 4 个 .so 作为 outs，其中 opp/op_impl/ai_core/tbe/op_tiling/liboptiling.so 只是 CANN custom opp 包中的一个文件；而 .bazelrc:421 的 ASCEND_CUSTOM_OPP_PATH="bazel-bin/3rdparty/aclnn_custom_ops/opp" 需要完整 opp 目录，其余文件属未声明输出，远端缓存命中时会得到不完整的 opp 包。genrule 的真实输入还包括本机 CANN 安装（未声明）且未加 no-remote-cache；cmd（:13）中 ascend950 为硬编码 SoC；cc_library.deps（:30-31）放了不提供 CcInfo 的 genrule，而 4 个 .so 已在 data 列出。.bazelrc:421 的 bazel-bin/... 是工作区相对路径而测试 CWD 为 runfiles 根；:417-418 的 `LD_LIBRARY
• [6.1] Quality — Commit 原子、message 与行为匹配 → issue Ascend pip lock 文件未随 PR 提交且存在三种不一致命名，顶层无条件 load 阻塞所有平台
deps/pip.bzl:88-89 的 pip_parse(name = "pip_ascend_torch") 要求 @rtp_deps//:requirements_lock_ascend_x86.txt，deps/BUILD:111 的 requirements_txt 同名；但已列举 deps/ 全量文件，只有输入文件 requirements_ascend_x86.txt，不存在任何 requirements_lock_ascend*.txt。PR 新增的仓库根 requirements_lock_ascend.txt 读取报 File does not exist（悬空 symlink），且缺 _x86 后缀，是第三个名字，全仓无 label 引用。而 WORKSPACE:80-81 与 arch_config/arch_select.bzl:8 都是与 --config 无关的顶层 eager load，故 cuda12/cuda12_9/cuda13/rocm/arm/cpu 全部配置都会在 WORKSPACE 求值
• [6.1] Quality — PR description 说明动机与设计 → issue Ascend pip lock 文件未随 PR 提交且存在三种不一致命名，顶层无条件 load 阻塞所有平台
deps/pip.bzl:88-89 的 pip_parse(name = "pip_ascend_torch") 要求 @rtp_deps//:requirements_lock_ascend_x86.txt，deps/BUILD:111 的 requirements_txt 同名；但已列举 deps/ 全量文件，只有输入文件 requirements_ascend_x86.txt，不存在任何 requirements_lock_ascend*.txt。PR 新增的仓库根 requirements_lock_ascend.txt 读取报 File does not exist（悬空 symlink），且缺 _x86 后缀，是第三个名字，全仓无 label 引用。而 WORKSPACE:80-81 与 arch_config/arch_select.bzl:8 都是与 --config 无关的顶层 eager load，故 cuda12/cuda12_9/cuda13/rocm/arm/cpu 全部配置都会在 WORKSPACE 求值
• [6.1] Quality — 逻辑变更未混入无关格式化 → issue 新增 Python/C++ 代码未过仓库既有 black / isort / flake8 / clang-format 约定
.flake8 仅 ignore W292、max-line-length=120，故 E501/F401 均生效：server_config_setup.py:439 的 logging 行约 146 字符、start_backend_server.py:148 约 168 字符均超限；server_config_setup.py:269/436/544 的 import torch_npu 属仅取副作用的未使用导入（F401），而本 PR 在 models_py/bindings/ascend/ops/tests/test_ascend_sampler.py:23 已给出正确的 # noqa: F401 写法，说明提交前未跑 pre-commit。.pre-commit-config.yaml 启用 black(88) 与 clang-format(120 列)：server_config_setup.py:437 与 start_backend_server.py:126 非 black 规范化结果（同函数 :421-424、`:4
• [6.1] Software Engineering — DRY：重复非平凡逻辑被抽取或显式复用 → issue _get_cuda_device_list 用 dir() 成员检测做控制流，条件恒为 False
:148 为 range(_dev_count if "_dev_count" in dir() else (torch.npu.device_count() if not torch.cuda.is_available() else torch.cuda.device_count()))。无参 dir() 返回当前函数的局部名字空间，而 _get_cuda_device_list 的局部变量只有 cuda_devices；_dev_count 是 _get_local_world_size(:126) 与 start_backend_server(:438) 的局部变量，不在本作用域内，条件恒为 False，作者跨函数复用局部变量的意图未实现。同时该表达式在「CUDA 不可用且未安装 torch_npu」时会抛 AttributeError: module 'torch' has no attribute 'npu'。该行约 168 字符，也超出 .flake8 的 max-line-length=120。
• [6.1] Software Engineering — KISS/YAGNI：无投机性抽象 → issue _ascend_excluded 与 whl_deps 的 Ascend 分支不符合既有约定
:26-36 的 _ascend_excluded 与紧邻的 _CUDA13_DEFERRED（:17-19）、_DSV4_PLATFORM_ONLY（:21-24）三点不一致：命名用小写而既有同类常量用 UPPER_SNAKE；完全没有注释说明每项为何排除，而既有两个常量都写清了原因；条目与实际解析结果不符——decord、xfastertransformer_devel、xfastertransformer_devel_icx 在 deps/requirements_base.txt:8/56/57 均带 platform_machine == "x86_64" 标记而 Ascend 目标是 x86_64，且 rtp_llm/BUILD:57-65 的 xft_dep 全部注释掉、:70/77 对 ascend 返回 []，这三个 target 并不在 Ascend 依赖图内。另 :52 ascend_deps = ... 用多空格对齐（buildifier 应为单空格）；:119 为 Ascend 输出
• [6.1] Software Engineering — SRP：模块/类职责单一 → issue ascend 分支把 cache_store 与 stack_tracer 引入 compute .so，依赖分层不对称且无说明
pybind/BUILD:94-99 为 th_compute_lib（librtp_compute_ops.so）的 ascend 分支加入 //rtp_llm/cpp/disaggregate/cache_store:cache_store_interface 与 @havenask//aios/autil:stack_tracer，而 using_cuda12（:88-92）与 using_rocm（:93）分支都不需要。compute .so 的职责是计算算子注册，cache_store 属传输/存储层，这种只在单平台出现的反向依赖通常是为消除未定义符号临时补上的，且缺少像 :67-70 那样的注释说明。同时「ExecOps.cc 只进 compute .so」这一不变量现在有三种写法：CUDA 用 srcs = exec_ops_srcs；ROCm 用组合目标 exec_ctx_rocm；Ascend（:70 + :94-99）用第三种 exec_ops_srcs + exec_ctx_ops，因而必
• [6.1] Tests — 分布式/跨平台变更有对应覆盖 → issue aclnn 自定义算子目标结构不规范、SoC 硬编码、测试环境 Python 前缀与仓库统一运行时冲突
  genrule（:3-20）只声明 4 个 .so 作为 outs，其中 opp/op_impl/ai_core/tbe/op_tiling/liboptiling.so 只是 CANN custom opp 包中的一个文件；而 .bazelrc:421 的 ASCEND_CUSTOM_OPP_PATH="bazel-bin/3rdparty/aclnn_custom_ops/opp" 需要完整 opp 目录，其余文件属未声明输出，远端缓存命中时会得到不完整的 opp 包。genrule 的真实输入还包括本机 CANN 安装（未声明）且未加 no-remote-cache；cmd（:13）中 ascend950 为硬编码 SoC；cc_library.deps（:30-31）放了不提供 CcInfo 的 genrule，而 4 个 .so 已在 data 列出。.bazelrc:421 的 bazel-bin/... 是工作区相对路径而测试 CWD 为 runfiles 根；:417-418 的 `LD_LIBRARY
• [6.1] Tests — 新逻辑有聚焦单测 + 相关集成/smoke 测试 → issue Ascend 默认 seq_size_per_block=128 缺少单测，且探测方式与相邻平台默认值不一致
:435-441 新增 Ascend 下 seq_size_per_block 默认 128，该值是 KV cache 分块与显存布局的关键输入，但 diff 未提供任何针对 setup_default_args 默认值决议链的单测。rtp_llm/config/test/server_config_setup_test.py:307 已具备现成手法（patch server_config_setup.torch.cuda.is_available），却未新增 Ascend 用例，文件内检索不到 seq_size_per_block 断言。此外该 try 块不受任何平台条件保护，在所有平台启动路径上都会执行：一旦 CI 基础镜像装入 torch_npu，:437 会把默认值由 64 改为 128，静默影响既有断言。判定方式也与相邻两个平台不一致——:415/:422 的 /dev/kfd 与 :430 的 /dev/alixpu 都是无副作用的设备文件探测，新分支改用 import torch_npu + `torch.npu.is_
• [6.1] Tests — 边界 case 覆盖（空、单元素、最大值） → issue Ascend 默认 seq_size_per_block=128 缺少单测，且探测方式与相邻平台默认值不一致
:435-441 新增 Ascend 下 seq_size_per_block 默认 128，该值是 KV cache 分块与显存布局的关键输入，但 diff 未提供任何针对 setup_default_args 默认值决议链的单测。rtp_llm/config/test/server_config_setup_test.py:307 已具备现成手法（patch server_config_setup.torch.cuda.is_available），却未新增 Ascend 用例，文件内检索不到 seq_size_per_block 断言。此外该 try 块不受任何平台条件保护，在所有平台启动路径上都会执行：一旦 CI 基础镜像装入 torch_npu，:437 会把默认值由 64 改为 128，静默影响既有断言。判定方式也与相邻两个平台不一致——:415/:422 的 /dev/kfd 与 :430 的 /dev/alixpu 都是无副作用的设备文件探测，新分支改用 import torch_npu + `torch.npu.is_

RTP-LLM Checklist

• [I] 代码质量 — 同一功能用统一工具函数 → issue 新增 Python/C++ 代码未过仓库既有 black / isort / flake8 / clang-format 约定
.flake8 仅 ignore W292、max-line-length=120，故 E501/F401 均生效：server_config_setup.py:439 的 logging 行约 146 字符、start_backend_server.py:148 约 168 字符均超限；server_config_setup.py:269/436/544 的 import torch_npu 属仅取副作用的未使用导入（F401），而本 PR 在 models_py/bindings/ascend/ops/tests/test_ascend_sampler.py:23 已给出正确的 # noqa: F401 写法，说明提交前未跑 pre-commit。.pre-commit-config.yaml 启用 black(88) 与 clang-format(120 列)：server_config_setup.py:437 与 start_backend_server.py:126 非 black 规范化结果（同函数 :421-424、`:4

Python Static-First Checklist

• [P.A] 静态结构与类型纪律 — 禁止 hasattr 做控制流分支 → issue _get_cuda_device_list 用 dir() 成员检测做控制流，条件恒为 False
:148 为 range(_dev_count if "_dev_count" in dir() else (torch.npu.device_count() if not torch.cuda.is_available() else torch.cuda.device_count()))。无参 dir() 返回当前函数的局部名字空间，而 _get_cuda_device_list 的局部变量只有 cuda_devices；_dev_count 是 _get_local_world_size(:126) 与 start_backend_server(:438) 的局部变量，不在本作用域内，条件恒为 False，作者跨函数复用局部变量的意图未实现。同时该表达式在「CUDA 不可用且未安装 torch_npu」时会抛 AttributeError: module 'torch' has no attribute 'npu'。该行约 168 字符，也超出 .flake8 的 max-line-length=120。
• [P.B] 错误处理 — 禁止 bare except 或静默吞异常 → issue torch_npu 探测逻辑多处内联重复，未复用同 PR 新增的 is_ascend()，且降级路径无日志
  同一份 try: import torch_npu / if torch.npu.is_available() 探测在本文件内联三次（:265-275 推导 local_world_size、:435-441 设 seq_size_per_block、:540-548 调 set_device），写法互不一致（一处 else 兜底、两处 except ImportError: pass）；model_loader/loader.py:594-600 是第四份，start_backend_server.py:126/148/438 又各自裸写。而本 PR 已在 device/device_type.py:17-49 提供规范实现 get_device_type()/is_ascend()（只依赖 os/enum/torch，叶子模块，无循环依赖风险）。异常语义也不安全：except ImportError 过窄（CANN 无法 dlopen 抛 OSError、版本错配抛 AttributeError 都会逃逸）；而

Strengths

• 平台扩展严格遵循仓库既有约定：新增 @//:using_ascend（BUILD:91）后一律以 select() 追加分支，未修改任何 cuda/rocm/arm 既有分支取值，绝大多数平台依赖图与行为保持不变，回滚面清晰。
• 设备来源被收敛到单一入口 getTorchCudaDevice()（ExecOps.cc:597-603），CUDA/ROCm 下仍返回 torch::kCUDA，Sampler.cc:123、ModelTypes.cc、cuda_graph_runner.h:55-61 的替换属零行为变更的安全抽象。
• MemoryLayoutStrategy.cc:3-8 用注释写清「NPU 上 torch::from_blob 丢失 ACL format tag 会破坏下游 FIA/scatter attention kernel」这一非显然约束，属高质量的 WHY 型注释，makeCacheTypedBlob() 的封装位置也正确。
• pybind 分层链接契约被严格遵守：pybind/BUILD:70 让 Ascend 经 exec_ops_srcs 把 ExecOps.cc 只编入 librtp_compute_ops.so，th_transformer_lib（:142-145）只依赖 exec_ctx_ops + exec_ops_hdr，与 :134-136 注释声明的「globalRuntime() 只存在一份」不变量一致。
• NormalExecutor.cc:203-207/216-220/320-324 采用「保留 sync()、只把 wait stream 换成 Ascend 默认流」的正确做法，保住了 host join 与 worker 异常回抛，是 MTP 侧应当对齐的范式。
• fuse_copy_kernel.h:15-19 用 #if USING_CUDA || USING_ROCM 收窄 invokeFusedCopy/invokeFusedStridedCopy 的声明范围，避免在无 cudaStream_t 的平台暴露不可用声明，与 FusedCopyOp.cc 的 aclrtMemcpyAsync 实现自洽。
• graphMemGetInfo（cuda_graph_device_shims.cc:178-180）与 ExecOps.cc:587-590 对 getDeviceMemoryInfo 返回的 (used, free) 语义处理正确（total = used + free）；BlockPool.cc:49 也补齐了 MEMORY_NPU 的 switch 分支，避免 -Wswitch -Werror 编译中断。
• PrefixToCandidateTokens.h:504-505 把 decltype(&fclose) 换成显式 lambda deleter，规避不同工具链下 fclose 声明差异，RAII 语义与所有权行为等价。
• 3rdparty/aclnn_custom_ops/BUILD:15-18 用 target_compatible_with 把自定义算子编译限制在 using_ascend；rtp_llm/BUILD:70/77 为 ascend 正确排除 decord/av 等依赖，与既有 arm 分支风格一致。
• server_config_setup.py:437 的 Ascend seq_size_per_block=128 带 == 0 前置判断，不覆盖 MI308X(16)/alixpu(256) 分支与用户显式配置，日志写明了 FIA v2 的性能理由，回滚路径清晰。

[LLLLKKKK]

📍 实际位置 rtp_llm/cpp/cuda_graph/cuda_graph_device_shims.cc:212（不在 diff 展示范围内，就近挂载）

[P0] ROCm 下 graphCaptureBegin 不再调用 capture_begin，既有平台的图捕获被静默破坏

.cc:212-219 中 graphCaptureBegin 仅在 #if USING_CUDA 下调用 graph.capture_begin(pool)，#else 只有 (void)pool; (void)graph;。.bazelrc:351-352 对 rocm 设 -DUSING_CUDA=0 -DUSING_ROCM=1，ROCm 必定走该空分支；header .h:172-173 为 USING_CUDA || USING_ROCM 声明的正是 at::cuda::CUDAGraph& 重载，且 cuda_graph/BUILD:89-92、models/BUILD:190、pybind/BUILD:140 的 rocm 分支均链接 cuda_graph_impl。唯一调用点 cuda_graph_runner.cc:1186 之后紧跟无条件 graph.capture_end()（:1197）：捕获从未开始却调用结束，图捕获与 replay 完全不可用。该改动对 Ascend 也无收益——...

建议: 改为三分支：#if USING_CUDA 调 graph.capture_begin(pool)；新增 #elif USING_ROCM 恢复 ROCm 原有的 capture_begin() 调用；仅在真正不支持的分支（Ascend/CPU）做空操作或 fail-fast。同时补一条 ROCm 侧图捕获+replay 的最小回归用例，避免同类 #else 合并再次吞掉既有平台行为；仓库已有 rocm 配置，可在 CI 中实际验证。

[LLLLKKKK]

[P0] Ascend pip lock 文件未随 PR 提交且存在三种不一致命名，顶层无条件 load 阻塞所有平台

deps/pip.bzl:88-89 的 pip_parse(name = "pip_ascend_torch") 要求 @rtp_deps//:requirements_lock_ascend_x86.txt，deps/BUILD:111 的 requirements_txt 同名；但已列举 deps/ 全量文件，只有输入文件 requirements_ascend_x86.txt，不存在任何 requirements_lock_ascend*.txt。PR 新增的仓库根 requirements_lock_ascend.txt 读取报 File does not exist（悬空 symlink），且缺 _x86 后缀，是第三个名字，全仓无 label 引用。而 WORKSPACE:80-81 与 arch_config/arch_select.bzl:8 都是与 --config 无关的顶层 eager load，故 cuda12/cuda12_9/cuda13/rocm/arm/cpu 全部配置都会在 WORKSPACE...

建议: 提交由 bazel run //deps:requirements_ascend_x86.update 生成的 deps/requirements_lock_ascend_x86.txt 并入库，统一 deps/pip.bzl:89、deps/BUILD:111 与根 symlink 三处命名（建议与 src 对齐用 _x86 后缀），并像 requirements_lock_rocm.txt 一样加入 deps/BUILD 的 exports_files。合并前请在不叠加 --override_repository @rtp_deps 的树上跑一次 bazel query //rtp_llm/...（任一非 Ascend config）确认 WORKSPACE 可正常求值——内部 override 会掩盖此问题使开源侧单向破损。若短期无法提供 lock 文件，应把 pip_ascend_torch 与两处 load 一并延后引入。另请修正 deps/pip.bzl:90 的制表符缩进。

Checklist: [6.1] Commit 原子、message 与行为匹配；[6.1] PR description 说明动机与设计

[LLLLKKKK]

[P1] Ascend 上事件 record 退化为空操作，transfer_done_event 的 D2H 完成不变量被静默破坏

.h:138-139 注释称 Ascend 下 record/block 对「degrade to stream-synchronous semantics」，但 graphRecordEvent（:141-143）与 graphBlockEvent（:150-152）实现是 (void)event; (void)stream; 纯空操作，toGraphStream/graphGetStreamFromPool/graphGetCurrentStream（:68-118）一律返回 nullptr，GraphStreamGuard（:57-59）为空构造，不存在任何降级同步。生产端 SpeculativeSampler.cc:255-257 是 .to(torch::kCPU, /*non_blocking=*/true) 后 graphRecordEvent，消费端未改动：MtpBatchStreamProcessor.cc:1278（及 :1043）仍 `transfer_done_event->synchro...

建议: Ascend 分支不要留空：graphRecordEvent 应调用 event.record(<当前 NPU stream>)（FusedCopyOp.cc:49 已在用 c10_npu::getCurrentNPUStream()），或在无法建立事件时退化为该 stream 的 synchronize()，使消费侧 synchronize() 的契约成立；同时把 .h:138-139 注释改成与实现一致的描述（「Ascend 当前只使用默认流，暂不提供跨流定序」）。若短期无法实现事件语义，请把 accept_*_cpu 的 D2H 改为 non_blocking=false 并在 shim 头部注明该退化，并补一条 Ascend MTP 解码的 smoke 用例锁定 accept 序列。

[LLLLKKKK]

AI Code Review - PR #1349 (non-blocking suggestions)

19 条 P2/P3 建议，不阻塞合并。阻塞判定与完整摘要见上一条 review。

[LLLLKKKK]

[P2] ExecOps.h 被 cpp 核心层新增引用但两个 Bazel 目标未声明该依赖，且属反向跨层依赖

cuda_graph_runner.h:3 新增 #include "rtp_llm/models_py/bindings/core/ExecOps.h"（为用 :56/:60 的 getTorchCudaDevice()），但暴露该头的 cc_library(name="cuda_graph_hdrs_lib")（cuda_graph/BUILD:45-57）deps 只有 //rtp_llm/models_py/bindings:op_defs，缺 //rtp_llm/models_py/bindings/core:exec_ops_hdr（只有 cuda_graph_impl 在 BUILD:78 声明了），且本 PR 未修改 cuda_graph/BUILD。同型问题见 CudaCopyUtil.cc:5：transfer/tcp/BUILD:9-13 的 deps 中 //:rtp_compute_ops 是 linkshared=1 的 cc_binary（不传播编译上下文）、`no_block_copy_h...

建议: 在 cuda_graph_hdrs_lib.deps 与 cuda_copy_util.deps 中显式加入 //rtp_llm/models_py/bindings/core:exec_ops_hdr（与 th_transformer_lib 的做法一致），并在 Ascend 与 CUDA 两个 config 下各验证一次这两个 target 的构建。更好的做法是把 getTorchCudaDevice() 这类纯设备查询下沉到 rtp_llm/cpp/utils 的小头文件，避免引擎层/cache 层头文件反向依赖 models_py/bindings 层。

[LLLLKKKK]

[P2] NPU 保 ACL format 的 from_blob 包装无法复用，KV 拷贝路径仍直接用 torch::from_blob

MemoryLayoutStrategy.cc:3-8 明确记录 NPU 上 torch::from_blob 会丢失 ACL format tag 并破坏下游 FIA/scatter kernel，故 :17-23 引入 makeCacheTypedBlob() 走 at_npu::native::from_blob。但该 helper 藏在 .cc 匿名 namespace 中无法复用，而本 PR 同时把 CudaCopyUtil.cc:11-14 的 wrapRawPtr()（:30-32/:55-57 传入 getTorchCudaDevice()，Ascend 下即 kPrivateUse1）与 KVCacheMemoryConnector.cc:2751-2762 的设备选择改成了 NPU，包装方式却仍是 :13/:2757/:2760 的原生 torch::from_blob。这些 tensor 随后进入 execNoBlockCopy（Ascend 走 dst.copy_(src)），而本 P...

建议: 把 makeCacheTypedBlob() 提升为共享工具（例如放到 rtp_llm/cpp/cache/ 公共头或 models_py/bindings/core/torch_utils/TypeConvert.h），在所有包装设备侧裸指针的位置统一使用：CudaCopyUtil.cc:13、KVCacheMemoryConnector.cc:2757/2760。若这两条 connector 路径在 Ascend 上确实不启用，请显式 fail-fast 或加注释说明未支持，而不是把设备改成 NPU 却保留 CUDA 时代的包装方式；若确认纯 memcpy 语义与 format tag 无关，请在 makeCacheTypedBlob 注释里写清适用边界。（BlockPool.cc:272 的 from_blob 位于 #if USING_CUDA 内，不受影响。）

[LLLLKKKK]

[P2] BlockInfo::is_cuda 语义被扩为"设备内存"，但字段名、注释与部分消费者未同步

:35-36 把判定改为 t.is_cuda() || t.is_privateuseone()（同一表达式被求值两遍），字段含义从「CUDA 显存」变为「设备内存」，但 cache/BlockInfo.h:14 的注释仍写 "Torch device of the backing storage (CPU/CUDA)"、字段名也未改。该字段直接流入 model_rpc/DecodeRpcServer.cc 的 addBlock(..., block.is_cuda, true) → MemoryUtil::regUserMr(buf, size, gpu, ...) 按 GPU 显存注册 MR；测试辅助仍以 b.is_cuda 为条件直调 CUDA API。全仓已 grep 到 17 处同形表达式（含 MemoryLayoutStrategy.cc:307、PyWrappedModel.cc:80/298/781/793、ModelTypes.cc:416、MtpExecutor.cc:135/150/1104），而 `KVCache...

建议: 抽一个共享 helper（如 isDeviceTensor(const torch::Tensor&)）放入 rtp_llm/cpp/utils 供全部调用点使用，并把 :36 的重复求值合并为一个局部 const bool on_device；把 BlockInfo::is_cuda 更名为 is_device/is_accelerator 并同步 BlockInfo.h:14 注释与 KVCacheMemoryConnector.cc:2751 的形参命名。同时确认 RDMA regUserMr 在收到 NPU 指针且 gpu=true 时行为可接受，并为以 is_cuda 为条件直调 CUDA API 的测试辅助函数补平台断言。

[LLLLKKKK]

[P2] BlockPool::where() 的 Ascend 分支丢失 PINNED 语义，NPU 设备池被按 host 内存注册 MR

新的 where()（:751-760）在 #if USING_ASCEND 下只判 is_privateuseone()，其余一律返回 MEMORY_CPU，丢掉了非 Ascend 分支保留的 is_pinned() ? MEMORY_CPU_PINNED : MEMORY_CPU 区分；initializeCacheBuffer() 对 HOST 池会调 pin_memory()，Ascend 上该池被误报为 MEMORY_CPU，memoryTypeName(where()) 日志随之失真。更实质的是 regUserMr/deregUserMr（:580、:614）用 where() == MEMORY_GPU 推导 gpu 标志，Ascend 设备池返回 MEMORY_NPU 不满足该条件，于是会把 NPU 设备内存按 host 内存注册；本 PR 已为 Ascend 链接 cache_store_interface（pybind/BUILD:95），该路径并非不可达。

建议: Ascend 分支保留 pinned 判定，例如统一为 if (is_privateuseone() || is_cuda()) return <device type>; return is_pinned() ? MEMORY_CPU_PINNED : MEMORY_CPU;（可直接复用同 PR 新增的 TypeConvert.h:92-96 torchDeviceToMemoryType()）。把 :580/:614 的判定改为「是否设备内存」语义——将 OpData.cc:80-82 的 isDeviceMemory()（MEMORY_GPU || MEMORY_NPU）从文件内 static 提升为公共 helper 并在此复用，避免 MEMORY_NPU 在各消费点被逐个漏判；KVCacheAllocator.cc:294 硬编码的 MEMORY_GPU 也应改用同 PR 新增的 allocationTypeToMemoryType(allocation_type_)。若 Ascend 暂不支持 cache_store，请在该路径显式 fail-fast。

[LLLLKKKK]

[P2] aclnn 自定义算子目标结构不规范、SoC 硬编码、测试环境 Python 前缀与仓库统一运行时冲突

genrule（:3-20）只声明 4 个 .so 作为 outs，其中 opp/op_impl/ai_core/tbe/op_tiling/liboptiling.so 只是 CANN custom opp 包中的一个文件；而 .bazelrc:421 的 ASCEND_CUSTOM_OPP_PATH="bazel-bin/3rdparty/aclnn_custom_ops/opp" 需要完整 opp 目录，其余文件属未声明输出，远端缓存命中时会得到不完整的 opp 包。genrule 的真实输入还包括本机 CANN 安装（未声明）且未加 no-remote-cache；cmd（:13）中 ascend950 为硬编码 SoC；cc_library.deps（:30-31）放了不提供 CcInfo 的 genrule，而 4 个 .so 已在 data 列出。.bazelrc:421 的 bazel-bin/... 是工作区相对路径而测试 CWD 为 runfiles 根；:417-418 的 `LD_LIBR...

建议: 1) 若 opp 是目录产物，改用 declare_directory 风格的自定义 rule，或把 opp 包打成单个 tar 作为唯一 out；2) 给 genrule 追加 no-remote-cache（或把 CANN 版本纳入输入）；3) 把 ascend950 提到 --define 或 config_setting；4) 从 cc_library.deps 移除 :build_aclnn_ops；5) ASCEND_CUSTOM_OPP_PATH 改为 runfiles 相对路径并验证 libcust_opapi.so 在测试中确实可被加载；6) ascend 统一使用与 --python_top 一致的 Python 运行时（或说明为何必须独立环境并在镜像中固化）；7) 逐项为 -Wno-error=*/-Wno-* 加注释说明触发文件与清理计划。

Checklist: [6.1] 错误语义：fail-fast/retry/fallback/silent 行为显式；[6.1] 分布式/跨平台变更有对应覆盖

[LLLLKKKK]

[P2] ascend 分支把 cache_store 与 stack_tracer 引入 compute .so，依赖分层不对称且无说明

pybind/BUILD:94-99 为 th_compute_lib（librtp_compute_ops.so）的 ascend 分支加入 //rtp_llm/cpp/disaggregate/cache_store:cache_store_interface 与 @havenask//aios/autil:stack_tracer，而 using_cuda12（:88-92）与 using_rocm（:93）分支都不需要。compute .so 的职责是计算算子注册，cache_store 属传输/存储层，这种只在单平台出现的反向依赖通常是为消除未定义符号临时补上的，且缺少像 :67-70 那样的注释说明。同时「ExecOps.cc 只进 compute .so」这一不变量现在有三种写法：CUDA 用 srcs = exec_ops_srcs；ROCm 用组合目标 exec_ctx_rocm；Ascend（:70 + :94-99）用第三种 exec_ops_srcs + exec_ctx_ops，...

建议: 定位 ascend 下实际缺失的符号来源，把依赖补在真正引用它的目标上（例如 ascend_host_utils 或 exec_ctx_ops），而不是在 pybind 汇聚层补洞；若确实无法避免，请像 :67-70 那样注明是哪个符号、来自哪个 TU。同时把 exec_ctx_rocm 泛化并改名（如 exec_ctx_combined），让 Ascend/CUDA/ROCm 统一走同一 deps 形状，由该目标自己维护 ExecOps.cc 的依赖清单；kv_cache_utils 若只是 ascend 需要也应放进 select 分支。

Checklist: [6.1] 依赖方向：无循环依赖/跨层惊喜；[6.1] 分层边界：新概念在正确层级，不泄漏内部；[6.1] SRP：模块/类职责单一

[LLLLKKKK]

[P3] 条件编译结构冗余与死代码，ROCm 的关键说明被埋入不可达分支

Sampler.cc:126-149 把原 ROCm 分支嵌进 #if USING_CUDA 内层：.bazelrc:351 对 rocm 设 USING_CUDA=0，故 :127 的 #if USING_ROCM 永不可达，:128-130 那段「hipMemcpyAsync from pageable memory 真异步、record_stream() 拒绝 at::hip 流」的关键说明成为死注释；ROCm 实际走 :146-148 的「Ascend/other devices」分支，行为等价（getTorchCudaDevice() 在 ROCm 返回 kCUDA）故无功能回归，但注释归属错误会误导维护者。同类噪声：ComputeInit.cc:11-12 是空的 #if/#endif 块本次又被追加一个平台宏；cuda_graph_device_shims.h:130-136 的 makeGraphEvent() 与 :9-10 的 GRAPH_DEVICE_TYPE 重复硬编码同一枚举，`:101-1...

建议: Sampler.cc 改为平铺的 #if USING_CUDA ... #elif USING_ROCM ... #else 三分支，把 ROCm 说明放回可达分支并把 #else 注释改为准确描述 Ascend/CPU 的阻塞拷贝；删除 ComputeInit.cc:11-12 的空块；makeGraphEvent() 统一 return torch::Event(GRAPH_DEVICE_TYPE); 并给 graphGetStreamFromPool 的 Ascend 分支补 (void)is_high_priority;；SpecLogitsVerifyRunner.cc 二选一（去掉外层 #if USING_CUDA 或恢复直接 record/block）；cuda_host_utils.h 的 Ascend 注释改为明确「该头不支持 Ascend，请使用 ascend_host_utils.h」；删除与 default 等价的 using_ascend 空分支。

[LLLLKKKK]

[P3] _ascend_excluded 与 whl_deps 的 Ascend 分支不符合既有约定

:26-36 的 _ascend_excluded 与紧邻的 _CUDA13_DEFERRED（:17-19）、_DSV4_PLATFORM_ONLY（:21-24）三点不一致：命名用小写而既有同类常量用 UPPER_SNAKE；完全没有注释说明每项为何排除，而既有两个常量都写清了原因；条目与实际解析结果不符——decord、xfastertransformer_devel、xfastertransformer_devel_icx 在 deps/requirements_base.txt:8/56/57 均带 platform_machine == "x86_64" 标记而 Ascend 目标是 x86_64，且 rtp_llm/BUILD:57-65 的 xft_dep 全部注释掉、:70/77 对 ascend 返回 []，这三个 target 并不在 Ascend 依赖图内。另 :52 ascend_deps = ... 用多空格对齐（buildifier 应为单空格）；:119 为 Ascend ...

建议: 改名为 _ASCEND_EXCLUDED 并补注释说明排除依据（哪些包在 Ascend lock 中不存在、为何不存在）；剔除 decord/xfastertransformer_devel/xfastertransformer_devel_icx 三项冗余条目，若排除确有意图则在注释中写明；:52 改为单空格。whl_deps() 的 Ascend 分支与 cuda13 保持一致改为直链形式，使产出 wheel 在无额外索引配置的环境也可安装；若团队约定依赖内部 index 解析 local version，请在该行补注释指明所需 index，并顺带记录 :108 using_cuda12 的相同遗留以便后续统一。lock 文件入库后请用 bazel query 在 ascend 配置下复核清单完整性。

Checklist: [6.1] 兼容性：外部 HTTP/RPC API、持久数据、配置、环境迁移安全；[6.1] KISS/YAGNI：无投机性抽象

LGTM:阻断已解除(rtpcli 自动清除旧红标)

[LLLLKKKK]

AI Code Review - PR #1349 (non-blocking suggestions)

22 条 P2/P3 建议，不阻塞合并。阻塞判定与完整摘要见上一条 review。

[LLLLKKKK]

📍 实际位置 rtp_llm/cpp/cache/BlockPool.cc:580（不在 diff 展示范围内，就近挂载）

[P2] where() 新增 MEMORY_NPU 后 regUserMr 仍按 MEMORY_GPU 判定，NPU 显存被当作 host 内存注册

where() 在 USING_ASCEND 下改为返回 MEMORY_NPU（设备）或 MEMORY_CPU（:751-760），但全仓仅剩的两处 == MemoryType::MEMORY_GPU 消费者（:580、:614）未同步：Ascend 上 gpu 恒为 false，NPU 显存块会以 host 内存语义传给 memory_util->regUserMr(...)／对应 dereg（:643、:670）。而同批 NPU 张量在 MemoryLayoutStrategy.cc 与 LayerBlockConverterImpl.h:35 已被显式视作加速器内存——同一 pool 在 BlockInfo 层是 device、在 MR 注册层是 host。Ascend 分支同时丢掉 MEMORY_CPU_PINNED 区分，host 池的日志（:212）与后续判定退化为 CPU；:203 的 is_cuda = cache_aligned_buffer_.is_cuda() 也未扩宽。

建议: 把 580/614 改为「是否设备内存」的统一判定（复用 OpData.cc:81 的 isDeviceMemory 语义，或 where() != MEMORY_CPU && where() != MEMORY_CPU_PINNED），与 BlockInfo.is_cuda 保持一致；Ascend 分支保留 pinned 判定：is_privateuseone() ? MEMORY_NPU : (is_pinned() ? MEMORY_CPU_PINNED : MEMORY_CPU)，并同步 :203 的谓词。若 Ascend 确认暂不走该注册路径，请加断言或注释明确不支持而非静默错配。新增枚举值时建议全仓搜索等值比较点一次性补齐。

[LLLLKKKK]

[P2] BlockInfo::is_cuda 语义被扩为「加速器内存」，字段名/注释未同步且设备判定三套并存

:35-36 把 info.is_cuda = t.is_cuda() 改为 t.is_cuda() || t.is_privateuseone()，device_index 同步扩宽，同一谓词在两行内重复书写；该表达式在本 PR 生产代码中共出现 18 处、分布 9 个文件，Torch_ext.h 只是把 CHECK_CUDA 宏放宽，未提供可复用谓词——任何遗漏（如 PyWrappedModel.cc:170）都会在 Ascend 上静默走错分支且编译期无法发现。BlockInfo.h 注释仍写 “CPU/CUDA”、字段名仍是 is_cuda，语义已漂移；getAllBuffers() 结果经 P2PConnectorWorker 送入 regMem。设备选择同样三套并存：getTorchCudaDevice()、手写宏链（BlockPool.cc:189-197、MtpExecutor.cc:2234-2238）、以及大量裸 torch::kCUDA/.cuda()（`cuda_graph_runner.cc...

建议: 在 ExecOps.h（与 getTorchCudaDevice() 同一头文件）提供 inline bool isAcceleratorTensor(const at::Tensor&) / isAcceleratorDevice(const at::Device&) 并全仓替换该谓词，本处先改为局部变量复用一次；把 BlockInfo::is_cuda 更名为 is_device/on_accelerator（内部结构体、非 proto，可安全改名）或至少更新注释为「CPU / accelerator device memory」，并在提交说明中列出所有 regMem 实现的核对结论。设备选择统一收敛到 getTorchCudaDevice()，删除手写宏链；cuda_graph_runner.cc 要么文件内统一，要么回退这几处并注明该文件仅支持 CUDA/ROCm。若 RDMA/ACCL 后端尚未支持 NPU 内存注册，应在 Ascend 上显式拒绝该 transfer backend。

[LLLLKKKK]

[P2] NPU 保 ACL format 的 from_blob 包装无法复用，KV 设备内存其他包装点仍用裸 torch::from_blob

MemoryLayoutStrategy.cc:3-23 明确写出不变量：「torch::from_blob 会丢失 NPU 张量的 ACL 格式标记，破坏下游 FIA/scatter attention kernel」，并为此提供 makeCacheTypedBlob()——但它位于匿名 namespace，仅本文件可用。同一 KV 内存的其他包装点全部仍是裸 torch::from_blob：本 PR 刚改为 getTorchCudaDevice() 的 CudaCopyUtil.cc:11-13（Ascend 下由此产出 PrivateUse1 张量）、KVCacheMemoryConnector.cc:2753-2760 的设备张量。即「必须用 at_npu::native::from_blob」这条不变量只在一个文件内成立，其余路径按注释所述会丢格式标记。

建议: 把该包装提升为可复用工具（放到 ExecOps.h/TypeConvert.h 旁的公共头，命名体现语义，例如 fromBlobPreserveDeviceFormat()），并替换 CudaCopyUtil.cc、KVCacheMemoryConnector.cc 中作用于 KV 设备内存的 torch::from_blob；若某些点确实只做字节级拷贝、无需格式标记，请就近注释说明为何可安全使用裸 from_blob，避免不变量只靠单点注释维持。

[LLLLKKKK]

[P2] Ascend 下 cuda_graph::setDevice 为空实现，且实参为 index=-1，worker 线程不绑定 NPU

setDevice 的 #if USING_ASCEND 分支只有 (void)rank;（:87-89），CUDA/ROCm 分支则 cudaSetDevice/hipSetDevice 并校验。唯一调用点 AsyncRunner.cc:12 在 worker 线程启动时调用，设备绑定是线程级的，Ascend 下 GraphStreamGuard 亦为空壳，故该线程内无任何绑定动作，而 NormalOutputDispatcher.cc:30-34 会在该线程做 pinned D2H。更关键的是实参本身失效：MtpExecutor.cc:552-557 的 AsyncRunner 由 torch::Stream(c10::Stream::DEFAULT, getTorchCudaDevice()) 构造，而 ExecOps.cc:603 返回不带 index 的设备，故 stream_.device_index() 为 -1，setDevice(-1) 即使日后实现也是错误参数；`NormalExecutor.cc:81-83...

建议: 在 USING_ASCEND 分支实现真实绑定（aclrtSetDevice(rank) 配合 torch_npu 的 set_device，按 ascend_host_utils 现有风格做错误校验），并把 MtpExecutor 的 AsyncRunner 构造对齐 NormalExecutor，传入携带 local_rank 的 torch::Device(kPrivateUse1, local_rank)，避免把 -1 传进 setDevice(int)。若本阶段不打算支持，请改为 RTP_LLM_CHECK_WITH_INFO 失败或至少 WARNING，把「不支持」显式化而非静默返回；并补一个 local_rank>0 的用例覆盖 worker 线程的设备绑定。

[LLLLKKKK]

[P2] Ascend 下所有 AsyncRunner 退化到默认流，异步重叠收益消失且顺序依赖隐式化

:549-557 把 4 个 AsyncRunner 与 collect_metrics_stream_ 全部构造为 torch::Stream(c10::Stream::DEFAULT, getTorchCudaDevice())，非 Ascend 分支（:559-564）则是 graphGetStreamFromPool(true) 的独立流。配合 AsyncRunner::workerLoop:86 的 GraphStreamGuard（Ascend 下为空壳，GraphStream = void*），worker 线程与主线程会向同一条默认流并发入队；launch() 与 sync() 之间主线程仍在同流入队，入队顺序取决于线程交错。同时 target_verify/draft_prefill 预处理与主流串行化，异步预处理的 overlap 收益归零。该不变量目前只在 AsyncRunner.h 注释里零散说明，无断言保护。

建议: 建议：(1) 若 Ascend 暂不支持多流，用开关在 Ascend 上直接关闭异步 runner 走同步路径，而不是保留 worker 线程但共享默认流，这样顺序语义与回滚手段都更明确；(2) 若保留，显式用 c10_npu::getCurrentNPUStream() 构造并在注释中写明「worker 与主线程共享默认流，仅靠 host join 保序」；(3) 在 PR 描述中记录 Ascend 上异步预处理不产生重叠这一预期，避免后续按 CUDA 结论做性能归因。

[LLLLKKKK]

[P3] Sampler.cc 中 ROCm 分支被外层 USING_CUDA 包成死代码，关键说明落在不可达处

:126-127 形成 #if USING_CUDA 内嵌 #if USING_ROCM。.bazelrc:351-352 显示 ROCm 构建定义 USING_CUDA=0、USING_ROCM=1，外层为真时内层恒为假，:128-131 的 ROCm 专用分支及其「hipMemcpyAsync from pageable memory / record_stream 拒绝 at::hip streams，故保留阻塞拷贝」注释永不可达。ROCm 实际落入 :146-148 的 #else（to(getTorchCudaDevice())，行为等价、无功能回归）。死分支带着关键说明存在，会误导后续维护者以为在此调整 ROCm 行为生效。

建议: 把嵌套宏改为并列的 #if USING_CUDA / #elif USING_ROCM / #else 三分支，让 ROCm 的阻塞拷贝说明落在真正生效处；或直接删除已被 #else 覆盖的 ROCm 分支，避免同一语义存在两处不一致的平台判定。

[LLLLKKKK]

[P3] Ascend 内存查询缺少返回码校验，另有冗余守卫、未使用 include 与空预处理块

1. :60-63 的 Ascend 分支不校验返回值，而 CUDA/ROCm 分支分别有 check_cuda_value/ROCM_CHECK；其调用的 ascend_host_utils.cc:48-53 同样忽略 aclrtGetMemInfo 返回码（该文件 :12-19 已有会抛异常的 check 模板未被使用），查询失败时 total 为 0，5% 下限静默退化并影响 KV cache 容量决策。2) MtpBatchStreamProcessor.cc:537 的 #if !USING_ASCEND 冗余：ready_event 创建整体在 #if USING_CUDA 内，Ascend 下恒为空，已有的 if (spec_logits_result.ready_event) 判断即可。3) KVCacheAllocator.cc:13 新增的 TypeConvert.h include 在该文件无使用点。4) ComputeInit.cc:11-12 是内部完全为空的 `#if USING_CUDA || US...

建议: 为 Ascend 内存查询补上与其他平台对等的返回码校验（复用 ascend_host_utils 的 check 模板，并覆盖 aclrtGetDevice/aclrtGetDeviceCount），失败时快速失败或明确告警；删除 #if !USING_ASCEND 冗余守卫、未使用的 include、ComputeInit.cc:11-12 的空块与多余空行。另建议统一 torch_npu/.../NPUStream.h 的引入方式：PyWrappedModel.h:27-32 用 pragma 抑制 -Wunused-function，而 PyWrappedModel.cc 的同一 include 没有，目前只是靠 include 顺序侥幸通过 -Werror。

Checklist: [6.1] 错误语义：fail-fast/retry/fallback/silent 行为显式

[LLLLKKKK]

[P3] GRAPH_DEVICE_TYPE 的 Ascend 定义永不生效，与 makeGraphEvent 分支重复

:10 已为 Ascend 定义 #define GRAPH_DEVICE_TYPE torch::kPrivateUse1，但该宏的唯一使用点 :142 位于 makeGraphEvent() 中 #if USING_ASCEND 的 #else 分支内，导致 Ascend 定义永不生效；:140 又在 Ascend 分支硬写一遍 torch::kPrivateUse1，同一常量出现两份来源。三个分支的宏值书写风格也不统一（torch::kPrivateUse1 与 c10::DeviceType::HIP/CUDA 混用）。

建议: 删除 makeGraphEvent() 中的 #if USING_ASCEND 分支，直接 return torch::Event(GRAPH_DEVICE_TYPE);（Ascend 下宏已展开为 torch::kPrivateUse1）；统一三个分支的宏值书写风格。

Checklist: [6.1] KISS/YAGNI：无投机性抽象

[LLLLKKKK]

📍 实际位置 rtp_llm/model_loader/loader.py:587（不在 diff 展示范围内，就近挂载）

[P3] Python 侧 Ascend 探测出现三套实现，且设备无关化后命名与重复探测未同步整理

loader.py:586-600 的 force_clean_cuda_memory 用 try: import torch_npu / torch.npu.is_available() 内联判定；weight_manager.py:120-122 改用函数内 from rtp_llm.device.device_type import get_device_type, DeviceType 比较枚举；而本 PR 新增的 device_type.py:44 的 is_ascend() 已被 server_config_setup.py:266/429/532 与 start_backend_server.py:147/183 使用——同一 PR 内三套探测并存。命名同样未同步：force_clean_cuda_memory 现在也清理 NPU，_get_cuda_device_list/cuda_device_list（start_backend_server.py:143/177）返回 NPU 列表，`setup_cuda_...

建议: loader.py 与 weight_manager.py 都改为调用 device_type.is_ascend()；force_clean_cuda_memory 改名为 force_clean_device_memory 或在 docstring 说明覆盖 NPU；把 _get_cuda_device_list/cuda_device_list/setup_cuda_device_and_accl_env 改为设备中性命名（并同步 ExecOps.cc:601 注释）；把 is_ascend() 结果提到循环外，或在 device_type.py 用 functools.lru_cache 缓存 get_device_type()（进程级常量，无泄漏风险）。若 weight_manager.py 的函数内 import 是为规避循环依赖，请补一行注释说明。

Checklist: [6.1] DRY：重复非平凡逻辑被抽取或显式复用；[P.D] lru_cache 注意内存泄漏风险；[I] 同一功能用统一工具函数

[LLLLKKKK]

[P3] 新增 Python/C++ 代码未过仓库既有 black / isort / flake8 / clang-format 约定

仓库 .pre-commit-config.yaml 启用 black + isort(--profile=black) + clang-format，flake8 max-line-length=120。但：(1) server_config_setup.py:431 的 logging.info("[Ascend] set SEQ_SIZE_PER_BLOCK 128 ...") 单行约 141 字符；(2) start_backend_server.py:20 把 from rtp_llm.device.device_type import ... 插在 rtp_llm.config.log_config 与 rtp_llm.config.py_config_modules 之间，违反 isort 排序；(3) :447-450 的多行 if (...) 折叠后仅约 66 字符，black 会收成一行；(4) MtpExecutorTest.cc:479、MtpBatchStreamProcessorTest.cc:148 单行约 1...

建议: 提交前对本 PR 改动文件执行 pre-commit run --files ...（black + isort + clang-format），避免后续格式化提交与逻辑提交混在一起、干扰 CI 二分定位。

Checklist: [6.1] 逻辑变更未混入无关格式化