From bb817df263e11be79da70f7773caeab920a6a608 Mon Sep 17 00:00:00 2001
From: jiqing-feng <jiqing.feng@intel.com>
Date: Thu, 7 May 2026 18:44:28 -0700
Subject: [PATCH] fix kernel

Signed-off-by: jiqing-feng <jiqing.feng@intel.com>
---
 csrc/xpu_kernels.h | 1 +
 1 file changed, 1 insertion(+)

diff --git a/csrc/xpu_kernels.h b/csrc/xpu_kernels.h
index caa7e6716..d3dac55b2 100644
--- a/csrc/xpu_kernels.h
+++ b/csrc/xpu_kernels.h
@@ -22,6 +22,7 @@ template <typename T, int TILE_SIZE, int NUM_PER_TH, int DATA_TYPE> class kDequa
 
 template <typename T, size_t GROUP_SIZE, size_t NUM_PER_THREAD, size_t SUBG_SIZE, int BITS> class kgemv_4bit_inference {
   public:
+    [[sycl::reqd_sub_group_size(SUBG_SIZE)]]
     SYCL_EXTERNAL void operator()(sycl::nd_item<1> item) const;
 
     kgemv_4bit_inference(