From 1f3d318734c61cf6f3b209726cdd3f9c300a782e Mon Sep 17 00:00:00 2001 From: "Jingxin (Philip) Li" Date: Tue, 1 Sep 2026 23:47:08 +0800 Subject: [PATCH] sycl : add Kronecker product FWHT support for sizes 384, 640, 768, 1280 (#28016) --- ggml/src/ggml-sycl/fwht.cpp | 172 ++++++++++++++++++++++++++++++++++++ tests/test-backend-ops.cpp | 124 +++++++++++++++++++++++--- 2 files changed, 283 insertions(+), 13 deletions(-) diff --git a/ggml/src/ggml-sycl/fwht.cpp b/ggml/src/ggml-sycl/fwht.cpp index 2312b3d131..39f273beaa 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 d61b379282..28fe48b06e 100644 --- a/tests/test-backend-ops.cpp +++ b/tests/test-backend-ops.cpp @@ -4660,6 +4660,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, @@ -4684,20 +4729,57 @@ 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); + }; + + const bool is_kronecker = + ((n_cols % 12 == 0) && is_pow2(n_cols / 12)) || ((n_cols % 20 == 0) && is_pow2(n_cols / 20)); + + if (is_kronecker) { + const int64_t B = (n_cols % 12 == 0 && is_pow2(n_cols / 12)) ? 12 : 20; + const int64_t M = n_cols / B; + 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)); @@ -9298,6 +9380,14 @@ static std::vector> make_test_cases_eval() { test_cases.emplace_back(new test_mul_mat_hadamard(GGML_TYPE_F32, GGML_TYPE_F32, 256, 512, 256)); // many rows 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) + 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) #if 0 // > 4GB A matrix. Too slow to be enabled by default. @@ -10500,7 +10590,15 @@ static std::vector> make_test_cases_perf() { test_cases.emplace_back(new test_mul_mat_hadamard(GGML_TYPE_F32, GGML_TYPE_F32, 128, 2048, 128)); test_cases.emplace_back(new test_mul_mat_hadamard(GGML_TYPE_F32, GGML_TYPE_F32, 256, 2048, 256)); test_cases.emplace_back(new test_mul_mat_hadamard(GGML_TYPE_F32, GGML_TYPE_F32, 512, 2048, 512)); - + 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) + 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