From 8739beb3063e1223ee68b2b6e17c4ca866adfbb5 Mon Sep 17 00:00:00 2001 From: "Jingxin (Philip) Li" Date: Sat, 5 Sep 2026 10:37:12 +0800 Subject: [PATCH 1/2] sycl : fix test-backend-ops CI break && restore Kronecker product FWHT support (#28016) (#28254) MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit * Reapply "sycl : add Kronecker product FWHT support for sizes 384, 640, 768, 12…" (#28184) This reverts commit c845263f8b7d60113e213a3bd2d5cc6472ccf204. * tests : fix unused variable M in test-backend-ops * tests: fix trailing space error and isolate kronecker tests for sycl backend only (cherry picked from commit 4d9176092d00586775af140581bb0b558ddc4389) --- ggml/src/ggml-sycl/fwht.cpp | 172 ++++++++++++++++++++++++++++++++++++ tests/test-backend-ops.cpp | 129 ++++++++++++++++++++++++--- 2 files changed, 288 insertions(+), 13 deletions(-) diff --git a/ggml/src/ggml-sycl/fwht.cpp b/ggml/src/ggml-sycl/fwht.cpp index 2312b3d131b7..39f273beaa9f 100644 --- a/ggml/src/ggml-sycl/fwht.cpp +++ b/ggml/src/ggml-sycl/fwht.cpp @@ -1,6 +1,50 @@ #include "fwht.hpp" #include +#define P 1.0f +#define N -1.0f + +// constant Hadamard matrix via Paley I construction +static constexpr float H12[12][12] = { + { P, P, P, P, P, P, P, P, P, P, P, P }, + { P, N, P, N, P, P, P, N, N, N, P, N }, + { P, N, N, P, N, P, P, P, N, N, N, P }, + { P, P, N, N, P, N, P, P, P, N, N, N }, + { P, N, P, N, N, P, N, P, P, P, N, N }, + { P, N, N, P, N, N, P, N, P, P, P, N }, + { P, N, N, N, P, N, N, P, N, P, P, P }, + { P, P, N, N, N, P, N, N, P, N, P, P }, + { P, P, P, N, N, N, P, N, N, P, N, P }, + { P, P, P, P, N, N, N, P, N, N, P, N }, + { P, N, P, P, P, N, N, N, P, N, N, P }, + { P, P, N, P, P, P, N, N, N, P, N, N } +}; + +static constexpr float H20[20][20] = { + { P, P, P, P, P, P, P, P, P, P, P, P, P, P, P, P, P, P, P, P }, + { P, N, P, N, N, P, P, P, P, N, P, N, P, N, N, N, N, P, P, N }, + { P, N, N, P, N, N, P, P, P, P, N, P, N, P, N, N, N, N, P, P }, + { P, P, N, N, P, N, N, P, P, P, P, N, P, N, P, N, N, N, N, P }, + { P, P, P, N, N, P, N, N, P, P, P, P, N, P, N, P, N, N, N, N }, + { P, N, P, P, N, N, P, N, N, P, P, P, P, N, P, N, P, N, N, N }, + { P, N, N, P, P, N, N, P, N, N, P, P, P, P, N, P, N, P, N, N }, + { P, N, N, N, P, P, N, N, P, N, N, P, P, P, P, N, P, N, P, N }, + { P, N, N, N, N, P, P, N, N, P, N, N, P, P, P, P, N, P, N, P }, + { P, P, N, N, N, N, P, P, N, N, P, N, N, P, P, P, P, N, P, N }, + { P, N, P, N, N, N, N, P, P, N, N, P, N, N, P, P, P, P, N, P }, + { P, P, N, P, N, N, N, N, P, P, N, N, P, N, N, P, P, P, P, N }, + { P, N, P, N, P, N, N, N, N, P, P, N, N, P, N, N, P, P, P, P }, + { P, P, N, P, N, P, N, N, N, N, P, P, N, N, P, N, N, P, P, P }, + { P, P, P, N, P, N, P, N, N, N, N, P, P, N, N, P, N, N, P, P }, + { P, P, P, P, N, P, N, P, N, N, N, N, P, P, N, N, P, N, N, P }, + { P, P, P, P, P, N, P, N, P, N, N, N, N, P, P, N, N, P, N, N }, + { P, N, P, P, P, P, N, P, N, P, N, N, N, N, P, P, N, N, P, N }, + { P, N, N, P, P, P, P, N, P, N, P, N, N, N, N, P, P, N, N, P }, + { P, P, N, N, P, P, P, P, N, P, N, P, N, N, N, N, P, P, N, N } +}; + +#undef P +#undef N template static void fwht_kernel(const float * __restrict__ src, float * __restrict__ dst, const int64_t n_rows, @@ -80,6 +124,122 @@ static void launch_fwht(const float * src, float * dst, const int64_t n_rows, co }); } +template +static void kronecker_kernel(const float * __restrict__ src, + float * __restrict__ dst, + const int64_t n_rows, + const float scale, + const sycl::nd_item<2> & item) { + static_assert(m == 12 || m == 20, "block size has to be 12 or 20."); + + const sycl::sub_group sg = item.get_sub_group(); + + const int64_t r = item.get_global_id(0); + if (r >= n_rows) { + return; + } + + src += r * N; + dst += r * N; + + constexpr int blocks_per_group = N / m; + constexpr int el_w = blocks_per_group / WARP_SIZE; + static_assert(el_w >= 1 && blocks_per_group % WARP_SIZE == 0, "blocks_per_group must be a multiple of WARP_SIZE"); + float reg[el_w * m]; + const int lane = sg.get_local_linear_id(); + +#pragma unroll + for (int i = 0; i < el_w; ++i) { + const int b_idx = i * WARP_SIZE + lane; + +#pragma unroll + for (int j = 0; j < m; ++j) { + reg[i * m + j] = src[b_idx * m + j] * scale; + } + } + +#pragma unroll + for (int b = 0; b < el_w; ++b) { + float z[m] = { 0.0f }; + +#pragma unroll + for (int i = 0; i < m; ++i) { +#pragma unroll + for (int j = 0; j < m; ++j) { + const float h = (m == 12 ? H12[j][i] : H20[j][i]); + z[i] += reg[b * m + j] * h; + } + } + +#pragma unroll + for (int i = 0; i < m; ++i) { + reg[b * m + i] = z[i]; + } + } + +#pragma unroll + for (int h = 1; h < WARP_SIZE; h *= 2) { +#pragma unroll + for (int j = 0; j < el_w; ++j) { +#pragma unroll + for (int k = 0; k < m; ++k) { + const float val = reg[j * m + k]; + const float val2 = dpct::permute_sub_group_by_xor(sg, val, h, WARP_SIZE); + + reg[j * m + k] = (lane & h) == 0 ? val + val2 : val2 - val; + } + } + } + +#pragma unroll + for (int h = WARP_SIZE; h < blocks_per_group; h *= 2) { + const int step = h / WARP_SIZE; +#pragma unroll + for (int j = 0; j < el_w; j += 2 * step) { +#pragma unroll + for (int s = 0; s < step; ++s) { +#pragma unroll + for (int k = 0; k < m; ++k) { + const float x = reg[(j + s) * m + k]; + const float y = reg[(j + s + step) * m + k]; + + reg[(j + s) * m + k] = x + y; + reg[(j + s + step) * m + k] = x - y; + } + } + } + } + +#pragma unroll + for (int i = 0; i < el_w; ++i) { + const int b_idx = i * WARP_SIZE + lane; +#pragma unroll + for (int k = 0; k < m; ++k) { + dst[b_idx * m + k] = reg[i * m + k]; + } + } +} + +template +static void launch_kronecker(const float * src, + float * dst, + const int64_t n_rows, + const float scale, + dpct::queue_ptr stream) { + constexpr int rows_per_block = 4; + + const int64_t num_blocks = (n_rows + rows_per_block - 1) / rows_per_block; + + // dim 1 is the fastest-varying, so a sub-group is exactly one row's WARP_SIZE lanes. + const sycl::range<2> global(num_blocks * rows_per_block, WARP_SIZE); + const sycl::range<2> local(rows_per_block, WARP_SIZE); + + stream->parallel_for(sycl::nd_range<2>(global, local), + [=](sycl::nd_item<2> item) [[sycl::reqd_sub_group_size(WARP_SIZE)]] { + kronecker_kernel(src, dst, n_rows, scale, item); + }); +} + bool ggml_sycl_op_fwht(ggml_backend_sycl_context & ctx, const ggml_tensor * src, ggml_tensor * dst) { if (src->type != GGML_TYPE_F32 || dst->type != GGML_TYPE_F32) { return false; @@ -113,6 +273,18 @@ bool ggml_sycl_op_fwht(ggml_backend_sycl_context & ctx, const ggml_tensor * src, case 512: launch_fwht<512>(src_d, dst_d, rows, scale, stream); return true; + case 384: + launch_kronecker<384, 12>(src_d, dst_d, rows, scale, stream); + return true; + case 768: + launch_kronecker<768, 12>(src_d, dst_d, rows, scale, stream); + return true; + case 640: + launch_kronecker<640, 20>(src_d, dst_d, rows, scale, stream); + return true; + case 1280: + launch_kronecker<1280, 20>(src_d, dst_d, rows, scale, stream); + return true; default: return false; } diff --git a/tests/test-backend-ops.cpp b/tests/test-backend-ops.cpp index 3e46715a2096..4a5fb0bd7270 100644 --- a/tests/test-backend-ops.cpp +++ b/tests/test-backend-ops.cpp @@ -4633,6 +4633,51 @@ struct test_mul_mat : public test_case { } }; +#define P 1.0f +#define N -1.0f + +// constant Hadamard matrix via Paley I construction +static constexpr float H12[12][12] = { + { P, P, P, P, P, P, P, P, P, P, P, P }, + { P, N, P, N, P, P, P, N, N, N, P, N }, + { P, N, N, P, N, P, P, P, N, N, N, P }, + { P, P, N, N, P, N, P, P, P, N, N, N }, + { P, N, P, N, N, P, N, P, P, P, N, N }, + { P, N, N, P, N, N, P, N, P, P, P, N }, + { P, N, N, N, P, N, N, P, N, P, P, P }, + { P, P, N, N, N, P, N, N, P, N, P, P }, + { P, P, P, N, N, N, P, N, N, P, N, P }, + { P, P, P, P, N, N, N, P, N, N, P, N }, + { P, N, P, P, P, N, N, N, P, N, N, P }, + { P, P, N, P, P, P, N, N, N, P, N, N } +}; + +static constexpr float H20[20][20] = { + { P, P, P, P, P, P, P, P, P, P, P, P, P, P, P, P, P, P, P, P }, + { P, N, P, N, N, P, P, P, P, N, P, N, P, N, N, N, N, P, P, N }, + { P, N, N, P, N, N, P, P, P, P, N, P, N, P, N, N, N, N, P, P }, + { P, P, N, N, P, N, N, P, P, P, P, N, P, N, P, N, N, N, N, P }, + { P, P, P, N, N, P, N, N, P, P, P, P, N, P, N, P, N, N, N, N }, + { P, N, P, P, N, N, P, N, N, P, P, P, P, N, P, N, P, N, N, N }, + { P, N, N, P, P, N, N, P, N, N, P, P, P, P, N, P, N, P, N, N }, + { P, N, N, N, P, P, N, N, P, N, N, P, P, P, P, N, P, N, P, N }, + { P, N, N, N, N, P, P, N, N, P, N, N, P, P, P, P, N, P, N, P }, + { P, P, N, N, N, N, P, P, N, N, P, N, N, P, P, P, P, N, P, N }, + { P, N, P, N, N, N, N, P, P, N, N, P, N, N, P, P, P, P, N, P }, + { P, P, N, P, N, N, N, N, P, P, N, N, P, N, N, P, P, P, P, N }, + { P, N, P, N, P, N, N, N, N, P, P, N, N, P, N, N, P, P, P, P }, + { P, P, N, P, N, P, N, N, N, N, P, P, N, N, P, N, N, P, P, P }, + { P, P, P, N, P, N, P, N, N, N, N, P, P, N, N, P, N, N, P, P }, + { P, P, P, P, N, P, N, P, N, N, N, N, P, P, N, N, P, N, N, P }, + { P, P, P, P, P, N, P, N, P, N, N, N, N, P, P, N, N, P, N, N }, + { P, N, P, P, P, P, N, P, N, P, N, N, N, N, P, P, N, N, P, N }, + { P, N, N, P, P, P, P, N, P, N, P, N, N, N, N, P, P, N, N, P }, + { P, P, N, N, P, P, P, P, N, P, N, P, N, N, N, N, P, P, N, N } +}; + +#undef P +#undef N + // GGML_HINT_SRC0_IS_HADAMARD struct test_mul_mat_hadamard : public test_mul_mat { test_mul_mat_hadamard(ggml_type type_a = GGML_TYPE_F32, ggml_type type_b = GGML_TYPE_F32, @@ -4657,20 +4702,58 @@ struct test_mul_mat_hadamard : public test_mul_mat { void initialize_tensors(ggml_context * ctx) override { for (ggml_tensor * t = ggml_get_first_tensor(ctx); t != NULL; t = ggml_get_next_tensor(ctx, t)) { if (strcmp(t->name, "a") == 0) { - const int64_t n_cols = t->ne[0]; - const int64_t n_rows = ggml_nrows(t); + const int64_t n_cols = t->ne[0]; + const int64_t n_rows = ggml_nrows(t); std::vector data(n_cols * n_rows); - float scale = 1.0f / sqrtf((float)n_cols); - for (int64_t r = 0; r < n_rows; r++) { - float * row_data = data.data() + r * n_cols; - for (int64_t i = 0; i < n_cols; i++) { - int pop = 0; - int64_t val = r & i; - while (val) { - pop += (val & 1); - val >>= 1; + float scale = 1.0f / sqrtf((float) n_cols); + + auto is_pow2 = [](const int64_t a) { + return (a > 0) && ((a & (a - 1)) == 0); + }; +#ifdef GGML_USE_SYCL + const bool is_kronecker = + ((n_cols % 12 == 0) && is_pow2(n_cols / 12)) || ((n_cols % 20 == 0) && is_pow2(n_cols / 20)); +#else + const bool is_kronecker = false; +#endif + if (is_kronecker) { + const int64_t B = (n_cols % 12 == 0 && is_pow2(n_cols / 12)) ? 12 : 20; + for (int64_t r = 0; r < n_rows; r++) { + float * row_data = data.data() + r * n_cols; + const int64_t r_mod = r % n_cols; + const int64_t r_b = r_mod / B; + const int64_t r_m = r_mod % B; + + for (int64_t i = 0; i < n_cols; i++) { + const int64_t c_b = i / B; + const int64_t c_m = i % B; + + int pop = 0; + int64_t val = r_b & c_b; + while (val) { + pop += (val & 1); + val >>= 1; + } + const float sign_m = (pop % 2 == 0) ? 1.0f : -1.0f; + const float sign_b = (B == 12) ? H12[c_m][r_m] : H20[c_m][r_m]; + + row_data[i] = scale * sign_b * sign_m; + } + } + } + + else if (is_pow2(n_cols)) { + for (int64_t r = 0; r < n_rows; r++) { + float * row_data = data.data() + r * n_cols; + for (int64_t i = 0; i < n_cols; i++) { + int pop_cnt = 0; + int64_t val = r & i; + while (val) { + pop_cnt += (val & 1); + val >>= 1; + } + row_data[i] = (pop_cnt % 2 == 0) ? scale : -scale; } - row_data[i] = (pop % 2 == 0) ? scale : -scale; } } ggml_backend_tensor_set(t, data.data(), 0, data.size() * sizeof(float)); @@ -9292,7 +9375,17 @@ static std::vector> make_test_cases_eval() { test_cases.emplace_back(new test_fwht_signed(8192, 8192, 2)); test_cases.emplace_back(new test_fwht_signed(8192, 8192, 1, GGML_TYPE_F16)); test_cases.emplace_back(new test_mul_mat_hadamard(GGML_TYPE_F32, GGML_TYPE_F32, 32, 1, 32)); // too small (N<64) - + test_cases.emplace_back(new test_mul_mat_hadamard(GGML_TYPE_F32, GGML_TYPE_F32, 1024, 1, 1024)); // too big (N>512) +#ifdef GGML_USE_SYCL + test_cases.emplace_back(new test_mul_mat_hadamard(GGML_TYPE_F32, GGML_TYPE_F32, 384, 1, 384)); // m=12 (N=384) + test_cases.emplace_back(new test_mul_mat_hadamard(GGML_TYPE_F32, GGML_TYPE_F32, 384, 32, 384)); // m=12 (batch) + test_cases.emplace_back( + new test_mul_mat_hadamard(GGML_TYPE_F32, GGML_TYPE_F32, 384, 4, 384, { 2, 3 })); // m=12 (multi-dim) + test_cases.emplace_back(new test_mul_mat_hadamard(GGML_TYPE_F32, GGML_TYPE_F32, 768, 1, 768)); // m=12 (N=768) + test_cases.emplace_back(new test_mul_mat_hadamard(GGML_TYPE_F32, GGML_TYPE_F32, 640, 1, 640)); // m=20 (N=640) + test_cases.emplace_back(new test_mul_mat_hadamard(GGML_TYPE_F32, GGML_TYPE_F32, 640, 32, 640)); // m=20 (batch) + test_cases.emplace_back(new test_mul_mat_hadamard(GGML_TYPE_F32, GGML_TYPE_F32, 1280, 1, 1280)); // m=20 (N=1280) +#endif #if 0 // > 4GB A matrix. Too slow to be enabled by default. test_cases.emplace_back(new test_mul_mat(GGML_TYPE_F16, GGML_TYPE_F16, 900000, 3, 2592, {1, 1}, {1, 1})); @@ -10578,6 +10671,16 @@ static std::vector> make_test_cases_perf() { } } +#ifdef GGML_USE_SYCL + test_cases.emplace_back(new test_mul_mat_hadamard(GGML_TYPE_F32, GGML_TYPE_F32, 384, 1, 384)); // m=12 (N=384) + test_cases.emplace_back(new test_mul_mat_hadamard(GGML_TYPE_F32, GGML_TYPE_F32, 384, 32, 384)); // m=12 (batch) + test_cases.emplace_back( + new test_mul_mat_hadamard(GGML_TYPE_F32, GGML_TYPE_F32, 384, 4, 384, { 2, 3 })); // m=12 (multi-dim) + test_cases.emplace_back(new test_mul_mat_hadamard(GGML_TYPE_F32, GGML_TYPE_F32, 768, 1, 768)); // m=12 (N=768) + test_cases.emplace_back(new test_mul_mat_hadamard(GGML_TYPE_F32, GGML_TYPE_F32, 640, 1, 640)); // m=20 (N=640) + test_cases.emplace_back(new test_mul_mat_hadamard(GGML_TYPE_F32, GGML_TYPE_F32, 640, 32, 640)); // m=20 (batch) + test_cases.emplace_back(new test_mul_mat_hadamard(GGML_TYPE_F32, GGML_TYPE_F32, 1280, 1, 1280)); // m=20 (N=1280) +#endif test_cases.emplace_back(new test_solve_tri(GGML_TYPE_F32, { 64, 64, 4, 4 }, { 32, 64, 4, 4 })); test_cases.emplace_back(new test_solve_tri(GGML_TYPE_F32, { 128, 128, 4, 2 }, { 32, 128, 4, 2 })); // qwen3next with CHUNK_SIZE 64 From c69d6745bdb086ecd41932aedf8c0cbeca814ec3 Mon Sep 17 00:00:00 2001 From: bri-prism <288398250+bri-prism@users.noreply.github.com> Date: Sun, 27 Sep 2026 03:08:19 -0700 Subject: [PATCH 2/2] sycl: FWHT kernels for block widths above 512 (#29243) The SYCL FWHT covers 64 to 512 via the standard butterfly network, plus 384/640/768/1280 via the Kronecker/Paley construction added separately in Hadamard hint can produce (1024, 2048, 4096, 8192); those still fall through to the default case and run as a dense GEMM against the materialized rotation tensor, correct but O(n^2) instead of O(n log n). fwht_kernel_wide runs one row per work-group instead of per sub-group, so each work-item keeps N/NT values rather than N/WARP_SIZE. Butterflies below the sub-group width still shuffle; those up to the work-group width go through work-group local memory; the rest stay in registers. Same butterfly and sign convention as the existing narrow kernel. ggml's SYCL backend registration (dpct::dev_mgr) unconditionally requires a GPU-labeled platform to exist and throws before any op-level test can run, so test-backend-ops could not be exercised on this box (a GPU-less pod) even via the CPU device. Verified instead with a standalone harness: the same kernel body run through a real SYCL CPU device (Intel oneAPI DPC++ 2026.1, OpenCL CPU backend), checked against an independent recursive-doubling Hadamard reference, cross-validated by first running the existing unmodified narrow kernel through the identical harness and confirming it passes (rules out a reference-convention bug before trusting a pass on the new code). Random-input results for all four widths, single- and multi-row: N=1024 NT=256 rows=1 max_abs_err=1.7e-07 max_rel_err=4.9e-04 PASS N=2048 NT=256 rows=1 max_abs_err=1.9e-07 max_rel_err=2.0e-04 PASS N=4096 NT=256 rows=1 max_abs_err=2.0e-07 max_rel_err=1.4e-04 PASS N=8192 NT=256 rows=1 max_abs_err=2.5e-07 max_rel_err=3.8e-03 PASS N=1024 NT=256 rows=7 max_abs_err=2.4e-07 max_rel_err=1.0e-03 PASS N=2048 NT=256 rows=5 max_abs_err=3.0e-07 max_rel_err=9.4e-04 PASS N=4096 NT=256 rows=3 max_abs_err=2.7e-07 max_rel_err=1.7e-03 PASS N=8192 NT=256 rows=2 max_abs_err=2.5e-07 max_rel_err=1.9e-03 PASS This covers the kernel algorithm itself; it does not exercise the ggml dispatch/supports_op integration end to end, which needs a real GPU (or a SYCL GPU plugin) to get past backend registration. test-backend-ops build is verified: fwht.cpp recompiles with zero warnings as part of ggml-sycl. (cherry picked from commit c8296709920f9c1ae168bfd5fe66f9f73637bd60) --- ggml/src/ggml-sycl/fwht.cpp | 113 ++++++++++++++++++++++++++++++++++++ 1 file changed, 113 insertions(+) diff --git a/ggml/src/ggml-sycl/fwht.cpp b/ggml/src/ggml-sycl/fwht.cpp index 39f273beaa9f..fb48d7fec70d 100644 --- a/ggml/src/ggml-sycl/fwht.cpp +++ b/ggml/src/ggml-sycl/fwht.cpp @@ -124,6 +124,107 @@ static void launch_fwht(const float * src, float * dst, const int64_t n_rows, co }); } +// Wide blocks: one row per work-group instead of per sub-group, so each work-item +// keeps N/NT values rather than N/WARP_SIZE. Butterflies below the sub-group width +// still shuffle; those up to NT go through work-group local memory; the rest stay +// in registers. +template +static void fwht_kernel_wide(const float * __restrict__ src, + float * __restrict__ dst, + const int64_t n_rows, + const float scale, + const sycl::nd_item<2> & item, + float * smem) { + const int64_t r = item.get_global_id(0); + if (r >= n_rows) { + return; + } + + src += r * N; + dst += r * N; + + constexpr int el_w = N / NT; + static_assert(el_w >= 1 && N % NT == 0, "row must be a whole number of work-group widths"); + + const int tid = item.get_local_id(1); + + float reg[el_w]; +#pragma unroll + for (int i = 0; i < el_w; ++i) { + reg[i] = src[i * NT + tid] * scale; + } + + const sycl::sub_group sg = item.get_sub_group(); + const int lane = sg.get_local_linear_id(); + + // Butterflies inside the sub-group, same pattern as the narrow kernel. +#pragma unroll + for (int h = 1; h < WARP_SIZE; h *= 2) { +#pragma unroll + for (int j = 0; j < el_w; ++j) { + const float val = reg[j]; + const float val2 = dpct::permute_sub_group_by_xor(sg, val, h, WARP_SIZE); + + reg[j] = (lane & h) == 0 ? val + val2 : val2 - val; + } + } + + // Butterflies from the sub-group width up to NT: the partner lane is outside + // this sub-group, so it goes through work-group local memory instead of a shuffle. + for (int h = WARP_SIZE; h < NT; h *= 2) { +#pragma unroll + for (int j = 0; j < el_w; ++j) { + smem[j * NT + tid] = reg[j]; + } + item.barrier(sycl::access::fence_space::local_space); +#pragma unroll + for (int j = 0; j < el_w; ++j) { + const float val = reg[j]; + const float val2 = smem[j * NT + (tid ^ h)]; + reg[j] = (tid & h) == 0 ? val + val2 : val2 - val; + } + item.barrier(sycl::access::fence_space::local_space); + } + + // Butterflies across registers: h is a multiple of NT, so the partner of element + // i*NT + tid lives in reg[i + h/NT] on the same work-item. + for (int h = NT; h < N; h *= 2) { + const int step = h / NT; + for (int j = 0; j < el_w; j += 2 * step) { + for (int k = 0; k < step; ++k) { + const float x = reg[j + k]; + const float y = reg[j + k + step]; + + reg[j + k] = x + y; + reg[j + k + step] = x - y; + } + } + } + +#pragma unroll + for (int i = 0; i < el_w; ++i) { + dst[i * NT + tid] = reg[i]; + } +} + +template +static void launch_fwht_wide(const float * src, + float * dst, + const int64_t n_rows, + const float scale, + dpct::queue_ptr stream) { + const sycl::range<2> global(n_rows, NT); + const sycl::range<2> local(1, NT); + + stream->submit([&](sycl::handler & cgh) { + sycl::local_accessor smem(sycl::range<1>(N), cgh); + cgh.parallel_for(sycl::nd_range<2>(global, local), + [=](sycl::nd_item<2> item) [[sycl::reqd_sub_group_size(WARP_SIZE)]] { + fwht_kernel_wide(src, dst, n_rows, scale, item, get_pointer(smem)); + }); + }); +} + template static void kronecker_kernel(const float * __restrict__ src, float * __restrict__ dst, @@ -285,6 +386,18 @@ bool ggml_sycl_op_fwht(ggml_backend_sycl_context & ctx, const ggml_tensor * src, case 1280: launch_kronecker<1280, 20>(src_d, dst_d, rows, scale, stream); return true; + case 1024: + launch_fwht_wide<1024, 256>(src_d, dst_d, rows, scale, stream); + return true; + case 2048: + launch_fwht_wide<2048, 256>(src_d, dst_d, rows, scale, stream); + return true; + case 4096: + launch_fwht_wide<4096, 256>(src_d, dst_d, rows, scale, stream); + return true; + case 8192: + launch_fwht_wide<8192, 256>(src_d, dst_d, rows, scale, stream); + return true; default: return false; }