feat(intel-xpu): support MiniMax-H3 Turbo-SLA inference

.gitignore

@@ -35,3 +35,4 @@ lightx2v_ros/build
 lightx2v_ros/install
 lightx2v_ros/log
 .gitnexus
+lightx2v_kernel_xpu/_cmake_build

configs/platforms/intel_xpu/minimax_h3_fl2v_turbo_sla_4step.json

@@ -0,0 +1,44 @@
+{
+  "infer_steps": 5,
+  "target_video_length": 362,
+  "target_height": 768,
+  "target_width": 1344,
+  "fps": 24,
+  "target_fps": 24,
+  "enable_cfg": false,
+  "cpu_offload": true,
+  "offload_granularity": "block",
+  "text_encoder_cpu_offload": true,
+  "text_encoder_offload_granularity": "block",
+  "text_encoder_host_pinned": false,
+  "text_encoder_release_block_offload_buffers": true,
+  "vae_cpu_offload": true,
+  "lazy_load": false,
+  "unload_modules": false,
+  "attn_type": "dynamic_sparse_attn",
+  "refiner_attn_type": "intel_xpu_cute_attn",
+  "dynamic_sparse_attn_setting": {
+    "sparsity_ratio": 0.85,
+    "operator": "intel_xpu_cute_attn"
+  },
+  "rms_type": "intel_xpu",
+  "rope_type": "minimax_h3_xpu_rope",
+  "feature_caching": "NoCaching",
+  "use_compile": false,
+  "video_flow_shift": 6.0,
+  "audio_flow_shift": 3.0,
+  "h3_step_update": "training_euler",
+  "vae_spatial_scale_factor": 16,
+  "audio_sampling_rate": 32000,
+  "audio_latents_per_second": 40,
+  "audio_channels": 2,
+  "keep_latents_dtype_in_scheduler": true,
+  "lora_dynamic_apply": false,
+  "lora_configs": [
+    {
+      "path": "/llm/models/Minimax-h3-Turbo-SLA/minimax_h3_fl2v_turbo_4step_v0.1_768p_sla_bf16.safetensors",
+      "strength": 1.0,
+      "alpha": 128
+    }
+  ]
+}

lightx2v/common/ops/attn/dynamic_sparse_attn.py

@@ -1,3 +1,5 @@
+from functools import partial
+
 import torch
 from loguru import logger
 
@@ -77,8 +79,9 @@ def __init__(self, config=None):
         if not 0.0 <= self.sparsity_ratio < 1.0:
             raise ValueError(f"dynamic sparse attention sparsity_ratio must be in [0, 1), got {self.sparsity_ratio}")
 
-        self.arch = get_cuda_arch(torch.cuda.current_device())
         self.topk = 1 - self.sparsity_ratio
+        self.arch = get_cuda_arch(torch.cuda.current_device()) if torch.cuda.is_available() else None
+
         if self.operator == "triton":
             self.BLKQ, self.BLKK = 64, 64
             self.apply_func = self.apply_triton
@@ -100,6 +103,16 @@ def __init__(self, config=None):
         elif self.operator == "magi":
             self.BLKQ, self.BLKK = 128, 128
             self.apply_func = self.apply_magi
+        elif self.operator == "intel_xpu_cute_attn":
+            self.BLKQ, self.BLKK = 128, 128
+            from lightx2v_platform.ops.attn.intel_xpu.xpu_sla_attn import apply_intel_xpu_cute_attn
+
+            self.apply_func = partial(
+                apply_intel_xpu_cute_attn,
+                keep_ratio=self.topk,
+                block_q=self.BLKQ,
+                block_k=self.BLKK,
+            )
         else:
             raise NotImplementedError(f"Not supported SLA operator: {self.operator}.")
 

lightx2v_kernel_xpu/CMakeLists.txt

@@ -284,6 +284,9 @@ if(ENABLE_CUTE_FMHA)
     set(H3_OVERLAY_ROOT "${CMAKE_BINARY_DIR}/minimax_h3_cute_overlay")
     set(H3_OVERLAY_DIR "${H3_OVERLAY_ROOT}/flash_attention_v2/collective")
     file(MAKE_DIRECTORY "${H3_OVERLAY_DIR}")
+    # Remove a stale sparse kernel overlay left by older build trees.  The
+    # dense target must resolve the upstream kernel header unchanged.
+    file(REMOVE "${H3_OVERLAY_ROOT}/flash_attention_v2/kernel/xe_fmha_fwd_kernel.hpp")
     set(H3_MAINLOOP_SOURCE
         "${CUTLASS_SYCL_ROOT}/applications/flash_attention_v2/collective/xe_fmha_fwd_mainloop.hpp")
     set(H3_MAINLOOP_PATCH
@@ -306,9 +309,69 @@ if(ENABLE_CUTE_FMHA)
         "${H3_OVERLAY_DIR}/xe_fmha_fwd_mainloop.hpp"
         COPYONLY)
     file(REMOVE "${H3_MAINLOOP_TEMP}")
+
+    set(H3_SPARSE_OVERLAY_ROOT "${CMAKE_BINARY_DIR}/minimax_h3_sparse_cute_overlay")
+    set(H3_SPARSE_OVERLAY_DIR "${H3_SPARSE_OVERLAY_ROOT}/flash_attention_v2/collective")
+    file(MAKE_DIRECTORY "${H3_SPARSE_OVERLAY_DIR}")
+    configure_file(
+        "${H3_OVERLAY_DIR}/xe_fmha_fwd_mainloop.hpp"
+        "${H3_SPARSE_OVERLAY_DIR}/xe_fmha_fwd_mainloop.hpp"
+        COPYONLY)
+
+    # Add block-LUT traversal after the long-sequence dataflow patch.  The
+    # sparse path keeps CUTE's DPAS/online-softmax implementation and changes
+    # only which physical K/V tiles each Q workgroup visits.
+    set(H3_SPARSE_MAINLOOP_PATCH
+        "${CMAKE_CURRENT_SOURCE_DIR}/cute/patches/minimax_h3_sparse_lut.patch")
+    execute_process(
+        COMMAND patch --batch --forward --fuzz=0
+            ${H3_SPARSE_OVERLAY_DIR}/xe_fmha_fwd_mainloop.hpp
+            ${H3_SPARSE_MAINLOOP_PATCH}
+        RESULT_VARIABLE H3_SPARSE_MAINLOOP_RESULT
+        OUTPUT_VARIABLE H3_SPARSE_MAINLOOP_STDOUT
+        ERROR_VARIABLE H3_SPARSE_MAINLOOP_STDERR)
+    if(NOT H3_SPARSE_MAINLOOP_RESULT EQUAL 0)
+        message(FATAL_ERROR
+            "Failed to apply MiniMax-H3 sparse mainloop patch:\n${H3_SPARSE_MAINLOOP_STDOUT}${H3_SPARSE_MAINLOOP_STDERR}")
+    endif()
+    set(H3_SPARSE_TAIL_PATCH
+        "${CMAKE_CURRENT_SOURCE_DIR}/cute/patches/minimax_h3_sparse_tail.patch")
+    execute_process(
+        COMMAND patch --batch --forward --fuzz=0
+            ${H3_SPARSE_OVERLAY_DIR}/xe_fmha_fwd_mainloop.hpp ${H3_SPARSE_TAIL_PATCH}
+        RESULT_VARIABLE H3_SPARSE_TAIL_RESULT
+        OUTPUT_VARIABLE H3_SPARSE_TAIL_STDOUT
+        ERROR_VARIABLE H3_SPARSE_TAIL_STDERR)
+    if(NOT H3_SPARSE_TAIL_RESULT EQUAL 0)
+        message(FATAL_ERROR
+            "Failed to apply MiniMax-H3 sparse tail patch:\n${H3_SPARSE_TAIL_STDOUT}${H3_SPARSE_TAIL_STDERR}")
+    endif()
+
+    set(H3_KERNEL_DIR "${H3_SPARSE_OVERLAY_ROOT}/flash_attention_v2/kernel")
+    file(MAKE_DIRECTORY "${H3_KERNEL_DIR}")
+    set(H3_KERNEL_SOURCE
+        "${CUTLASS_SYCL_ROOT}/applications/flash_attention_v2/kernel/xe_fmha_fwd_kernel.hpp")
+    set(H3_KERNEL_PATCH
+        "${CMAKE_CURRENT_SOURCE_DIR}/cute/patches/minimax_h3_sparse_kernel.patch")
+    set(H3_KERNEL_TEMP "${H3_KERNEL_DIR}/xe_fmha_fwd_kernel.hpp.new")
+    execute_process(
+        COMMAND patch --batch --forward --fuzz=0
+            --output=${H3_KERNEL_TEMP} ${H3_KERNEL_SOURCE} ${H3_KERNEL_PATCH}
+        RESULT_VARIABLE H3_KERNEL_PATCH_RESULT
+        OUTPUT_VARIABLE H3_KERNEL_PATCH_STDOUT
+        ERROR_VARIABLE H3_KERNEL_PATCH_STDERR)
+    if(NOT H3_KERNEL_PATCH_RESULT EQUAL 0)
+        message(FATAL_ERROR
+            "Failed to apply MiniMax-H3 sparse kernel patch:\n${H3_KERNEL_PATCH_STDOUT}${H3_KERNEL_PATCH_STDERR}")
+    endif()
+    configure_file("${H3_KERNEL_TEMP}" "${H3_KERNEL_DIR}/xe_fmha_fwd_kernel.hpp" COPYONLY)
+    file(REMOVE "${H3_KERNEL_TEMP}")
     file(COPY
         "${CUTLASS_SYCL_ROOT}/applications/flash_attention_v2/collective/fmha_fusion.hpp"
         DESTINATION "${H3_OVERLAY_DIR}")
+    file(COPY
+        "${CUTLASS_SYCL_ROOT}/applications/flash_attention_v2/collective/fmha_fusion.hpp"
+        DESTINATION "${H3_SPARSE_OVERLAY_DIR}")
 
     add_library(cute_fmha_minimax_h3_torch SHARED cute/cute_fmha_torch.cpp)
     set_target_properties(cute_fmha_minimax_h3_torch PROPERTIES
@@ -341,6 +404,39 @@ if(ENABLE_CUTE_FMHA)
         "-spirv-ext=+SPV_INTEL_split_barrier,+SPV_INTEL_2d_block_io,+SPV_INTEL_subgroup_matrix_multiply_accumulate")
     target_link_libraries(cute_fmha_minimax_h3_torch PRIVATE
         ${TORCH_LIBRARIES} ${TORCH_PYTHON_LIBRARY})
+
+    add_library(cute_fmha_minimax_h3_sparse_torch SHARED cute/cute_fmha_torch.cpp)
+    set_target_properties(cute_fmha_minimax_h3_sparse_torch PROPERTIES
+        PREFIX ""
+        BUILD_RPATH "${TORCH_INSTALL_PREFIX}/lib"
+        INSTALL_RPATH "\$ORIGIN;${TORCH_INSTALL_PREFIX}/lib")
+    target_compile_features(cute_fmha_minimax_h3_sparse_torch PRIVATE cxx_std_17)
+    target_compile_definitions(cute_fmha_minimax_h3_sparse_torch PRIVATE
+        CUTLASS_ENABLE_SYCL SYCL_INTEL_TARGET ${XPU_ARCH_DEFINE}
+        CUTE_FMHA_TORCH_LIBRARY=sycl_kernels_cute_minimax_h3_sparse
+        CUTE_FMHA_SPARSE=1)
+    target_include_directories(cute_fmha_minimax_h3_sparse_torch PRIVATE
+        "${CUTLASS_SYCL_ROOT}/include"
+        "${CUTLASS_SYCL_ROOT}/tools/util/include"
+        "${CUTLASS_SYCL_ROOT}/examples/common"
+        "${H3_SPARSE_OVERLAY_ROOT}"
+        "${CUTLASS_SYCL_ROOT}/applications"
+        "${CMAKE_CURRENT_SOURCE_DIR}/cute"
+        ${TORCH_INCLUDE_DIRS})
+    target_compile_options(cute_fmha_minimax_h3_sparse_torch PRIVATE
+        -O3 -fsycl -fsycl-targets=spir64_gen
+        -Xsycl-target-backend "-device ${XPU_TARGET}"
+        -fno-sycl-instrument-device-code -Wno-unknown-pragmas
+        -Wno-unused-variable -Wno-unused-but-set-variable
+        -Wno-unused-local-typedef -Wno-uninitialized -Wno-reorder-ctor
+        -Wno-logical-op-parentheses -Wno-unused-function -Wno-deprecated-copy)
+    target_link_options(cute_fmha_minimax_h3_sparse_torch PRIVATE
+        -fsycl -fsycl-targets=spir64_gen
+        -Xsycl-target-backend "-device ${XPU_TARGET}"
+        -Xspirv-translator
+        "-spirv-ext=+SPV_INTEL_split_barrier,+SPV_INTEL_2d_block_io,+SPV_INTEL_subgroup_matrix_multiply_accumulate")
+    target_link_libraries(cute_fmha_minimax_h3_sparse_torch PRIVATE
+        ${TORCH_LIBRARIES} ${TORCH_PYTHON_LIBRARY})
 endif()
 
 # ── Install targets for wheel assembly ────────────────────────────────────────
@@ -352,4 +448,5 @@ install(FILES "${ESIMD_RUNTIME}" DESTINATION sycl_kernels)
 if(ENABLE_CUTE_FMHA)
     install(TARGETS cute_fmha_torch LIBRARY DESTINATION sycl_kernels)
     install(TARGETS cute_fmha_minimax_h3_torch LIBRARY DESTINATION sycl_kernels)
+    install(TARGETS cute_fmha_minimax_h3_sparse_torch LIBRARY DESTINATION sycl_kernels)
 endif()

lightx2v_kernel_xpu/README.md

@@ -274,3 +274,27 @@ version `0.0.1` because CUTE FMHA is disabled there.
 CUTE FMHA is disabled on Windows because the current sycl-tla kernel produces
 incorrect attention results there. `build.bat` builds only the existing
 ESIMD/oneDNN extension, and `sycl_kernels.has_cute_fmha()` returns `False`.
+
+### MiniMax-H3 SLA block-sparse attention
+
+The Linux BMG build also exposes an SLA forward path for BF16 MiniMax-H3
+self-attention in `[B,L,H,128]` layout:
+
+```python
+lut = sycl_kernels.sla_block_map(q, k, keep_ratio=0.15, block_q=128, block_k=128)
+out = sycl_kernels.sparse_block_attention(q, k, v, lut)
+# Or route and execute in one call:
+out = sycl_kernels.sla_sparse_attention(q, k, v, keep_ratio=0.15)
+```
+
+The CUTE kernel fuses sparse QK, online softmax, and PV and never materializes
+the score matrix. The optimized contract is forward-only, non-causal, B=1,
+BF16, D=128, equal Q/K/V head counts, and 128x128 SLA blocks. Other supported
+64/128 block or GQA shapes currently use a slower Triton-XPU fallback.
+
+Benchmark the kernel, router, and dense CUTE baseline separately with:
+
+```bash
+ONEAPI_DEVICE_SELECTOR=level_zero:0 PYTHONPATH=python \
+  python test/bench_sla_sparse_attention.py
+```

lightx2v_kernel_xpu/build.sh

@@ -71,7 +71,7 @@ mkdir -p python/sycl_kernels
 find _cmake_build -maxdepth 1 -type f -name '_ext*.so' -exec cp -f {} python/sycl_kernels/ \;
 find _cmake_build -maxdepth 1 -type f -name 'rms_norm_torch*.so' -exec cp -f {} python/sycl_kernels/ \;
 find _cmake_build -maxdepth 1 -type f -name 'minimax_h3_rope_torch*.so' -exec cp -f {} python/sycl_kernels/ \;
-find _cmake_build -maxdepth 1 -type f -name 'cute_fmha_torch*.so' -exec cp -f {} python/sycl_kernels/ \;
+find _cmake_build -maxdepth 1 -type f -name 'cute_fmha*.so' -exec cp -f {} python/sycl_kernels/ \;
 cp -f lgrf_uni/libesimd.unify.lgrf.so python/sycl_kernels/
 
 if [[ "${SKIP_TESTS:-0}" != "1" ]]; then

lightx2v_kernel_xpu/cute/cute_fmha_torch.cpp

@@ -210,6 +210,8 @@ template <
 void run_d128_tile(
     const void* q_ptr, const void* k_ptr, const void* v_ptr, void* o_ptr,
     int B, int H, int Lq, int Lkv, int D, float scale,
+    const int* block_lut = nullptr, int lut_q_blocks = 0,
+    int lut_topk = 0, int lut_block_tiles = 0,
     int64_t q_stride_seq = -1, int64_t q_stride_head = -1,
     int64_t q_stride_batch = -1, int64_t k_stride_seq = -1,
     int64_t k_stride_head = -1, int64_t k_stride_batch = -1,
@@ -294,7 +296,11 @@ void run_d128_tile(
           nullptr, stride_K,   // k_cache
           nullptr, stride_V,   // v_cache
       },
-      {scale, nullptr, 0, nullptr},
+      {scale, nullptr, 0, nullptr,
+#if defined(CUTE_FMHA_SPARSE)
+       block_lut, H, lut_q_blocks, lut_topk, lut_block_tiles
+#endif
+      },
       {},
       hw_info};
 
@@ -378,12 +384,63 @@ at::Tensor sdp(const at::Tensor& q, const at::Tensor& k, const at::Tensor& v) {
   return o;
 }
 
+#if defined(CUTE_FMHA_SPARSE)
+at::Tensor sparse_sdp(
+    const at::Tensor& q, const at::Tensor& k, const at::Tensor& v,
+    const at::Tensor& block_lut) {
+  TORCH_CHECK(q.dim() == 4 && k.dim() == 4 && v.dim() == 4,
+              "sycl_kernels sparse CUTE FMHA: expect q/k/v [B,L,H,D]");
+  TORCH_CHECK(block_lut.dim() == 4,
+              "sycl_kernels sparse CUTE FMHA: expect LUT [B,H,Qblocks,topk]");
+  TORCH_CHECK(q.device().is_xpu() && k.device().is_xpu() && v.device().is_xpu() &&
+                  block_lut.device().is_xpu(),
+              "sycl_kernels sparse CUTE FMHA: all tensors must be on XPU");
+  TORCH_CHECK(q.scalar_type() == at::kBFloat16 && k.scalar_type() == at::kBFloat16 &&
+                  v.scalar_type() == at::kBFloat16,
+              "sycl_kernels sparse CUTE FMHA: q/k/v must be BF16");
+  TORCH_CHECK(block_lut.scalar_type() == at::kInt,
+              "sycl_kernels sparse CUTE FMHA: LUT must be int32");
+
+  const int B = checked_int(q.size(0), "batch");
+  const int L = checked_int(q.size(1), "sequence length");
+  const int H = checked_int(q.size(2), "head count");
+  const int D = checked_int(q.size(3), "head dimension");
+  constexpr int BlockQ = 128;
+  constexpr int BlockK = 128;
+  constexpr int CuteKvTile = 32;
+  const int q_blocks = (L + BlockQ - 1) / BlockQ;
+  TORCH_CHECK(B == 1 && D == 128,
+              "sycl_kernels sparse CUTE FMHA: only B=1,D=128 are supported");
+  TORCH_CHECK(k.sizes() == q.sizes() && v.sizes() == q.sizes(),
+              "sycl_kernels sparse CUTE FMHA: q/k/v shapes must match");
+  TORCH_CHECK(block_lut.size(0) == B && block_lut.size(1) == H &&
+                  block_lut.size(2) == q_blocks && block_lut.size(3) > 0,
+              "sycl_kernels sparse CUTE FMHA: invalid LUT shape");
+
+  auto qc = q.contiguous(), kc = k.contiguous(), vc = v.contiguous();
+  auto lut = block_lut.contiguous();
+  auto output = at::empty_like(qc);
+  const float scale = 1.0f / std::sqrt(static_cast<float>(D));
+  run_d128_tile<cutlass::bfloat16_t, 0, BlockQ>(
+      qc.data_ptr(), kc.data_ptr(), vc.data_ptr(), output.data_ptr(),
+      B, H, L, L, D, scale, lut.const_data_ptr<int>(), q_blocks,
+      checked_int(lut.size(3), "LUT topk"), BlockK / CuteKvTile);
+  return output;
+}
+#endif
+
 }  // namespace
 
 TORCH_LIBRARY(CUTE_FMHA_TORCH_LIBRARY, m) {
   m.def("sdp(Tensor q, Tensor k, Tensor v) -> Tensor");
+#if defined(CUTE_FMHA_SPARSE)
+  m.def("sparse_sdp(Tensor q, Tensor k, Tensor v, Tensor block_lut) -> Tensor");
+#endif
 }
 
 TORCH_LIBRARY_IMPL(CUTE_FMHA_TORCH_LIBRARY, XPU, m) {
   m.impl("sdp", &sdp);
+#if defined(CUTE_FMHA_SPARSE)
+  m.impl("sparse_sdp", &sparse_sdp);
+#endif
 }

lightx2v_kernel_xpu/cute/patches/minimax_h3_sparse_kernel.patch

@@ -0,0 +1,21 @@
+--- xe_fmha_fwd_kernel.hpp
++++ xe_fmha_fwd_kernel.hpp
+@@ -226,7 +226,9 @@
+       if (CollectiveMainloop::CausalMask && seq_coord < discard_seq_coord) continue;
+       const int seq_len_new = CollectiveMainloop::CausalMask ? full_tile_offset + cute::min(seq_len_kv, seq_coord - discard_seq_coord) + q_sg_tile : seq_len_kv;
+       const int seq_len = seq_len_new + seq_len_kv_cache;
+-      const int k_blocks = cute::ceil_div(seq_len, get<1>(TileShapeQK{}));
++      const int k_blocks = params.mainloop.ptr_block_lut
++          ? params.mainloop.lut_topk * params.mainloop.lut_block_tiles
++          : cute::ceil_div(seq_len, get<1>(TileShapeQK{}));
+
+       int offset_q = 0, offset_k = 0, offset_v = 0, offset_o = 0;
+       int offset_k_cache = 0, offset_v_cache = 0;
+@@ -287,7 +289,7 @@
+                tArA, tA_max, tA_sum,
+                blk_qv, 0, k_blocks, k_blocks,
+-               thr_id, seq_len, seq_len_kv_cache, idx_b,
++               thr_id, seq_len, seq_len_kv_cache, idx_b, head_q,
+                full_tile_offset, discard_seq_coord,
+                K_cache(_,_,head,l_coord),
+                V_cache(_,_,head,l_coord));