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.

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;
     }