To build for CRI and update sycl-tla SHA-ID by mkumargarg · Pull Request #390 · sgl-project/sgl-kernel-xpu

mkumargarg · 2026-08-13T09:06:05Z

No description provided.

Copilot

Pull request overview

This PR updates the SYCL/XPU build to support CRI (Xe35) alongside existing BMG (Xe20) builds, and adjusts tests to compute references on CPU while running kernels on XPU.

Changes:

Added Xe35 (CRI) build buckets and target-specific SYCL offline compilation paths.
Introduced Level Zero–based device IP auto-detection to select build target when not explicitly provided.
Updated multiple tests to create inputs / references on CPU and compare against XPU outputs.

Reviewed changes

Copilot reviewed 33 out of 33 changed files in this pull request and generated 12 comments.

Show a summary per file

File	Description
tests/test_swiglu_with_alpha_limit.py	Generate reference on CPU; run kernel on XPU and compare on CPU.
tests/test_per_tensor_quant_fp8.py	Refactors FP8 quant test to mix CPU/XPU paths.
tests/test_moe_gemm.py	Switches test tensor creation to CPU helpers.
tests/test_merge_state_v2.py	Moves Torch baseline tensors to CPU and adds divide-by-zero guard in perf ratio.
tests/test_fused_qk_norm_rope.py	Moves reference copies and comparisons to CPU; adjusts some tensor creation to CPU.
tests/test_activation.py	Computes references on CPU and compares against XPU outputs moved back to CPU.
src/torch_extension_sycl.cc	Gates Xe20-only kernel registrations behind `SYCL_INTEL_TARGET == 20`.
src/sycl/kernels/mla/collective/xe_mla_epilogue.hpp	Replaces `cute::transform` reductions with explicit unrolled loops.
src/sycl/kernels/mla_sparse/collective/xe_mla_sparse_epilogue.hpp	Same reduction refactor for sparse MLA epilogue.
src/sycl/kernels/mla_sparse/collective/xe_mla_sparse_2stage_epilogue.hpp	Same reduction refactor for sparse 2-stage epilogue.
src/sycl/kernels/flash_attention_v2/collective/xe_fmha_fwd_epilogue.hpp	Same reduction refactor for FMHA epilogues.
src/SGEMMLoraBFwdXe20.cmake	Adds generated LoRA-B sources to Xe35 bucket.
src/SGEMMLoraAFwdXe20.cmake	Adds generated LoRA-A sources to Xe35 bucket.
src/MlaSparsePrefillXe20.cmake	Adds generated sparse MLA prefill sources to Xe35 bucket.
src/MlaSparseDecodeXe20.cmake	Adds generated sparse MLA decode sources to Xe35 bucket.
src/MlaPrefillXe20.cmake	Adds generated MLA prefill sources to Xe35 bucket.
src/MlaDecodeXe20.cmake	Adds generated MLA decode sources to Xe35 bucket.
src/GdnAttnXe20.cmake	Avoids duplicate compilation by removing Xe20 sources from common bucket before appending to Xe20.
src/FMHAPrefillXe20.cmake	Adds generated FMHA prefill sources to Xe35 bucket.
src/FMHADecodeXe20.cmake	Adds generated FMHA decode sources to Xe35 bucket.
src/CMakeLists.txt	Adds recursive SYCL source discovery and splits sources into common/Xe20/Xe35 lists.
src/BuildOnLinux.cmake	Builds Xe20-only and Xe35-only SYCL libs conditionally based on `DPCPP_SYCL_TARGET`.
python/sgl_kernel/init.py	Makes Triton import optional; exports `fp8_paged_mqa_logits_triton = None` when unavailable.
CMakeLists.txt	Adds CRI target selection, sets `SYCL_INTEL_TARGET`, and updates sycl-tla commit SHA.
cmake/Modules/FindSYCL.cmake	Adjusts device link flags handling in `SYCL_LINK_DEVICE_OBJECTS`.
cmake/DeviceDetection.cmake	Adds Level Zero device IP version detection helper.
cmake/BuildFlags.cmake	Wires in device detection and maps `DPCPP_SYCL_TARGET` to AOT flags/defines.

Suppressed comments (1)

tests/test_moe_gemm.py:39

create_random_cpu_tensor is defined twice; the second definition overwrites the first and drops the (already incorrect) docstring. This makes the helper harder to maintain and can hide future edits to the first implementation.

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

    assert torch.allclose(
-        output_ref, output, atol=atol, rtol=rtol
+        output_ref, output.to("cpu"), atol=atol, rtol=rtol
    ), f"dtype = {dtype}Output mismatch: max_diff={torch.max(torch.abs(output_ref - output))}"


+    x = torch.rand((num_tokens, hidden_dim), dtype=dtype)

-    sglang_out, sglang_scale = sglang_scaled_fp8_quant(x)
-    torch_out = torch_scaled_fp8_quant(x, sglang_scale)
+    x_hpu = x.to(device)
+    sglang_out, sglang_scale = sglang_scaled_fp8_quant(x_hpu)
+    torch_out = torch_scaled_fp8_quant(x, sglang_scale.cpu())


  target_include_directories(${lib} PUBLIC ${TORCH_XPU_OPS_INCLUDE_DIRS})
  target_include_directories(${lib} PUBLIC ${ATen_XPU_INCLUDE_DIRS})
  target_include_directories(${lib} PUBLIC ${SYCL_INCLUDE_DIR})
  target_include_directories(${lib} PRIVATE ${Python3_INCLUDE_DIRS})
-  target_include_directories(${lib} PRIVATE ${CMAKE_CURRENT_SOURCE_DIR})
  target_link_libraries(${lib} PRIVATE ${Python3_LIBRARIES})


        list(APPEND device_cpp_xe20 ${GENERATED_FILE})
+        list(APPEND device_cpp_xe35 ${GENERATED_FILE})


        list(APPEND device_cpp_xe20 ${GENERATED_FILE})
+        list(APPEND device_cpp_xe35 ${GENERATED_FILE})


                list(APPEND device_cpp_xe20 ${GENERATED_FILE_2STAGE})
+                list(APPEND device_cpp_xe35 ${GENERATED_FILE_2STAGE})


            "${CMAKE_CURRENT_BINARY_DIR}/sycl/mla_sparse_decode_kernel_${ELEM_TAG}_128.cpp")
        configure_file(${MLA_SPARSE_DECODE_TEMPLATE} ${GENERATED_FILE} @ONLY)
        list(APPEND device_cpp_xe20 ${GENERATED_FILE})
+        list(APPEND device_cpp_xe35 ${GENERATED_FILE})
    endif()


                set(GENERATED_FILE
                    "${CMAKE_CURRENT_BINARY_DIR}/sycl/xe_fmha_fwd_decode_page_${QG_SZ}_${HEAD_DIM}_${PAGE_SIZE}_${DT16}.cpp")
                configure_file(${FMHA_DECODE_TEMPLATE} ${GENERATED_FILE} @ONLY)
                list(APPEND device_cpp_xe20 ${GENERATED_FILE})
+                list(APPEND device_cpp_xe35 ${GENERATED_FILE})


        set(GENERATED_FILE
            "${CMAKE_CURRENT_BINARY_DIR}/sycl/xe_fmha_fwd_prefill_page_${HEAD_DIM}_${DT16}.cpp")
        configure_file(${FMHA_PREFILL_TEMPLATE} ${GENERATED_FILE} @ONLY)
        list(APPEND device_cpp_xe20 ${GENERATED_FILE})
+        list(APPEND device_cpp_xe35 ${GENERATED_FILE})


    # Build the generated file and dependency file ##########################
-    # Only pass -Xs when there are offline compiler flags (AOT targets).
-    # For JIT (spir64) targets SYCL_OFFLINE_COMPILER_FLAGS is empty and
-    # passing a bare -Xs causes icx to consume the following -o flag as
-    # its argument, producing "no such file or directory" errors.
-    if(SYCL_OFFLINE_COMPILER_FLAGS)
-      set(_sycl_xs_flags -Xs ${SYCL_OFFLINE_COMPILER_FLAGS})
-    else()
-      set(_sycl_xs_flags)
-    endif()
-
    add_custom_command(
      OUTPUT ${output_file}
      DEPENDS ${object_files}


sunjiweiswift · 2026-08-18T02:56:38Z

pls rebase main

These are the changes required to build/enable for cri. Authored-by: Suryaprakash Shanmugam <suryaprakash.shanmugam@intel.com> Meng, Hengyu <hengyu.meng@intel.com> Signed-off-by: Suryaprakash Shanmugam <suryaprakash.shanmugam@intel.com> Meng, Hengyu <hengyu.meng@intel.com> Manoj Kumar <manoj1.kumar@intel.com>

sunjiweiswift · 2026-08-21T05:55:15Z

@yuankuns

The flash_attention CI passes after the commit(1d453dd). Please check the content related to FMHA_PREFILL_ENABLE_SCORE_BLOCK2D_512.

Copilot AI lite review requested due to automatic review settings August 13, 2026 09:06

mkumargarg requested review from airMeng and kareemshaik80 as code owners August 13, 2026 09:06

Copilot started reviewing on behalf of mkumargarg August 13, 2026 09:06 View session

mkumargarg mentioned this pull request Aug 13, 2026

[Internal/CRI] Changes to build/enable CRI - WIP #292

Closed

mkumargarg added the run-ci label Aug 13, 2026

Copilot AI reviewed Aug 13, 2026

View reviewed changes

mkumargarg requested review from sspintel and sunjiweiswift August 18, 2026 04:55

sunjiweiswift approved these changes Aug 18, 2026

View reviewed changes

mkumargarg force-pushed the my_branch_upstream_cri_kernels_2 branch from 911f519 to 19b9c12 Compare August 19, 2026 08:47

mkumargarg force-pushed the my_branch_upstream_cri_kernels_2 branch 4 times, most recently from a14130f to 3c66921 Compare August 23, 2026 06:38

try FMHA_PREFILL_ENABLE_SCORE_BLOCK2D_512 OFF

3c66921

Navigation Menu

Uh oh!

To build for CRI and update sycl-tla SHA-ID - #390

To build for CRI and update sycl-tla SHA-ID#390
mkumargarg wants to merge 2 commits into
mainfrom
my_branch_upstream_cri_kernels_2

mkumargarg commented Aug 13, 2026

Uh oh!

Copilot AI left a comment

Uh oh!

sunjiweiswift commented Aug 18, 2026 •

edited

Loading

Uh oh!

sunjiweiswift commented Aug 21, 2026

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

3 participants

		list(APPEND device_cpp_xe20 ${GENERATED_FILE})
		list(APPEND device_cpp_xe35 ${GENERATED_FILE})

Uh oh!

Conversation

mkumargarg commented Aug 13, 2026

Uh oh!

Copilot AI left a comment

Choose a reason for hiding this comment

Pull request overview

Reviewed changes

Uh oh!

sunjiweiswift commented Aug 18, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

sunjiweiswift commented Aug 21, 2026

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

3 participants

sunjiweiswift commented Aug 18, 2026 •

edited

Loading