From f95de9776b5b90dd993f36d2bd66a3eee21c887f Mon Sep 17 00:00:00 2001 From: Nick Lafleur <55208706+nicklafleur@users.noreply.github.com> Date: Tue, 28 Jul 2026 03:44:06 -0400 Subject: [PATCH] ggml-metal: FWHT kernel for metal backend (#25924) * metal fwht wip * shape guard and formatting * formatting * Formatting and typos Co-authored-by: YiChen Lv <63285796+forforever73@users.noreply.github.com> * fix narrowing issue Co-authored-by: YiChen Lv <63285796+forforever73@users.noreply.github.com> * cont : minor style --------- Co-authored-by: YiChen Lv <63285796+forforever73@users.noreply.github.com> Co-authored-by: Georgi Gerganov --- ggml/src/ggml-metal/ggml-metal-device.cpp | 15 ++++++ ggml/src/ggml-metal/ggml-metal-device.h | 1 + ggml/src/ggml-metal/ggml-metal-impl.h | 4 ++ ggml/src/ggml-metal/ggml-metal-ops.cpp | 52 ++++++++++++++++++ ggml/src/ggml-metal/ggml-metal-ops.h | 1 + ggml/src/ggml-metal/ggml-metal.metal | 64 ++++++++++++++++++++++- tests/test-backend-ops.cpp | 7 +++ 7 files changed, 143 insertions(+), 1 deletion(-) diff --git a/ggml/src/ggml-metal/ggml-metal-device.cpp b/ggml/src/ggml-metal/ggml-metal-device.cpp index 270c1411a0..16e98eb519 100644 --- a/ggml/src/ggml-metal/ggml-metal-device.cpp +++ b/ggml/src/ggml-metal/ggml-metal-device.cpp @@ -1252,6 +1252,21 @@ ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_argsort_merge(gg return res; } +ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_fwht(ggml_metal_library_t lib, int n) { + char base[256]; + char name[256]; + + snprintf(base, 256, "kernel_fwht_f32_%d", n); + snprintf(name, 256, "%s", base); + + ggml_metal_pipeline_with_params res = ggml_metal_library_get_pipeline(lib, name); + if (!res.pipeline) { + res = ggml_metal_library_compile_pipeline(lib, base, name, nullptr); + } + + return res; +} + // note: reuse the argsort kernel for top_k ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_top_k(ggml_metal_library_t lib, const ggml_tensor * op) { assert(op->op == GGML_OP_TOP_K); diff --git a/ggml/src/ggml-metal/ggml-metal-device.h b/ggml/src/ggml-metal/ggml-metal-device.h index b36fa8110b..d0956df506 100644 --- a/ggml/src/ggml-metal/ggml-metal-device.h +++ b/ggml/src/ggml-metal/ggml-metal-device.h @@ -139,6 +139,7 @@ struct ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_mul_mv_id struct ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_argmax (ggml_metal_library_t lib, const struct ggml_tensor * op); struct ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_argsort (ggml_metal_library_t lib, const struct ggml_tensor * op); struct ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_argsort_merge (ggml_metal_library_t lib, const struct ggml_tensor * op); +struct ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_fwht (ggml_metal_library_t lib, int n); struct ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_top_k (ggml_metal_library_t lib, const struct ggml_tensor * op); struct ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_top_k_merge (ggml_metal_library_t lib, const struct ggml_tensor * op); struct ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_bin (ggml_metal_library_t lib, const struct ggml_tensor * op, int32_t n_fuse ); diff --git a/ggml/src/ggml-metal/ggml-metal-impl.h b/ggml/src/ggml-metal/ggml-metal-impl.h index 330278d003..9f350aad5b 100644 --- a/ggml/src/ggml-metal/ggml-metal-impl.h +++ b/ggml/src/ggml-metal/ggml-metal-impl.h @@ -1157,6 +1157,10 @@ typedef struct { int32_t len; } ggml_metal_kargs_argsort_merge; +typedef struct { + int32_t nrows; +} ggml_metal_kargs_fwht; + typedef struct { int64_t ne0; float start; diff --git a/ggml/src/ggml-metal/ggml-metal-ops.cpp b/ggml/src/ggml-metal/ggml-metal-ops.cpp index c716f118f6..76626a4518 100644 --- a/ggml/src/ggml-metal/ggml-metal-ops.cpp +++ b/ggml/src/ggml-metal/ggml-metal-ops.cpp @@ -1979,6 +1979,46 @@ int ggml_metal_op_pool_1d(ggml_metal_op_t ctx, int idx) { return 1; } +// supported FWHT sizes, must stay in sync with the +// kernel_fwht_f32_ templates in ggml-metal.metal +static bool ggml_metal_fwht_supported_size(int64_t n) { + return n == 64 || n == 128 || n == 256 || n == 512; +} + +int ggml_metal_op_fwht(ggml_metal_op_t ctx, int idx) { + ggml_tensor * op = ctx->node(idx); + + ggml_metal_library_t lib = ctx->lib; + ggml_metal_encoder_t enc = ctx->enc; + + ggml_tensor * src1 = op->src[1]; + + const int64_t n = src1->ne[0]; + const int64_t nrows = ggml_nrows(src1); + + ggml_metal_kargs_fwht args = { + /*.nrows = */ (int32_t) nrows, + }; + + auto pipeline = ggml_metal_library_get_pipeline_fwht(lib, n); + + ggml_metal_encoder_set_pipeline(enc, pipeline); + ggml_metal_encoder_set_bytes(enc, &args, sizeof(args), 0); + ggml_metal_encoder_set_buffer(enc, ggml_metal_get_buffer_id(src1), 1); + ggml_metal_encoder_set_buffer(enc, ggml_metal_get_buffer_id(op), 2); + + const int th_max = ggml_metal_pipeline_max_theads_per_threadgroup(pipeline); + const int simd_size = 32; + + int sg_per_tg = 2; + sg_per_tg = std::min(sg_per_tg, th_max/simd_size); + sg_per_tg = std::max(sg_per_tg, 1); + + const int64_t n_tg = (nrows + sg_per_tg - 1) / sg_per_tg; + ggml_metal_encoder_dispatch_threadgroups(enc, n_tg, 1, 1, 32*sg_per_tg, 1, 1); + + return 1; +} int ggml_metal_op_pool_2d(ggml_metal_op_t ctx, int idx) { ggml_tensor * op = ctx->node(idx); @@ -2046,6 +2086,18 @@ int ggml_metal_op_mul_mat(ggml_metal_op_t ctx, int idx) { ggml_metal_library_t lib = ctx->lib; ggml_metal_encoder_t enc = ctx->enc; + const int32_t hint = ggml_get_op_params_i32(op, 1); + + if (hint == GGML_HINT_SRC0_IS_HADAMARD) { + if (op->src[1]->type == GGML_TYPE_F32 && + op->type == GGML_TYPE_F32 && + ggml_is_contiguous(op->src[1]) && + ggml_is_contiguous(op) && + ggml_are_same_shape(op->src[1], op) && + ggml_metal_fwht_supported_size(op->src[1]->ne[0])) { + return ggml_metal_op_fwht(ctx, idx); + } + } const ggml_metal_device_props * props_dev = ggml_metal_device_get_props(ctx->dev); GGML_TENSOR_LOCALS( int32_t, ne0, op->src[0], ne); diff --git a/ggml/src/ggml-metal/ggml-metal-ops.h b/ggml/src/ggml-metal/ggml-metal-ops.h index 89a6ad82f1..2783ecb8b6 100644 --- a/ggml/src/ggml-metal/ggml-metal-ops.h +++ b/ggml/src/ggml-metal/ggml-metal-ops.h @@ -64,6 +64,7 @@ int ggml_metal_op_set (ggml_metal_op_t ctx, int idx); int ggml_metal_op_cpy (ggml_metal_op_t ctx, int idx); int ggml_metal_op_pool_1d (ggml_metal_op_t ctx, int idx); int ggml_metal_op_pool_2d (ggml_metal_op_t ctx, int idx); +int ggml_metal_op_fwht (ggml_metal_op_t ctx, int idx); int ggml_metal_op_mul_mat (ggml_metal_op_t ctx, int idx); int ggml_metal_op_mul_mat_id (ggml_metal_op_t ctx, int idx); int ggml_metal_op_add_id (ggml_metal_op_t ctx, int idx); diff --git a/ggml/src/ggml-metal/ggml-metal.metal b/ggml/src/ggml-metal/ggml-metal.metal index 969fddfa5b..f14ee0792b 100644 --- a/ggml/src/ggml-metal/ggml-metal.metal +++ b/ggml/src/ggml-metal/ggml-metal.metal @@ -5762,7 +5762,7 @@ kernel void kernel_upscale_bicubic_f32( const float w_y2 = bicubic_weight1(1.0f - fd1); const float w_y3 = bicubic_weight2(2.0f - fd1); - const device const char * src_slice = src0 + i03 * args.nb03 + i02 * args.nb02; + const device char * src_slice = src0 + i03 * args.nb03 + i02 * args.nb02; device float * dst_ptr = (device float *)(dst + i3 * args.nb3 + i2 * args.nb2 + i1 * args.nb1); @@ -6172,6 +6172,68 @@ kernel void kernel_argsort_merge_f32_i32( template [[host_name("kernel_argsort_merge_f32_i32_asc")]] kernel argsort_merge_t kernel_argsort_merge_f32_i32; template [[host_name("kernel_argsort_merge_f32_i32_desc")]] kernel argsort_merge_t kernel_argsort_merge_f32_i32; +template +kernel void kernel_fwht_f32( + constant ggml_metal_kargs_fwht & args, + device const float * src, + device float * dst, + uint3 tgpig[[threadgroup_position_in_grid]], + ushort sgitg[[simdgroup_index_in_threadgroup]], + ushort tiisg[[thread_index_in_simdgroup]], + ushort3 ntg[[threads_per_threadgroup]]) { + + constexpr int NW = N_SIMDWIDTH; + constexpr int NE = N / NW; + + const float scale = 1.0f / sqrt((float) N); + + const int sg_per_tg = ntg.x / NW; + const int64_t r = tgpig.x * sg_per_tg + sgitg; + if (r >= args.nrows) { + return; + } + + src += r * N; + dst += r * N; + + const int lane = tiisg; + + float reg[NE]; + for (int i = 0; i < NE; i++) { + reg[i] = src[i*NW + lane]*scale; + } + for (int i = 1; i < NW; i *= 2) { + for (int j = 0; j < NE; j++) { + const float val = reg[j]; + const float val2 = simd_shuffle_xor(val, i); + reg[j] = (lane & i) == 0 ? val2 + val : val2 - val; + } + } + + for (int i = NW; i < N; i *= 2) { + const int step = i / NW; + for (int j = 0; j < NE; 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; + } + } + } + + for (int i = 0; i < NE; i++) { + dst[i*NW + lane] = reg[i]; + } +} + +typedef decltype(kernel_fwht_f32<64>) kernel_fwht_t; + +template [[host_name("kernel_fwht_f32_64")]] kernel kernel_fwht_t kernel_fwht_f32<64>; +template [[host_name("kernel_fwht_f32_128")]] kernel kernel_fwht_t kernel_fwht_f32<128>; +template [[host_name("kernel_fwht_f32_256")]] kernel kernel_fwht_t kernel_fwht_f32<256>; +template [[host_name("kernel_fwht_f32_512")]] kernel kernel_fwht_t kernel_fwht_f32<512>; + constant bool FC_flash_attn_ext_pad_has_mask [[function_constant(FC_FLASH_ATTN_EXT_PAD + 0)]]; constant int32_t FC_flash_attn_ext_pad_ncpsg [[function_constant(FC_FLASH_ATTN_EXT_PAD + 25)]]; diff --git a/tests/test-backend-ops.cpp b/tests/test-backend-ops.cpp index e7cd6d0cb6..b4061e35a0 100644 --- a/tests/test-backend-ops.cpp +++ b/tests/test-backend-ops.cpp @@ -8793,6 +8793,9 @@ static std::vector> make_test_cases_eval() { test_cases.emplace_back(new test_mul_mat_hadamard(GGML_TYPE_F32, GGML_TYPE_F32, 512, 1, 512)); test_cases.emplace_back(new test_mul_mat_hadamard(GGML_TYPE_F32, GGML_TYPE_F32, 128, 32, 128)); test_cases.emplace_back(new test_mul_mat_hadamard(GGML_TYPE_F32, GGML_TYPE_F32, 128, 4, 128, {2, 3})); + 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) #if 0 // > 4GB A matrix. Too slow to be enabled by default. @@ -9803,6 +9806,10 @@ static std::vector> make_test_cases_perf() { test_cases.emplace_back(new test_mul_mat_hadamard(GGML_TYPE_F32, GGML_TYPE_F32, 64, 1, 64)); test_cases.emplace_back(new test_mul_mat_hadamard(GGML_TYPE_F32, GGML_TYPE_F32, 256, 1, 256)); test_cases.emplace_back(new test_mul_mat_hadamard(GGML_TYPE_F32, GGML_TYPE_F32, 128, 32, 128)); + test_cases.emplace_back(new test_mul_mat_hadamard(GGML_TYPE_F32, GGML_TYPE_F32, 64, 2048, 64)); + 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_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 }));