diff --git a/vllm/custom-esimd-kernels-vllm/csrc/xpu/esimd_kernel.sycl b/vllm/custom-esimd-kernels-vllm/csrc/xpu/esimd_kernel.sycl index 3f5821c8..9c694617 100644 --- a/vllm/custom-esimd-kernels-vllm/csrc/xpu/esimd_kernel.sycl +++ b/vllm/custom-esimd-kernels-vllm/csrc/xpu/esimd_kernel.sycl @@ -160,10 +160,14 @@ at::Tensor esimd_gemv_int4( at::Tensor output) { int64_t N = weight.size(0); int64_t K = weight.size(1) * 2; // packed: K/2 bytes → K elements + // Infer group size from scale shape [N, K/group_size]: 128 (legacy GGML) + // or 32 (vLLM 0.21 sym_int4 / q4_0 32-element blocks). + int64_t n_groups = weight_scale.size(weight_scale.dim() - 1); + int group_size = (n_groups > 0) ? (int)(K / n_groups) : 128; EXTRACT_PTR(p_in, input); EXTRACT_PTR(p_w, weight); EXTRACT_PTR(p_sc, weight_scale); EXTRACT_PTR(p_out, output); auto& dpcpp_queue = get_device_queue(input); - GEMV_int4_host(p_in, p_w, p_sc, p_out, N, K, dpcpp_queue); + GEMV_int4_host(p_in, p_w, p_sc, p_out, N, K, dpcpp_queue, group_size); return output; } @@ -175,6 +179,8 @@ at::Tensor esimd_gemv_int4_fused2( at::Tensor w1, at::Tensor s1, at::Tensor o1) { int64_t K = w0.size(1) * 2; // packed: K/2 bytes → K elements int64_t N0 = w0.size(0), N1 = w1.size(0); + int64_t ng0 = s0.size(s0.dim() - 1); + int group_size = (ng0 > 0) ? (int)(K / ng0) : 128; EXTRACT_PTR(p_in, input); EXTRACT_PTR(pw0, w0); EXTRACT_PTR(ps0, s0); EXTRACT_PTR(po0, o0); EXTRACT_PTR(pw1, w1); EXTRACT_PTR(ps1, s1); EXTRACT_PTR(po1, o1); @@ -183,7 +189,7 @@ at::Tensor esimd_gemv_int4_fused2( uint8_t* sp[2] = {ps0, ps1}; uint8_t* op[2] = {po0, po1}; uint32_t ns[2] = {(uint32_t)N0, (uint32_t)N1}; - GEMV_int4_fused_host<2>(p_in, wp, sp, op, ns, (uint32_t)K, dpcpp_queue); + GEMV_int4_fused_host<2>(p_in, wp, sp, op, ns, (uint32_t)K, dpcpp_queue, group_size); return o0; } diff --git a/vllm/custom-esimd-kernels-vllm/csrc/xpu/esimd_kernel_gemm.sycl b/vllm/custom-esimd-kernels-vllm/csrc/xpu/esimd_kernel_gemm.sycl index 1e16cda0..2ec9330a 100644 --- a/vllm/custom-esimd-kernels-vllm/csrc/xpu/esimd_kernel_gemm.sycl +++ b/vllm/custom-esimd-kernels-vllm/csrc/xpu/esimd_kernel_gemm.sycl @@ -47,7 +47,8 @@ at::Tensor esimd_gemm_fp8_pert( // ---- INT4 GEMM per-group scale: DPAS kernel for M>=2 ---- // input: [M, K] fp16, weight: [N, K/2] uint8 packed int4, -// weight_scale: [N, K/128] fp16 per-group, output: [M, N] fp16 pre-allocated. +// weight_scale: [N, K/group] fp16 per-group (group=128 GGML or 32 q4_0), +// output: [M, N] fp16 pre-allocated. // Complementary to esimd_gemv_int4 (M=1); for M>=2 use this. at::Tensor esimd_gemm_int4_pgrp( at::Tensor input, at::Tensor weight, at::Tensor weight_scale, @@ -66,12 +67,20 @@ at::Tensor esimd_gemm_int4_pgrp( "esimd_gemm_int4_pgrp: output must be fp16"); TORCH_CHECK(N % 16 == 0, "esimd_gemm_int4_pgrp: N must be a multiple of 16 (DPAS N tile), got N=", N); - TORCH_CHECK(K % INT4_GEMM_GROUP_SIZE == 0, - "esimd_gemm_int4_pgrp: K must be a multiple of 128 (group_size), got K=", K); - TORCH_CHECK(weight_scale.size(0) == N && weight_scale.size(1) == K / INT4_GEMM_GROUP_SIZE, + TORCH_CHECK(K % 128 == 0, + "esimd_gemm_int4_pgrp: K must be a multiple of 128 (DPAS K_LOAD), got K=", K); + // Infer group size from scale shape [N, K/group_size]: 128 (legacy GGML) + // or 32 (vLLM 0.21 sym_int4 / q4_0 32-element blocks). + int64_t n_groups = weight_scale.size(weight_scale.dim() - 1); + int group_size = (n_groups > 0) ? (int)(K / n_groups) : INT4_GEMM_GROUP_SIZE; + TORCH_CHECK(group_size == 128 || group_size == 32, + "esimd_gemm_int4_pgrp: unsupported group_size=", group_size, + " (expected 128 or 32); scale shape [", weight_scale.size(0), + ", ", weight_scale.size(weight_scale.dim()-1), "], K=", K); + TORCH_CHECK(weight_scale.size(0) == N && (int64_t)group_size * n_groups == K, "esimd_gemm_int4_pgrp: weight_scale shape mismatch; expected [", N, - ", ", K / INT4_GEMM_GROUP_SIZE, "], got [", weight_scale.size(0), - ", ", weight_scale.size(1), "]"); + ", ", K / group_size, "], got [", weight_scale.size(0), + ", ", weight_scale.size(weight_scale.dim()-1), "]"); TORCH_CHECK(output.size(0) == M && output.size(1) == N, "esimd_gemm_int4_pgrp: output shape mismatch; expected [", M, ", ", N, "]"); @@ -83,6 +92,6 @@ at::Tensor esimd_gemm_int4_pgrp( reinterpret_cast(p_w), reinterpret_cast(p_sc), reinterpret_cast(p_out), - (uint32_t)M, (uint32_t)N, (uint32_t)K, dpcpp_queue); + (uint32_t)M, (uint32_t)N, (uint32_t)K, dpcpp_queue, group_size); return output; } diff --git a/vllm/custom-esimd-kernels-vllm/csrc/xpu/esimd_kernels/int4_GEMM.h b/vllm/custom-esimd-kernels-vllm/csrc/xpu/esimd_kernels/int4_GEMM.h index 08abb9df..098471c8 100644 --- a/vllm/custom-esimd-kernels-vllm/csrc/xpu/esimd_kernels/int4_GEMM.h +++ b/vllm/custom-esimd-kernels-vllm/csrc/xpu/esimd_kernels/int4_GEMM.h @@ -11,7 +11,10 @@ namespace xesimd = sycl::ext::intel::experimental::esimd; using fp16 = sycl::half; // ============================================================================ -// INT4 GEMM via DPAS (XMX matrix engine), per-group scale (group_size = 128). +// INT4 GEMM via DPAS (XMX matrix engine), per-group scale. +// Supports GROUP_SIZE = 128 (legacy GGML) and 32 (vLLM 0.21 sym_int4 / q4_0). +// GROUP_SIZE is a template param (default 128); host dispatches on the scale +// last dim, so the op signature is unchanged and group128 callers are intact. // // Mirrors the FP8 DPAS V9 kernel (fp8_GEMM_pert.h:FP8_GEMM_DPAS_V9) and // adapts three things that are INT4-specific: @@ -42,8 +45,10 @@ using fp16 = sycl::half; // // Requirements: // N % 16 == 0 -// K % (K_THREADS * 128) == 0 (each K-thread owns whole scale groups) -// input fp16, weight uint8 [N, K/2], weight_scale fp16 [N, K/128], +// K % (K_THREADS * K_LOAD) == 0 (each K-thread owns whole 128-wide loads; +// K_LOAD=128 is a multiple of both 128 and 32, so this also keeps each +// thread owning whole scale groups for either GROUP_SIZE) +// input fp16, weight uint8 [N, K/2], weight_scale fp16 [N, K/GROUP_SIZE], // output fp16 [M, N] pre-allocated. // ============================================================================ @@ -78,11 +83,11 @@ int4_pair_to_vnni_scaled(simd byte_u32, return interleaved.template bit_cast_view(); } -template +template struct GEMM_int4_pgrp_kernel { const fp16* input; // [M, K] const uint8_t* weight; // [N, K/2] packed int4 - const fp16* scale; // [N, K/GROUP_SIZE] + const fp16* scale; // [N, K/GROUP_SIZE] (per-group scales) fp16* output; // [M, N] int M, N, K; int n_groups; // K / GROUP_SIZE @@ -90,7 +95,12 @@ struct GEMM_int4_pgrp_kernel { void operator()(sycl::nd_item<1> item) const SYCL_ESIMD_KERNEL { constexpr int N_TILE = 16; constexpr int M_TILE = 8; - constexpr int K_LOAD = INT4_GEMM_GROUP_SIZE; // 128 + // K_LOAD stays 128 for DPAS efficiency regardless of GROUP_SIZE. + constexpr int K_LOAD = 128; + // How many scale groups a single K_LOAD (=128 K) spans: 1 for + // GROUP_SIZE=128, 4 for GROUP_SIZE=32. K_SUB(=16) divides GROUP_SIZE + // for both, so every K_SUB falls entirely inside one scale group. + constexpr int SUBS_PER_GROUP = GROUP_SIZE / 16; // 8 (g128) or 2 (g32) constexpr int K_SUB = 16; constexpr int SUBS_PER_KLOAD = K_LOAD / K_SUB; // 8 constexpr int SLM_PER_THREAD = M_TILES * 128 * 4; @@ -131,18 +141,27 @@ struct GEMM_int4_pgrp_kernel { const fp16* s_base = scale + (size_t)n_start * n_groups; for (int k_base = k_start; k_base < k_end; k_base += K_LOAD) { - int group_idx = k_base / INT4_GEMM_GROUP_SIZE; - simd scales_f16; - #pragma unroll - for (int n = 0; n < N_TILE; n++) { - scales_f16[n] = s_base[n * n_groups + group_idx]; - } - simd scale_m8_f16 = scales_f16 * fp16(-8.0f); + // Base group index of this K_LOAD. For GROUP_SIZE=128 this is the + // single group; for 32 it is the first of SUBS_PER_KLOAD/SUBS_PER_GROUP + // (=4) groups covered by the 128-wide load. + int base_group = k_base / GROUP_SIZE; #pragma unroll for (int sub = 0; sub < SUBS_PER_KLOAD; sub++) { int k_sub = k_base + sub * K_SUB; + // Which scale group this K_SUB belongs to. Each K_SUB (=16) + // lies entirely in one group (K_SUB divides GROUP_SIZE). For + // GROUP_SIZE=128, SUBS_PER_GROUP=8 so all subs share base_group; + // for 32, SUBS_PER_GROUP=2 so the group advances every 2 subs. + int group_idx = base_group + sub / SUBS_PER_GROUP; + simd scales_f16; + #pragma unroll + for (int n = 0; n < N_TILE; n++) { + scales_f16[n] = s_base[n * n_groups + group_idx]; + } + simd scale_m8_f16 = scales_f16 * fp16(-8.0f); + // set_x is a uint32-element offset; one uint32 covers 8 int4 // K-elements, so stepping by k_sub K means set_x(k_sub / 8). payB_t.set_x((uint32_t)(k_sub / 8)); @@ -269,17 +288,17 @@ struct GEMM_int4_pgrp_kernel { } }; -template +template inline void gemm_int4_pgrp_host_impl( const fp16* input, const uint8_t* weight, const fp16* scale, fp16* output, uint32_t M, uint32_t N, uint32_t K, sycl::queue& q) { - int n_groups = (int)K / INT4_GEMM_GROUP_SIZE; + int n_groups = (int)K / GROUP_SIZE; int num_wg = ((int)N + 15) / 16; q.submit([&](sycl::handler& h) { h.parallel_for( sycl::nd_range<1>({(size_t)(num_wg * K_THREADS)}, {(size_t)K_THREADS}), - GEMM_int4_pgrp_kernel{ + GEMM_int4_pgrp_kernel{ input, weight, scale, output, (int)M, (int)N, (int)K, n_groups}); }); @@ -291,30 +310,36 @@ inline void gemm_int4_pgrp_host_impl( inline void GEMM_int4_pgrp_host( const fp16* input, const uint8_t* weight, const fp16* scale, fp16* output, uint32_t M, uint32_t N, uint32_t K, - sycl::queue& q) { + sycl::queue& q, int group_size = INT4_GEMM_GROUP_SIZE) { int m_tiles = ((int)M + 7) / 8; int n_wgs = ((int)N + 15) / 16; int k_threads = std::max(1, std::min(4, 640 / std::max(n_wgs, 1))); - while (k_threads > 1 && ((int)K % (k_threads * INT4_GEMM_GROUP_SIZE) != 0)) k_threads--; + // Each K-thread must own whole 128-wide K_LOADs (K_LOAD=128 is a multiple + // of both group sizes), so the per-thread K span must be a multiple of 128. + // This is independent of group_size. + while (k_threads > 1 && ((int)K % (k_threads * 128) != 0)) k_threads--; if (k_threads == 3) k_threads = 2; - #define DISPATCH(KT, MT) gemm_int4_pgrp_host_impl( \ + #define DISPATCH(KT, MT, G) gemm_int4_pgrp_host_impl( \ input, weight, scale, output, M, N, K, q) - if (k_threads >= 4) { - if (m_tiles <= 1) DISPATCH(4, 1); - else if (m_tiles <= 2) DISPATCH(4, 2); - else if (m_tiles <= 4) DISPATCH(4, 4); - else DISPATCH(4, 8); - } else if (k_threads >= 2) { - if (m_tiles <= 1) DISPATCH(2, 1); - else if (m_tiles <= 2) DISPATCH(2, 2); - else if (m_tiles <= 4) DISPATCH(2, 4); - else DISPATCH(2, 8); - } else { - if (m_tiles <= 1) DISPATCH(1, 1); - else if (m_tiles <= 2) DISPATCH(1, 2); - else if (m_tiles <= 4) DISPATCH(1, 4); - else DISPATCH(1, 8); - } + #define DISPATCH_MT(KT, G) \ + do { \ + if (m_tiles <= 1) DISPATCH(KT, 1, G); \ + else if (m_tiles <= 2) DISPATCH(KT, 2, G); \ + else if (m_tiles <= 4) DISPATCH(KT, 4, G); \ + else DISPATCH(KT, 8, G); \ + } while (0) + #define DISPATCH_KT(G) \ + do { \ + if (k_threads >= 4) DISPATCH_MT(4, G); \ + else if (k_threads >= 2) DISPATCH_MT(2, G); \ + else DISPATCH_MT(1, G); \ + } while (0) + + if (group_size == 32) { DISPATCH_KT(32); } + else { DISPATCH_KT(128); } + + #undef DISPATCH_KT + #undef DISPATCH_MT #undef DISPATCH } diff --git a/vllm/custom-esimd-kernels-vllm/csrc/xpu/esimd_kernels/int4_GEMV.h b/vllm/custom-esimd-kernels-vllm/csrc/xpu/esimd_kernels/int4_GEMV.h index bb71356f..e526d58f 100644 --- a/vllm/custom-esimd-kernels-vllm/csrc/xpu/esimd_kernels/int4_GEMV.h +++ b/vllm/custom-esimd-kernels-vllm/csrc/xpu/esimd_kernels/int4_GEMV.h @@ -66,8 +66,41 @@ // across threads. // ============================================================================ +// Default / legacy group size (GGML q4_0 with 128-element blocks). static constexpr int INT4_GROUP_SIZE = 128; +// ---------------------------------------------------------------------------- +// Per-lane scale construction for arbitrary GROUP_SIZE <= VL. +// +// VL=128 is kept fixed for bandwidth, but the scale group can be smaller +// (e.g. 32 for vLLM 0.21 sym_int4 / ggml q4_0 with 32-element blocks). +// One K-loop iteration of VL=128 elements then spans VL/GROUP_SIZE groups. +// +// The kernel deinterleaves K into even (K=0,2,4,...) and odd (K=1,3,5,...) +// halves, each VL/2 lanes. Even lane i -> K position 2*i, so it belongs to +// group (2*i) / GROUP_SIZE. Odd lane i -> K position 2*i+1, group (2*i+1)/GS. +// Because GROUP_SIZE is even and lanes are contiguous, both even and odd +// halves partition into segments of (GROUP_SIZE/2) lanes per group: +// lanes [g*(GS/2) : (g+1)*(GS/2)) use group scale g. +// +// build_lane_scale fills a simd by broadcasting each of the +// (VL/GROUP_SIZE) group scales onto its (GROUP_SIZE/2)-lane segment. +// When GROUP_SIZE == VL this degenerates to a single scalar broadcast. +// ---------------------------------------------------------------------------- +template +SYCL_ESIMD_FUNCTION inline simd build_lane_scale( + const fp16* s_row, int group_idx) { + constexpr int N_GROUPS_PER_ITER = VL / GROUP_SIZE; // e.g. 128/32 = 4 + constexpr int LANES_PER_GROUP = GROUP_SIZE / 2; // e.g. 32/2 = 16 + simd sv; + #pragma unroll + for (int g = 0; g < N_GROUPS_PER_ITER; ++g) { + float gs = static_cast(s_row[group_idx + g]); + sv.template select(g * LANES_PER_GROUP) = gs; + } + return sv; +} + // ============================================================================ // INT4 unpacking: uint8[VL/2] → two float[VL/2] vectors @@ -130,10 +163,14 @@ inline void select_vl_ks_int4(uint32_t N, uint32_t K, int& vl, int& ks) { if (N <= 128 && K >= 2048) { ks = 8; } else if (N <= 512 && K >= 2048) { ks = 4; } - // Enforce: kp = K / ks must be a multiple of 128 (GROUP_SIZE). - // If not, halve ks until it is. + // Enforce: kp = K / ks must be a multiple of VL (=128) so the K-loop + // (step VL) never over-reads and never splits the VL-block across threads. + // This is stricter than (and therefore satisfies) the scale-group + // alignment for both GROUP_SIZE=128 and GROUP_SIZE=32 (128 % 32 == 0). + // If not aligned, halve ks until it is. + constexpr int KP_ALIGN = 128; int kp = K / ks; - while (kp % INT4_GROUP_SIZE != 0 && ks > 1) { + while (kp % KP_ALIGN != 0 && ks > 1) { ks /= 2; kp = K / ks; } @@ -162,7 +199,7 @@ inline void select_vl_ks_int4(uint32_t N, uint32_t K, int& vl, int& ks) { // 5. After loop: horizontal reduce acc_even + acc_odd → scalar output. // ============================================================================ -template +template struct GEMV_int4_kernel { const fp16* input; // [1, K] fp16 — input activation vector const uint8_t* weight; // [N, K/2] uint8 — packed INT4 weights @@ -194,8 +231,8 @@ struct GEMV_int4_kernel { const uint8_t* w_row = weight + (size_t)n * (K / 2); // packed weight const fp16* s_row = scale + (size_t)n * n_groups; // group scales - // Track which group we're in (advances by VL/GROUP_SIZE per iteration). - int group_idx = ks / INT4_GROUP_SIZE; + // Track which group we are in (advances by VL/GROUP_SIZE per iteration). + int group_idx = ks / GROUP_SIZE; for (int k = ks; k < ks + kp; k += VL) { // --- Load input: VL fp16 values (VL*2 = 256 bytes) --- @@ -217,16 +254,17 @@ struct GEMV_int4_kernel { simd wf_even, wf_odd; int4_dequant(raw, wf_even, wf_odd); - // --- Load per-group scale --- - // With VL=128 and GROUP_SIZE=128, exactly 1 group per iteration. - // Scale is fp16, convert to float for FMA precision. - float gs = static_cast(s_row[group_idx]); - group_idx += VL / INT4_GROUP_SIZE; // +1 when VL=GROUP_SIZE + // --- Load per-group scale(s) --- + // VL elements span VL/GROUP_SIZE groups (1 when GROUP_SIZE==VL). + // Build per-lane scale vectors for the even/odd halves; both halves + // share the same segment layout (see build_lane_scale comment). + simd sv = + build_lane_scale(s_row, group_idx); + group_idx += VL / GROUP_SIZE; // --- FMA: accumulate input * weight * scale --- - // Apply scale to weight first (scalar broadcast), then multiply by input. - wf_even *= gs; - wf_odd *= gs; + wf_even *= sv; + wf_odd *= sv; acc_even += in_even * wf_even; acc_odd += in_odd * wf_odd; } @@ -271,14 +309,15 @@ inline void GEMV_int4_host( uint8_t* output_data, uint32_t N, uint32_t K, - sycl::queue& q) { + sycl::queue& q, + int group_size = INT4_GROUP_SIZE) { auto* p_in = reinterpret_cast(input_data); auto* p_w = reinterpret_cast(weight_data); auto* p_sc = reinterpret_cast(scale_data); auto* p_out = reinterpret_cast(output_data); - int n_groups = K / INT4_GROUP_SIZE; + int n_groups = K / group_size; int vl, ks; select_vl_ks_int4(N, K, vl, ks); @@ -286,20 +325,25 @@ inline void GEMV_int4_host( int global = N * ks; // total threads = N workgroups × K_SPLIT threads/WG int local = ks; // threads per workgroup - // Phase 1: VL is always 128. Only K_SPLIT varies. - #define LAUNCH_INT4(S) \ + // VL is always 128. K_SPLIT (S) varies; GROUP_SIZE (G) is 128 or 32. + #define LAUNCH_INT4(S, G) \ q.submit([&](sycl::handler& h) { \ h.parallel_for(sycl::nd_range<1>(global, local), \ - GEMV_int4_kernel<128, S>{ \ + GEMV_int4_kernel<128, S, G>{ \ p_in, p_w, p_sc, p_out, (int)N, (int)K, n_groups}); \ }); - if (ks == 1) { LAUNCH_INT4(1) } - else if (ks == 2) { LAUNCH_INT4(2) } - else if (ks == 4) { LAUNCH_INT4(4) } - else if (ks == 8) { LAUNCH_INT4(8) } - else { LAUNCH_INT4(1) } + #define DISPATCH_KS(G) \ + if (ks == 1) { LAUNCH_INT4(1, G) } \ + else if (ks == 2) { LAUNCH_INT4(2, G) } \ + else if (ks == 4) { LAUNCH_INT4(4, G) } \ + else if (ks == 8) { LAUNCH_INT4(8, G) } \ + else { LAUNCH_INT4(1, G) } + if (group_size == 32) { DISPATCH_KS(32) } + else { DISPATCH_KS(128) } + + #undef DISPATCH_KS #undef LAUNCH_INT4 } @@ -318,7 +362,7 @@ inline void GEMV_int4_host( // it belongs to via cumulative N sums (same pattern as fp8_GEMV_v2.h). // ============================================================================ -template +template struct GEMV_int4_fused_kernel { const fp16* input; // [1, K] fp16 — shared input const uint8_t* weights[GEMV_COUNT]; // packed INT4 weight per matrix @@ -361,7 +405,7 @@ struct GEMV_int4_fused_kernel { simd acc_even = 0.0f; simd acc_odd = 0.0f; - int group_idx = ks / INT4_GROUP_SIZE; + int group_idx = ks / GROUP_SIZE; for (int k = ks; k < ks + kp; k += VL) { // Load + deinterleave input. @@ -376,12 +420,13 @@ struct GEMV_int4_fused_kernel { simd wf_even, wf_odd; int4_dequant(raw, wf_even, wf_odd); - // Per-group scale + FMA. - float gs = static_cast(s_row[group_idx]); - group_idx += VL / INT4_GROUP_SIZE; + // Per-group scale(s) + FMA (VL spans VL/GROUP_SIZE groups). + simd sv = + build_lane_scale(s_row, group_idx); + group_idx += VL / GROUP_SIZE; - wf_even *= gs; - wf_odd *= gs; + wf_even *= sv; + wf_odd *= sv; acc_even += in_even * wf_even; acc_odd += in_odd * wf_odd; } @@ -419,14 +464,15 @@ inline void GEMV_int4_fused_host( uint8_t* output_ptrs[GEMV_COUNT], uint32_t Ns[GEMV_COUNT], uint32_t K, - sycl::queue& q) { + sycl::queue& q, + int group_size = INT4_GROUP_SIZE) { auto* p_in = reinterpret_cast(input_data); uint32_t total_N = 0; for (int i = 0; i < GEMV_COUNT; i++) total_N += Ns[i]; - int n_groups = K / INT4_GROUP_SIZE; + int n_groups = K / group_size; int vl, ks; select_vl_ks_int4(total_N, K, vl, ks); @@ -434,9 +480,9 @@ inline void GEMV_int4_fused_host( int global = total_N * ks; int local = ks; - #define LAUNCH_INT4_FUSED(S) \ + #define LAUNCH_INT4_FUSED(S, G) \ q.submit([&](sycl::handler& h) { \ - GEMV_int4_fused_kernel<128, S, GEMV_COUNT> kern; \ + GEMV_int4_fused_kernel<128, S, GEMV_COUNT, G> kern; \ kern.input = p_in; \ kern.K = (int)K; \ kern.n_groups = n_groups; \ @@ -452,11 +498,16 @@ inline void GEMV_int4_fused_host( h.parallel_for(sycl::nd_range<1>(global, local), kern); \ }); - if (ks == 1) { LAUNCH_INT4_FUSED(1) } - else if (ks == 2) { LAUNCH_INT4_FUSED(2) } - else if (ks == 4) { LAUNCH_INT4_FUSED(4) } - else if (ks == 8) { LAUNCH_INT4_FUSED(8) } - else { LAUNCH_INT4_FUSED(1) } + #define DISPATCH_KS_FUSED(G) \ + if (ks == 1) { LAUNCH_INT4_FUSED(1, G) } \ + else if (ks == 2) { LAUNCH_INT4_FUSED(2, G) } \ + else if (ks == 4) { LAUNCH_INT4_FUSED(4, G) } \ + else if (ks == 8) { LAUNCH_INT4_FUSED(8, G) } \ + else { LAUNCH_INT4_FUSED(1, G) } + + if (group_size == 32) { DISPATCH_KS_FUSED(32) } + else { DISPATCH_KS_FUSED(128) } + #undef DISPATCH_KS_FUSED #undef LAUNCH_INT4_FUSED }