Commit c82967099 for llama.cpp
commit c8296709920f9c1ae168bfd5fe66f9f73637bd60
Author: bri-prism <288398250+bri-prism@users.noreply.github.com>
Date: Sun Sep 27 03:08:19 2026 -0700
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.
diff --git a/ggml/src/ggml-sycl/fwht.cpp b/ggml/src/ggml-sycl/fwht.cpp
index 39f273bea..fb48d7fec 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 <int N, int NT>
+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 <int N, int NT>
+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<float, 1> 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<N, NT>(src, dst, n_rows, scale, item, get_pointer(smem));
+ });
+ });
+}
+
template <int N, int m>
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;
}