metal: add F16 input to the FWHT (#29094)

* metal: add F16 input to the FWHT

The Metal FWHT kernel accepts F32 input only. This change makes the source
type a template parameter, so the kernel reads an F16 source directly instead
of requiring a converted copy. The F32 instantiations are unchanged.

The pipeline name now carries the source type, and supports_op accepts an F16
src1 for the Hadamard hint at the four sizes the kernels cover. Every other
F16 src1 path still goes through ggml_metal_supports_mul_mat_op.

These are the test cases mentioned in #27779.

test-backend-ops on M5 Pro: MUL_MAT_HADAMARD 16/16, MUL_MAT 1265/1265.

* metal: ask the same FWHT question in supports_op and the dispatch

supports_op admitted an F16 src1 on the type, the hint and the width alone, but the
dispatch also requires src1 and dst to be contiguous and the same shape. A Hadamard
hinted MUL_MAT that passed the first and failed the second reached the generic path,
which has no F32 src0 by F16 src1 kernel, and aborted on a nil pipeline:

  kernel not found in any metal library: base = 'kernel_mul_mv_f32_f16_4'
  ggml_metal_encoder_set_pipeline: nil Metal pipeline

ggml_metal_use_fwht now holds the whole condition and both callers use it, so they
cannot drift apart again. The added test case has src1 and dst of different shapes,
which aborted before this change and is declined by the Metal backend after it.

* metal: branchless butterfly select in the FWHT simdgroup kernel

Review suggestion. Replaces the ternary in the shuffle stages with
val2 - val + 2*((lane & i) == 0)*val, which is the same value without the
select.

Measured on M5 Pro, interleaved A/B, five rounds, first discarded, on a
Hadamard matmul with block 512 and 65536 rows so the kernel rather than the
launch dominates: 1324.6 us before, 1285.0 us after, a 3.0% gain, and faster
in every round. At the shapes already in the perf suite the op runs 1.6 to
3.9 us against a 1.6 us launch floor, so the difference is not visible there.

FOR_UNROLL on the same loops was also measured and made no difference, the
delta changing sign between rounds, so it is not included.

* metal: move the FWHT dispatch predicates to ggml-metal-common

Review feedback. ggml_metal_use_fwht and ggml_metal_fwht_supported_size were
static inline in ggml-metal-device.h. They now follow the
ggml_metal_op_mul_mat_use_mm pattern: declared in ggml-metal-common.h and
implemented in ggml-metal-common.cpp, which is already the home for helpers
shared between supports_op and the op dispatch. The predicate is named
ggml_metal_op_mul_mat_use_fwht to sit alongside the _use_mm pair it parallels.

This also fixes the macos-latest-arm64 build. The header needed ggml-impl.h
for ggml_get_op_params_i32, but ggml-metal-device.h is reached from
tools/tuning through ggml-metal-tuning.h, and that target does not have
ggml/src on its include path. ggml-metal-common.cpp already includes
ggml-impl.h, so the accessor is used normally there and the header goes back
to needing nothing extra.

* metal: keep the FWHT size check internal and group the dispatch helpers

Applies the patch from the review. ggml_metal_fwht_supported_size becomes
static in ggml-metal-common.cpp since nothing outside it needs the size list,
which also drops stdint.h from the header again, and
ggml_metal_op_mul_mat_use_fwht joins the existing _use_mm declarations under
their shared comment instead of carrying its own block.

* tests: drop the mismatched-shape Hadamard case

I added a case with m != k to cover an abort, but the hint is a promise that
src0 is a Hadamard matrix, so src0 is square and dst has the same shape as
src1. Every other case in the suite holds to that. The case was not a valid
op, and on CPU it compared the FWHT against a real matmul of a non-square
src0, which cannot agree.

The supports_op and dispatch conditions still come from one predicate, which
is what keeps them from disagreeing on contiguity.

ggml/src/ggml-metal/ggml-metal-common.cpp

@@ -7,6 +7,24 @@
 
 #include <vector>
 
+// must stay in sync with the kernel_fwht_<type>_<N> templates in misc.metal
+static bool ggml_metal_fwht_supported_size(int64_t n) {
+    return n == 64 || n == 128 || n == 256 || n == 512;
+}
+
+// the FWHT kernels handle a Hadamard-hinted MUL_MAT only under these conditions. supports_op
+// and the dispatch must ask the same question: an F16 src1 that is admitted but then falls
+// through reaches the generic path, which has no F32 src0 by F16 src1 kernel.
+bool ggml_metal_op_mul_mat_use_fwht(const struct ggml_tensor * op) {
+    return ggml_get_op_params_i32(op, 1) == GGML_HINT_SRC0_IS_HADAMARD &&
+           op->type == GGML_TYPE_F32 &&
+           (op->src[1]->type == GGML_TYPE_F32 || op->src[1]->type == GGML_TYPE_F16) &&
+           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]);
+}
+
 bool ggml_metal_op_mul_mat_use_mm(const struct ggml_tensor * op, bool has_simdgroup_mm) {
     const int64_t ne00 = op->src[0]->ne[0];
     const int64_t ne11 = op->src[1]->ne[1];

ggml/src/ggml-metal/ggml-metal-common.h

@@ -48,6 +48,7 @@ bool ggml_mem_ranges_check(ggml_mem_ranges_t mrs, const struct ggml_tensor * ten
 void ggml_graph_optimize(struct ggml_cgraph * gf);
 
 // mat-mat vs mat-vec dispatch; used by both supports_op and ggml_metal_op_mul_mat*
+bool ggml_metal_op_mul_mat_use_fwht (const struct ggml_tensor * op);
 bool ggml_metal_op_mul_mat_use_mm   (const struct ggml_tensor * op, bool has_simdgroup_mm);
 bool ggml_metal_op_mul_mat_id_use_mm(const struct ggml_tensor * op, bool has_simdgroup_mm);
 

ggml/src/ggml-metal/ggml-metal-device.cpp

@@ -1472,11 +1472,11 @@ 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) {
+ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_fwht(ggml_metal_library_t lib, int n, ggml_type tsrc) {
     char base[256];
     char name[256];
 
-    snprintf(base, 256, "kernel_fwht_f32_%d", n);
+    snprintf(base, 256, "kernel_fwht_%s_%d", ggml_type_name(tsrc), n);
     snprintf(name, 256, "%s", base);
 
     ggml_metal_pipeline_with_params res = ggml_metal_library_get_pipeline(lib, name);

ggml/src/ggml-metal/ggml-metal-device.h

@@ -145,7 +145,8 @@ 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_fwht              (ggml_metal_library_t lib, int n, enum ggml_type tsrc);
+
 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_radix       (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);

ggml/src/ggml-metal/ggml-metal-device.m

@@ -1836,6 +1836,12 @@ bool ggml_metal_device_supports_op(ggml_metal_device_t dev, const struct ggml_te
         case GGML_OP_SOLVE_TRI:
             return has_simdgroup_reduction && op->src[0]->type == GGML_TYPE_F32;
         case GGML_OP_MUL_MAT:
+            // the FWHT kernels read an F16 source directly; every other F16 src1 path
+            // still goes through ggml_metal_supports_mul_mat_op
+            if (op->src[0]->type == GGML_TYPE_F32 && op->src[1]->type == GGML_TYPE_F16 &&
+                ggml_metal_op_mul_mat_use_fwht(op)) {
+                return has_simdgroup_reduction;
+            }
             return ggml_metal_supports_mul_mat_op(
                     has_simdgroup_reduction, op, true,
                     ggml_metal_op_mul_mat_use_mm(op, has_simdgroup_mm));

ggml/src/ggml-metal/ggml-metal-ops.cpp

@@ -2319,12 +2319,6 @@ 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_<N> 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);
 
@@ -2340,7 +2334,7 @@ int ggml_metal_op_fwht(ggml_metal_op_t ctx, int idx) {
         /*.nrows = */ (int32_t) nrows,
     };
 
-    auto pipeline = ggml_metal_library_get_pipeline_fwht(lib, n);
+    auto pipeline = ggml_metal_library_get_pipeline_fwht(lib, n, src1->type);
 
     ggml_metal_encoder_set_pipeline(enc, pipeline);
     ggml_metal_encoder_set_bytes(enc, &args, sizeof(args), 0);
@@ -2426,17 +2420,8 @@ 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);
-        }
+    if (ggml_metal_op_mul_mat_use_fwht(op)) {
+        return ggml_metal_op_fwht(ctx, idx);
     }
     const ggml_metal_device_props * props_dev = ggml_metal_device_get_props(ctx->dev);
 

ggml/src/ggml-metal/kernels/misc.metal

@@ -374,10 +374,10 @@ template [[host_name("kernel_snake_f16")]]  kernel void kernel_snake<half>(const
 template [[host_name("kernel_snake_bf16")]] kernel void kernel_snake<bfloat>(constant ggml_metal_kargs_snake &, device const bfloat *, device const float *, device const float *, device bfloat *, uint, uint, uint);
 #endif
 
-template<int N>
-kernel void kernel_fwht_f32(
+template<int N, typename src_t>
+kernel void kernel_fwht(
         constant ggml_metal_kargs_fwht & args,
-        device const float * src,
+        device const src_t * src,
         device float * dst,
         uint3  tgpig[[threadgroup_position_in_grid]],
         ushort sgitg[[simdgroup_index_in_threadgroup]],
@@ -402,13 +402,13 @@ kernel void kernel_fwht_f32(
 
     float reg[NE];
     for (int i = 0; i < NE; i++) {
-        reg[i] = src[i*NW + lane]*scale;
+        reg[i] = float(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;
+            reg[j] = val2 - val + 2*((lane & i) == 0)*val;
         }
     }
 
@@ -429,12 +429,18 @@ kernel void kernel_fwht_f32(
     }
 }
 
-typedef decltype(kernel_fwht_f32<64>) kernel_fwht_t;
+typedef decltype(kernel_fwht<64, float>) kernel_fwht_f32_t;
+typedef decltype(kernel_fwht<64, half>)  kernel_fwht_f16_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>;
+template [[host_name("kernel_fwht_f32_64")]]  kernel kernel_fwht_f32_t kernel_fwht<64,  float>;
+template [[host_name("kernel_fwht_f32_128")]] kernel kernel_fwht_f32_t kernel_fwht<128, float>;
+template [[host_name("kernel_fwht_f32_256")]] kernel kernel_fwht_f32_t kernel_fwht<256, float>;
+template [[host_name("kernel_fwht_f32_512")]] kernel kernel_fwht_f32_t kernel_fwht<512, float>;
+
+template [[host_name("kernel_fwht_f16_64")]]  kernel kernel_fwht_f16_t kernel_fwht<64,  half>;
+template [[host_name("kernel_fwht_f16_128")]] kernel kernel_fwht_f16_t kernel_fwht<128, half>;
+template [[host_name("kernel_fwht_f16_256")]] kernel kernel_fwht_f16_t kernel_fwht<256, half>;
+template [[host_name("kernel_fwht_f16_512")]] kernel kernel_fwht_f16_t kernel_fwht<512, half>;
 
 kernel void kernel_dsv4_hc_comb_f32(
         constant ggml_metal_kargs_dsv4_hc_comb & args,

tests/test-backend-ops.cpp

@@ -9847,6 +9847,13 @@ static std::vector<std::unique_ptr<test_case>> 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_F16, 64, 1, 64));
+    test_cases.emplace_back(new test_mul_mat_hadamard(GGML_TYPE_F32, GGML_TYPE_F16, 128, 1, 128));
+    test_cases.emplace_back(new test_mul_mat_hadamard(GGML_TYPE_F32, GGML_TYPE_F16, 256, 1, 256));
+    test_cases.emplace_back(new test_mul_mat_hadamard(GGML_TYPE_F32, GGML_TYPE_F16, 512, 1, 512));
+    test_cases.emplace_back(new test_mul_mat_hadamard(GGML_TYPE_F32, GGML_TYPE_F16, 128, 32, 128));
+    test_cases.emplace_back(new test_mul_mat_hadamard(GGML_TYPE_F32, GGML_TYPE_F16, 128, 4, 128, {2, 3}));
+    test_cases.emplace_back(new test_mul_mat_hadamard(GGML_TYPE_F32, GGML_TYPE_F16, 256, 512, 256)); // many rows
 
 #if 0
     // > 4GB A matrix. Too slow to be enabled by default.