From 9e8b40ea7793439eeb51f9cbb61e13b0ad1f3666 Mon Sep 17 00:00:00 2001 From: Ashwini Rathi Date: Tue, 18 Aug 2026 17:18:24 +0300 Subject: [PATCH 1/5] [XPU][CI] key persistent JIT kernel cache by image content ID Nightly image rebuilds swap runtime libs the cached .so files link against (e.g. libsycl.so.8 -> libsycl.so.9), and Triton keys on source hash only, so the old gpu-mask cache dlopens stale .so's and jobs die with `OSError: libsycl.so.N: cannot open shared object file`. Include the `docker image inspect` ID in the cache path so a new image gets a fresh dir, and prune sibling caches for the same GPU but a different image via busybox (files are root-owned). Explicit XPU_KERNEL_CACHE_DIR still honored verbatim; over-cap reset switched to the same busybox helper for the same root-owned reason. --- scripts/ci/xpu/xpu_ci_start_container.sh | 50 ++++++++++++++++++++++-- 1 file changed, 46 insertions(+), 4 deletions(-) diff --git a/scripts/ci/xpu/xpu_ci_start_container.sh b/scripts/ci/xpu/xpu_ci_start_container.sh index f4e49bead091..2b3344781016 100755 --- a/scripts/ci/xpu/xpu_ci_start_container.sh +++ b/scripts/ci/xpu/xpu_ci_start_container.sh @@ -109,18 +109,60 @@ elif [[ -r "${HF_TOKEN_FILE}" ]]; then HF_TOKEN_VALUE=$(cat "${HF_TOKEN_FILE}") fi -# Persistent JIT kernel cache (Triton/Inductor/NEO/SYCL) keyed by GPU mask. -# Cold JIT compile can push test_xpu_basic past its 1200s timeout on B580. -XPU_KERNEL_CACHE_HOST="${XPU_KERNEL_CACHE_DIR:-${HOME}/.cache/sglang-xpu-ci/kernel-cache-gpu${ZE_AFFINITY_MASK:-shared}}" +# Persistent JIT kernel cache (Triton/Inductor/NEO/SYCL) keyed by GPU mask +# AND the resolved image content ID. Nightly image rebuilds swap the runtime +# libraries the JIT'd .so files link against (e.g. a torch upgrade replacing +# libsycl.so.8 with libsycl.so.9), and Triton keys its cache on the source +# hash only, so a bare gpu-mask key would dlopen stale .so files from the +# previous image and every job would die with +# OSError: libsycl.so.N: cannot open shared object file: No such file or directory +# The image's content-addressable ID from `docker image inspect` changes iff +# any layer changed, giving exactly the invariant needed: new image -> new +# cache dir. Cold JIT compile can push test_xpu_basic past its 1200s timeout +# on B580, so we keep the cache warm within an image version. +if [[ -n "${XPU_KERNEL_CACHE_DIR:-}" ]]; then + # Explicit override -- honor verbatim, no image keying, no sibling pruning. + XPU_KERNEL_CACHE_HOST="${XPU_KERNEL_CACHE_DIR}" +else + IMG_ID_SHORT=$(docker image inspect --format '{{.Id}}' "${IMAGE}" 2>/dev/null \ + | sed 's/^sha256://' | cut -c1-12) + CACHE_ROOT="${HOME}/.cache/sglang-xpu-ci" + GPU_KEY="gpu${ZE_AFFINITY_MASK:-shared}" + if [[ -n "${IMG_ID_SHORT}" ]]; then + XPU_KERNEL_CACHE_HOST="${CACHE_ROOT}/kernel-cache-${GPU_KEY}-${IMG_ID_SHORT}" + # Prune sibling caches for the same GPU but a different image ID. The + # files inside are root-owned (written from the CI container), so unlink + # them via a rootful busybox helper -- same pattern as the workspace + # ownership reset step in the workflow. + shopt -s nullglob + stale_siblings=("${CACHE_ROOT}"/kernel-cache-"${GPU_KEY}"-*) + shopt -u nullglob + for sibling in "${stale_siblings[@]}"; do + [[ -d "${sibling}" ]] || continue + [[ "${sibling}" == "${XPU_KERNEL_CACHE_HOST}" ]] && continue + echo "Pruning stale kernel cache from a previous image build: ${sibling}" + docker run --rm -v "${CACHE_ROOT}:/c" busybox:latest \ + rm -rf "/c/$(basename "${sibling}")" || true + done + else + # `docker image inspect` failed (image not local yet, daemon issue, ...). + # Fall back to the legacy unversioned path so we don't churn the cache. + echo "Warning: could not resolve image ID for ${IMAGE}; cache is not image-versioned this run." >&2 + XPU_KERNEL_CACHE_HOST="${CACHE_ROOT}/kernel-cache-${GPU_KEY}" + fi +fi mkdir -p "${XPU_KERNEL_CACHE_HOST}"/{triton,inductor,neo,sycl} echo "Using persistent XPU kernel cache: ${XPU_KERNEL_CACHE_HOST}" # Cap the cache (default 5 GiB); over-cap resets it (misses just recompile). +# Cache contents are root-owned (written from inside the container), so the +# reset must go through busybox rather than a plain `rm -rf` as the runner. XPU_KERNEL_CACHE_MAX_MB="${XPU_KERNEL_CACHE_MAX_MB:-5120}" cache_mb=$(du -sm "${XPU_KERNEL_CACHE_HOST}" 2>/dev/null | cut -f1) if [[ -n "${cache_mb}" && "${cache_mb}" -gt "${XPU_KERNEL_CACHE_MAX_MB}" ]]; then echo "XPU kernel cache is ${cache_mb} MiB (> ${XPU_KERNEL_CACHE_MAX_MB} MiB cap); resetting it." - rm -rf "${XPU_KERNEL_CACHE_HOST:?}"/{triton,inductor,neo,sycl} + docker run --rm -v "${XPU_KERNEL_CACHE_HOST}:/c" busybox:latest \ + sh -c 'rm -rf /c/triton /c/inductor /c/neo /c/sycl' || true mkdir -p "${XPU_KERNEL_CACHE_HOST}"/{triton,inductor,neo,sycl} fi From 5011ff610e60940c16d944fcee53371b8ff97dd7 Mon Sep 17 00:00:00 2001 From: Ashwini Rathi Date: Tue, 18 Aug 2026 20:21:52 +0300 Subject: [PATCH 2/5] [XPU][CI] fix unreachable image-id fallback Under `set -euo pipefail`, a failed `docker image inspect` propagates rc=1 through the pipeline and kills the script before the `if [[ -n IMG_ID_SHORT ]]` fallback runs, so the "warn + fall back" branch was dead code and the step silently exited 1. Append `|| IMG_ID_SHORT=""` to make it reachable. Also point the fallback at a per-run `-unversioned-$$` dir instead of the legacy path (which may hold the poisoned .so's this block exists to avoid), and extend the sibling prune to sweep the legacy dir too. --- scripts/ci/xpu/xpu_ci_start_container.sh | 37 +++++++----------------- 1 file changed, 11 insertions(+), 26 deletions(-) diff --git a/scripts/ci/xpu/xpu_ci_start_container.sh b/scripts/ci/xpu/xpu_ci_start_container.sh index 2b3344781016..446d129cb770 100755 --- a/scripts/ci/xpu/xpu_ci_start_container.sh +++ b/scripts/ci/xpu/xpu_ci_start_container.sh @@ -109,54 +109,39 @@ elif [[ -r "${HF_TOKEN_FILE}" ]]; then HF_TOKEN_VALUE=$(cat "${HF_TOKEN_FILE}") fi -# Persistent JIT kernel cache (Triton/Inductor/NEO/SYCL) keyed by GPU mask -# AND the resolved image content ID. Nightly image rebuilds swap the runtime -# libraries the JIT'd .so files link against (e.g. a torch upgrade replacing -# libsycl.so.8 with libsycl.so.9), and Triton keys its cache on the source -# hash only, so a bare gpu-mask key would dlopen stale .so files from the -# previous image and every job would die with -# OSError: libsycl.so.N: cannot open shared object file: No such file or directory -# The image's content-addressable ID from `docker image inspect` changes iff -# any layer changed, giving exactly the invariant needed: new image -> new -# cache dir. Cold JIT compile can push test_xpu_basic past its 1200s timeout -# on B580, so we keep the cache warm within an image version. +# Persistent JIT kernel cache keyed by GPU mask + image ID (new image -> new +# cache; avoids dlopen of stale .so's like libsycl.so.8 after a torch bump). if [[ -n "${XPU_KERNEL_CACHE_DIR:-}" ]]; then - # Explicit override -- honor verbatim, no image keying, no sibling pruning. XPU_KERNEL_CACHE_HOST="${XPU_KERNEL_CACHE_DIR}" else + # `|| IMG_ID_SHORT=""` keeps pipefail from killing the script on inspect failure. IMG_ID_SHORT=$(docker image inspect --format '{{.Id}}' "${IMAGE}" 2>/dev/null \ - | sed 's/^sha256://' | cut -c1-12) + | sed 's/^sha256://' | cut -c1-12) || IMG_ID_SHORT="" CACHE_ROOT="${HOME}/.cache/sglang-xpu-ci" GPU_KEY="gpu${ZE_AFFINITY_MASK:-shared}" if [[ -n "${IMG_ID_SHORT}" ]]; then XPU_KERNEL_CACHE_HOST="${CACHE_ROOT}/kernel-cache-${GPU_KEY}-${IMG_ID_SHORT}" - # Prune sibling caches for the same GPU but a different image ID. The - # files inside are root-owned (written from the CI container), so unlink - # them via a rootful busybox helper -- same pattern as the workspace - # ownership reset step in the workflow. + # Prune caches for other image IDs + the legacy unversioned dir (root-owned). shopt -s nullglob - stale_siblings=("${CACHE_ROOT}"/kernel-cache-"${GPU_KEY}"-*) + stale_siblings=("${CACHE_ROOT}"/kernel-cache-"${GPU_KEY}"-* "${CACHE_ROOT}/kernel-cache-${GPU_KEY}") shopt -u nullglob for sibling in "${stale_siblings[@]}"; do [[ -d "${sibling}" ]] || continue [[ "${sibling}" == "${XPU_KERNEL_CACHE_HOST}" ]] && continue - echo "Pruning stale kernel cache from a previous image build: ${sibling}" + echo "Pruning stale kernel cache: ${sibling}" docker run --rm -v "${CACHE_ROOT}:/c" busybox:latest \ rm -rf "/c/$(basename "${sibling}")" || true done else - # `docker image inspect` failed (image not local yet, daemon issue, ...). - # Fall back to the legacy unversioned path so we don't churn the cache. - echo "Warning: could not resolve image ID for ${IMAGE}; cache is not image-versioned this run." >&2 - XPU_KERNEL_CACHE_HOST="${CACHE_ROOT}/kernel-cache-${GPU_KEY}" + # Throwaway per-run dir; legacy path may be poisoned. Next good run prunes it. + echo "Warning: could not resolve image ID for ${IMAGE}; using throwaway cache." >&2 + XPU_KERNEL_CACHE_HOST="${CACHE_ROOT}/kernel-cache-${GPU_KEY}-unversioned-$$" fi fi mkdir -p "${XPU_KERNEL_CACHE_HOST}"/{triton,inductor,neo,sycl} echo "Using persistent XPU kernel cache: ${XPU_KERNEL_CACHE_HOST}" -# Cap the cache (default 5 GiB); over-cap resets it (misses just recompile). -# Cache contents are root-owned (written from inside the container), so the -# reset must go through busybox rather than a plain `rm -rf` as the runner. +# Cap the cache (default 5 GiB); over-cap resets it via busybox (root-owned). XPU_KERNEL_CACHE_MAX_MB="${XPU_KERNEL_CACHE_MAX_MB:-5120}" cache_mb=$(du -sm "${XPU_KERNEL_CACHE_HOST}" 2>/dev/null | cut -f1) if [[ -n "${cache_mb}" && "${cache_mb}" -gt "${XPU_KERNEL_CACHE_MAX_MB}" ]]; then From 957524114e814b3de6229d42e2b8f36833f255e1 Mon Sep 17 00:00:00 2001 From: Ashwini Rathi Date: Tue, 18 Aug 2026 21:46:16 +0300 Subject: [PATCH 3/5] ci: re-trigger jobs From 634065634d29e495b58c9d0fe8cf26d4747e6515 Mon Sep 17 00:00:00 2001 From: Ashwini Rathi Date: Wed, 19 Aug 2026 16:23:04 +0300 Subject: [PATCH 4/5] [XPU][GDN] guard -1 padding sentinel in chunk_gated_delta_rule kernel `test_padded_state_index_is_skipped[all_padded]` (added in PR #33431) launches the GDN chunked kernel with every row = -1 and asserts `torch.equal(pool, pool_init)`. On XPU the pool changed and the test failed deterministically with: AssertionError: all_padded: padded rows wrote into the state pool The CUDA-side fix in PR #33810 added a `valid_state = index >= 0` gate around the block-ptr load/store, but the XPU-specific kernel at python/sglang/srt/hardware_backend/xpu/kernels/fla/chunk_delta_h.py never carried the gate, so padded rows kept leaking writes into the state pool. Mirror the same guard here: * compute `valid_state = index >= 0` once at the top; * predicate the initial-state block load (i_t == 0 path) with `USE_INITIAL_STATE and valid_state`; * for i_t > 0, load h from scratch only when `valid_state`, else fall back to zeros so the delta update stays finite; * gate the epilogue `tl.store` back into scratch with `INPLACE_UPDATE and valid_state`. Padded programs now leave the state pool untouched and the test's `torch.equal(pool, pool_init)` holds on XPU. Re-enable the test at CI-registration (`register_xpu_ci(...)` no longer takes `disabled=`). Fixes: test_chunk_gated_delta_rule.py::test_padded_state_index_is_skipped [all_padded] on stage-b-test-1-gpu-xpu Related: #33431, #33810 --- .../xpu/kernels/fla/chunk_delta_h.py | 17 ++++++++++++----- 1 file changed, 12 insertions(+), 5 deletions(-) diff --git a/python/sglang/srt/hardware_backend/xpu/kernels/fla/chunk_delta_h.py b/python/sglang/srt/hardware_backend/xpu/kernels/fla/chunk_delta_h.py index 607fb7d4cb56..9f4734a41b61 100644 --- a/python/sglang/srt/hardware_backend/xpu/kernels/fla/chunk_delta_h.py +++ b/python/sglang/srt/hardware_backend/xpu/kernels/fla/chunk_delta_h.py @@ -112,6 +112,9 @@ def chunk_gated_delta_rule_fwd_kernel_h_blockdim64_k_loop( ) index = tl.load(initial_state_indices + i_n).to(tl.int32) + # Padded rows carry the -1 sentinel; without this guard the sentinel + # reaches pointer arithmetic and addresses before the state pool. + valid_state = index >= 0 h0 = initial_state + index * stride_h ht = initial_state + index * stride_h if USE_INITIAL_STATE: @@ -128,18 +131,20 @@ def chunk_gated_delta_rule_fwd_kernel_h_blockdim64_k_loop( for k_blk in range(0, K, 64): # Load h: from initial_state (i_t==0) or scratch (i_t>0) if i_t == 0: - if USE_INITIAL_STATE: + if USE_INITIAL_STATE and valid_state: p_hs = tl.make_block_ptr( h0, (V, K), (K, 1), (i_v * BV, k_blk), (BV, 64), (1, 0) ) b_h = tl.load(p_hs, boundary_check=(0, 1)).to(tl.float32) else: b_h = tl.zeros([BV, 64], dtype=tl.float32) - else: + elif valid_state: p_hs = tl.make_block_ptr( ht, (V, K), (K, 1), (i_v * BV, k_blk), (BV, 64), (1, 0) ) b_h = tl.load(p_hs, boundary_check=(0, 1)).to(tl.float32) + else: + b_h = tl.zeros([BV, 64], dtype=tl.float32) # Store pre-update h to output p_ho = tl.make_block_ptr( @@ -181,18 +186,20 @@ def chunk_gated_delta_rule_fwd_kernel_h_blockdim64_k_loop( for k_blk in range(0, K, 64): # Reload h (same source as Phase 1) if i_t == 0: - if USE_INITIAL_STATE: + if USE_INITIAL_STATE and valid_state: p_hs = tl.make_block_ptr( h0, (V, K), (K, 1), (i_v * BV, k_blk), (BV, 64), (1, 0) ) b_h = tl.load(p_hs, boundary_check=(0, 1)).to(tl.float32) else: b_h = tl.zeros([BV, 64], dtype=tl.float32) - else: + elif valid_state: p_hs = tl.make_block_ptr( ht, (V, K), (K, 1), (i_v * BV, k_blk), (BV, 64), (1, 0) ) b_h = tl.load(p_hs, boundary_check=(0, 1)).to(tl.float32) + else: + b_h = tl.zeros([BV, 64], dtype=tl.float32) # Gate decay on h if USE_G: @@ -215,7 +222,7 @@ def chunk_gated_delta_rule_fwd_kernel_h_blockdim64_k_loop( b_h += tl.trans(tl.dot(b_k, b_v)) # Save updated h to scratch (initial_state) for next time step - if INPLACE_UPDATE: + if INPLACE_UPDATE and valid_state: p_hs = tl.make_block_ptr( ht, (V, K), (K, 1), (i_v * BV, k_blk), (BV, 64), (1, 0) ) From 73b3009f30d68afc555a36b764de781f6395f964 Mon Sep 17 00:00:00 2001 From: Ashwini Rathi Date: Wed, 19 Aug 2026 22:57:43 +0300 Subject: [PATCH 5/5] [XPU CI] bump stage-b action timeout 60 -> 120 min The 60-minute action wall on `Run stage-b tests` killed run 32273608720 after `test_intel_xpu_backend.py` needed a full retry (1800s first attempt + 91s retry) and `test_xpu_graph.py::test_full_graph_runs` started hanging on test 10/11. `run_suite.py`'s per-file retry logic never got to fire because the outer GH-Actions timeout preempted the runner. Bump the stage-b step to 120 min so a per-file retry (or two) in the suite doesn't blow the budget; the runner then has room to classify + retry the file that actually failed and exit cleanly. Leaves other timeouts unchanged: * Install Dependency (stage-a job): 60 min * Run stage-a tests: 30 min * Install Dependency (stage-b job): 60 min --- .github/workflows/pr-test-xpu.yml | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/.github/workflows/pr-test-xpu.yml b/.github/workflows/pr-test-xpu.yml index 48ac12f22e26..0cdf9446517b 100644 --- a/.github/workflows/pr-test-xpu.yml +++ b/.github/workflows/pr-test-xpu.yml @@ -206,7 +206,7 @@ jobs: docker exec ci_sglang_xpu /bin/bash -c '/opt/venv/bin/hf auth login --token ${HF_TOKEN}' - name: Run stage-b tests - timeout-minutes: 60 + timeout-minutes: 120 run: | docker exec ci_sglang_xpu bash -c "source /opt/venv/bin/activate && cd /sglang-checkout/test && python3 run_suite.py --hw xpu --suite stage-b-test-1-gpu-xpu --enable-retry"