Commit b86d2f075 for llama.cpp
commit b86d2f07542b29ab099aed34fd6b6d1b2fd4b81c
Author: bri-prism <288398250+bri-prism@users.noreply.github.com>
Date: Wed Oct 7 11:56:18 2026 -0700
cuda: FWHT kernels for block widths above 512 (#29100)
The CUDA FWHT covers widths 64 to 512. It runs one row per warp and keeps N/32
values per lane, so wider blocks need more registers per lane than that layout
allows.
fwht_cuda_block runs one row per thread block with 256 threads, so each thread
keeps N/256 values. Stages below the warp width still shuffle, those up to the
block width go through shared memory, and the rest stay in registers. Same
butterfly and sign convention as the warp kernel.
Widths 64 to 512 keep the warp kernel. 1024 through 8192 use the new one, for
both F32 and F16 sources. ggml_cuda_op_mul_mat_use_fwht (the shared
supports_op/dispatch predicate added in #29096) does not check width, so it
needed no change here: any width it admits that ggml_cuda_op_fwht can't serve
already falls through correctly to the cuBLAS path.
Rebased onto current master with #29096's F16 commit underneath it, since this
depends on the same F16 template infrastructure; that commit applied cleanly,
the only conflict was in test-backend-ops.cpp where an unrelated intervening
commit's own test additions landed near this block.
test-backend-ops on an A10 (lambdalabs): MUL_MAT 1297/1297, including all
FWHT/Hadamard cases (18 existing, 4 new F32 wide, 4 new F16 wide, 2 new
many-rows, 1 too-big boundary moved to 16384).
diff --git a/ggml/src/ggml-cuda/fwht.cu b/ggml/src/ggml-cuda/fwht.cu
index 67eea594d..d69e2ed81 100644
--- a/ggml/src/ggml-cuda/fwht.cu
+++ b/ggml/src/ggml-cuda/fwht.cu
@@ -2,6 +2,9 @@
#include "convert.cuh"
#include "fwht.cuh"
+// wide FWHT blocks use one row per thread block with this many threads
+#define GGML_CUDA_FWHT_BLOCK_NT 256
+
template <int N, typename T>
__launch_bounds__(4*ggml_cuda_get_physical_warp_size(), 1)
__global__ void fwht_cuda(const T * src, float * dst, const int64_t n_rows, const float scale) {
@@ -59,6 +62,87 @@ __global__ void fwht_cuda(const T * src, float * dst, const int64_t n_rows, cons
}
}
+// Wide blocks: one row per thread block instead of per warp, so each thread keeps N/NT
+// values rather than N/32. Stages below the warp width still shuffle, those up to the
+// block width go through shared memory, and the rest stay in registers.
+template <int N, int NT, typename T>
+__launch_bounds__(NT, 1)
+__global__ void fwht_cuda_block(const T * src, float * dst, const int64_t n_rows, const float scale) {
+ constexpr int warp_size = ggml_cuda_get_physical_warp_size();
+ constexpr int NE = N / NT;
+ static_assert(NE >= 1 && N % NT == 0 && NT % warp_size == 0, "bad FWHT block shape");
+
+ __shared__ float s[N];
+
+ const int64_t r = blockIdx.x;
+ if (r >= n_rows) {
+ return;
+ }
+
+ src += r * N;
+ dst += r * N;
+
+ const int tid = threadIdx.x;
+ const int lane = tid % warp_size;
+
+ ggml_cuda_pdl_sync();
+
+ float reg[NE];
+#pragma unroll
+ for (int i = 0; i < NE; ++i) {
+ reg[i] = ggml_cuda_cast<float>(src[i * NT + tid]) * scale;
+ }
+
+ // stages within a warp: partner differs in the lane bits
+#pragma unroll
+ for (int h = 1; h < warp_size; h *= 2) {
+#pragma unroll
+ for (int j = 0; j < NE; j++) {
+ const float val = reg[j];
+ const float val2 = __shfl_xor_sync(0xFFFFFFFF, val, h, warp_size);
+ reg[j] = (lane & h) == 0 ? val + val2 : val2 - val;
+ }
+ }
+
+ // stages across warps: partner differs in the thread-index bits above the lane
+#pragma unroll
+ for (int h = warp_size; h < NT; h *= 2) {
+#pragma unroll
+ for (int j = 0; j < NE; j++) {
+ s[j * NT + tid] = reg[j];
+ }
+ __syncthreads();
+#pragma unroll
+ for (int j = 0; j < NE; j++) {
+ const float val = reg[j];
+ const float val2 = s[j * NT + (tid ^ h)];
+ reg[j] = (tid & h) == 0 ? val + val2 : val2 - val;
+ }
+ __syncthreads();
+ }
+
+ // stages above the block width: partner is another register of the same thread
+#pragma unroll
+ for (int h = NT; h < N; h *= 2) {
+ const int step = h / NT;
+#pragma unroll
+ for (int j = 0; j < NE; j += 2 * step) {
+#pragma unroll
+ 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 < NE; ++i) {
+ dst[i * NT + tid] = reg[i];
+ }
+}
+
template <typename T>
static bool ggml_cuda_op_fwht_impl(ggml_backend_cuda_context & ctx, const ggml_tensor * src, ggml_tensor * dst) {
const int n = src->ne[0];
@@ -94,7 +178,34 @@ static bool ggml_cuda_op_fwht_impl(ggml_backend_cuda_context & ctx, const ggml_t
ggml_cuda_kernel_launch(fwht_cuda<512, T>, launch_params, src_d, dst_d, rows, scale);
return true;
default:
- return false;
+ break;
+ }
+
+ // wide blocks: one row per thread block
+ {
+ constexpr int nt = GGML_CUDA_FWHT_BLOCK_NT;
+
+ dim3 grid_dims_w(rows, 1, 1);
+ dim3 block_dims_w(nt, 1, 1);
+ const ggml_cuda_kernel_launch_params launch_params_w =
+ ggml_cuda_kernel_launch_params(grid_dims_w, block_dims_w, 0, stream);
+
+ switch (n) {
+ case 1024:
+ ggml_cuda_kernel_launch(fwht_cuda_block<1024, nt, T>, launch_params_w, src_d, dst_d, rows, scale);
+ return true;
+ case 2048:
+ ggml_cuda_kernel_launch(fwht_cuda_block<2048, nt, T>, launch_params_w, src_d, dst_d, rows, scale);
+ return true;
+ case 4096:
+ ggml_cuda_kernel_launch(fwht_cuda_block<4096, nt, T>, launch_params_w, src_d, dst_d, rows, scale);
+ return true;
+ case 8192:
+ ggml_cuda_kernel_launch(fwht_cuda_block<8192, nt, T>, launch_params_w, src_d, dst_d, rows, scale);
+ return true;
+ default:
+ return false;
+ }
}
}