[TENT] Add Intel XPU platform to the Transfer Engine by XinyuYe-Intel · Pull Request #4030 · kvcache-ai/Mooncake

XinyuYe-Intel · 2026-09-11T06:09:44Z

Description

Add an XpuPlatform to the TENT runtime so Mooncake can stage transfers through Intel XPU (GPU) device memory. XpuPlatform extends CpuPlatform: host DRAM allocation, NUMA probing, and host<->host copies are inherited unchanged, and only the XPU-device-aware operations are overridden.

Implementation is a native, direct-link SYCL build: the platform sources include <sycl/sycl.hpp> and link libsycl directly via a private XpuSyclBackend (device enumeration with GPU-preferred/any-device fallback, sycl::malloc_device USM allocation, an interior-pointer registry so chunked staging addresses classify correctly, and VRAM<->host memcpy). Because SYCL requires the Intel DPC++ compiler, USE_XPU=ON now requires USE_TENT=ON and fails fast unless the build is configured with icpx (CXX=icpx); -fsycl is applied to platform_xpu and propagated to everything that links it.

runtime: MTYPE_XPU memory type + USE_XPU loader branch
platform: XpuPlatform (probe/allocate/free/copy/getMemoryType/ getLocation) backed by XpuSyclBackend
build: USE_XPU option/gate in common.cmake; platform_xpu with -fsycl
test: tent_xpu_platform_test drives the real backend (alloc -> H2D ->
D2H -> byte-equality -> free, host/interior pointer classification), skipping when no SYCL device is present; runs against the OpenCL CPU device on GPU-less hosts
docker/xpu.Dockerfile: build image based on intel/pytorch:xpu with the oneAPI DPC++ compiler added, so it both builds and runs the platform

Test on Intel B60 platform.

Module

Type of Change

How Has This Been Tested?

Built with the Intel DPC++ compiler and ran the XPU acceptance test on an Intel
B60 platform. The test performs a full alloc -> H2D -> D2H -> byte-equality ->
free cycle against the real oneAPI SYCL backend and validates host/interior
pointer classification; it skips gracefully when no SYCL device is visible and
runs against the OpenCL CPU device on GPU-less hosts.

Test commands:

# Configure with the Intel DPC++ compiler and enable the XPU platform
CXX=icpx cmake -S . -B build-xpu -DUSE_TENT=ON -DUSE_XPU=ON
cmake --build build-xpu --target tent_xpu_platform_test

# Run the acceptance test
ctest --test-dir build-xpu -R tent_xpu_platform_test --output-on-failure
# or directly:
./build-xpu/mooncake-transfer-engine/tent/tests/tent_xpu_platform_test

Test results:

Unit tests pass
Integration tests pass (if applicable)
Manual testing done (alloc -> H2D -> D2H -> byte-equality -> free and
pointer classification verified on Intel B60)

Hardware validation on Intel Arc Pro B60 (Level-Zero)

CI runners have no Intel GPU, so tent_xpu_platform_test there skips the device path (or falls back to the OpenCL CPU device). I validated the real GPU path on an Intel Arc Pro B60 node.

Environment

Node with Intel(R) Arc(TM) Pro B60 Graphics, oneAPI DPC++ compiler icpx 2026.1.1.
Base image intel/pytorch:xpu + oneAPI compiler (as in docker/xpu.Dockerfile).

sycl-ls on the node:

[level_zero:gpu][level_zero:0] Intel(R) Arc(TM) Pro B60 Graphics 20.1.0
[opencl:gpu][opencl:1]         Intel(R) Arc(TM) Pro B60 Graphics OpenCL 3.0 NEO
[opencl:cpu][opencl:0]         Intel(R) Xeon(R) Platinum 8480+

Build

CXX=icpx cmake -S . -B build-xpu -G Ninja \
  -DUSE_TENT=ON -DUSE_XPU=ON -DBUILD_UNIT_TESTS=ON \
  -DWITH_STORE=OFF -DWITH_STORE_RUST=OFF -DWITH_P2P_STORE=OFF
cmake --build build-xpu --target tent_xpu_platform_test   # 88/88, exit 0

Test — full suite (default selector): 5/5 PASSED

[==========] 5 tests from XpuPlatformTest
[  PASSED  ] 5 tests.

Test — pinned to the discrete GPU so the CPU device cannot be chosen, proving real VRAM sycl::malloc_device + H2D/D2H + byte-equality round-trip and interior-pointer classification run on the B60:

ONEAPI_DEVICE_SELECTOR=level_zero:gpu \
  ./build-xpu/.../tent_xpu_platform_test \
  --gtest_filter=XpuPlatformTest.DeviceAllocCopyFreeRoundTrip:XpuPlatformTest.InteriorPointerClassifiesAsXpu
# [  PASSED  ] 2 tests.

Notes for reviewers

CI coverage: tent_xpu_platform_test is a no-op/skip on GPU-less CI runners (it GTEST_SKIPs when no SYCL device is visible, or runs against the OpenCL CPU device). Real GPU coverage is the B60 run above — a green CI run does not by itself imply GPU-path coverage.
Build requires icpx for the whole tree: because -fsycl propagates from platform_xpu to everything that links transfer_engine, USE_XPU=ON requires the entire build to be configured with CXX=icpx. This is already enforced by the gate in common.cmake; don't mix g++ for the rest of the tree.

Checklist

I have performed a self-review of my own code
I have formatted my code using ./scripts/code_format.sh
I have run pre-commit on the files changed in this PR and all hooks pass
I have updated the documentation (if applicable)
I have added tests to prove my changes are effective
For changes >500 LOC: I have filed an RFC issue

AI Assistance Disclosure

AI tools were used (specify below)

GitHub Copilot was used to assist with scaffolding the XpuPlatform /
XpuSyclBackend sources, CMake wiring, and the acceptance test. The human
submitter has reviewed every changed line and can defend the change end-to-end.

Copilot

🟡 Changes recommended

Critical XPU routing, staging, and location-handling issues remain unresolved.

Once you've addressed the issues Copilot identified, you can request another Copilot review.

Pull request overview

Adds Intel XPU support to TENT through a direct-link SYCL backend, build integration, runtime support, tests, and Docker tooling.

Changes:

Adds XPU allocation, copying, pointer classification, and topology support.
Adds USE_XPU/DPC++ build wiring and runtime integration.
Adds acceptance tests and an Intel XPU build image.

File summaries

File	Change
`mooncake-transfer-engine/tent/tests/xpu/xpu_platform_test.cpp`	Tests XPU allocation, transfers, and pointer classification.
`mooncake-transfer-engine/tent/tests/CMakeLists.txt`	Registers the XPU test target.
`mooncake-transfer-engine/tent/src/runtime/platform.cpp`	Integrates XPU platform loading.
`mooncake-transfer-engine/tent/src/platform/xpu/xpu_sycl_backend.h`	Implements SYCL device, USM, and transfer handling.
`mooncake-transfer-engine/tent/src/platform/xpu/xpu_platform.cpp`	Implements XPU platform operations.
`mooncake-transfer-engine/tent/src/platform/xpu/CMakeLists.txt`	Configures SYCL compilation and linking.
`mooncake-transfer-engine/tent/src/platform/CMakeLists.txt`	Registers the XPU platform target.
`mooncake-transfer-engine/tent/include/tent/runtime/platform.h`	Adds XPU memory-type support.
`mooncake-transfer-engine/tent/include/tent/platform/xpu.h`	Declares the XPU platform interface.
`mooncake-common/common.cmake`	Adds XPU options and compiler validation.
`docker/xpu.Dockerfile`	Provides an Intel XPU build/runtime environment.

Review details

Suppressed comments (4)

mooncake-transfer-engine/tent/src/platform/xpu/xpu_platform.cpp:120

getMemoryType recognizes only allocations recorded by XpuSyclBackend. A normal TENT registerLocalMemory call for framework-owned Intel USM (for example, a oneAPI/PyTorch XPU buffer) is therefore reported as CPU, and copy falls through to host memcpy instead of performing H2D/D2H. Add an external-USM classification/registration path that also tracks interior ranges, or explicitly reject/document unsupported external buffers.

    if (backend().isDevicePtr(addr)) return MTYPE_XPU;

mooncake-transfer-engine/tent/src/platform/xpu/xpu_platform.cpp:74

LocationParser also uses index -1 for malformed values such as xpu:foo; this conditional then treats that error exactly like an omitted ordinal and silently allocates on device 0. Distinguish the exact bare xpu case from an invalid xpu:<ordinal> and return InvalidArgument for the latter.

        int device_index = location.index() >= 0 ? location.index() : 0;

mooncake-transfer-engine/tent/src/platform/xpu/xpu_sycl_backend.h:171

The singleton mutex is held across the blocking queue.memcpy(...).wait(), so all H2D/D2H copies serialize globally, including copies targeting different XPU devices. This prevents concurrent TENT staging workers from overlapping device transfers and can become a major throughput bottleneck; narrow the critical section or use per-device synchronization while preserving allocation lifetime safety.

    int copy(void *dst, const void *src, size_t len, bool to_host) {
        std::lock_guard<std::mutex> lock(mu_);
        const void *dev = to_host ? src : dst;
        const Alloc *a = findLocked(dev);
        if (!a) return 1;  // device side must be a known (interior) allocation
        const uintptr_t start = reinterpret_cast<uintptr_t>(dev);
        if (start + len > a->base + a->size) return 1;  // would run past end
        try {
            queues_[a->device].memcpy(dst, src, len).wait();

mooncake-transfer-engine/tent/src/platform/xpu/xpu_sycl_backend.h:65

The fallback enumerates every SYCL device, but sycl::malloc_device is only supported when the device has sycl::aspect::usm_device_allocations. A visible host/plugin device without that aspect will make probe() report an XPU and then fail every allocation (the test will not skip), despite the documented any-device fallback. Filter candidates by the required USM aspect, or mark unsupported devices unavailable before exposing them through deviceCount().

            for (const auto &d : sycl::device::get_devices()) {

Files reviewed: 11/11 changed files
Comments generated: 6
Review effort level: Lite

💡 Add a code-review agent skill or configure MCP servers for context-aware, tailored reviews. Learn more in the docs.

XinyuYe-Intel · 2026-09-11T08:59:00Z

+            initialized_ = true;
+            return 0;
+        } catch (const sycl::exception &) {
+            return 1;


Fixed. The catch block now calls queues_.clear() before returning failure, so a throw partway through enumeration never leaves a partially-populated queue vector. In the same area, device selection (both the GPU loop and the fallback) now requires sycl::aspect::usm_device_allocations, so we never pick a device that cannot service the USM allocations this backend depends on. Fixed in 1d61971.

staryxchen

LGTM overall. please address the AI's comments.

alogfans · 2026-09-11T07:27:14Z

It seems that end-to-end verification cannot be proceed with this PR only. The runtime needs to add support of proposed transport.

Add an XpuPlatform to the TENT runtime so Mooncake can stage transfers through Intel XPU (GPU) device memory. XpuPlatform extends CpuPlatform: host DRAM allocation, NUMA probing, and host<->host copies are inherited unchanged, and only the XPU-device-aware operations are overridden. Implementation is a native, direct-link SYCL build: the platform sources include <sycl/sycl.hpp> and link libsycl directly via a private XpuSyclBackend (device enumeration with GPU-preferred/any-device fallback, sycl::malloc_device USM allocation, an interior-pointer registry so chunked staging addresses classify correctly, and VRAM<->host memcpy). Because SYCL requires the Intel DPC++ compiler, USE_XPU=ON now requires USE_TENT=ON and fails fast unless the build is configured with icpx (CXX=icpx); -fsycl is applied to platform_xpu and propagated to everything that links it. - runtime: MTYPE_XPU memory type + USE_XPU loader branch - platform: XpuPlatform (probe/allocate/free/copy/getMemoryType/ getLocation) backed by XpuSyclBackend - build: USE_XPU option/gate in common.cmake; platform_xpu with -fsycl - test: tent_xpu_platform_test drives the real backend (alloc -> H2D -> D2H -> byte-equality -> free, host/interior pointer classification), skipping when no SYCL device is present; runs against the OpenCL CPU device on GPU-less hosts - docker/xpu.Dockerfile: build image based on intel/pytorch:xpu with the oneAPI DPC++ compiler added, so it both builds and runs the platform Test on Intel B60 platform.

XinyuYe-Intel · 2026-09-11T09:03:05Z

LGTM overall. please address the AI's comments.

All addressed, thanks.

XinyuYe-Intel · 2026-09-11T09:06:03Z

It seems that end-to-end verification cannot be proceed with this PR only. The runtime needs to add support of proposed transport.

Yes, you are right, XpuTransport will be added in the following PR as in this RFC.

stmatengss · 2026-09-12T09:22:01Z

Please resolve the format issues: https://github.com/kvcache-ai/Mooncake/actions/runs/34581790883/job/103206817476?pr=4030 Thanks!

Copilot AI lite review requested due to automatic review settings September 11, 2026 06:09

XinyuYe-Intel requested review from 00fish0, alogfans, chestnut-Q, doujiang24, dtcccc and staryxchen as code owners September 11, 2026 06:09

github-actions Bot added run-ci Transfer Engine Common labels Sep 11, 2026

Copilot started reviewing on behalf of XinyuYe-Intel September 11, 2026 06:17 View session

Copilot AI reviewed Sep 11, 2026

View reviewed changes

staryxchen reviewed Sep 11, 2026

View reviewed changes

XinyuYe-Intel force-pushed the feat/xpu-sycl-platform branch from 1093c00 to 1d61971 Compare September 11, 2026 08:58

Navigation Menu

Uh oh!

[TENT] Add Intel XPU platform to the Transfer Engine - #4030

[TENT] Add Intel XPU platform to the Transfer Engine#4030
XinyuYe-Intel wants to merge 1 commit into
kvcache-ai:mainfrom
XinyuYe-Intel:feat/xpu-sycl-platform

XinyuYe-Intel commented Sep 11, 2026 •

edited

Loading

Uh oh!

Copilot AI left a comment

Uh oh!

Uh oh!

Uh oh!

Uh oh!

XinyuYe-Intel Sep 11, 2026

Uh oh!

Uh oh!

Uh oh!

staryxchen left a comment

Uh oh!

alogfans commented Sep 11, 2026

Uh oh!

XinyuYe-Intel commented Sep 11, 2026

Uh oh!

XinyuYe-Intel commented Sep 11, 2026

Uh oh!

stmatengss commented Sep 12, 2026

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

5 participants

Uh oh!

Conversation

XinyuYe-Intel commented Sep 11, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Description

Module

Type of Change

How Has This Been Tested?

Hardware validation on Intel Arc Pro B60 (Level-Zero)

Checklist

AI Assistance Disclosure

Uh oh!

Copilot AI left a comment

Choose a reason for hiding this comment

🟡 Changes recommended

Suppressed comments (4)

Uh oh!

Uh oh!

Uh oh!

Uh oh!

XinyuYe-Intel Sep 11, 2026

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Uh oh!

staryxchen left a comment

Choose a reason for hiding this comment

Uh oh!

alogfans commented Sep 11, 2026

Uh oh!

XinyuYe-Intel commented Sep 11, 2026

Uh oh!

XinyuYe-Intel commented Sep 11, 2026

Uh oh!

stmatengss commented Sep 12, 2026

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

5 participants

XinyuYe-Intel commented Sep 11, 2026 •

edited

Loading