feat(xpu): add Intel GPU (XPU) support#1167
Open
aslanxie wants to merge 1 commit into
Open
Conversation
Collaborator
AI Code Review - PR #1167Status: BLOCKING Summary: P0/0 · P1/0 · P2/2 · P3/1 Non-blocking SuggestionsP2
P3
Checklist Violations (2 fail / 56 total)General Principles Checklist
Strengths
|
Contributor
There was a problem hiding this comment.
Pull request overview
Note
Copilot couldn't run its full agentic review because no GitHub Actions runner was available. Make sure your repository has a runner available to run Copilot's review, or add a copilot-setup-steps.yml file specifying one with the runs-on attribute. See the docs for more details.
Adds Intel GPU (XPU) inference support to RTP-LLM while keeping existing CUDA/ROCm flows intact by gating behavior behind --config=xpu and device-type resolution.
Changes:
- Introduces XPU device detection/selection (Python) and XPU runtime support (C++), including memory/accounting, sampling, and beam-search fallbacks.
- Adds XPU attention + KV-cache layout adaptation and registers XPU-specific module factories/ops (with CPU-runnable unit tests).
- Extends Bazel build + dependency plumbing for SYCL/oneAPI toolchains, torch-xpu discovery, and XPU-specific wheel/requirements handling.
Reviewed changes
Copilot reviewed 93 out of 97 changed files in this pull request and generated 1 comment.
Show a summary per file
| File | Description |
|---|---|
| rtp_llm/start_backend_server.py | Uses device-agnostic GPU helpers and XPU-aware env binding / startup guards |
| rtp_llm/ops/init.py | Adds XPU-specific note around libpython linkage behavior |
| rtp_llm/models_py/utils/arch.py | Adds CUDA-only guard for get_sm() |
| rtp_llm/models_py/standalone/auto_model.py | Binds per-rank device early; XPU MHA override; XPU KV cache layout + pinning tweak |
| rtp_llm/models_py/modules/hybrid/causal_attention.py | Imports XPU fused norm implementation |
| rtp_llm/models_py/modules/factory/linear/impl/xpu/f16_linear.py | Adds XPU F16/BF16 linear strategy (PyTorch fallback) |
| rtp_llm/models_py/modules/factory/linear/impl/xpu/init.py | Registers XPU linear strategy |
| rtp_llm/models_py/modules/factory/linear/init.py | Loads XPU linear registry based on device type |
| rtp_llm/models_py/modules/factory/fused_moe/impl/xpu/init.py | Adds XPU MoE placeholder package |
| rtp_llm/models_py/modules/factory/fused_moe/init.py | XPU uses batched triton MoE fallback only |
| rtp_llm/models_py/modules/factory/attention/xpu_impl/test/test_no_rope.py | Adds CPU-runnable test ensuring RoPE isn’t applied for RopeStyle.No |
| rtp_llm/models_py/modules/factory/attention/xpu_impl/test/test_kv_cache_layout.py | Adds CPU-runnable test guarding XPU NSHD KV-cache layout contract |
| rtp_llm/models_py/modules/factory/attention/xpu_impl/test/BUILD | Adds Bazel py_tests for XPU attention helpers |
| rtp_llm/models_py/modules/factory/attention/common.py | Makes write_cache_store no-op when kv_cache is None |
| rtp_llm/models_py/modules/factory/attention/init.py | Registers XPU vLLM flash-attn implementations |
| rtp_llm/models_py/modules/base/xpu/vllm_xpu_ops.py | Wraps vllm-xpu-kernels with PyTorch fallbacks |
| rtp_llm/models_py/modules/base/xpu/select_topk.py | Adds XPU MoE top-k selection fallback |
| rtp_llm/models_py/modules/base/xpu/not_implemented_ops.py | Adds explicit XPU stubs for not-yet-supported ops |
| rtp_llm/models_py/modules/base/xpu/norm.py | Adds XPU norm implementations (vllm kernels + fallbacks) |
| rtp_llm/models_py/modules/base/xpu/moe_gating.py | Adds XPU MoE gating fallback |
| rtp_llm/models_py/modules/base/xpu/activation.py | Adds XPU fused activation fallback |
| rtp_llm/models_py/modules/base/common/embedding.py | Adds fallback/fast-fail behavior when compiled embedding op missing |
| rtp_llm/models_py/modules/base/init.py | Adds XPU base op wiring (norm/activation/moe stubs) |
| rtp_llm/models_py/model_desc/disaggregate_qwen3.py | Selects block map per layer during micro-batch processing |
| rtp_llm/models_py/bindings/xpu/XpuTorchExt.h | Adds XPU-only torch extension include wrapper |
| rtp_llm/models_py/bindings/xpu/RegisterXpuOps.cc | Registers XPU pybind ops |
| rtp_llm/models_py/bindings/xpu/BUILD | Adds Bazel target for XPU bindings registration |
| rtp_llm/models_py/bindings/core/ExecOps.h | Introduces getTorchDevice() + maybePinMemory() helpers |
| rtp_llm/models_py/bindings/core/ExecOps.cc | Adds XPU runtime sync/event/device/memory accounting support |
| rtp_llm/models_py/bindings/core/CudaSampleOp.cc | Adds XPU sampling fallback; disables speculative/rejection sampling on XPU |
| rtp_llm/models_py/bindings/core/CudaOps.cc | Adds XPU copy + mask-logits fallbacks |
| rtp_llm/models_py/bindings/core/CudaBeamSearchOp.cc | Adds XPU beam-search fallback implemented via PyTorch ops |
| rtp_llm/models_py/bindings/core/BUILD | Updates BUILD selects/features and sampling deps across CUDA/ROCm/XPU |
| rtp_llm/models_py/bindings/common/kernels/BUILD | Avoids compiling CUDA fuse-copy kernel in XPU builds |
| rtp_llm/models_py/bindings/common/FusedCopyOp.cc | Adds SYCL memcpy fallback for fused copy ops on XPU |
| rtp_llm/models_py/bindings/common/BUILD | Adjusts bindings compilation/linking for XPU; adds xpu_sycl_compile feature |
| rtp_llm/models_py/bindings/OpDefs.h | Adds XPU KV-cache reshaping (NSHD) and XPU-only attention inputs field |
| rtp_llm/models_py/bindings/OpDefs.cc | Exposes XPU-only position_ids to Python via pybind |
| rtp_llm/models/base_model.py | Uses resolved device string instead of always cuda |
| rtp_llm/model_loader/weight_manager.py | Avoids CUDA streams + CUDA IPC on XPU; adds XPU sync path |
| rtp_llm/model_loader/loader.py | Uses XPU-aware memory cleanup and XPU memory stats logging |
| rtp_llm/frontend/frontend_app.py | Adds uvicorn import fallback for loop setup API changes |
| rtp_llm/device/device_type.py | Adds DeviceType.Xpu, override mechanism, and cached resolution |
| rtp_llm/device/device_impl.py | Adds XpuImpl and device-agnostic helpers (count/set/name/masks) |
| rtp_llm/device/init.py | Wires XpuImpl into device factory |
| rtp_llm/cpp/utils/TensorDebugUtils.h | Treats XPU tensors as “device tensors” for debug dump behavior |
| rtp_llm/cpp/utils/ErrorCode.h | Adds missing include for XPU toolchain |
| rtp_llm/cpp/pybind/th_utils.h | Fixes CPU tensor checks and makes CHECK_TH_CUDA XPU-aware |
| rtp_llm/cpp/pybind/ComputeInit.cc | Enables exec ctx bindings for XPU builds |
| rtp_llm/cpp/pybind/BUILD | Adds XPU selects for pybind linking and exec ops |
| rtp_llm/cpp/normal_engine/speculative/SpeculativeSampler.cc | Routes device tensors through getTorchDevice(); handles XPU D2H |
| rtp_llm/cpp/normal_engine/speculative/MtpExecutor.cc | Switches to getTorchDevice() and uses maybePinMemory() |
| rtp_llm/cpp/normal_engine/speculative/MtpBatchStreamProcessor.cc | Adds XPU-safe pinning + device routing |
| rtp_llm/cpp/normal_engine/NormalSamplerInputGatherer.cc | Allocates sampler tensors on resolved device |
| rtp_llm/cpp/normal_engine/NormalOutputDispatcher.cc | Moves label tensor to resolved device for loss computation |
| rtp_llm/cpp/normal_engine/NormalModelInputGatherer.cc | Moves multimodal tensors via resolved device routing |
| rtp_llm/cpp/normal_engine/NormalEngine.cc | Adds XPU ctor guard + warmup + emptyCache path |
| rtp_llm/cpp/models/logits_processor/MultiSeqLogitsProcessor.cc | Moves vocab mask to resolved device |
| rtp_llm/cpp/models/logits_processor/BaseLogitsProcessor.cc | Moves vocab mask to resolved device |
| rtp_llm/cpp/models/eplb/ExpertBalancer.cc | Uses getTorchDevice() and maybePinMemory() for buffers |
| rtp_llm/cpp/models/Sampler.cc | Routes tensors to resolved device; adds XPU-specific success handling |
| rtp_llm/cpp/models/PyWrappedModel.h | Adds XPU stream include and disables prefill-CP on XPU |
| rtp_llm/cpp/models/PyWrappedModel.cc | Makes host pinning conditional; uses getTorchDevice() throughout |
| rtp_llm/cpp/models/ModelTypes.cc | Uses maybePinMemory() and routes device tensors via getTorchDevice() |
| rtp_llm/cpp/models/BUILD | Adjusts deps for XPU build (no CUDA graph impl on XPU) |
| rtp_llm/cpp/engine_base/stream/GenerateStream.cc | Adds XPU generator support and treats XPU tensors like CUDA for D2H |
| rtp_llm/cpp/engine_base/WeightsConverter.cc | Copies tensors to resolved device, not hardcoded CUDA |
| rtp_llm/cpp/engine_base/TorchProfiler.h | Enables profiling activities for XPU (Kineto) |
| rtp_llm/cpp/cache/connector/p2p/transfer/tcp/CudaCopyUtil.cc | Routes wrapRawPtr device to getTorchDevice() |
| rtp_llm/cpp/cache/connector/p2p/LayerBlockConverterImpl.h | Treats XPU tensors like device tensors in block info |
| rtp_llm/cpp/cache/connector/memory/KVCacheMemoryConnector.cc | Uses getTorchDevice() for device buffers |
| rtp_llm/cpp/cache/MemoryLayoutStrategy.cc | Treats XPU tensors like device tensors in block info |
| rtp_llm/cpp/cache/MemoryEvaluationHelper.cc | Adds XPU memory accounting via XPUCachingAllocator and device props |
| rtp_llm/cpp/cache/KVCacheManager.cc | Uses resolved device and maybePinMemory() for sync buffers |
| rtp_llm/cpp/cache/BlockPool.cc | Avoids pin_memory on XPU; uses resolved device for GPU buffers |
| rtp_llm/cpp/cache/BlockInfo.h | Documents is_cuda semantics for XPU and notes rename TODO |
| rtp_llm/config/server_config_setup.py | XPU seq_size defaults, device binding via gpu_set_device, and XPU speculative guard |
| rtp_llm/BUILD | Filters wheel requirements for XPU to avoid CUDA-only deps |
| deps/requirements_xpu.txt | Adds dedicated XPU requirements set (torch +xpu pinned) |
| deps/pip.bzl | Adds pip_parse for XPU lockfile and TF_NEED_XPU-gated requirement exports |
| deps/BUILD | Adds lockfile compilation rule for XPU requirements |
| bazel/device_defs.bzl | Adds XPU test env marker |
| arch_config/arch_select.bzl | Adds XPU requirement resolver, filtering/remapping for XPU wheel metadata |
| WORKSPACE | Adds XPU and torch-xpu configure rules; gates XPU wheel install_deps |
| BUILD.pytorch | Adds using_xpu config + links torch XPU libs + XPU python headers |
| BUILD | Adds using_xpu config_setting |
| 3rdparty/gpus/xpu_configure.bzl | Adds oneAPI/SYCL toolchain auto-config + XPU python runtime wiring |
| 3rdparty/gpus/xpu/BUILD.tpl | Adds XPU runtime/header Bazel targets template |
| 3rdparty/gpus/torch_xpu_configure.bzl | Adds torch-xpu site-packages auto-detection with TF_NEED_XPU enforcement |
| 3rdparty/gpus/crosstool/xpu_cc_toolchain_config.bzl.tpl | Defines Bazel cc_toolchain_config for SYCL (xpu_sycl_compile feature) |
| 3rdparty/gpus/crosstool/clang/bin/crosstool_wrapper_driver_xpu.tpl | Adds icx/icpx wrapper that filters unsupported flags + rewrites @params |
| .bazelrc | Adds --config=xpu with toolchain, defines, env, and SYCL target settings |
Comments suppressed due to low confidence (5)
rtp_llm/start_backend_server.py:1
- On XPU,
cuda_device_listcomes fromget_visible_device_list(), which strips sub-device selectors (e.g.0.0) down to0. Writing this back intoZE_AFFINITY_MASKcan silently change device visibility/partitioning on systems using sub-devices. Prefer not rewritingZE_AFFINITY_MASKat all when it is already set, or preserve the original (raw) entries (e.g. keep0.0,1.0) when computing the per-rank mask.
rtp_llm/models_py/standalone/auto_model.py:1 - This condition disables pinned-memory on ROCm as well (since
get_device_string()returns\"cuda\"for both CUDA and ROCm). If pinning is still intended on ROCm, gate this on the resolved device type (e.g.not _is_xpu_device()/maybePinMemory(...)) rather thanself.device == \"cuda\".
rtp_llm/models_py/modules/base/xpu/vllm_xpu_ops.py:1 - Using
assertfor runtime validation is unsafe because asserts can be stripped with Python optimizations (-O), which would turn a required precondition into unchecked behavior. Raise aValueError/RuntimeErrorinstead so the failure mode is consistent in production.
rtp_llm/models_py/bindings/core/CudaSampleOp.cc:1 - The XPU sampling path mutates
params.top_kin-place (viadata_ptr) even though it is conceptually an input. If the caller reuses the sametop_ktensor across steps/requests, this will persistently change the configuration and can lead to incorrect sampling behavior. Prefer computing an internaltop_k_effective(clone/copy) used only inside this function, leavingparams.top_kuntouched.
rtp_llm/start_backend_server.py:1 - The function name
_get_cuda_device_list()no longer matches its behavior (it now returns a generic GPU/XPU-visible device list). Renaming it (e.g._get_gpu_device_list) would avoid confusion and reduce the chance of future CUDA-only assumptions creeping back in.
💡 Add Copilot custom instructions for smarter, more guided reviews. Learn how to get started.
Comment on lines
+76
to
+82
| pip_parse( | ||
| name = "pip_xpu_torch", | ||
| requirements_lock = "@rtp_deps//:requirements_lock_xpu.txt", | ||
| python_interpreter = "/opt/conda310/bin/python3", | ||
| extra_pip_args = PIP_EXTRA_ARGS + ["--extra-index-url=https://download.pytorch.org/whl/xpu"], | ||
| timeout = 3600, | ||
| ) |
Collaborator
AI Code Review - PR #1167Status: BLOCKING Summary: P0/0 · P1/1 · P2/0 · P3/0 Blocking IssuesP1
Checklist Violations (2 fail / 134 total)General Principles Checklist
Strengths
|
6420b3c to
ab6345f
Compare
Contributor
There was a problem hiding this comment.
Pull request overview
Copilot reviewed 93 out of 97 changed files in this pull request and generated 1 comment.
Comments suppressed due to low confidence (8)
rtp_llm/models_py/bindings/core/CudaSampleOp.cc:1
- Reinterpreting an
int32_t*fromparams.top_kasuint32_t*is undefined behavior under C++ strict-aliasing rules and can miscompile under optimization. Useint32_t*(orauto* = params.top_k.data_ptr<int32_t>()) and compare against integer constants without changing the pointer type (cast per-value if you need unsigned semantics).
rtp_llm/models_py/bindings/core/CudaSampleOp.cc:1 - Reinterpreting an
int32_t*fromparams.top_kasuint32_t*is undefined behavior under C++ strict-aliasing rules and can miscompile under optimization. Useint32_t*(orauto* = params.top_k.data_ptr<int32_t>()) and compare against integer constants without changing the pointer type (cast per-value if you need unsigned semantics).
rtp_llm/models_py/modules/base/xpu/vllm_xpu_ops.py:1 - This is a runtime correctness check but uses
assert, which is stripped when Python runs with optimizations (-O), potentially turning a clear error into a later, harder-to-debug failure. Replace with an explicit conditional that raisesValueError/RuntimeErrorwhen eithercu_seqlens_qorcu_seqlens_kis missing.
rtp_llm/models_py/modules/base/xpu/moe_gating.py:1 - These input validations use
assert, which can be disabled withpython -O, removing safety checks and potentially causing silent misbehavior or confusing tensor errors. Prefer raisingValueError(orRuntimeError) with the same message so the checks are always enforced.
rtp_llm/cpp/models/PyWrappedModel.cc:1 - This reintroduces a local pin-memory flag that duplicates the new shared helpers in
ExecOps.h(kPinHostMemory/maybePinMemory). To reduce drift and future backend conditionals, consider switching these call sites to use the shared helpers instead of defining a second compile-time flag in this translation unit.
rtp_llm/device/device_impl.py:1 - Stripping the sub-device portion from
ZE_AFFINITY_MASKchanges semantics when users intentionally target tiles/sub-devices (e.g.0.0). This can cause the parent process to widen affinity unintentionally once it re-exportsZE_AFFINITY_MASK. A safer approach is to preserve the raw mask tokens for re-exporting and only derive a separate 'device index list' fortorch.xpu.set_device()mapping.
rtp_llm/start_backend_server.py:1 - This error message is hard to parse (grammar, line continuation, and the meaning of 'result'). Since this block is now device-agnostic, consider raising
RuntimeErrorwith a clearer message such as: world_size must be <= device_count or divisible by device_count, and include the resolved backend (CUDA/ROCm/XPU) and the visible-device mask env var for easier debugging.
rtp_llm/model_loader/loader.py:1 - This helper now handles both CUDA and XPU, but the name
force_clean_cuda_memorysuggests CUDA-only behavior. Consider renaming to something backend-neutral likeforce_clean_gpu_memory(and optionally keepingforce_clean_cuda_memoryas a backward-compatible alias) to avoid misleading call sites as more backends are added.
Comment on lines
+96
to
+99
| elif normalized in _XPU_PACKAGE_REMAP: | ||
| xpu_reqs.append(_XPU_PACKAGE_REMAP[normalized]) | ||
| else: | ||
| xpu_reqs.append(req) |
Collaborator
AI Code Review - PR #1167Status: BLOCKING Summary: P0/0 · P1/2 · P2/0 · P3/0 Blocking IssuesP1
Checklist Violations (5 fail / 56 total)General Principles Checklist
Strengths
|
ab6345f to
bc6a7d1
Compare
Contributor
There was a problem hiding this comment.
Pull request overview
Copilot reviewed 97 out of 101 changed files in this pull request and generated no new comments.
Comments suppressed due to low confidence (8)
rtp_llm/models_py/bindings/core/CudaOps.cc:1
c10::xpu::getCurrentXPUStream()is typically an XPUStream object, not asycl::queue&. This is likely to fail to compile (or bind to the wrong type) on some PyTorch XPU versions. Prefer retrieving the SYCL queue from the XPU stream explicitly (e.g., get the current XPU stream, then access its underlying queue) and include the appropriate SYCL headers if needed.
rtp_llm/models_py/bindings/common/FusedCopyOp.cc:1- Same issue as in
CudaOps.cc: bindingc10::xpu::getCurrentXPUStream()directly tosycl::queue&is likely incorrect (PyTorch commonly exposes an XPUStream wrapper). Please fetch the queue from the stream explicitly (and ensure required SYCL headers are present) so this code compiles reliably across supported torch-xpu versions.
rtp_llm/device/device_impl.py:1 - Stripping sub-device (tile) selectors from
ZE_AFFINITY_MASKchanges semantics for users who intentionally specify sub-devices (e.g.0.0,0.1). This becomes more problematic becausestart_backend_server.pylater writesZE_AFFINITY_MASKback from this list, effectively overwriting the user’s original mask with the stripped version. Consider returning the raw entries for XPU (preserve0.0) and only deriving a separate 'device indices' list where strictly required fortorch.xpu.set_device()indexing.
rtp_llm/start_backend_server.py:1 - Overwriting
ZE_AFFINITY_MASKin the parent process can unintentionally discard a user-provided affinity (including sub-device selection), and can also affect sibling ranks because the parent env is shared during process spawn. A safer pattern is: (1) only set the mask if it is not already set, or (2) avoid rewritingZE_AFFINITY_MASKaltogether and rely on per-ranktorch.xpu.set_device()in the child, or (3) if you need per-rank masking, set per-rankZE_AFFINITY_MASKto a single entry (not the full list) derived from the rank.
rtp_llm/models_py/utils/arch.py:1 is_xpuis imported but not used in this module (based on the shown file content). Please remove the unused import to avoid lint/type-check noise.
rtp_llm/models_py/bindings/core/CudaSampleOp.cc:1- Reinterpreting an
int32_ttensor buffer asuint32_t*is brittle and makes sentinel/negativetop_kvalues harder to reason about. Prefer keepingtop_kasint32_t*and applying forcing logic without type-punning; also consider using a local copy oftop_krather than mutatingparams.top_kin place if callers expect it to remain an input.
rtp_llm/models_py/bindings/core/CudaBeamSearchOp.cc:1 token_ids_outis allocated and then immediately replaced by the result ofgather, which discards the original allocation. This causes an avoidable allocation. Prefer constructingtoken_ids_outdirectly from thegatherresult (orcopy_into the preallocated buffer if you need to control memory reuse).
rtp_llm/models_py/bindings/core/CudaBeamSearchOp.cc:1token_ids_outis allocated and then immediately replaced by the result ofgather, which discards the original allocation. This causes an avoidable allocation. Prefer constructingtoken_ids_outdirectly from thegatherresult (orcopy_into the preallocated buffer if you need to control memory reuse).
Collaborator
AI Code Review - PR #1167Status: BLOCKING Summary: P0/0 · P1/2 · P2/0 · P3/0 Blocking IssuesP1
Checklist Violations (4 fail / 55 total)General Principles Checklist
Strengths
|
bc6a7d1 to
fe77bda
Compare
Contributor
There was a problem hiding this comment.
Pull request overview
Copilot reviewed 98 out of 102 changed files in this pull request and generated no new comments.
Comments suppressed due to low confidence (6)
rtp_llm/models_py/bindings/core/CudaSampleOp.cc:1
params.logitsis overwritten with softmax probabilities (lines 665-667), but later the “original logits” argmax fallback (line 731) uses the mutated tensor. This breaks the documented intent and can change behavior for degenerate rows (NaN/Inf), where softmax can further distort values. Keep anorig_logitstensor (or avoid mutatingparams.logits) and compute the fallback argmax from the pre-softmax logits.
rtp_llm/models_py/bindings/core/CudaSampleOp.cc:1params.logitsis overwritten with softmax probabilities (lines 665-667), but later the “original logits” argmax fallback (line 731) uses the mutated tensor. This breaks the documented intent and can change behavior for degenerate rows (NaN/Inf), where softmax can further distort values. Keep anorig_logitstensor (or avoid mutatingparams.logits) and compute the fallback argmax from the pre-softmax logits.
rtp_llm/start_backend_server.py:1- The raised exception type (
Exception) and message are hard to interpret (“result: … not support … local gpu”, plus a backslash continuation). PreferValueError/RuntimeErrorand a clearer message that explains the constraint (e.g., “WORLD_SIZE must be <= device_count or divisible by device_count”) and includesworld_size+device_countonce.
rtp_llm/start_backend_server.py:1 _get_cuda_device_list()now returns a generic GPU-visible device list for CUDA/ROCm/XPU. Renaming it to_get_gpu_device_list()(and updating callers accordingly) would prevent confusion now that this function is no longer CUDA-specific.
rtp_llm/models_py/modules/factory/attention/xpu_impl/test/BUILD:1- This comment conflicts with
test_kv_cache_layout.py, which explicitly imports compiled bindings (rtp_llm.ops.compute_ops/KVCache) and only skips at runtime if the import fails. Please clarify the comment (e.g., that tests are “pure Python” but may require the compiledrtp_llm.opsextension at runtime, and will SkipTest otherwise) to avoid misleading future maintainers.
rtp_llm/models_py/bindings/core/CudaSampleOp.cc:1 - The new XPU
sampleGreedypath implements complex sampling semantics (temperature=0 fast-path, per-row greedy viado_sample, top-k/top-p filters, degenerate probability handling, generator handling) but there’s no targeted unit test coverage validating output tokens /successfor these cases. Adding focused C++ tests (similar to existing Sampler tests) for at least temperature==0, top_k==1, top_p filtering, and NaN/Inf/all-zero row behavior would significantly reduce regression risk.
Collaborator
AI Code Review - PR #1167Status: BLOCKING Summary: P0/0 · P1/1 · P2/0 · P3/0 Blocking IssuesP1
Checklist Violations (2 fail / 55 total)General Principles Checklist
Strengths
|
Collaborator
AI Code Review - PR #1167Status: BLOCKING Summary: P0/0 · P1/1 · P2/0 · P3/0 Blocking IssuesP1
Checklist Violations (2 fail / 55 total)General Principles Checklist
Strengths
|
Contributor
There was a problem hiding this comment.
Pull request overview
Copilot reviewed 98 out of 102 changed files in this pull request and generated no new comments.
Comments suppressed due to low confidence (7)
rtp_llm/start_backend_server.py:1
- On XPU this overwrites
ZE_AFFINITY_MASKusingcuda_device_list, which is derived fromget_visible_device_list()and strips sub-device/tile selectors (e.g.0.0). This can silently change the affinity mask and bind the child process to the wrong subdevice(s). A safer approach is to (a) not overwriteZE_AFFINITY_MASKif it is already set, or (b) preserve and re-use the rawZE_AFFINITY_MASKentries when re-exporting, while keeping a separate list of integer device indices fortorch.xpu.set_device().
rtp_llm/device/device_impl.py:1 CUDA_VISIBLE_DEVICEScan legally be set to an empty/whitespace string to indicate 'no devices'. In that casesplit(',')returns[''], and downstreamint('')conversions (or list indexing) will fail in non-obvious places. Consider treating empty/whitespace-only values as an empty device list (or falling back togpu_device_count()==0) by checkingcuda_devices.strip()before splitting.
rtp_llm/models_py/bindings/core/CudaSampleOp.cc:1- Reinterpreting the
int32_ttop_kbuffer asuint32_t*breaks the intended<= 0semantics: negativetop_kvalues (often used to mean 'disabled' / 'no top-k') become large unsigned values, causinghas_top_kto evaluate incorrectly and changing behavior vs CUDA/ROCm. It also introduces strict-aliasing/UB risk. Use anint32_t*(orauto* top_k_ptr = params.top_k.data_ptr<int32_t>();) and perform signed comparisons/clamping onint64_t kderived from that signed value.
rtp_llm/cpp/models/PyWrappedModel.cc:1 - This introduces a second pinned-memory compile-time flag (
kPinHostMem) with duplicated logic, whilertp_llm/models_py/bindings/core/ExecOps.halready defineskPinHostMemoryandmaybePinMemory(...). To prevent the pinned-memory contract from drifting across files/backends, prefer reusingkPinHostMemoryand/ormaybePinMemory(t)here instead of maintaining a separate local preprocessor-based flag.
rtp_llm/cpp/models/PyWrappedModel.cc:1 - This introduces a second pinned-memory compile-time flag (
kPinHostMem) with duplicated logic, whilertp_llm/models_py/bindings/core/ExecOps.halready defineskPinHostMemoryandmaybePinMemory(...). To prevent the pinned-memory contract from drifting across files/backends, prefer reusingkPinHostMemoryand/ormaybePinMemory(t)here instead of maintaining a separate local preprocessor-based flag.
rtp_llm/start_backend_server.py:1 - This error message is hard to interpret ("result:" and "not support"), and it still hard-codes "local gpu" even when the selected backend is XPU. Consider raising a
ValueError/RuntimeErrorwith a clearer description like: world_size must be <= device_count or divisible by device_count, and include the resolved device type (cuda/xpu) plus the relevant visible-device env var name (CUDA_VISIBLE_DEVICES vs ZE_AFFINITY_MASK).
rtp_llm/models_py/bindings/core/CudaSampleOp.cc:1 - The XPU sampling path is a new, substantial behavioral fork from the CUDA/ROCm kernels (temperature==0 fast path, top_k/top_p filtering, do_sample overrides, and degenerate-row handling). Please add unit coverage that compares XPU sampling behavior against the existing CUDA/ROCm expectations for a few representative cases (e.g., negative/zero top_k semantics, top_p edge values, mixed temperature rows including 0, and do_sample=false rows). This will also catch regressions like the signed/unsigned
top_khandling issue.
Contributor
There was a problem hiding this comment.
Pull request overview
Copilot reviewed 98 out of 102 changed files in this pull request and generated 1 comment.
Comments suppressed due to low confidence (11)
rtp_llm/device/device_impl.py:1
gpu_is_available()checkstorch.xpu.device_count()>0for XPU but does not checktorch.cuda.device_count()>0for CUDA/ROCm. WithCUDA_VISIBLE_DEVICESmasking all GPUs,torch.cuda.is_available()can still be true whiledevice_count()==0, which makes call sites take GPU paths and later fail. Make CUDA availability symmetric by including atorch.cuda.device_count() > 0check (and optionally treat empty/invalid masks as unavailable).
rtp_llm/device/device_impl.py:1get_visible_device_list()can return invalid/empty entries whenZE_AFFINITY_MASKorCUDA_VISIBLE_DEVICEScontains whitespace or is set to an empty string (e.g.CUDA_VISIBLE_DEVICES=''produces['']). This can later crash when converting toint(e.g. inXpuImpl.get_device_id()), or produce confusing device pinning. Strip whitespace and filter out empty tokens for both masks before returning the list.
rtp_llm/start_backend_server.py:1- The helper is now GPU-backend agnostic (CUDA/ROCm/XPU) but is still named
_get_cuda_device_list(). Rename it (and its local variable usages likecuda_device_list) to something backend-neutral (e.g._get_gpu_device_list) to avoid confusion and reduce the chance of future CUDA-only assumptions creeping back in.
rtp_llm/start_backend_server.py:1 - This error message is hard to interpret ("result:") and is now used for non-CUDA backends as well, but still says "local gpu". Consider raising a
ValueError/RuntimeErrorwith a clearer explanation, e.g. that WORLD_SIZE must be either <= device_count or a multiple of device_count, and include the resolved device type (cuda/xpu) in the message. Also avoid the\\line continuation inside the f-string to prevent odd whitespace/newlines in the final message.
rtp_llm/models_py/modules/base/xpu/vllm_xpu_ops.py:1 - This uses a Python
assertfor required runtime validation. Asserts can be disabled withpython -O, which would turn this into undefined behavior later in the function. Replace this with an explicit check that raises a deterministic exception (e.g.ValueError/RuntimeError) whencu_seqlens_qorcu_seqlens_kis missing.
rtp_llm/models_py/bindings/core/CudaSampleOp.cc:1 - Reinterpreting
int32_t*asuint32_t*can violate C++ strict-aliasing rules and is unnecessary here since you only compare/assign small positive values. Prefer keepingint32_t*and doing comparisons/assignments in signed integers (or copy into astd::vector<int32_t>if you need a host-side view). This avoids undefined behavior on compilers with aggressive aliasing optimizations.
rtp_llm/models_py/bindings/core/CudaSampleOp.cc:1 - Reinterpreting
int32_t*asuint32_t*can violate C++ strict-aliasing rules and is unnecessary here since you only compare/assign small positive values. Prefer keepingint32_t*and doing comparisons/assignments in signed integers (or copy into astd::vector<int32_t>if you need a host-side view). This avoids undefined behavior on compilers with aggressive aliasing optimizations.
rtp_llm/models_py/bindings/core/CudaBeamSearchOp.cc:1 - The XPU beam search fallback introduces substantial new decoding logic (topk flattening, history gather, scatter write at
sequence_lengths, overflow masking). There should be unit tests covering at least: (1) correct token history propagation across steps, (2) correctbeam_indices_outmapping, (3)sequence_lengthsincrement behavior, and (4) overflow (sequence_lengths >= max_seq_len) handling. Adding a small deterministic test (CPU/XPU-tensor based) would help prevent subtle regressions.
rtp_llm/models_py/bindings/core/CudaBeamSearchOp.cc:1 - The XPU beam search fallback introduces substantial new decoding logic (topk flattening, history gather, scatter write at
sequence_lengths, overflow masking). There should be unit tests covering at least: (1) correct token history propagation across steps, (2) correctbeam_indices_outmapping, (3)sequence_lengthsincrement behavior, and (4) overflow (sequence_lengths >= max_seq_len) handling. Adding a small deterministic test (CPU/XPU-tensor based) would help prevent subtle regressions.
rtp_llm/cpp/pybind/ComputeInit.cc:1 - This preprocessor block is empty and has no effect. It should be removed to reduce noise and avoid confusion about missing conditional code.
rtp_llm/cpp/models/PyWrappedModel.cc:1 - This duplicates the project-wide pinned-memory policy that was already introduced in
ExecOps.h(kPinHostMemoryandmaybePinMemory). To keep pinning behavior consistent across the codebase, prefer reusingkPinHostMemory/maybePinMemoryhere instead of defining a localkPinHostMemthat could drift from the canonical definition.
Comment on lines
+19
to
+36
| def _find_repo_root(start): | ||
| cur = os.path.abspath(start) | ||
| while True: | ||
| if os.path.exists(os.path.join(cur, "arch_config", "arch_select.bzl")) and \ | ||
| os.path.exists(os.path.join(cur, "deps", "requirements_lock_xpu.txt")): | ||
| return cur | ||
| parent = os.path.dirname(cur) | ||
| if parent == cur: | ||
| raise RuntimeError( | ||
| "Could not locate repo root containing arch_config/arch_select.bzl " | ||
| "and deps/requirements_lock_xpu.txt" | ||
| ) | ||
| cur = parent | ||
|
|
||
|
|
||
| _REPO_ROOT = _find_repo_root(os.path.dirname(__file__)) | ||
| _BZL_PATH = os.path.join(_REPO_ROOT, "arch_config", "arch_select.bzl") | ||
| _LOCK_PATH = os.path.join(_REPO_ROOT, "deps", "requirements_lock_xpu.txt") |
Collaborator
AI Code Review - PR #1167Status: BLOCKING Summary: P0/0 · P1/1 · P2/0 · P3/0 Blocking IssuesP1
Checklist Violations (2 fail / 55 total)General Principles Checklist
Strengths
|
aslanxie
added a commit
to aslanxie/rtp-llm-xpu
that referenced
this pull request
Jul 11, 2026
Under RTP_LLM_DEVICE_TYPE=xpu, a missing rtp_llm.ops / xpu consumer import (broken build or pybind config) is a real regression, not an environment gap. Add _skip_or_fail_on_missing_import() helper that fails the test on XPU when the package-level import fails, and only skips when not targeting XPU. Addresses PR alibaba#1167 review comment.
Comment on lines
+228
to
+232
| auto safe_mask = (write_pos < max_seq_len).squeeze(-1); // [batch, beam_out] | ||
| write_pos = write_pos.clamp(0, max_seq_len - 1); | ||
| // Write tokens only for beams that have not overflowed (device-side where). | ||
| auto safe_ids = output_ids.where(safe_mask, token_ids_out.gather(2, write_pos).squeeze(-1)); | ||
| token_ids_out.scatter_(2, write_pos, safe_ids.unsqueeze(-1)); |
Comment on lines
+548
to
+554
| auto top_k_force = reinterpret_cast<uint32_t*>(params.top_k.data_ptr<int32_t>()); | ||
| for (int64_t b = 0; b < batch_size; b++) { | ||
| if (params.temperature.data_ptr<float>()[b] == 0.0f) { | ||
| top_k_force[b] = 1; | ||
| } | ||
| } | ||
| } |
Comment on lines
+648
to
+651
| // 3. Top-k=1 fast path (greedy argmax) | ||
| auto top_k_ptr = reinterpret_cast<uint32_t*>(params.top_k.data_ptr<int32_t>()); | ||
| if (std::all_of(top_k_ptr, top_k_ptr + batch_size, [](uint32_t t) { return t == 1; }) | ||
| && !params.output_all_probs.has_value()) { |
Comment on lines
+673
to
+675
| auto filtered_probs = probs_t.clone(); | ||
| bool has_top_k = !std::all_of(top_k_ptr, top_k_ptr + batch_size, [](uint32_t t) { return t <= 0; }); | ||
| if (has_top_k) { |
Comment on lines
+69
to
+73
| # XPU wheel installation is gated on TF_NEED_XPU (set by --config=xpu) via the | ||
| # xpu_pip_gate repo declared in pip_deps(). On non-XPU containers and plain | ||
| # `bazel sync`, install_deps() is a no-op, so the pip_xpu_torch_* whl_library | ||
| # repos are never declared and XPU-only pins (e.g. scikit-learn==1.8.0, which | ||
| # Requires-Python>=3.11) are never resolved under Python 3.10. |
Collaborator
AI Code Review - PR #1167Status: BLOCKING Summary: P0/0 · P1/1 · P2/0 · P3/0 Blocking IssuesP1
Checklist Violations (2 fail / 68 total)General Principles Checklist
Strengths
|
aslanxie
added a commit
to aslanxie/rtp-llm-xpu
that referenced
this pull request
Jul 12, 2026
The previous fix (skip-vs-fail based on RTP_LLM_DEVICE_TYPE=xpu) never triggered in real CI: the three xpu-tagged py_test targets never set that env var, so _ON_XPU stayed false and import failures kept being skipped. - BUILD: wire env = device_test_envs() (//bazel:device_defs.bzl) into the xpu-tagged py_test targets, injecting the repo's existing TEST_USING_DEVICE=XPU convention (select on @//:using_xpu). - Extract the skip-vs-fail policy into a shared, stdlib-only _import_guard.py (is_xpu_test_env() checks both TEST_USING_DEVICE=XPU and RTP_LLM_DEVICE_TYPE=xpu; skip_or_fail_on_missing_import() fails on XPU, skips elsewhere), used by all three existing test files instead of duplicated per-file logic. - Add test_import_guard.py: focused regression tests for the helper itself (no rtp_llm/torch import, so it always runs) covering all four fail/skip x xpu/non-xpu combinations end-to-end. Addresses PR alibaba#1167 review comment.
Comment on lines
48
to
+50
| #define CHECK_TH_CUDA(x) TORCH_CHECK(x.is_cuda(), #x " must be a CUDA tensor") | ||
| #define CHECK_CPU(x) TORCH_CHECK(!x.is_cuda(), #x " must be a CPU tensor") | ||
| #endif |
Comment on lines
+517
to
+519
| if (params.do_sample.has_value()) { | ||
| auto top_k_ptr = params.top_k.data_ptr<int32_t>(); | ||
| bool any_greedy = false; |
81dfd62 to
178fff6
Compare
Comment on lines
+1202
to
+1207
| if _is_xpu_device(): | ||
| xpu_mask = os.environ.get("ZE_AFFINITY_MASK", None) | ||
| if xpu_mask is not None and xpu_mask.strip(): | ||
| # ZE_AFFINITY_MASK may use dotted format "0.0,1.0" for sub-devices; | ||
| # extract only the device portion (before the dot) for indexing. | ||
| return [e.split(".")[0] for e in xpu_mask.split(",")] |
| params.logits[b].copy_(saved_logits[b]); | ||
| } | ||
| } | ||
| saved_logits.reset(); |
| // are 0, so mark those rows top_k=1 to make the per-row top_k filter pick | ||
| // the argmax token deterministically (instead of softmax/multinomial). | ||
| { | ||
| auto top_k_force = reinterpret_cast<uint32_t*>(params.top_k.data_ptr<int32_t>()); |
| } | ||
|
|
||
| // 3. Top-k=1 fast path (greedy argmax) | ||
| auto top_k_ptr = reinterpret_cast<uint32_t*>(params.top_k.data_ptr<int32_t>()); |
| // which must remain the original (pre-filter) softmax source for | ||
| // return_original_all_probs and cum_log_probs. | ||
| auto filtered_probs = probs_t.clone(); | ||
| bool has_top_k = !std::all_of(top_k_ptr, top_k_ptr + batch_size, [](uint32_t t) { return t <= 0; }); |
- Add --config=xpu to .bazelrc; new xpu_configure.bzl (479 lines) that auto-detects the Intel oneAPI/SYCL toolchain and emits a Bazel C++ toolchain (xpu_cc_toolchain_config.bzl.tpl, crosstool_wrapper_driver_xpu.tpl) - Add torch_xpu_configure.bzl to locate PyTorch XPU headers/libs - Extend arch_select.bzl to select XPU platform; update BUILD, BUILD.pytorch, bazel/defs.bzl, bazel/device_defs.bzl with USING_XPU compile-time flag - Add pip_xpu_torch Bazel repo rule for XPU Python deps (requirements_xpu.txt, requirements_lock_xpu.txt); gate the rule behind TF_NEED_XPU so that `bazel sync` on non-XPU Python 3.10 containers never resolves XPU-only pins (e.g. scikit-learn==1.8.0 requires Python>=3.11) - NormalEngine: include XPUCachingAllocator; add DeviceGuard for ctor thread; enable prefill warm-up and emptyCache() via XPU caching allocator - PyWrappedModel: replace all hard-coded torch::kCUDA / .cuda() calls with getTorchDevice(); add kPinHostMem=false on XPU (pinned memory unsupported); guard kv_cache_block_id_host and block_id pin_memory() calls accordingly - KVCacheManager, BlockPool, MemoryEvaluationHelper: add XPU memory evaluation and KV cache layout support - Sampler, ModelTypes, ComputeInit: add XPU device branching - New bindings/xpu/: RegisterXpuBaseBindings.hpp registers ~40 ops as pure-PyTorch fallbacks on XPU (RMSNorm, LayerNorm, attention, sampling, KV cache ops, write_cache_store) - New modules/base/xpu/: activation.py, norm.py (RMSNorm/LayerNorm), moe_gating.py, select_topk.py, vllm_xpu_ops.py (flash_attn_varlen wrapper with SDPA fallback when vllm-xpu-kernels FA2 is absent) - New modules/factory/attention/xpu_impl/vllm_flash_attn.py (906 lines): full paged-KV-cache attention for prefill and decode using flash_attn_varlen; rejects unsupported RoPE styles (Yarn/Llama3/Mrope/Glm2/DynamicNTK); hoist decode hash computation to step boundary (layer_idx wrap-around detection) instead of recomputing per layer - New modules/factory/linear/impl/xpu/f16_linear.py and modules/factory/fused_moe/impl/xpu/ stubs - Add unit tests: test_kv_cache_layout.py, test_no_rope.py - Add XpuImpl class extending GpuImpl: get_device_id() with fallback to LOCAL_RANK + ZE_AFFINITY_MASK; _get_mem_info() via torch.xpu.mem_get_info; guard empty ZE_AFFINITY_MASK to avoid spurious device-0 targeting - device_type.py: add XPU enum value and auto-detection logic - server_config_setup.py: device selection via ZE_AFFINITY_MASK - start_backend_server.py: XPU startup path - model_loader/loader.py, weight_manager.py: XPU weight loading; use gpu_is_available() + XPU memory APIs for fastsafetensor progress log - frontend_app.py, base_model.py: XPU path plumbing
178fff6 to
ad1c03b
Compare
Comment on lines
+1203
to
+1207
| xpu_mask = os.environ.get("ZE_AFFINITY_MASK", None) | ||
| if xpu_mask is not None and xpu_mask.strip(): | ||
| # ZE_AFFINITY_MASK may use dotted format "0.0,1.0" for sub-devices; | ||
| # extract only the device portion (before the dot) for indexing. | ||
| return [e.split(".")[0] for e in xpu_mask.split(",")] |
| import torch | ||
|
|
||
| from rtp_llm.device.device_type import DeviceType, get_device_type, is_cuda, is_hip | ||
| from rtp_llm.device.device_type import DeviceType, get_device_type, is_cuda, is_hip, is_xpu |
Collaborator
AI Code Review - PR #1167Status: BLOCKING Summary: P0/0 · P1/4 · P2/0 · P3/0 Blocking IssuesP1
Checklist Violations (10 fail / 137 total)General Principles Checklist
RTP-LLM Checklist
Strengths
|
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Overview
Add Intel GPU (XPU) inference support to RTP-LLM, reusing vllm-xpu-kernels to optimize performance on Intel GPU.
Guiding principles:
--config=xpu— DO NOT break existing code logicChanges
1. Build Infrastructure
xpu_configure.bzl,xpu_cc_toolchain_config.bzl.tpl,crosstool_wrapper_driver_xpu.tpl)xpu_configure.bzl(analogous tocuda_configure), plustorch_xpu_configure.bzlto locate PyTorch XPU headers/libs.bazelrc --config=xpupreset with-DUSING_XPU=1and oneAPI compiler flagspip_xpu_torchBazel repo rule gated behindTF_NEED_XPU, sobazel syncon non-XPU Python 3.10 containers never resolves XPU-only pinsrequirements_xpu.txt) and dependency lockfile (requirements_lock_xpu.txt)2. C++ Device Generalization
NormalEngine,KVCacheManager,BlockPool,MemoryEvaluationHelper,Sampler,ModelTypes,ComputeInitselect()branches in BUILD, BUILD.pytorch,bazel/defs.bzl,bazel/device_defs.bzlXPUCachingAllocatorintegration withDeviceGuardfor ctor thread; prefill warm-up andemptyCache()support3. Python Device & Attention
XpuImplextendsGpuImpl,device_type.py)RegisterXpuBaseBindings.hpp(RMSNorm, LayerNorm, attention, sampling, KV cache ops, write_cache_store)4. Module Factories & Server Integration
f16_linear.py, fused MoE stubs)torch::kCUDA/.cuda()calls replaced withgetTorchDevice(); pinned memory disabled on XPU (kPinHostMem=false)server_config_setup.py,start_backend_server.py) — device selection viaZE_AFFINITY_MASKloader.py,weight_manager.py,frontend_app.py,base_model.py)test_kv_cache_layout.py,test_no_rope.pyTest Environment
How to Build