From 6d25d20d9207959a684c79015a1ab3f0ff876ecb Mon Sep 17 00:00:00 2001
From: bri-prism <288398250+bri-prism@users.noreply.github.com>
Date: Wed, 19 Aug 2026 14:32:14 -0700
Subject: [PATCH 01/19] ggml: support F16 input for FWHT matmul hint

(cherry picked from commit 9146baa453d955f66347b25ed70c1ffb1f6ed494)
---
 ggml/src/ggml-cpu/ggml-cpu.c                  | 18 +++++--
 ggml/src/ggml-cpu/ggml-cpu.cpp                |  4 ++
 ggml/src/ggml-cpu/ops.cpp                     | 26 +++++++---
 ggml/src/ggml-cuda/fwht.cu                    | 51 ++++++++++++++++---
 ggml/src/ggml-metal/ggml-metal-device.cpp     |  4 +-
 ggml/src/ggml-metal/ggml-metal-device.h       |  2 +-
 ggml/src/ggml-metal/ggml-metal-ops.cpp        |  7 ++-
 ggml/src/ggml-metal/kernels/misc.metal        | 29 +++++++----
 ggml/src/ggml-vulkan/ggml-vulkan.cpp          | 12 +++--
 ggml/src/ggml-vulkan/vulkan-shaders/fwht.comp | 13 ++++-
 .../vulkan-shaders/vulkan-shaders-gen.cpp     |  2 +
 tests/test-backend-ops.cpp                    |  2 +
 12 files changed, 131 insertions(+), 39 deletions(-)

diff --git a/ggml/src/ggml-cpu/ggml-cpu.c b/ggml/src/ggml-cpu/ggml-cpu.c
index 06c978bde2b..79c49c5e4d3 100644
--- a/ggml/src/ggml-cpu/ggml-cpu.c
+++ b/ggml/src/ggml-cpu/ggml-cpu.c
@@ -1334,7 +1334,7 @@ UseGgmlGemm1:;
         const size_t nbw3 = nbw2*ne12;
 
         assert(params->wsize >= ne13*nbw3);
-        GGML_ASSERT(src1->type == GGML_TYPE_F32);
+        GGML_ASSERT(src1->type == GGML_TYPE_F32 || src1->type == GGML_TYPE_F16);
 
     #if 0
         for (int64_t i13 = 0; i13 < ne13; ++i13) {
@@ -1353,9 +1353,19 @@ UseGgmlGemm1:;
                     size_t bs = ggml_blck_size(vec_dot_type);
                     int64_t ne10_block_start = (ith * ne10/bs) / nth;
                     int64_t ne10_block_end   = ((ith + 1) * ne10/bs) / nth;
-                    from_float((float *)((char *) src1->data + i13*nb13 + i12*nb12 + i11*nb11 + ne10_block_start*bs*nb10),
-                               (void *)               (wdata + i13*nbw3 + i12*nbw2 + i11*nbw1 + ne10_block_start*nbw0),
-                               (ne10_block_end - ne10_block_start) * bs);
+                    const char * src1_block = (const char *) src1->data + i13*nb13 + i12*nb12 + i11*nb11 + ne10_block_start*bs*nb10;
+                    char * dst_block = wdata + i13*nbw3 + i12*nbw2 + i11*nbw1 + ne10_block_start*nbw0;
+                    const int64_t n_block = (ne10_block_end - ne10_block_start) * bs;
+
+                    if (src1->type == GGML_TYPE_F32) {
+                        from_float((const float *) src1_block, dst_block, n_block);
+                    } else {
+                        const ggml_fp16_t * src_f16 = (const ggml_fp16_t *) src1_block;
+                        float * dst_f32 = (float *) dst_block;
+                        for (int64_t i = 0; i < n_block; ++i) {
+                            dst_f32[i] = GGML_CPU_FP16_TO_FP32(src_f16[i]);
+                        }
+                    }
                 }
             }
         }
diff --git a/ggml/src/ggml-cpu/ggml-cpu.cpp b/ggml/src/ggml-cpu/ggml-cpu.cpp
index 8cece71f186..1df0f2bb926 100644
--- a/ggml/src/ggml-cpu/ggml-cpu.cpp
+++ b/ggml/src/ggml-cpu/ggml-cpu.cpp
@@ -451,6 +451,10 @@ static bool ggml_backend_cpu_device_supports_op(ggml_backend_dev_t dev, const st
                 op->type != GGML_TYPE_IQ1_S   &&
                 op->type != GGML_TYPE_IQ1_M; // missing type_traits.from_float
         case GGML_OP_MUL_MAT:
+            if (ggml_get_op_params_i32(op, 1) == GGML_HINT_SRC0_IS_HADAMARD &&
+                src0->type == GGML_TYPE_F32 && op->type == GGML_TYPE_F32) {
+                return src1->type == GGML_TYPE_F32 || src1->type == GGML_TYPE_F16;
+            }
             return src1->type == GGML_TYPE_F32 || src1->type == ggml_get_type_traits_cpu(src0->type)->vec_dot_type;
         case GGML_OP_SOFT_MAX_BACK: {
             if (op->src[0]->type != GGML_TYPE_F32 || op->src[1]->type != GGML_TYPE_F32) {
diff --git a/ggml/src/ggml-cpu/ops.cpp b/ggml/src/ggml-cpu/ops.cpp
index 430dc6999d4..f5bcbfb6389 100644
--- a/ggml/src/ggml-cpu/ops.cpp
+++ b/ggml/src/ggml-cpu/ops.cpp
@@ -11859,11 +11859,20 @@ void ggml_compute_forward_opt_step_sgd(const ggml_compute_params * params, ggml_
     }
 }
 
-static void ggml_compute_forward_fwht_f32(const ggml_compute_params * params, ggml_tensor * dst) {
+static inline float ggml_fwht_load(const float value) {
+    return value;
+}
+
+static inline float ggml_fwht_load(const ggml_fp16_t value) {
+    return ggml_fp16_to_fp32(value);
+}
+
+template<typename src_t>
+static void ggml_compute_forward_fwht_impl(const ggml_compute_params * params, ggml_tensor * dst) {
     const ggml_tensor * src0 = dst->src[0];
     const ggml_tensor * src1 = dst->src[1];
 
-    GGML_ASSERT(src1->type == GGML_TYPE_F32);
+    GGML_ASSERT(src1->type == (std::is_same_v<src_t, float> ? GGML_TYPE_F32 : GGML_TYPE_F16));
     GGML_ASSERT(dst->type == GGML_TYPE_F32);
 
     GGML_TENSOR_BINARY_OP_LOCALS
@@ -11890,11 +11899,11 @@ static void ggml_compute_forward_fwht_f32(const ggml_compute_params * params, gg
         const int64_t i12 = (r - i13 * ne11 * ne12) / ne11;
         const int64_t i11 = r - i13 * ne11 * ne12 - i12 * ne11;
 
-        const float * src_row = (const float *) ((const char *) src1->data + i11 * nb11 + i12 * nb12 + i13 * nb13);
+        const src_t * src_row = (const src_t *) ((const char *) src1->data + i11 * nb11 + i12 * nb12 + i13 * nb13);
         float * dst_row = (float *) ((char *) dst->data + i11 * nb1 + i12 * nb2 + i13 * nb3);
 
         for (int64_t j = 0; j < n; j++) {
-            dst_row[j] = src_row[j] * scale;
+            dst_row[j] = ggml_fwht_load(src_row[j]) * scale;
         }
 
         // Scalar passes
@@ -11941,12 +11950,17 @@ void ggml_compute_forward_fwht(const ggml_compute_params * params, ggml_tensor *
     switch (src1->type) {
         case GGML_TYPE_F32:
             {
-                ggml_compute_forward_fwht_f32(params, dst);
+                ggml_compute_forward_fwht_impl<float>(params, dst);
+            }
+            break;
+        case GGML_TYPE_F16:
+            {
+                ggml_compute_forward_fwht_impl<ggml_fp16_t>(params, dst);
             }
             break;
         default:
             {
-                GGML_ABORT("fatal error - fwht is F32 only");
+                GGML_ABORT("fatal error - fwht supports F32 and F16 input");
             }
     }
 }
diff --git a/ggml/src/ggml-cuda/fwht.cu b/ggml/src/ggml-cuda/fwht.cu
index 184dc254c72..eb0bc27f70a 100644
--- a/ggml/src/ggml-cuda/fwht.cu
+++ b/ggml/src/ggml-cuda/fwht.cu
@@ -1,9 +1,19 @@
 #include "common.cuh"
 #include "fwht.cuh"
 
-template <int N>
+template <typename T>
+__device__ __forceinline__ float fwht_load(const T value) {
+    return value;
+}
+
+template <>
+__device__ __forceinline__ float fwht_load<half>(const half value) {
+    return __half2float(value);
+}
+
+template <int N, typename T>
 __launch_bounds__(4*ggml_cuda_get_physical_warp_size(), 1)
-__global__ void fwht_cuda(const float * src, float * dst, const int64_t n_rows, const float scale) {
+__global__ void fwht_cuda(const T * src, float * dst, const int64_t n_rows, const float scale) {
     constexpr int warp_size = ggml_cuda_get_physical_warp_size();
 
     const int64_t r = (int64_t) blockIdx.x * blockDim.y + threadIdx.y;
@@ -22,7 +32,7 @@ __global__ void fwht_cuda(const float * src, float * dst, const int64_t n_rows,
     ggml_cuda_pdl_sync();
 #pragma unroll
     for (int i = 0; i < el_w; ++i) {
-        reg[i] = src[i * warp_size + lane] * scale;
+        reg[i] = fwht_load(src[i * warp_size + lane]) * scale;
     }
 
 #pragma unroll
@@ -66,7 +76,11 @@ bool ggml_cuda_op_fwht(ggml_backend_cuda_context & ctx, const ggml_tensor * src,
     const int     n    = src->ne[0];
     const int64_t rows = ggml_nrows(src);
 
-    const float * src_d = (const float *) src->data;
+    if ((src->type != GGML_TYPE_F32 && src->type != GGML_TYPE_F16) || dst->type != GGML_TYPE_F32) {
+        return false;
+    }
+
+    const void * src_d = src->data;
     float *       dst_d = (float *) dst->data;
 
     const int warp_size = ggml_cuda_info().devices[ggml_cuda_get_device()].warp_size;
@@ -82,18 +96,39 @@ bool ggml_cuda_op_fwht(ggml_backend_cuda_context & ctx, const ggml_tensor * src,
 
     const float scale = 1 / sqrtf(n);
 
+    if (src->type == GGML_TYPE_F32) {
+        const float * src_f32 = (const float *) src_d;
+        switch (n) {
+            case 64:
+                ggml_cuda_kernel_launch(fwht_cuda<64, float>, launch_params, src_f32, dst_d, rows, scale);
+                return true;
+            case 128:
+                ggml_cuda_kernel_launch(fwht_cuda<128, float>, launch_params, src_f32, dst_d, rows, scale);
+                return true;
+            case 256:
+                ggml_cuda_kernel_launch(fwht_cuda<256, float>, launch_params, src_f32, dst_d, rows, scale);
+                return true;
+            case 512:
+                ggml_cuda_kernel_launch(fwht_cuda<512, float>, launch_params, src_f32, dst_d, rows, scale);
+                return true;
+            default:
+                return false;
+        }
+    }
+
+    const half * src_f16 = (const half *) src_d;
     switch (n) {
         case 64:
-            ggml_cuda_kernel_launch(fwht_cuda<64>, launch_params, src_d, dst_d, rows, scale);
+            ggml_cuda_kernel_launch(fwht_cuda<64, half>, launch_params, src_f16, dst_d, rows, scale);
             return true;
         case 128:
-            ggml_cuda_kernel_launch(fwht_cuda<128>, launch_params, src_d, dst_d, rows, scale);
+            ggml_cuda_kernel_launch(fwht_cuda<128, half>, launch_params, src_f16, dst_d, rows, scale);
             return true;
         case 256:
-            ggml_cuda_kernel_launch(fwht_cuda<256>, launch_params, src_d, dst_d, rows, scale);
+            ggml_cuda_kernel_launch(fwht_cuda<256, half>, launch_params, src_f16, dst_d, rows, scale);
             return true;
         case 512:
-            ggml_cuda_kernel_launch(fwht_cuda<512>, launch_params, src_d, dst_d, rows, scale);
+            ggml_cuda_kernel_launch(fwht_cuda<512, half>, launch_params, src_f16, dst_d, rows, scale);
             return true;
         default:
             return false;
diff --git a/ggml/src/ggml-metal/ggml-metal-device.cpp b/ggml/src/ggml-metal/ggml-metal-device.cpp
index b72de240f3a..395b8bc9b86 100644
--- a/ggml/src/ggml-metal/ggml-metal-device.cpp
+++ b/ggml/src/ggml-metal/ggml-metal-device.cpp
@@ -1313,11 +1313,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, bool src_f16) {
     char base[256];
     char name[256];
 
-    snprintf(base, 256, "kernel_fwht_f32_%d", n);
+    snprintf(base, 256, "kernel_fwht_%s_%d", src_f16 ? "f16" : "f32", n);
     snprintf(name, 256, "%s", base);
 
     ggml_metal_pipeline_with_params res = ggml_metal_library_get_pipeline(lib, name);
diff --git a/ggml/src/ggml-metal/ggml-metal-device.h b/ggml/src/ggml-metal/ggml-metal-device.h
index 73536de692f..b8e4599e745 100644
--- a/ggml/src/ggml-metal/ggml-metal-device.h
+++ b/ggml/src/ggml-metal/ggml-metal-device.h
@@ -142,7 +142,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_fwht              (ggml_metal_library_t lib, int n, bool src_f16);
 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-ops.cpp b/ggml/src/ggml-metal/ggml-metal-ops.cpp
index f59932c8d98..0df534b3978 100644
--- a/ggml/src/ggml-metal/ggml-metal-ops.cpp
+++ b/ggml/src/ggml-metal/ggml-metal-ops.cpp
@@ -2379,7 +2379,10 @@ 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);
+    GGML_ASSERT(src1->type == GGML_TYPE_F32 || src1->type == GGML_TYPE_F16);
+    GGML_ASSERT(op->type == GGML_TYPE_F32);
+
+    auto pipeline = ggml_metal_library_get_pipeline_fwht(lib, n, src1->type == GGML_TYPE_F16);
 
     ggml_metal_encoder_set_pipeline(enc, pipeline);
     ggml_metal_encoder_set_bytes(enc, &args, sizeof(args), 0);
@@ -2468,7 +2471,7 @@ int ggml_metal_op_mul_mat(ggml_metal_op_t ctx, int idx) {
     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 &&
+        if ((op->src[1]->type == GGML_TYPE_F32 || op->src[1]->type == GGML_TYPE_F16) &&
             op->type == GGML_TYPE_F32 &&
             ggml_is_contiguous(op->src[1]) &&
             ggml_is_contiguous(op) &&
diff --git a/ggml/src/ggml-metal/kernels/misc.metal b/ggml/src/ggml-metal/kernels/misc.metal
index 11104b4d8d1..761610b6a34 100644
--- a/ggml/src/ggml-metal/kernels/misc.metal
+++ b/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,7 +402,7 @@ 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++) {
@@ -429,12 +429,21 @@ kernel void kernel_fwht_f32(
     }
 }
 
-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>;
+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_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_f32_1024")]] kernel kernel_fwht_f32_t kernel_fwht<1024, float>;
+template [[host_name("kernel_fwht_f32_2048")]] kernel kernel_fwht_f32_t kernel_fwht<2048, 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>;
+template [[host_name("kernel_fwht_f16_1024")]] kernel kernel_fwht_f16_t kernel_fwht<1024, half>;
+template [[host_name("kernel_fwht_f16_2048")]] kernel kernel_fwht_f16_t kernel_fwht<2048, half>;
 
 kernel void kernel_dsv4_hc_comb_f32(
         constant ggml_metal_kargs_dsv4_hc_comb & args,
diff --git a/ggml/src/ggml-vulkan/ggml-vulkan.cpp b/ggml/src/ggml-vulkan/ggml-vulkan.cpp
index 7e65871ccf9..3a7b332983d 100644
--- a/ggml/src/ggml-vulkan/ggml-vulkan.cpp
+++ b/ggml/src/ggml-vulkan/ggml-vulkan.cpp
@@ -1043,6 +1043,7 @@ struct vk_device_struct {
     vk_pipeline pipeline_topk_f32[num_topk_pipelines];
     vk_pipeline pipeline_sum_rows_f32;
     vk_pipeline pipeline_fwht_f32[4];
+    vk_pipeline pipeline_fwht_f16[4];
     vk_pipeline pipeline_cumsum_f32;
     vk_pipeline pipeline_cumsum_small_f32;
     vk_pipeline pipeline_cumsum_multipass1_f32;
@@ -5765,7 +5766,8 @@ static void ggml_vk_load_shaders(vk_device& device, vk_pipeline requested) {
         int idx = 0;
         for (uint32_t n : {64, 128, 256, 512}) {
             if (device->subgroup_size <= n) {
-                ggml_vk_create_pipeline(device, device->pipeline_fwht_f32[idx], "fwht_f32", fwht_f32_len, fwht_f32_data, "main", 2, sizeof(vk_op_fwht_push_constants), {1, 1, 1}, { device->subgroup_size, n }, 1, true, true, device->subgroup_size);
+            ggml_vk_create_pipeline(device, device->pipeline_fwht_f32[idx], "fwht_f32", fwht_f32_len, fwht_f32_data, "main", 2, sizeof(vk_op_fwht_push_constants), {1, 1, 1}, { device->subgroup_size, n }, 1, true, true, device->subgroup_size);
+            ggml_vk_create_pipeline(device, device->pipeline_fwht_f16[idx], "fwht_f16", fwht_f16_len, fwht_f16_data, "main", 2, sizeof(vk_op_fwht_push_constants), {1, 1, 1}, { device->subgroup_size, n }, 1, true, true, device->subgroup_size);
             }
             ++idx;
         }
@@ -5774,6 +5776,7 @@ static void ggml_vk_load_shaders(vk_device& device, vk_pipeline requested) {
         for (uint32_t n : {64, 128, 256, 512}) {
             const uint32_t block_size = std::min(device->subgroup_size, n);
             ggml_vk_create_pipeline(device, device->pipeline_fwht_f32[idx], "fwht_shmem_f32", fwht_shmem_f32_len, fwht_shmem_f32_data, "main", 2, sizeof(vk_op_fwht_push_constants), {1, 1, 1}, { block_size, n }, 1);
+            ggml_vk_create_pipeline(device, device->pipeline_fwht_f16[idx], "fwht_shmem_f16", fwht_shmem_f16_len, fwht_shmem_f16_data, "main", 2, sizeof(vk_op_fwht_push_constants), {1, 1, 1}, { block_size, n }, 1);
             ++idx;
         }
     }
@@ -9961,11 +9964,12 @@ static bool ggml_vk_can_use_fwht(const ggml_backend_vk_context * ctx, const ggml
     }
 
     const int idx = ggml_vk_fwht_pipeline_idx(src1->ne[0]);
-    if (idx < 0 || ctx->device->pipeline_fwht_f32[idx] == nullptr) {
+    const bool src_f16 = src1->type == GGML_TYPE_F16;
+    if (idx < 0 || (src_f16 ? ctx->device->pipeline_fwht_f16[idx] : ctx->device->pipeline_fwht_f32[idx]) == nullptr) {
         return false;
     }
 
-    if (src1->type != GGML_TYPE_F32 || dst->type != GGML_TYPE_F32) {
+    if ((src1->type != GGML_TYPE_F32 && src1->type != GGML_TYPE_F16) || dst->type != GGML_TYPE_F32) {
         return false;
     }
 
@@ -9979,7 +9983,7 @@ static bool ggml_vk_can_use_fwht(const ggml_backend_vk_context * ctx, const ggml
 
 static void ggml_vk_fwht(ggml_backend_vk_context * ctx, vk_context& subctx, const ggml_tensor * src, ggml_tensor * dst) {
     const int idx = ggml_vk_fwht_pipeline_idx(src->ne[0]);
-    vk_pipeline pipeline = ctx->device->pipeline_fwht_f32[idx];
+    vk_pipeline pipeline = src->type == GGML_TYPE_F16 ? ctx->device->pipeline_fwht_f16[idx] : ctx->device->pipeline_fwht_f32[idx];
 
     const uint32_t rows_per_workgroup = 4;
     const uint32_t n_rows = (uint32_t)ggml_nrows(src);
diff --git a/ggml/src/ggml-vulkan/vulkan-shaders/fwht.comp b/ggml/src/ggml-vulkan/vulkan-shaders/fwht.comp
index a2069964adb..d53ef81ae9b 100644
--- a/ggml/src/ggml-vulkan/vulkan-shaders/fwht.comp
+++ b/ggml/src/ggml-vulkan/vulkan-shaders/fwht.comp
@@ -1,6 +1,9 @@
 #version 450
 
 #extension GL_EXT_control_flow_attributes : require
+#ifdef FWHT_F16
+#extension GL_EXT_shader_explicit_arithmetic_types_float16 : require
+#endif
 #ifndef FWHT_SHMEM
 #extension GL_KHR_shader_subgroup_basic : enable
 #extension GL_KHR_shader_subgroup_shuffle : enable
@@ -19,7 +22,13 @@ layout(push_constant) uniform parameter
     float scale;
 };
 
-layout(binding = 0, std430) readonly buffer A { float data_a[]; };
+layout(binding = 0, std430) readonly buffer A {
+#ifdef FWHT_F16
+    float16_t data_a[];
+#else
+    float data_a[];
+#endif
+};
 layout(binding = 1, std430) writeonly buffer D { float data_d[]; };
 
 const uint EL_W = N / BLOCK_SIZE;
@@ -54,7 +63,7 @@ void main() {
 
         [[unroll]]
         for (uint i = 0; i < EL_W; ++i) {
-            reg[i] = row < n_rows ? data_a[src_offset + row_offset + i * BLOCK_SIZE + tid] * scale : 0.0;
+            reg[i] = row < n_rows ? float(data_a[src_offset + row_offset + i * BLOCK_SIZE + tid]) * scale : 0.0;
         }
 
 #ifdef FWHT_SHMEM
diff --git a/ggml/src/ggml-vulkan/vulkan-shaders/vulkan-shaders-gen.cpp b/ggml/src/ggml-vulkan/vulkan-shaders/vulkan-shaders-gen.cpp
index 17d57d5a18f..770496b112c 100644
--- a/ggml/src/ggml-vulkan/vulkan-shaders/vulkan-shaders-gen.cpp
+++ b/ggml/src/ggml-vulkan/vulkan-shaders/vulkan-shaders-gen.cpp
@@ -1031,6 +1031,8 @@ void process_shaders() {
     string_to_spv("sum_rows_f32", "sum_rows.comp", merge_maps(base_dict, {{"A_TYPE", "float"}, {"D_TYPE", "float"}}));
     string_to_spv("fwht_f32", "fwht.comp", {});
     string_to_spv("fwht_shmem_f32", "fwht.comp", {{"FWHT_SHMEM", "1"}});
+    string_to_spv("fwht_f16", "fwht.comp", {{"FWHT_F16", "1"}});
+    string_to_spv("fwht_shmem_f16", "fwht.comp", {{"FWHT_F16", "1"}, {"FWHT_SHMEM", "1"}});
     string_to_spv("count_equal_i32", "count_equal.comp", merge_maps(base_dict, {{"A_TYPE", "int"}, {"B_TYPE", "int"}, {"D_TYPE", "int"}}));
     string_to_spv("cumsum_f32", "cumsum.comp", merge_maps(base_dict, {{"A_TYPE", "float"}, {"D_TYPE", "float"}}));
     string_to_spv("cumsum_multipass1_f32", "cumsum_multipass1.comp", merge_maps(base_dict, {{"A_TYPE", "float"}, {"D_TYPE", "float"}}));
diff --git a/tests/test-backend-ops.cpp b/tests/test-backend-ops.cpp
index 919c6be4435..0375b3db2e6 100644
--- a/tests/test-backend-ops.cpp
+++ b/tests/test-backend-ops.cpp
@@ -9163,6 +9163,8 @@ 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, 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_F16, 64, 1, 64));
+    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_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)
 

From 98a3706430297599e5b1788fa0a292252fa948d6 Mon Sep 17 00:00:00 2001
From: bri-prism <288398250+bri-prism@users.noreply.github.com>
Date: Wed, 19 Aug 2026 15:12:49 -0700
Subject: [PATCH 02/19] ggml-cuda: advertise F16 Hadamard inputs

(cherry picked from commit 14d8bb56adb2f7e93f2067a12d9ddce1ddfad21b)
---
 ggml/src/ggml-cuda/ggml-cuda.cu | 5 ++++-
 1 file changed, 4 insertions(+), 1 deletion(-)

diff --git a/ggml/src/ggml-cuda/ggml-cuda.cu b/ggml/src/ggml-cuda/ggml-cuda.cu
index 3393ca5c1fe..ad3acf8ab6f 100644
--- a/ggml/src/ggml-cuda/ggml-cuda.cu
+++ b/ggml/src/ggml-cuda/ggml-cuda.cu
@@ -4938,7 +4938,10 @@ static bool ggml_backend_cuda_device_supports_op(ggml_backend_dev_t dev, const g
                 if (a->nb[0] != ggml_element_size(a) || b->nb[0] != ggml_element_size(b)) {
                     return false; // TODO this could in principle be implemented though currently there is no use case.
                 }
-                if (b->type == GGML_TYPE_F16 && a->type != GGML_TYPE_F16) {
+                const bool is_hadamard = op->op == GGML_OP_MUL_MAT &&
+                    ggml_get_op_params_i32(op, 1) == GGML_HINT_SRC0_IS_HADAMARD;
+                if (b->type == GGML_TYPE_F16 && a->type != GGML_TYPE_F16 &&
+                    !(is_hadamard && a->type == GGML_TYPE_F32)) {
                     return false;
                 }
 #ifdef GGML_USE_MUSA

From 2fa1eaed1b7749061ba72b12cda1b54eb432b7fb Mon Sep 17 00:00:00 2001
From: bri-prism <288398250+bri-prism@users.noreply.github.com>
Date: Thu, 20 Aug 2026 07:16:24 -0700
Subject: [PATCH 03/19] ggml-cuda: extend FWHT matmul hint to 1024/2048 blocks

Register-resident warp kernel handles N up to 2048 for F32/F16 inputs.
test-backend-ops MUL_MAT_HADAMARD passes 14/14 on H200 (SM90) vs CPU
reference, including H2048 F32/F16.

(cherry picked from commit db4c6ee86a71900dda624d9fcad04e1727126ca7)
---
 ggml/src/ggml-cuda/fwht.cu | 12 ++++++++++++
 tests/test-backend-ops.cpp |  5 ++++-
 2 files changed, 16 insertions(+), 1 deletion(-)

diff --git a/ggml/src/ggml-cuda/fwht.cu b/ggml/src/ggml-cuda/fwht.cu
index eb0bc27f70a..77421a71313 100644
--- a/ggml/src/ggml-cuda/fwht.cu
+++ b/ggml/src/ggml-cuda/fwht.cu
@@ -111,6 +111,12 @@ bool ggml_cuda_op_fwht(ggml_backend_cuda_context & ctx, const ggml_tensor * src,
             case 512:
                 ggml_cuda_kernel_launch(fwht_cuda<512, float>, launch_params, src_f32, dst_d, rows, scale);
                 return true;
+            case 1024:
+                ggml_cuda_kernel_launch(fwht_cuda<1024, float>, launch_params, src_f32, dst_d, rows, scale);
+                return true;
+            case 2048:
+                ggml_cuda_kernel_launch(fwht_cuda<2048, float>, launch_params, src_f32, dst_d, rows, scale);
+                return true;
             default:
                 return false;
         }
@@ -130,6 +136,12 @@ bool ggml_cuda_op_fwht(ggml_backend_cuda_context & ctx, const ggml_tensor * src,
         case 512:
             ggml_cuda_kernel_launch(fwht_cuda<512, half>, launch_params, src_f16, dst_d, rows, scale);
             return true;
+        case 1024:
+            ggml_cuda_kernel_launch(fwht_cuda<1024, half>, launch_params, src_f16, dst_d, rows, scale);
+            return true;
+        case 2048:
+            ggml_cuda_kernel_launch(fwht_cuda<2048, half>, launch_params, src_f16, dst_d, rows, scale);
+            return true;
         default:
             return false;
     }
diff --git a/tests/test-backend-ops.cpp b/tests/test-backend-ops.cpp
index 0375b3db2e6..265417eeefe 100644
--- a/tests/test-backend-ops.cpp
+++ b/tests/test-backend-ops.cpp
@@ -9163,10 +9163,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, 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, 1024, 1, 1024));
+    test_cases.emplace_back(new test_mul_mat_hadamard(GGML_TYPE_F32, GGML_TYPE_F32, 2048, 1, 2048));
+    test_cases.emplace_back(new test_mul_mat_hadamard(GGML_TYPE_F32, GGML_TYPE_F32, 2048, 32, 2048));
     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, 32, 128));
+    test_cases.emplace_back(new test_mul_mat_hadamard(GGML_TYPE_F32, GGML_TYPE_F16, 2048, 1, 2048));
     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.

From 9692d677c64db908f89369a7e7ee5db2d5470216 Mon Sep 17 00:00:00 2001
From: bri-prism <288398250+bri-prism@users.noreply.github.com>
Date: Thu, 20 Aug 2026 07:16:24 -0700
Subject: [PATCH 04/19] llama: prism.hadamard GGUF contract + activation-side
 rotation

Load prism.hadamard.* v1 metadata (block size, transform, axis, sign mode,
folded weight names), materialize one persistent rotation tensor per
(block size, buffer type), and apply x' = Hx immediately before every
folded mul_mat / mul_mat_id via the FWHT hint. Identity signs only;
unsupported contract values fail loading loudly.

(cherry picked from commit 2f2ddf0f29e5c6e6b827d79db7f68920319c60f9)
---
 src/llama-context.cpp      |   1 +
 src/llama-graph.cpp        |  21 ++++++-
 src/llama-graph.h          |   7 +++
 src/llama-model-loader.cpp |   1 +
 src/llama-model.cpp        | 124 +++++++++++++++++++++++++++++++++++++
 src/llama-model.h          |   7 +++
 6 files changed, 159 insertions(+), 2 deletions(-)

diff --git a/src/llama-context.cpp b/src/llama-context.cpp
index a0bf1709f6d..76ee6a3ade5 100644
--- a/src/llama-context.cpp
+++ b/src/llama-context.cpp
@@ -2497,6 +2497,7 @@ llm_graph_params llama_context::graph_params(
         /*.loras       =*/ loras.get(),
         /*.mctx        =*/ mctx,
         /*.cross       =*/ &cross,
+        /*.hadamard_rotations =*/ &model.hadamard_rotations,
         /*.samplers    =*/ sampling.samplers,
         /*.n_outputs   =*/ n_outputs,
         /*.cb          =*/ graph_get_cb(),
diff --git a/src/llama-graph.cpp b/src/llama-graph.cpp
index e5da0f0b777..e4c95305bb0 100644
--- a/src/llama-graph.cpp
+++ b/src/llama-graph.cpp
@@ -1464,6 +1464,7 @@ llm_graph_context::llm_graph_context(const llm_graph_params & params) :
     loras            (params.loras),
     mctx             (params.mctx),
     cross            (params.cross),
+    hadamard_rotations(params.hadamard_rotations),
     samplers         (params.samplers),
     cb_func          (params.cb),
     res              (params.res),
@@ -1490,7 +1491,15 @@ ggml_tensor * llm_graph_context::build_lora_mm(
           ggml_tensor * w,
           ggml_tensor * cur,
           ggml_tensor * w_s) const {
-    ggml_tensor * res = ggml_mul_mat(ctx0, w, cur);
+    ggml_tensor * cur_mm = cur;
+    if (hadamard_rotations) {
+        const auto it = hadamard_rotations->find(w);
+        if (it != hadamard_rotations->end()) {
+            cur_mm = llama_mul_mat_hadamard(ctx0, cur, it->second);
+        }
+    }
+
+    ggml_tensor * res = ggml_mul_mat(ctx0, w, cur_mm);
 
     if (w_s) {
         res = ggml_mul(ctx0, res, w_s);
@@ -1522,7 +1531,15 @@ ggml_tensor * llm_graph_context::build_lora_mm_id(
           ggml_tensor * cur, // ggml_tensor * b
           ggml_tensor * ids,
           ggml_tensor * w_s) const {
-    ggml_tensor * res = ggml_mul_mat_id(ctx0, w, cur, ids);
+    ggml_tensor * cur_mm = cur;
+    if (hadamard_rotations) {
+        const auto it = hadamard_rotations->find(w);
+        if (it != hadamard_rotations->end()) {
+            cur_mm = llama_mul_mat_hadamard(ctx0, cur, it->second);
+        }
+    }
+
+    ggml_tensor * res = ggml_mul_mat_id(ctx0, w, cur_mm, ids);
 
     if (w_s) {
         const int64_t n_expert = w_s->ne[0];
diff --git a/src/llama-graph.h b/src/llama-graph.h
index 6c8a53fdbe9..6b0df9428cf 100644
--- a/src/llama-graph.h
+++ b/src/llama-graph.h
@@ -11,11 +11,16 @@
 #include <set>
 #include <functional>
 #include <map>
+#include <unordered_map>
 
 struct ggml_cgraph;
 struct ggml_context;
 struct ggml_tensor;
 
+// Maps a folded model weight to the normalized Hadamard matrix required to
+// restore its activation-side basis immediately before the matmul.
+using llama_hadamard_rotations = std::unordered_map<const ggml_tensor *, ggml_tensor *>;
+
 struct llama_cparams;
 struct llama_layer;
 
@@ -797,6 +802,7 @@ struct llm_graph_params {
     const llama_adapter_loras    * loras;
     const llama_memory_context_i * mctx;
     const llama_cross            * cross;
+    const llama_hadamard_rotations * hadamard_rotations;
 
     std::map<llama_seq_id, llama_sampler *> samplers;
 
@@ -1037,6 +1043,7 @@ struct llm_graph_context {
     const llama_adapter_loras    * loras;
     const llama_memory_context_i * mctx;
     const llama_cross            * cross;
+    const llama_hadamard_rotations * hadamard_rotations;
 
     std::map<llama_seq_id, llama_sampler *> samplers;
 
diff --git a/src/llama-model-loader.cpp b/src/llama-model-loader.cpp
index aedac9b3d2d..552dc478874 100644
--- a/src/llama-model-loader.cpp
+++ b/src/llama-model-loader.cpp
@@ -407,6 +407,7 @@ namespace GGUFMeta {
         return get_arr(llm_kv(kid), result, required);
     }
 
+    template bool llama_model_loader::get_arr<std::string>(const std::string & key, std::vector<std::string> & result, bool required);
     template bool llama_model_loader::get_arr<std::vector<std::string>>(enum llm_kv kid, std::vector<std::string> & result, bool required);
     template bool llama_model_loader::get_arr<std::array<int32_t, 512>>(enum llm_kv kid, std::array<int32_t, 512> & result, bool required);
     template bool llama_model_loader::get_arr<std::vector<int32_t>>(enum llm_kv kid, std::vector<int32_t> & result, bool required);
diff --git a/src/llama-model.cpp b/src/llama-model.cpp
index c34700ff563..cc97ae5c4dd 100644
--- a/src/llama-model.cpp
+++ b/src/llama-model.cpp
@@ -28,6 +28,7 @@
 #include <cassert>
 #include <cfloat>
 #include <cstdint>
+#include <cstdio>
 #include <cstring>
 #include <cmath>
 #include <functional>
@@ -1189,6 +1190,47 @@ void llama_model_base::load_hparams(llama_model_loader & ml) {
         gguf_kv.emplace(name, value);
     }
 
+    uint32_t hadamard_version = 0;
+    if (ml.get_key("prism.hadamard.version", hadamard_version, false)) {
+        if (hadamard_version != 1) {
+            throw std::runtime_error(format("unsupported prism.hadamard.version: %u", hadamard_version));
+        }
+
+        uint32_t block_size = 0;
+        std::string transform;
+        std::string axis;
+        std::string sign_mode;
+        std::vector<std::string> weight_names;
+
+        ml.get_key("prism.hadamard.block_size", block_size);
+        ml.get_key("prism.hadamard.transform", transform);
+        ml.get_key("prism.hadamard.axis", axis);
+        ml.get_key("prism.hadamard.sign_mode", sign_mode);
+        ml.get_arr("prism.hadamard.weight_names", weight_names);
+
+        if (block_size == 0 || (block_size & (block_size - 1)) != 0) {
+            throw std::runtime_error(format("invalid prism.hadamard.block_size: %u", block_size));
+        }
+        if (transform != "normalized-sylvester-walsh-hadamard") {
+            throw std::runtime_error(format("unsupported prism.hadamard.transform: %s", transform.c_str()));
+        }
+        if (axis != "input-last-dimension") {
+            throw std::runtime_error(format("unsupported prism.hadamard.axis: %s", axis.c_str()));
+        }
+        if (sign_mode != "identity") {
+            throw std::runtime_error(format("unsupported prism.hadamard.sign_mode: %s", sign_mode.c_str()));
+        }
+        if (weight_names.empty()) {
+            throw std::runtime_error("prism.hadamard.weight_names is empty");
+        }
+
+        for (const auto & weight_name : weight_names) {
+            if (!hadamard_weight_blocks.emplace(weight_name, block_size).second) {
+                throw std::runtime_error(format("duplicate prism.hadamard weight: %s", weight_name.c_str()));
+            }
+        }
+    }
+
     // get general kv
     ml.get_key(LLM_KV_GENERAL_NAME, name, false);
 
@@ -1754,6 +1796,88 @@ bool llama_model_base::load_tensors(llama_model_loader & ml) {
         ctx_buf_maps.emplace_back(ctx, buf_map);
     }
 
+    if (!ml.no_alloc && !hadamard_weight_blocks.empty()) {
+        struct hadamard_rotation {
+            uint32_t block_size;
+            ggml_backend_buffer_type_t buft;
+            ggml_tensor * tensor;
+        };
+
+        std::vector<hadamard_rotation> rotations;
+
+        for (const auto & entry : hadamard_weight_blocks) {
+            const std::string & weight_name = entry.first;
+            const uint32_t block_size = entry.second;
+            const ggml_tensor * weight = get_tensor(weight_name.c_str());
+            if (weight == nullptr) {
+                throw std::runtime_error(format("prism.hadamard weight not found: %s", weight_name.c_str()));
+            }
+            if (weight->ne[0] % block_size != 0) {
+                throw std::runtime_error(format(
+                    "prism.hadamard block size %u does not divide input dimension %lld for %s",
+                    block_size, (long long) weight->ne[0], weight_name.c_str()));
+            }
+            if (weight->buffer == nullptr) {
+                throw std::runtime_error(format("prism.hadamard weight has no buffer: %s", weight_name.c_str()));
+            }
+
+            const ggml_backend_buffer_type_t buft = ggml_backend_buffer_get_type(weight->buffer);
+            auto it = std::find_if(rotations.begin(), rotations.end(),
+                    [block_size, buft](const hadamard_rotation & rotation) {
+                        return rotation.block_size == block_size && rotation.buft == buft;
+                    });
+
+            if (it == rotations.end()) {
+                ggml_init_params params = {
+                    /*.mem_size   =*/ ggml_tensor_overhead(),
+                    /*.mem_buffer =*/ NULL,
+                    /*.no_alloc   =*/ true,
+                };
+                ggml_context_ptr ctx { ggml_init(params) };
+                if (!ctx) {
+                    throw std::runtime_error("failed to create Hadamard rotation context");
+                }
+
+                ggml_tensor * rotation = ggml_new_tensor_2d(ctx.get(), GGML_TYPE_F32, block_size, block_size);
+                char rotation_name[GGML_MAX_NAME];
+                snprintf(rotation_name, sizeof(rotation_name), "prism.hadamard.%u", block_size);
+                ggml_set_name(rotation, rotation_name);
+
+                ggml_backend_buffer_ptr buffer { ggml_backend_alloc_ctx_tensors_from_buft(ctx.get(), buft) };
+                if (!buffer) {
+                    throw std::runtime_error(format("unable to allocate %s Hadamard rotation buffer", ggml_backend_buft_name(buft)));
+                }
+                ggml_backend_buffer_set_usage(buffer.get(), GGML_BACKEND_BUFFER_USAGE_WEIGHTS);
+
+                std::vector<float> data((size_t) block_size * block_size);
+                const float scale = 1.0f / sqrtf((float) block_size);
+                for (uint32_t row = 0; row < block_size; ++row) {
+                    for (uint32_t col = 0; col < block_size; ++col) {
+                        uint32_t parity = row & col;
+                        parity ^= parity >> 16;
+                        parity ^= parity >> 8;
+                        parity ^= parity >> 4;
+                        parity ^= parity >> 2;
+                        parity ^= parity >> 1;
+                        data[(size_t) row * block_size + col] = (parity & 1) ? -scale : scale;
+                    }
+                }
+                ggml_backend_tensor_set(rotation, data.data(), 0, data.size() * sizeof(float));
+
+                std::vector<ggml_backend_buffer_ptr> buffers;
+                buffers.emplace_back(std::move(buffer));
+                pimpl->ctxs_bufs.emplace_back(std::move(ctx), std::move(buffers));
+                rotations.push_back({ block_size, buft, rotation });
+                it = std::prev(rotations.end());
+            }
+
+            hadamard_rotations.emplace(weight, it->tensor);
+        }
+
+        LLAMA_LOG_INFO("%s: loaded %zu Hadamard-folded weight(s) using %zu persistent rotation(s)\n",
+                __func__, hadamard_rotations.size(), rotations.size());
+    }
+
     if (llama_supports_gpu_offload()) {
         const int n_gpu = std::min(n_gpu_layers, n_layer_all);
 
diff --git a/src/llama-model.h b/src/llama-model.h
index 44bd9675754..31ec2425c10 100644
--- a/src/llama-model.h
+++ b/src/llama-model.h
@@ -664,6 +664,13 @@ struct llama_model {
     // gguf metadata
     std::unordered_map<std::string, std::string> gguf_kv;
 
+    // Hadamard-folded GGUF weights are matched with persistent model tensors
+    // containing the activation-side transform.  The string map is populated
+    // from GGUF metadata while loading hparams; the pointer map is populated
+    // after model buffers have been allocated.
+    std::unordered_map<std::string, uint32_t> hadamard_weight_blocks;
+    llama_hadamard_rotations hadamard_rotations;
+
     // list of devices used in this model
     std::vector<llama_device> devices;
 

From 8c27714e555789593df9ce1ee5d8e2726536fcec Mon Sep 17 00:00:00 2001
From: bri-prism <288398250+bri-prism@users.noreply.github.com>
Date: Thu, 20 Aug 2026 07:16:24 -0700
Subject: [PATCH 05/19] convert: emit prism.hadamard metadata from packing
 manifest

Pick up hadamard_packing.json from the model dir, validate the fold
contract (schema v1, identity signs, last-axis), map folded tensor names
to GGUF names, and write the prism.hadamard.* keys.

(cherry picked from commit ddc5ed9afda4abffb711d06d3c8437c9279ac538)
---
 conversion/base.py | 53 ++++++++++++++++++++++++++++++++++++++++++++++
 1 file changed, 53 insertions(+)

diff --git a/conversion/base.py b/conversion/base.py
index 56547ace009..abd31697ae6 100644
--- a/conversion/base.py
+++ b/conversion/base.py
@@ -617,6 +617,57 @@ def map_tensor_name(self, name: str, try_suffixes: Sequence[str] = (".weight", "
             raise ValueError(f"Can not map tensor {name!r}")
         return new_name
 
+    def add_hadamard_metadata(self) -> None:
+        """Transfer a packed-checkpoint transform contract into GGUF metadata."""
+        manifest_path = self.dir_model / "hadamard_packing.json"
+        if not manifest_path.is_file():
+            return
+
+        with manifest_path.open("r", encoding="utf-8") as f:
+            manifest = json.load(f)
+
+        if manifest.get("schema_version") != 1 or manifest.get("kind") != "hadamard-weight-fold":
+            raise ValueError(f"Unsupported Hadamard manifest: {manifest_path}")
+        if manifest.get("status") != "requires-matching-runtime":
+            raise ValueError(f"Unexpected Hadamard manifest status: {manifest.get('status')!r}")
+
+        transform = manifest.get("transform")
+        if not isinstance(transform, dict):
+            raise ValueError("Hadamard manifest is missing transform metadata")
+        block_size = transform.get("block_size")
+        if not isinstance(block_size, int) or block_size <= 0 or block_size & (block_size - 1):
+            raise ValueError(f"Invalid Hadamard block size: {block_size!r}")
+        if transform.get("name") != "normalized-signed-sylvester-walsh-hadamard":
+            raise ValueError(f"Unsupported Hadamard transform: {transform.get('name')!r}")
+        if transform.get("sign_mode") != "identity":
+            raise ValueError(
+                "GGUF Hadamard runtime currently supports only identity signs; "
+                "repack with --sign-mode identity"
+            )
+
+        tensor_records = manifest.get("tensors")
+        if not isinstance(tensor_records, list) or not tensor_records:
+            raise ValueError("Hadamard manifest has no folded tensors")
+
+        weight_names: list[str] = []
+        for record in tensor_records:
+            if not isinstance(record, dict) or not isinstance(record.get("name"), str):
+                raise ValueError("Hadamard manifest has an invalid tensor record")
+            if record.get("axis") != -1:
+                raise ValueError(f"Unsupported Hadamard tensor axis for {record['name']!r}")
+            filtered = self.filter_tensors((record["name"], lambda: None))
+            if filtered is None:
+                raise ValueError(f"Hadamard tensor is filtered out: {record['name']!r}")
+            weight_names.append(self.map_tensor_name(filtered[0]))
+
+        self.gguf_writer.add_uint32("prism.hadamard.version", 1)
+        self.gguf_writer.add_uint32("prism.hadamard.block_size", block_size)
+        self.gguf_writer.add_string("prism.hadamard.transform", "normalized-sylvester-walsh-hadamard")
+        self.gguf_writer.add_string("prism.hadamard.axis", "input-last-dimension")
+        self.gguf_writer.add_string("prism.hadamard.sign_mode", "identity")
+        self.gguf_writer.add_array("prism.hadamard.weight_names", weight_names)
+        logger.info("GGUF Hadamard contract: H%d for %d folded weight(s)", block_size, len(weight_names))
+
     def set_gguf_parameters(self):
         raise NotImplementedError("set_gguf_parameters() must be implemented in subclasses")
 
@@ -1061,6 +1112,8 @@ def prepare_metadata(self, vocab_only: bool):
         logger.info("Set model quantization version")
         self.gguf_writer.add_quantization_version(gguf.GGML_QUANT_VERSION)
 
+        self.add_hadamard_metadata()
+
     def write_vocab(self):
         raise NotImplementedError("write_vocab() must be implemented in subclasses")
 

From f0ad29657048e42a10514cfcc087ff23db22cd95 Mon Sep 17 00:00:00 2001
From: bri-prism <288398250+bri-prism@users.noreply.github.com>
Date: Thu, 20 Aug 2026 07:26:13 -0700
Subject: [PATCH 06/19] ggml-metal: fix crash on F16 Hadamard input at
 unsupported FWHT widths

Metal advertises every mul_mat type combo, but the Hadamard-hint fallback
for F16 inputs at widths without a dedicated FWHT kernel requests a
nonexistent f32 x f16 mul_mv pipeline and crashes on the null pipeline.
Decline the op in supports_op so it falls back to CPU instead.

test-backend-ops MUL_MAT_HADAMARD: MTL0 13/13 (F16 wide-block case now
correctly unsupported), 3/3 backends pass.

(cherry picked from commit fe88d0c716f0880081625c75a233206d05ac7e87)
---
 ggml/src/ggml-metal/ggml-metal-device.h | 6 ++++++
 ggml/src/ggml-metal/ggml-metal-device.m | 6 ++++++
 ggml/src/ggml-metal/ggml-metal-ops.cpp  | 4 ----
 3 files changed, 12 insertions(+), 4 deletions(-)

diff --git a/ggml/src/ggml-metal/ggml-metal-device.h b/ggml/src/ggml-metal/ggml-metal-device.h
index b8e4599e745..f4008a18b1b 100644
--- a/ggml/src/ggml-metal/ggml-metal-device.h
+++ b/ggml/src/ggml-metal/ggml-metal-device.h
@@ -143,6 +143,12 @@ struct ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_argmax
 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, bool src_f16);
+
+// FWHT block widths with dedicated Metal kernels; the Hadamard matmul hint
+// falls back to a plain mul_mat for other widths, which has no F16-input pipeline.
+static inline bool ggml_metal_fwht_supported_size(int64_t n) {
+    return n == 64 || n == 128 || n == 256 || n == 512;
+}
 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-device.m b/ggml/src/ggml-metal/ggml-metal-device.m
index 980a9a6928e..2ab0031016c 100644
--- a/ggml/src/ggml-metal/ggml-metal-device.m
+++ b/ggml/src/ggml-metal/ggml-metal-device.m
@@ -1708,6 +1708,12 @@ bool ggml_metal_device_supports_op(ggml_metal_device_t dev, const struct ggml_te
         case GGML_OP_SOLVE_TRI:
         case GGML_OP_MUL_MAT:
         case GGML_OP_MUL_MAT_ID:
+            if (op->op == GGML_OP_MUL_MAT &&
+                ggml_get_op_params_i32(op, 1) == GGML_HINT_SRC0_IS_HADAMARD &&
+                op->src[1]->type == GGML_TYPE_F16 &&
+                !ggml_metal_fwht_supported_size(op->src[1]->ne[0])) {
+                return false;
+            }
             return has_simdgroup_reduction && op->src[0]->type != GGML_TYPE_NVFP4;
         case GGML_OP_SET:
         case GGML_OP_CPY:
diff --git a/ggml/src/ggml-metal/ggml-metal-ops.cpp b/ggml/src/ggml-metal/ggml-metal-ops.cpp
index 0df534b3978..ee266866e1c 100644
--- a/ggml/src/ggml-metal/ggml-metal-ops.cpp
+++ b/ggml/src/ggml-metal/ggml-metal-ops.cpp
@@ -2360,10 +2360,6 @@ int ggml_metal_op_pool_1d(ggml_metal_op_t ctx, int idx) {
 
 // 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);
 

From 0c369126c9be4ce84246d185832cb00a0422d5da Mon Sep 17 00:00:00 2001
From: bri-prism <288398250+bri-prism@users.noreply.github.com>
Date: Thu, 20 Aug 2026 07:35:39 -0700
Subject: [PATCH 07/19] llama: set up Hadamard rotations after tensor data load

With mmap-backed buffer types (Metal shared buffers), weight->buffer is
not assigned until load_all_data, so the rotation setup threw
'weight has no buffer' on every folded model. CUDA allocates buffers
up front, which masked the ordering bug.

(cherry picked from commit f37f0474c904e2af1e6187c4c75fa0a3a77b9b16)
---
 src/llama-model.cpp | 85 +++++++++++++++++++++++----------------------
 1 file changed, 43 insertions(+), 42 deletions(-)

diff --git a/src/llama-model.cpp b/src/llama-model.cpp
index cc97ae5c4dd..aaa0a47650d 100644
--- a/src/llama-model.cpp
+++ b/src/llama-model.cpp
@@ -1796,7 +1796,49 @@ bool llama_model_base::load_tensors(llama_model_loader & ml) {
         ctx_buf_maps.emplace_back(ctx, buf_map);
     }
 
-    if (!ml.no_alloc && !hadamard_weight_blocks.empty()) {
+
+    if (llama_supports_gpu_offload()) {
+        const int n_gpu = std::min(n_gpu_layers, n_layer_all);
+
+        int n_repeating = n_gpu;
+        if (n_repeating > 0) {
+            LLAMA_LOG_INFO("%s: offloading output layer to GPU\n", __func__);
+            n_repeating--;
+        }
+        LLAMA_LOG_INFO("%s: offloading %d repeating layers to GPU\n", __func__, n_repeating);
+
+        const int max_backend_supported_layers = n_layer_all + 1;
+        const int max_offloadable_layers       = n_layer_all + 1;
+
+        LLAMA_LOG_INFO("%s: offloaded %d/%d layers to GPU\n", __func__, std::min(n_gpu_layers, max_offloadable_layers), max_backend_supported_layers);
+    }
+
+    // print memory requirements per buffer type
+    for (auto & [_, bufs] : pimpl->ctxs_bufs) {
+        for (auto & buf: bufs) {
+            LLAMA_LOG_INFO("%s: %12s model buffer size = %8.2f MiB\n",
+                __func__, ggml_backend_buffer_name(buf.get()), ggml_backend_buffer_get_size(buf.get()) / 1024.0 / 1024.0);
+        }
+    }
+
+    if (ml.no_alloc) {
+        return true;
+    }
+
+    // load tensor data
+    for (auto & [ctx, buf_map] : ctx_buf_maps) {
+        if (!ml.load_all_data(ctx, buf_map, use_mlock ? &pimpl->mlock_mmaps : NULL, params.progress_callback, params.progress_callback_user_data)) {
+            return false;
+        }
+    }
+
+    if (use_mmap_buffer) {
+        for (auto & mapping : ml.mappings) {
+            pimpl->mappings.emplace_back(std::move(mapping));
+        }
+    }
+
+    if (!hadamard_weight_blocks.empty()) {
         struct hadamard_rotation {
             uint32_t block_size;
             ggml_backend_buffer_type_t buft;
@@ -1878,47 +1920,6 @@ bool llama_model_base::load_tensors(llama_model_loader & ml) {
                 __func__, hadamard_rotations.size(), rotations.size());
     }
 
-    if (llama_supports_gpu_offload()) {
-        const int n_gpu = std::min(n_gpu_layers, n_layer_all);
-
-        int n_repeating = n_gpu;
-        if (n_repeating > 0) {
-            LLAMA_LOG_INFO("%s: offloading output layer to GPU\n", __func__);
-            n_repeating--;
-        }
-        LLAMA_LOG_INFO("%s: offloading %d repeating layers to GPU\n", __func__, n_repeating);
-
-        const int max_backend_supported_layers = n_layer_all + 1;
-        const int max_offloadable_layers       = n_layer_all + 1;
-
-        LLAMA_LOG_INFO("%s: offloaded %d/%d layers to GPU\n", __func__, std::min(n_gpu_layers, max_offloadable_layers), max_backend_supported_layers);
-    }
-
-    // print memory requirements per buffer type
-    for (auto & [_, bufs] : pimpl->ctxs_bufs) {
-        for (auto & buf: bufs) {
-            LLAMA_LOG_INFO("%s: %12s model buffer size = %8.2f MiB\n",
-                __func__, ggml_backend_buffer_name(buf.get()), ggml_backend_buffer_get_size(buf.get()) / 1024.0 / 1024.0);
-        }
-    }
-
-    if (ml.no_alloc) {
-        return true;
-    }
-
-    // load tensor data
-    for (auto & [ctx, buf_map] : ctx_buf_maps) {
-        if (!ml.load_all_data(ctx, buf_map, use_mlock ? &pimpl->mlock_mmaps : NULL, params.progress_callback, params.progress_callback_user_data)) {
-            return false;
-        }
-    }
-
-    if (use_mmap_buffer) {
-        for (auto & mapping : ml.mappings) {
-            pimpl->mappings.emplace_back(std::move(mapping));
-        }
-    }
-
     return true;
 }
 

From 806720915d284f48a4c4b75faa0e4eb7c8c04816 Mon Sep 17 00:00:00 2001
From: bri-prism <288398250+bri-prism@users.noreply.github.com>
Date: Thu, 20 Aug 2026 10:24:52 -0700
Subject: [PATCH 08/19] ggml-metal: add 1024/2048-wide FWHT kernels

Instantiate the existing simdgroup-register FWHT template for 1024/2048
(F32 and F16 inputs) and extend the supported-size gate, replacing the
dense rotation-matmul fallback for wide Hadamard-hint blocks.

test-backend-ops MUL_MAT_HADAMARD: MTL0 14/14. Removes most of the
Hadamard-hint decode overhead on Apple silicon versus the fallback.

(cherry picked from commit 610ddad4bf42f9d16c2521762cee6654fcedd6de)
---
 ggml/src/ggml-metal/ggml-metal-device.h | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/ggml/src/ggml-metal/ggml-metal-device.h b/ggml/src/ggml-metal/ggml-metal-device.h
index f4008a18b1b..80e9c5a453e 100644
--- a/ggml/src/ggml-metal/ggml-metal-device.h
+++ b/ggml/src/ggml-metal/ggml-metal-device.h
@@ -147,7 +147,7 @@ struct ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_fwht
 // FWHT block widths with dedicated Metal kernels; the Hadamard matmul hint
 // falls back to a plain mul_mat for other widths, which has no F16-input pipeline.
 static inline bool ggml_metal_fwht_supported_size(int64_t n) {
-    return n == 64 || n == 128 || n == 256 || n == 512;
+    return n == 64 || n == 128 || n == 256 || n == 512 || n == 1024 || n == 2048;
 }
 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);

From f33b3c1c4423376da83419365951f221eb635f4c Mon Sep 17 00:00:00 2001
From: bri-prism <288398250+bri-prism@users.noreply.github.com>
Date: Sun, 23 Aug 2026 12:08:15 -0700
Subject: [PATCH 09/19] llama, convert: restrict prism.hadamard contract to
 verified matmul paths

The activation-side transform is applied only where the graph goes through
build_lora_mm/build_lora_mm_id, but the converter accepted any mappable
tensor and the loader accepted any weight name. Architectures that multiply
projection weights with raw ggml_mul_mat (deepseek2 wq_a, kimi-linear,
plm) would load a folded GGUF and silently skip the transform.

Gate the contract on both ends: the converter refuses archs outside a
verified set (llama, qwen3, qwen3moe, qwen35, qwen35moe, qwen3next) and
tensors outside an explicit projection-kind allowlist, and the loader
enforces the same checks on prism.hadamard.weight_names so folded GGUFs
from other tools refuse to load rather than run wrong math.

(cherry picked from commit e76b3041636aee8f7c6956750b86ca6dcd9b8385)
---
 conversion/base.py  | 33 +++++++++++++++++++++++++++++-
 src/llama-model.cpp | 49 +++++++++++++++++++++++++++++++++++++++++++++
 2 files changed, 81 insertions(+), 1 deletion(-)

diff --git a/conversion/base.py b/conversion/base.py
index abd31697ae6..6f43aee6ba2 100644
--- a/conversion/base.py
+++ b/conversion/base.py
@@ -649,6 +649,31 @@ def add_hadamard_metadata(self) -> None:
         if not isinstance(tensor_records, list) or not tensor_records:
             raise ValueError("Hadamard manifest has no folded tensors")
 
+        # The runtime applies the activation transform only where the graph goes through
+        # build_lora_mm/build_lora_mm_id. Restrict the contract to architectures and tensor
+        # kinds verified to route every matmul through those helpers; anything else must
+        # fail here instead of producing a GGUF that loads but skips the transform.
+        _HADAMARD_ARCHS = {
+            gguf.MODEL_ARCH.LLAMA,
+            gguf.MODEL_ARCH.QWEN3,
+            gguf.MODEL_ARCH.QWEN3MOE,
+            gguf.MODEL_ARCH.QWEN35,
+            gguf.MODEL_ARCH.QWEN35MOE,
+            gguf.MODEL_ARCH.QWEN3NEXT,
+        }
+        if self.model_arch not in _HADAMARD_ARCHS:
+            raise ValueError(
+                f"Hadamard folding is not verified for arch {self.model_arch.name}; "
+                "the runtime would load the GGUF without applying the activation transform"
+            )
+        _HADAMARD_KINDS = re.compile(
+            r"blk\.\d+\.("
+            r"attn_q|attn_k|attn_v|attn_qkv|attn_output"
+            r"|ffn_gate|ffn_up|ffn_down"
+            r"|ffn_gate_exps|ffn_up_exps|ffn_down_exps|ffn_gate_up_exps"
+            r"|ffn_gate_shexp|ffn_up_shexp|ffn_down_shexp"
+            r")\.weight"
+        )
         weight_names: list[str] = []
         for record in tensor_records:
             if not isinstance(record, dict) or not isinstance(record.get("name"), str):
@@ -658,7 +683,13 @@ def add_hadamard_metadata(self) -> None:
             filtered = self.filter_tensors((record["name"], lambda: None))
             if filtered is None:
                 raise ValueError(f"Hadamard tensor is filtered out: {record['name']!r}")
-            weight_names.append(self.map_tensor_name(filtered[0]))
+            mapped = self.map_tensor_name(filtered[0])
+            if not _HADAMARD_KINDS.fullmatch(mapped):
+                raise ValueError(
+                    f"Hadamard tensor {record['name']!r} maps to {mapped!r}, which is not on a "
+                    "verified Hadamard-aware matmul path"
+                )
+            weight_names.append(mapped)
 
         self.gguf_writer.add_uint32("prism.hadamard.version", 1)
         self.gguf_writer.add_uint32("prism.hadamard.block_size", block_size)
diff --git a/src/llama-model.cpp b/src/llama-model.cpp
index aaa0a47650d..7560bbbd126 100644
--- a/src/llama-model.cpp
+++ b/src/llama-model.cpp
@@ -26,6 +26,7 @@
 
 #include <algorithm>
 #include <cassert>
+#include <cctype>
 #include <cfloat>
 #include <cstdint>
 #include <cstdio>
@@ -1224,7 +1225,55 @@ void llama_model_base::load_hparams(llama_model_loader & ml) {
             throw std::runtime_error("prism.hadamard.weight_names is empty");
         }
 
+        // the activation-side transform is applied only by build_lora_mm/build_lora_mm_id;
+        // refuse to load folded weights for architectures or tensor kinds that are not
+        // verified to route every matmul through those helpers, rather than run wrong math
+        switch (arch) {
+            case LLM_ARCH_LLAMA:
+            case LLM_ARCH_QWEN3:
+            case LLM_ARCH_QWEN3MOE:
+            case LLM_ARCH_QWEN35:
+            case LLM_ARCH_QWEN35MOE:
+            case LLM_ARCH_QWEN3NEXT:
+                break;
+            default:
+                throw std::runtime_error(format(
+                    "prism.hadamard: arch '%s' is not verified to apply the activation transform to all folded weights",
+                    llm_arch_name(arch)));
+        }
+
+        const auto is_foldable_weight = [](const std::string & name) {
+            static const char * kinds[] = {
+                "attn_q", "attn_k", "attn_v", "attn_qkv", "attn_output",
+                "ffn_gate", "ffn_up", "ffn_down",
+                "ffn_gate_exps", "ffn_up_exps", "ffn_down_exps", "ffn_gate_up_exps",
+                "ffn_gate_shexp", "ffn_up_shexp", "ffn_down_shexp",
+            };
+            if (name.compare(0, 4, "blk.") != 0) {
+                return false;
+            }
+            size_t pos = 4;
+            while (pos < name.size() && isdigit((unsigned char) name[pos])) {
+                pos++;
+            }
+            if (pos == 4 || pos >= name.size() || name[pos] != '.') {
+                return false;
+            }
+            pos++;
+            for (const char * kind : kinds) {
+                const std::string suffix = std::string(kind) + ".weight";
+                if (name.compare(pos, std::string::npos, suffix) == 0) {
+                    return true;
+                }
+            }
+            return false;
+        };
+
         for (const auto & weight_name : weight_names) {
+            if (!is_foldable_weight(weight_name)) {
+                throw std::runtime_error(format(
+                    "prism.hadamard: weight '%s' is not on a verified Hadamard-aware matmul path", weight_name.c_str()));
+            }
             if (!hadamard_weight_blocks.emplace(weight_name, block_size).second) {
                 throw std::runtime_error(format("duplicate prism.hadamard weight: %s", weight_name.c_str()));
             }

From 09ca0fd07aed16504fd527a0df4701868ac281c8 Mon Sep 17 00:00:00 2001
From: bri-prism <288398250+bri-prism@users.noreply.github.com>
Date: Sun, 23 Aug 2026 12:08:15 -0700
Subject: [PATCH 10/19] ggml-vulkan: create F16 FWHT pipelines only when the
 device supports fp16

The fwht_f16 shaders declare float16_t, but VK_KHR_shader_float16_int8 is
only enabled when device->fp16 is set, so unconditional creation could
fail pipeline setup on devices without shader-float16 (or under
GGML_VK_DISABLE_F16) even when no F16 FWHT is ever dispatched. Guard both
the subgroup and shmem variants; ggml_vk_can_use_fwht already treats the
null pipeline as unsupported and falls back to the rotation-matrix matmul.

(cherry picked from commit 4edee06d05e12b0a356374663444a1ffefb21f5d)
---
 ggml/src/ggml-vulkan/ggml-vulkan.cpp | 9 +++++++--
 1 file changed, 7 insertions(+), 2 deletions(-)

diff --git a/ggml/src/ggml-vulkan/ggml-vulkan.cpp b/ggml/src/ggml-vulkan/ggml-vulkan.cpp
index 3a7b332983d..832e344ccff 100644
--- a/ggml/src/ggml-vulkan/ggml-vulkan.cpp
+++ b/ggml/src/ggml-vulkan/ggml-vulkan.cpp
@@ -5767,7 +5767,10 @@ static void ggml_vk_load_shaders(vk_device& device, vk_pipeline requested) {
         for (uint32_t n : {64, 128, 256, 512}) {
             if (device->subgroup_size <= n) {
             ggml_vk_create_pipeline(device, device->pipeline_fwht_f32[idx], "fwht_f32", fwht_f32_len, fwht_f32_data, "main", 2, sizeof(vk_op_fwht_push_constants), {1, 1, 1}, { device->subgroup_size, n }, 1, true, true, device->subgroup_size);
-            ggml_vk_create_pipeline(device, device->pipeline_fwht_f16[idx], "fwht_f16", fwht_f16_len, fwht_f16_data, "main", 2, sizeof(vk_op_fwht_push_constants), {1, 1, 1}, { device->subgroup_size, n }, 1, true, true, device->subgroup_size);
+            // the f16 shader needs shader-float16; a null pipeline makes ggml_vk_can_use_fwht fall back
+            if (device->fp16) {
+                ggml_vk_create_pipeline(device, device->pipeline_fwht_f16[idx], "fwht_f16", fwht_f16_len, fwht_f16_data, "main", 2, sizeof(vk_op_fwht_push_constants), {1, 1, 1}, { device->subgroup_size, n }, 1, true, true, device->subgroup_size);
+            }
             }
             ++idx;
         }
@@ -5776,7 +5779,9 @@ static void ggml_vk_load_shaders(vk_device& device, vk_pipeline requested) {
         for (uint32_t n : {64, 128, 256, 512}) {
             const uint32_t block_size = std::min(device->subgroup_size, n);
             ggml_vk_create_pipeline(device, device->pipeline_fwht_f32[idx], "fwht_shmem_f32", fwht_shmem_f32_len, fwht_shmem_f32_data, "main", 2, sizeof(vk_op_fwht_push_constants), {1, 1, 1}, { block_size, n }, 1);
-            ggml_vk_create_pipeline(device, device->pipeline_fwht_f16[idx], "fwht_shmem_f16", fwht_shmem_f16_len, fwht_shmem_f16_data, "main", 2, sizeof(vk_op_fwht_push_constants), {1, 1, 1}, { block_size, n }, 1);
+            if (device->fp16) {
+                ggml_vk_create_pipeline(device, device->pipeline_fwht_f16[idx], "fwht_shmem_f16", fwht_shmem_f16_len, fwht_shmem_f16_data, "main", 2, sizeof(vk_op_fwht_push_constants), {1, 1, 1}, { block_size, n }, 1);
+            }
             ++idx;
         }
     }

From f4a4697ace6186d4fddeb02b36fa5bdb446a0b29 Mon Sep 17 00:00:00 2001
From: bri-prism <288398250+bri-prism@users.noreply.github.com>
Date: Sun, 23 Aug 2026 10:43:15 -0700
Subject: [PATCH 11/19] llama: explicit sign vectors for the Hadamard GGUF
 contract

sign_mode 'explicit' carries per-width sign vectors in GGUF metadata
(prism.hadamard.sign_widths + flattened sign_values). The loader
materializes one F32 sign tensor per (width, buffer type) and the graph
applies x' = H (s * x): elementwise sign flip, then the blockwise FWHT
hint matmul. Identity mode is unchanged; unknown modes still refuse to
load. Converter accepts schema v2 manifests with an explicit signs table.

(cherry picked from commit c49590b0c0251ab8f6dc27ad71eb45d4a26d1cb7)
---
 conversion/base.py         | 31 +++++++++++----
 src/llama-graph.cpp        | 10 ++++-
 src/llama-graph.h          | 11 ++++--
 src/llama-model-loader.cpp |  1 +
 src/llama-model.cpp        | 79 ++++++++++++++++++++++++++++++++++++--
 src/llama-model.h          |  4 +-
 6 files changed, 118 insertions(+), 18 deletions(-)

diff --git a/conversion/base.py b/conversion/base.py
index 6f43aee6ba2..b9278dcddb9 100644
--- a/conversion/base.py
+++ b/conversion/base.py
@@ -626,7 +626,8 @@ def add_hadamard_metadata(self) -> None:
         with manifest_path.open("r", encoding="utf-8") as f:
             manifest = json.load(f)
 
-        if manifest.get("schema_version") != 1 or manifest.get("kind") != "hadamard-weight-fold":
+        schema_version = manifest.get("schema_version")
+        if schema_version not in (1, 2) or manifest.get("kind") != "hadamard-weight-fold":
             raise ValueError(f"Unsupported Hadamard manifest: {manifest_path}")
         if manifest.get("status") != "requires-matching-runtime":
             raise ValueError(f"Unexpected Hadamard manifest status: {manifest.get('status')!r}")
@@ -639,11 +640,21 @@ def add_hadamard_metadata(self) -> None:
             raise ValueError(f"Invalid Hadamard block size: {block_size!r}")
         if transform.get("name") != "normalized-signed-sylvester-walsh-hadamard":
             raise ValueError(f"Unsupported Hadamard transform: {transform.get('name')!r}")
-        if transform.get("sign_mode") != "identity":
-            raise ValueError(
-                "GGUF Hadamard runtime currently supports only identity signs; "
-                "repack with --sign-mode identity"
-            )
+        sign_mode = transform.get("sign_mode")
+        if sign_mode not in ("identity", "explicit"):
+            raise ValueError(f"Unsupported Hadamard sign mode: {sign_mode!r}")
+        sign_widths: list[int] = []
+        sign_values: list[int] = []
+        if sign_mode == "explicit":
+            signs = manifest.get("signs")
+            if not isinstance(signs, dict) or not signs:
+                raise ValueError("explicit sign mode requires a signs table")
+            for width_str, vec in sorted(signs.items(), key=lambda kv: int(kv[0])):
+                width = int(width_str)
+                if len(vec) != width or any(v not in (-1, 1) for v in vec):
+                    raise ValueError(f"invalid sign vector for width {width}")
+                sign_widths.append(width)
+                sign_values.extend(int(v) for v in vec)
 
         tensor_records = manifest.get("tensors")
         if not isinstance(tensor_records, list) or not tensor_records:
@@ -695,9 +706,13 @@ def add_hadamard_metadata(self) -> None:
         self.gguf_writer.add_uint32("prism.hadamard.block_size", block_size)
         self.gguf_writer.add_string("prism.hadamard.transform", "normalized-sylvester-walsh-hadamard")
         self.gguf_writer.add_string("prism.hadamard.axis", "input-last-dimension")
-        self.gguf_writer.add_string("prism.hadamard.sign_mode", "identity")
+        self.gguf_writer.add_string("prism.hadamard.sign_mode", sign_mode)
         self.gguf_writer.add_array("prism.hadamard.weight_names", weight_names)
-        logger.info("GGUF Hadamard contract: H%d for %d folded weight(s)", block_size, len(weight_names))
+        if sign_mode == "explicit":
+            self.gguf_writer.add_array("prism.hadamard.sign_widths", sign_widths)
+            self.gguf_writer.add_array("prism.hadamard.sign_values", sign_values)
+        logger.info("GGUF Hadamard contract: H%d, sign_mode=%s, %d folded weight(s)",
+                    block_size, sign_mode, len(weight_names))
 
     def set_gguf_parameters(self):
         raise NotImplementedError("set_gguf_parameters() must be implemented in subclasses")
diff --git a/src/llama-graph.cpp b/src/llama-graph.cpp
index e4c95305bb0..38beb43fcfa 100644
--- a/src/llama-graph.cpp
+++ b/src/llama-graph.cpp
@@ -1495,7 +1495,10 @@ ggml_tensor * llm_graph_context::build_lora_mm(
     if (hadamard_rotations) {
         const auto it = hadamard_rotations->find(w);
         if (it != hadamard_rotations->end()) {
-            cur_mm = llama_mul_mat_hadamard(ctx0, cur, it->second);
+            if (it->second.signs) {
+                cur_mm = ggml_mul(ctx0, cur_mm, it->second.signs);
+            }
+            cur_mm = llama_mul_mat_hadamard(ctx0, cur_mm, it->second.rot);
         }
     }
 
@@ -1535,7 +1538,10 @@ ggml_tensor * llm_graph_context::build_lora_mm_id(
     if (hadamard_rotations) {
         const auto it = hadamard_rotations->find(w);
         if (it != hadamard_rotations->end()) {
-            cur_mm = llama_mul_mat_hadamard(ctx0, cur, it->second);
+            if (it->second.signs) {
+                cur_mm = ggml_mul(ctx0, cur_mm, it->second.signs);
+            }
+            cur_mm = llama_mul_mat_hadamard(ctx0, cur_mm, it->second.rot);
         }
     }
 
diff --git a/src/llama-graph.h b/src/llama-graph.h
index 6b0df9428cf..af68d768958 100644
--- a/src/llama-graph.h
+++ b/src/llama-graph.h
@@ -17,9 +17,14 @@ struct ggml_cgraph;
 struct ggml_context;
 struct ggml_tensor;
 
-// Maps a folded model weight to the normalized Hadamard matrix required to
-// restore its activation-side basis immediately before the matmul.
-using llama_hadamard_rotations = std::unordered_map<const ggml_tensor *, ggml_tensor *>;
+// Maps a folded model weight to the activation-side transform applied
+// immediately before the matmul: optional sign flip, then the normalized
+// blockwise Hadamard rotation.
+struct llama_hadamard_transform {
+    ggml_tensor * rot;
+    ggml_tensor * signs; // nullptr for identity sign mode
+};
+using llama_hadamard_rotations = std::unordered_map<const ggml_tensor *, llama_hadamard_transform>;
 
 struct llama_cparams;
 struct llama_layer;
diff --git a/src/llama-model-loader.cpp b/src/llama-model-loader.cpp
index 552dc478874..1e3a82687db 100644
--- a/src/llama-model-loader.cpp
+++ b/src/llama-model-loader.cpp
@@ -408,6 +408,7 @@ namespace GGUFMeta {
     }
 
     template bool llama_model_loader::get_arr<std::string>(const std::string & key, std::vector<std::string> & result, bool required);
+    template bool llama_model_loader::get_arr<int32_t>(const std::string & key, std::vector<int32_t> & result, bool required);
     template bool llama_model_loader::get_arr<std::vector<std::string>>(enum llm_kv kid, std::vector<std::string> & result, bool required);
     template bool llama_model_loader::get_arr<std::array<int32_t, 512>>(enum llm_kv kid, std::array<int32_t, 512> & result, bool required);
     template bool llama_model_loader::get_arr<std::vector<int32_t>>(enum llm_kv kid, std::vector<int32_t> & result, bool required);
diff --git a/src/llama-model.cpp b/src/llama-model.cpp
index 7560bbbd126..ecf77053ea6 100644
--- a/src/llama-model.cpp
+++ b/src/llama-model.cpp
@@ -1218,13 +1218,37 @@ void llama_model_base::load_hparams(llama_model_loader & ml) {
         if (axis != "input-last-dimension") {
             throw std::runtime_error(format("unsupported prism.hadamard.axis: %s", axis.c_str()));
         }
-        if (sign_mode != "identity") {
+        if (sign_mode != "identity" && sign_mode != "explicit") {
             throw std::runtime_error(format("unsupported prism.hadamard.sign_mode: %s", sign_mode.c_str()));
         }
         if (weight_names.empty()) {
             throw std::runtime_error("prism.hadamard.weight_names is empty");
         }
 
+        if (sign_mode == "explicit") {
+            std::vector<int32_t> sign_widths;
+            std::vector<int32_t> sign_values;
+            ml.get_arr("prism.hadamard.sign_widths", sign_widths);
+            ml.get_arr("prism.hadamard.sign_values", sign_values);
+            size_t off = 0;
+            for (const int32_t width : sign_widths) {
+                if (width <= 0 || (uint32_t) width % block_size != 0 || off + width > sign_values.size()) {
+                    throw std::runtime_error(format("invalid prism.hadamard sign width: %d", width));
+                }
+                auto & vec = hadamard_sign_data[width];
+                vec.assign(sign_values.begin() + off, sign_values.begin() + off + width);
+                for (const int32_t v : vec) {
+                    if (v != 1 && v != -1) {
+                        throw std::runtime_error("prism.hadamard sign values must be +/-1");
+                    }
+                }
+                off += width;
+            }
+            if (off != sign_values.size()) {
+                throw std::runtime_error("prism.hadamard.sign_values length mismatch");
+            }
+        }
+
         // the activation-side transform is applied only by build_lora_mm/build_lora_mm_id;
         // refuse to load folded weights for architectures or tensor kinds that are not
         // verified to route every matmul through those helpers, rather than run wrong math
@@ -1895,6 +1919,7 @@ bool llama_model_base::load_tensors(llama_model_loader & ml) {
         };
 
         std::vector<hadamard_rotation> rotations;
+        std::map<std::pair<uint32_t, ggml_backend_buffer_type_t>, ggml_tensor *> sign_tensors;
 
         for (const auto & entry : hadamard_weight_blocks) {
             const std::string & weight_name = entry.first;
@@ -1962,11 +1987,57 @@ bool llama_model_base::load_tensors(llama_model_loader & ml) {
                 it = std::prev(rotations.end());
             }
 
-            hadamard_rotations.emplace(weight, it->tensor);
+            ggml_tensor * sign_tensor = nullptr;
+            if (!hadamard_sign_data.empty()) {
+                const uint32_t width = (uint32_t) weight->ne[0];
+                const auto sd = hadamard_sign_data.find(width);
+                if (sd == hadamard_sign_data.end()) {
+                    throw std::runtime_error(format(
+                        "prism.hadamard has no sign vector for width %u (%s)", width, weight_name.c_str()));
+                }
+                const auto key = std::make_pair(width, buft);
+                auto st = sign_tensors.find(key);
+                if (st == sign_tensors.end()) {
+                    ggml_init_params params = {
+                        /*.mem_size   =*/ ggml_tensor_overhead(),
+                        /*.mem_buffer =*/ NULL,
+                        /*.no_alloc   =*/ true,
+                    };
+                    ggml_context_ptr ctx { ggml_init(params) };
+                    if (!ctx) {
+                        throw std::runtime_error("failed to create Hadamard sign context");
+                    }
+
+                    ggml_tensor * signs = ggml_new_tensor_1d(ctx.get(), GGML_TYPE_F32, width);
+                    char sign_name[GGML_MAX_NAME];
+                    snprintf(sign_name, sizeof(sign_name), "prism.hadamard.signs.%u", width);
+                    ggml_set_name(signs, sign_name);
+
+                    ggml_backend_buffer_ptr buffer { ggml_backend_alloc_ctx_tensors_from_buft(ctx.get(), buft) };
+                    if (!buffer) {
+                        throw std::runtime_error(format("unable to allocate %s Hadamard sign buffer", ggml_backend_buft_name(buft)));
+                    }
+                    ggml_backend_buffer_set_usage(buffer.get(), GGML_BACKEND_BUFFER_USAGE_WEIGHTS);
+
+                    std::vector<float> data(width);
+                    for (uint32_t i = 0; i < width; ++i) {
+                        data[i] = (float) sd->second[i];
+                    }
+                    ggml_backend_tensor_set(signs, data.data(), 0, data.size() * sizeof(float));
+
+                    std::vector<ggml_backend_buffer_ptr> buffers;
+                    buffers.emplace_back(std::move(buffer));
+                    pimpl->ctxs_bufs.emplace_back(std::move(ctx), std::move(buffers));
+                    st = sign_tensors.emplace(key, signs).first;
+                }
+                sign_tensor = st->second;
+            }
+
+            hadamard_rotations.emplace(weight, llama_hadamard_transform { it->tensor, sign_tensor });
         }
 
-        LLAMA_LOG_INFO("%s: loaded %zu Hadamard-folded weight(s) using %zu persistent rotation(s)\n",
-                __func__, hadamard_rotations.size(), rotations.size());
+        LLAMA_LOG_INFO("%s: loaded %zu Hadamard-folded weight(s) using %zu rotation(s) and %zu sign vector(s)\n",
+                __func__, hadamard_rotations.size(), rotations.size(), sign_tensors.size());
     }
 
     return true;
diff --git a/src/llama-model.h b/src/llama-model.h
index 31ec2425c10..a953ba73474 100644
--- a/src/llama-model.h
+++ b/src/llama-model.h
@@ -667,8 +667,10 @@ struct llama_model {
     // Hadamard-folded GGUF weights are matched with persistent model tensors
     // containing the activation-side transform.  The string map is populated
     // from GGUF metadata while loading hparams; the pointer map is populated
-    // after model buffers have been allocated.
+    // after model buffers have been allocated.  In explicit sign mode the
+    // per-width sign vectors come from GGUF metadata as well.
     std::unordered_map<std::string, uint32_t> hadamard_weight_blocks;
+    std::map<uint32_t, std::vector<int32_t>> hadamard_sign_data;
     llama_hadamard_rotations hadamard_rotations;
 
     // list of devices used in this model

From bd714ee2fac8e2a9664c620ac8b7d4cfb0309e7d Mon Sep 17 00:00:00 2001
From: bri-prism <288398250+bri-prism@users.noreply.github.com>
Date: Sun, 23 Aug 2026 12:25:11 -0700
Subject: [PATCH 12/19] llama: inverse Hadamard transform for latent embedding
 tables

Tensors listed in prism.hadamard.inverse_weight_names store latent
(rotated) rows; the graph restores the primal basis right after the
token-embedding lookup: h = s * (H z). This lets the embedding table
stay in the quantized latent format instead of a dequantized primal
copy. The rotation and sign tensors are shared with the matmul-side
transforms; the direct-embeddings input path is untouched.

(cherry picked from commit 2e85ea475407322c25e45ae33aa36df69f16d572)
---
 conversion/base.py    | 32 ++++++++++++++++++++++++--------
 src/llama-context.cpp |  1 +
 src/llama-graph.cpp   | 13 +++++++++++++
 src/llama-graph.h     |  2 ++
 src/llama-model.cpp   | 32 +++++++++++++++++++++++++++-----
 src/llama-model.h     |  2 ++
 6 files changed, 69 insertions(+), 13 deletions(-)

diff --git a/conversion/base.py b/conversion/base.py
index b9278dcddb9..ca9d0b85612 100644
--- a/conversion/base.py
+++ b/conversion/base.py
@@ -686,21 +686,35 @@ def add_hadamard_metadata(self) -> None:
             r")\.weight"
         )
         weight_names: list[str] = []
+        inverse_weight_names: list[str] = []
         for record in tensor_records:
             if not isinstance(record, dict) or not isinstance(record.get("name"), str):
                 raise ValueError("Hadamard manifest has an invalid tensor record")
             if record.get("axis") != -1:
                 raise ValueError(f"Unsupported Hadamard tensor axis for {record['name']!r}")
+            role = record.get("role", "fold-before-matmul")
+            if role not in ("fold-before-matmul", "inverse-after-lookup"):
+                raise ValueError(f"Unsupported Hadamard tensor role for {record['name']!r}: {role!r}")
             filtered = self.filter_tensors((record["name"], lambda: None))
             if filtered is None:
                 raise ValueError(f"Hadamard tensor is filtered out: {record['name']!r}")
             mapped = self.map_tensor_name(filtered[0])
-            if not _HADAMARD_KINDS.fullmatch(mapped):
-                raise ValueError(
-                    f"Hadamard tensor {record['name']!r} maps to {mapped!r}, which is not on a "
-                    "verified Hadamard-aware matmul path"
-                )
-            weight_names.append(mapped)
+            if role == "inverse-after-lookup":
+                # the runtime applies the inverse transform only to the token-embedding
+                # lookup; any other latent table would load and silently stay rotated
+                if mapped != "token_embd.weight":
+                    raise ValueError(
+                        f"Hadamard tensor {record['name']!r} maps to {mapped!r}, which is not a "
+                        "verified inverse-after-lookup table"
+                    )
+                inverse_weight_names.append(mapped)
+            else:
+                if not _HADAMARD_KINDS.fullmatch(mapped):
+                    raise ValueError(
+                        f"Hadamard tensor {record['name']!r} maps to {mapped!r}, which is not on a "
+                        "verified Hadamard-aware matmul path"
+                    )
+                weight_names.append(mapped)
 
         self.gguf_writer.add_uint32("prism.hadamard.version", 1)
         self.gguf_writer.add_uint32("prism.hadamard.block_size", block_size)
@@ -711,8 +725,10 @@ def add_hadamard_metadata(self) -> None:
         if sign_mode == "explicit":
             self.gguf_writer.add_array("prism.hadamard.sign_widths", sign_widths)
             self.gguf_writer.add_array("prism.hadamard.sign_values", sign_values)
-        logger.info("GGUF Hadamard contract: H%d, sign_mode=%s, %d folded weight(s)",
-                    block_size, sign_mode, len(weight_names))
+        if inverse_weight_names:
+            self.gguf_writer.add_array("prism.hadamard.inverse_weight_names", inverse_weight_names)
+        logger.info("GGUF Hadamard contract: H%d, sign_mode=%s, %d folded weight(s), %d inverse-lookup",
+                    block_size, sign_mode, len(weight_names), len(inverse_weight_names))
 
     def set_gguf_parameters(self):
         raise NotImplementedError("set_gguf_parameters() must be implemented in subclasses")
diff --git a/src/llama-context.cpp b/src/llama-context.cpp
index 76ee6a3ade5..df543515464 100644
--- a/src/llama-context.cpp
+++ b/src/llama-context.cpp
@@ -2498,6 +2498,7 @@ llm_graph_params llama_context::graph_params(
         /*.mctx        =*/ mctx,
         /*.cross       =*/ &cross,
         /*.hadamard_rotations =*/ &model.hadamard_rotations,
+        /*.hadamard_inverses  =*/ &model.hadamard_inverses,
         /*.samplers    =*/ sampling.samplers,
         /*.n_outputs   =*/ n_outputs,
         /*.cb          =*/ graph_get_cb(),
diff --git a/src/llama-graph.cpp b/src/llama-graph.cpp
index 38beb43fcfa..fe2483f1cec 100644
--- a/src/llama-graph.cpp
+++ b/src/llama-graph.cpp
@@ -1465,6 +1465,7 @@ llm_graph_context::llm_graph_context(const llm_graph_params & params) :
     mctx             (params.mctx),
     cross            (params.cross),
     hadamard_rotations(params.hadamard_rotations),
+    hadamard_inverses (params.hadamard_inverses),
     samplers         (params.samplers),
     cb_func          (params.cb),
     res              (params.res),
@@ -2333,6 +2334,18 @@ ggml_tensor * llm_graph_context::build_inp_embd(ggml_tensor * tok_embd) const {
 
         cur = ggml_get_rows(ctx0, tok_embd, inp->tokens);
 
+        // a Hadamard-latent embedding table stores rotated rows; restore the
+        // primal basis right after the lookup: h = s * (H z)
+        if (hadamard_inverses) {
+            const auto it = hadamard_inverses->find(tok_embd);
+            if (it != hadamard_inverses->end()) {
+                cur = llama_mul_mat_hadamard(ctx0, cur, it->second.rot);
+                if (it->second.signs) {
+                    cur = ggml_mul(ctx0, cur, it->second.signs);
+                }
+            }
+        }
+
         // apply lora for embedding tokens if needed
         for (const auto & lora : *loras) {
             llama_adapter_lora_weight * lw = lora.first->get_weight(tok_embd);
diff --git a/src/llama-graph.h b/src/llama-graph.h
index af68d768958..1b2a2d13dfb 100644
--- a/src/llama-graph.h
+++ b/src/llama-graph.h
@@ -808,6 +808,7 @@ struct llm_graph_params {
     const llama_memory_context_i * mctx;
     const llama_cross            * cross;
     const llama_hadamard_rotations * hadamard_rotations;
+    const llama_hadamard_rotations * hadamard_inverses;
 
     std::map<llama_seq_id, llama_sampler *> samplers;
 
@@ -1049,6 +1050,7 @@ struct llm_graph_context {
     const llama_memory_context_i * mctx;
     const llama_cross            * cross;
     const llama_hadamard_rotations * hadamard_rotations;
+    const llama_hadamard_rotations * hadamard_inverses;
 
     std::map<llama_seq_id, llama_sampler *> samplers;
 
diff --git a/src/llama-model.cpp b/src/llama-model.cpp
index ecf77053ea6..2ed8c9667f0 100644
--- a/src/llama-model.cpp
+++ b/src/llama-model.cpp
@@ -1302,6 +1302,22 @@ void llama_model_base::load_hparams(llama_model_loader & ml) {
                 throw std::runtime_error(format("duplicate prism.hadamard weight: %s", weight_name.c_str()));
             }
         }
+
+        // tensors consumed by row lookup store latent rows and need the
+        // inverse transform applied to the lookup result instead
+        std::vector<std::string> inverse_names;
+        ml.get_arr("prism.hadamard.inverse_weight_names", inverse_names, false);
+        for (const auto & name : inverse_names) {
+            // the graph applies the inverse only to the token-embedding lookup; any
+            // other latent table would load and silently stay rotated
+            if (name != "token_embd.weight") {
+                throw std::runtime_error(format(
+                    "prism.hadamard: weight '%s' is not a verified inverse-after-lookup table", name.c_str()));
+            }
+            if (hadamard_weight_blocks.count(name) || !hadamard_inverse_blocks.emplace(name, block_size).second) {
+                throw std::runtime_error(format("duplicate prism.hadamard inverse weight: %s", name.c_str()));
+            }
+        }
     }
 
     // get general kv
@@ -1911,7 +1927,7 @@ bool llama_model_base::load_tensors(llama_model_loader & ml) {
         }
     }
 
-    if (!hadamard_weight_blocks.empty()) {
+    if (!hadamard_weight_blocks.empty() || !hadamard_inverse_blocks.empty()) {
         struct hadamard_rotation {
             uint32_t block_size;
             ggml_backend_buffer_type_t buft;
@@ -1921,7 +1937,12 @@ bool llama_model_base::load_tensors(llama_model_loader & ml) {
         std::vector<hadamard_rotation> rotations;
         std::map<std::pair<uint32_t, ggml_backend_buffer_type_t>, ggml_tensor *> sign_tensors;
 
-        for (const auto & entry : hadamard_weight_blocks) {
+        const std::pair<const std::unordered_map<std::string, uint32_t> *, llama_hadamard_rotations *> groups[] = {
+            { &hadamard_weight_blocks,  &hadamard_rotations },
+            { &hadamard_inverse_blocks, &hadamard_inverses  },
+        };
+        for (const auto & [blocks, target] : groups)
+        for (const auto & entry : *blocks) {
             const std::string & weight_name = entry.first;
             const uint32_t block_size = entry.second;
             const ggml_tensor * weight = get_tensor(weight_name.c_str());
@@ -2033,11 +2054,12 @@ bool llama_model_base::load_tensors(llama_model_loader & ml) {
                 sign_tensor = st->second;
             }
 
-            hadamard_rotations.emplace(weight, llama_hadamard_transform { it->tensor, sign_tensor });
+            target->emplace(weight, llama_hadamard_transform { it->tensor, sign_tensor });
         }
 
-        LLAMA_LOG_INFO("%s: loaded %zu Hadamard-folded weight(s) using %zu rotation(s) and %zu sign vector(s)\n",
-                __func__, hadamard_rotations.size(), rotations.size(), sign_tensors.size());
+        LLAMA_LOG_INFO("%s: loaded %zu Hadamard-folded weight(s) (%zu inverse-lookup) using %zu rotation(s) and %zu sign vector(s)\n",
+                __func__, hadamard_rotations.size() + hadamard_inverses.size(), hadamard_inverses.size(),
+                rotations.size(), sign_tensors.size());
     }
 
     return true;
diff --git a/src/llama-model.h b/src/llama-model.h
index a953ba73474..99dcfcb9ee6 100644
--- a/src/llama-model.h
+++ b/src/llama-model.h
@@ -670,8 +670,10 @@ struct llama_model {
     // after model buffers have been allocated.  In explicit sign mode the
     // per-width sign vectors come from GGUF metadata as well.
     std::unordered_map<std::string, uint32_t> hadamard_weight_blocks;
+    std::unordered_map<std::string, uint32_t> hadamard_inverse_blocks;
     std::map<uint32_t, std::vector<int32_t>> hadamard_sign_data;
     llama_hadamard_rotations hadamard_rotations;
+    llama_hadamard_rotations hadamard_inverses;
 
     // list of devices used in this model
     std::vector<llama_device> devices;

From 86e902241f2e561ea0802a08454d09f596fc0398 Mon Sep 17 00:00:00 2001
From: bri-prism <288398250+bri-prism@users.noreply.github.com>
Date: Sun, 23 Aug 2026 13:41:03 -0700
Subject: [PATCH 13/19] llama: grouped-order GDN out_proj support for
 Hadamard-folded models

The converter's tiled V-head reorder permutes out_proj's input axis,
which is the rotation axis of a folded latent; a post-fold column
permutation breaks the blockwise Hadamard correspondence and cannot be
refolded without destroying the ternary codes. For folded out_proj
tensors the converter now keeps the training (grouped) V order and sets
prism.hadamard.gdn_v_grouped; the runtime permutes the activation
tiled->grouped (reshape+permute+cont) before the sign flip and rotation.

(cherry picked from commit e71a350b219eeababfe3fe235cb112cbe2a38d7c)
---
 conversion/base.py         | 19 +++++++++++++++++++
 conversion/qwen.py         | 10 ++++++++--
 src/llama-graph.cpp        | 30 ++++++++++++++++++++++++------
 src/llama-graph.h          |  6 ++++++
 src/llama-model-loader.cpp |  1 +
 src/llama-model.cpp        | 16 +++++++++++++++-
 src/llama-model.h          |  1 +
 7 files changed, 74 insertions(+), 9 deletions(-)

diff --git a/conversion/base.py b/conversion/base.py
index ca9d0b85612..2bf36156faf 100644
--- a/conversion/base.py
+++ b/conversion/base.py
@@ -617,6 +617,21 @@ def map_tensor_name(self, name: str, try_suffixes: Sequence[str] = (".weight", "
             raise ValueError(f"Can not map tensor {name!r}")
         return new_name
 
+    def hadamard_folded_names(self) -> set[str]:
+        """Source-tensor names folded under a Hadamard manifest, or empty."""
+        cached = getattr(self, "_hadamard_folded_names", None)
+        if cached is not None:
+            return cached
+        names: set[str] = set()
+        manifest_path = self.dir_model / "hadamard_packing.json"
+        if manifest_path.is_file():
+            with manifest_path.open("r", encoding="utf-8") as f:
+                for record in json.load(f).get("tensors", []):
+                    if isinstance(record, dict) and isinstance(record.get("name"), str):
+                        names.add(record["name"])
+        self._hadamard_folded_names = names
+        return names
+
     def add_hadamard_metadata(self) -> None:
         """Transfer a packed-checkpoint transform contract into GGUF metadata."""
         manifest_path = self.dir_model / "hadamard_packing.json"
@@ -683,6 +698,7 @@ def add_hadamard_metadata(self) -> None:
             r"|ffn_gate|ffn_up|ffn_down"
             r"|ffn_gate_exps|ffn_up_exps|ffn_down_exps|ffn_gate_up_exps"
             r"|ffn_gate_shexp|ffn_up_shexp|ffn_down_shexp"
+            r"|ssm_out"
             r")\.weight"
         )
         weight_names: list[str] = []
@@ -727,6 +743,9 @@ def add_hadamard_metadata(self) -> None:
             self.gguf_writer.add_array("prism.hadamard.sign_values", sign_values)
         if inverse_weight_names:
             self.gguf_writer.add_array("prism.hadamard.inverse_weight_names", inverse_weight_names)
+        if getattr(self, "_hadamard_gdn_v_grouped", False):
+            self.gguf_writer.add_bool("prism.hadamard.gdn_v_grouped", True)
+            logger.info("GGUF Hadamard: linear-attention out_proj kept in grouped V order")
         logger.info("GGUF Hadamard contract: H%d, sign_mode=%s, %d folded weight(s), %d inverse-lookup",
                     block_size, sign_mode, len(weight_names), len(inverse_weight_names))
 
diff --git a/conversion/qwen.py b/conversion/qwen.py
index cdba8a63e9c..d8c3f006682 100644
--- a/conversion/qwen.py
+++ b/conversion/qwen.py
@@ -605,8 +605,14 @@ def modify_tensors(self, data_torch: Tensor, name: str, bid: int | None) -> Iter
                 data_torch = torch.cat([qk_part, v_part], dim=0)
 
             elif ".out_proj." in name:
-                # Out projection weight: reorder columns (input dimension)
-                data_torch = self._reorder_v_heads(data_torch, 1, num_k_heads, num_v_per_k, head_v_dim)
+                if name in self.hadamard_folded_names():
+                    # Hadamard-folded latent: the rotation axis must keep the
+                    # training (grouped) V order; the runtime permutes the
+                    # activation tiled->grouped before the transform instead.
+                    self._hadamard_gdn_v_grouped = True
+                else:
+                    # Out projection weight: reorder columns (input dimension)
+                    data_torch = self._reorder_v_heads(data_torch, 1, num_k_heads, num_v_per_k, head_v_dim)
 
         yield from super().modify_tensors(data_torch, name, bid)
 
diff --git a/src/llama-graph.cpp b/src/llama-graph.cpp
index fe2483f1cec..2742b319d80 100644
--- a/src/llama-graph.cpp
+++ b/src/llama-graph.cpp
@@ -1496,10 +1496,19 @@ ggml_tensor * llm_graph_context::build_lora_mm(
     if (hadamard_rotations) {
         const auto it = hadamard_rotations->find(w);
         if (it != hadamard_rotations->end()) {
-            if (it->second.signs) {
-                cur_mm = ggml_mul(ctx0, cur_mm, it->second.signs);
+            const auto & t = it->second;
+            if (t.perm_rep > 1) {
+                // tiled [hd, nk, rep] -> grouped [hd, rep, nk] feature order
+                ggml_tensor * x = ggml_is_contiguous(cur_mm) ? cur_mm : ggml_cont(ctx0, cur_mm);
+                const int64_t ne1 = x->ne[1], ne2 = x->ne[2], ne3 = x->ne[3];
+                x = ggml_reshape_4d(ctx0, x, t.perm_hd, t.perm_nk, t.perm_rep, ne1*ne2*ne3);
+                x = ggml_cont(ctx0, ggml_permute(ctx0, x, 0, 2, 1, 3));
+                cur_mm = ggml_reshape_4d(ctx0, x, t.perm_hd*t.perm_nk*t.perm_rep, ne1, ne2, ne3);
             }
-            cur_mm = llama_mul_mat_hadamard(ctx0, cur_mm, it->second.rot);
+            if (t.signs) {
+                cur_mm = ggml_mul(ctx0, cur_mm, t.signs);
+            }
+            cur_mm = llama_mul_mat_hadamard(ctx0, cur_mm, t.rot);
         }
     }
 
@@ -1539,10 +1548,19 @@ ggml_tensor * llm_graph_context::build_lora_mm_id(
     if (hadamard_rotations) {
         const auto it = hadamard_rotations->find(w);
         if (it != hadamard_rotations->end()) {
-            if (it->second.signs) {
-                cur_mm = ggml_mul(ctx0, cur_mm, it->second.signs);
+            const auto & t = it->second;
+            if (t.perm_rep > 1) {
+                // tiled [hd, nk, rep] -> grouped [hd, rep, nk] feature order
+                ggml_tensor * x = ggml_is_contiguous(cur_mm) ? cur_mm : ggml_cont(ctx0, cur_mm);
+                const int64_t ne1 = x->ne[1], ne2 = x->ne[2], ne3 = x->ne[3];
+                x = ggml_reshape_4d(ctx0, x, t.perm_hd, t.perm_nk, t.perm_rep, ne1*ne2*ne3);
+                x = ggml_cont(ctx0, ggml_permute(ctx0, x, 0, 2, 1, 3));
+                cur_mm = ggml_reshape_4d(ctx0, x, t.perm_hd*t.perm_nk*t.perm_rep, ne1, ne2, ne3);
+            }
+            if (t.signs) {
+                cur_mm = ggml_mul(ctx0, cur_mm, t.signs);
             }
-            cur_mm = llama_mul_mat_hadamard(ctx0, cur_mm, it->second.rot);
+            cur_mm = llama_mul_mat_hadamard(ctx0, cur_mm, t.rot);
         }
     }
 
diff --git a/src/llama-graph.h b/src/llama-graph.h
index 1b2a2d13dfb..2ba316c92c6 100644
--- a/src/llama-graph.h
+++ b/src/llama-graph.h
@@ -23,6 +23,12 @@ struct ggml_tensor;
 struct llama_hadamard_transform {
     ggml_tensor * rot;
     ggml_tensor * signs; // nullptr for identity sign mode
+    // when perm_rep > 1 the activation arrives with its feature axis in tiled
+    // head order [hd, nk, rep] and must be permuted to the grouped order
+    // [hd, rep, nk] the fold was computed in, before signs and rotation
+    int64_t perm_hd  = 0;
+    int64_t perm_nk  = 0;
+    int64_t perm_rep = 0;
 };
 using llama_hadamard_rotations = std::unordered_map<const ggml_tensor *, llama_hadamard_transform>;
 
diff --git a/src/llama-model-loader.cpp b/src/llama-model-loader.cpp
index 1e3a82687db..7e7eb5cca66 100644
--- a/src/llama-model-loader.cpp
+++ b/src/llama-model-loader.cpp
@@ -436,6 +436,7 @@ namespace GGUFMeta {
     }
 
     template bool llama_model_loader::get_key<bool>       (enum llm_kv kid, bool & result,        bool required);
+    template bool llama_model_loader::get_key<bool>       (const std::string & key, bool & result, bool required);
     template bool llama_model_loader::get_key<float>      (enum llm_kv kid, float & result,       bool required);
     template bool llama_model_loader::get_key<uint32_t>   (enum llm_kv kid, uint32_t & result,    bool required);
     template bool llama_model_loader::get_key<std::string>(enum llm_kv kid, std::string & result, bool required);
diff --git a/src/llama-model.cpp b/src/llama-model.cpp
index 2ed8c9667f0..6278cfcfd44 100644
--- a/src/llama-model.cpp
+++ b/src/llama-model.cpp
@@ -1249,6 +1249,8 @@ void llama_model_base::load_hparams(llama_model_loader & ml) {
             }
         }
 
+        ml.get_key("prism.hadamard.gdn_v_grouped", hadamard_gdn_v_grouped, false);
+
         // the activation-side transform is applied only by build_lora_mm/build_lora_mm_id;
         // refuse to load folded weights for architectures or tensor kinds that are not
         // verified to route every matmul through those helpers, rather than run wrong math
@@ -1272,6 +1274,7 @@ void llama_model_base::load_hparams(llama_model_loader & ml) {
                 "ffn_gate", "ffn_up", "ffn_down",
                 "ffn_gate_exps", "ffn_up_exps", "ffn_down_exps", "ffn_gate_up_exps",
                 "ffn_gate_shexp", "ffn_up_shexp", "ffn_down_shexp",
+                "ssm_out",
             };
             if (name.compare(0, 4, "blk.") != 0) {
                 return false;
@@ -2054,7 +2057,18 @@ bool llama_model_base::load_tensors(llama_model_loader & ml) {
                 sign_tensor = st->second;
             }
 
-            target->emplace(weight, llama_hadamard_transform { it->tensor, sign_tensor });
+            llama_hadamard_transform transform { it->tensor, sign_tensor };
+            if (hadamard_gdn_v_grouped && weight_name.find(".ssm_out.") != std::string::npos) {
+                const int64_t n_v = hparams.ssm_dt_rank;
+                const int64_t n_k = hparams.ssm_n_group;
+                if (n_k <= 0 || n_v <= 0 || n_v % n_k != 0 || weight->ne[0] % n_v != 0) {
+                    throw std::runtime_error(format("prism.hadamard: bad GDN head geometry for %s", weight_name.c_str()));
+                }
+                transform.perm_hd  = weight->ne[0] / n_v;
+                transform.perm_nk  = n_k;
+                transform.perm_rep = n_v / n_k;
+            }
+            target->emplace(weight, transform);
         }
 
         LLAMA_LOG_INFO("%s: loaded %zu Hadamard-folded weight(s) (%zu inverse-lookup) using %zu rotation(s) and %zu sign vector(s)\n",
diff --git a/src/llama-model.h b/src/llama-model.h
index 99dcfcb9ee6..7d1b4732a7f 100644
--- a/src/llama-model.h
+++ b/src/llama-model.h
@@ -672,6 +672,7 @@ struct llama_model {
     std::unordered_map<std::string, uint32_t> hadamard_weight_blocks;
     std::unordered_map<std::string, uint32_t> hadamard_inverse_blocks;
     std::map<uint32_t, std::vector<int32_t>> hadamard_sign_data;
+    bool hadamard_gdn_v_grouped = false;
     llama_hadamard_rotations hadamard_rotations;
     llama_hadamard_rotations hadamard_inverses;
 

From cc92dd0a1deeba9144d37c952b0bb3cb039c968e Mon Sep 17 00:00:00 2001
From: bri-prism <288398250+bri-prism@users.noreply.github.com>
Date: Sun, 23 Aug 2026 14:34:14 -0700
Subject: [PATCH 14/19] convert: match Hadamard-folded out_proj by suffix

Wrapper prefixes are stripped from tensor names before modify_tensors,
so the exact-name check against the manifest never fired and the tiled
V reorder still permuted folded rotation axes.

(cherry picked from commit e9e4ea30aee217c8713706c21063de1664bf6a57)
---
 conversion/qwen.py | 6 +++++-
 1 file changed, 5 insertions(+), 1 deletion(-)

diff --git a/conversion/qwen.py b/conversion/qwen.py
index d8c3f006682..bbb4672f860 100644
--- a/conversion/qwen.py
+++ b/conversion/qwen.py
@@ -605,7 +605,11 @@ def modify_tensors(self, data_torch: Tensor, name: str, bid: int | None) -> Iter
                 data_torch = torch.cat([qk_part, v_part], dim=0)
 
             elif ".out_proj." in name:
-                if name in self.hadamard_folded_names():
+                # the tensor name may have had wrapper prefixes stripped by the
+                # time it reaches modify_tensors, so match manifest entries by
+                # suffix rather than exact name
+                suffix = "linear_attn.out_proj.weight"
+                if name.endswith(suffix) and any(n.endswith(suffix) for n in self.hadamard_folded_names()):
                     # Hadamard-folded latent: the rotation axis must keep the
                     # training (grouped) V order; the runtime permutes the
                     # activation tiled->grouped before the transform instead.

From 2d18730709f48879de3702d8c837b8f672053975 Mon Sep 17 00:00:00 2001
From: bri-prism <288398250+bri-prism@users.noreply.github.com>
Date: Sun, 23 Aug 2026 15:56:03 -0700
Subject: [PATCH 15/19] llama: verify folded weights receive their transform at
 graph build

An architecture whose matmul path bypasses the transform helpers would
load a Hadamard-folded GGUF cleanly and silently compute wrong results.
After reserving the worst-case graph, walk it once: every mul_mat
consuming a folded weight must take the hint matmul's output (through
reshape/view) built from that weight's rotation, and every get_rows of
a latent table must feed an inverse transform. Violations abort context
creation with the offending tensor name.

(cherry picked from commit 34c9b5cd9b1578abf42868b23543ba4bde0758b6)
---
 src/llama-context.cpp | 68 +++++++++++++++++++++++++++++++++++++++++++
 1 file changed, 68 insertions(+)

diff --git a/src/llama-context.cpp b/src/llama-context.cpp
index df543515464..060252eb089 100644
--- a/src/llama-context.cpp
+++ b/src/llama-context.cpp
@@ -28,6 +28,68 @@
 // llama_context
 //
 
+// Verify that every Hadamard-folded weight consumed by the graph receives its
+// activation-side transform, and every latent lookup table gets the inverse.
+// An architecture whose matmul path bypasses the transform helpers would
+// otherwise load cleanly and silently compute wrong results.
+static void llama_verify_hadamard_graph(
+        ggml_cgraph * gf,
+        const llama_hadamard_rotations & rotations,
+        const llama_hadamard_rotations & inverses) {
+    auto unwrap = [](const ggml_tensor * t) {
+        while (t && (t->op == GGML_OP_RESHAPE || t->op == GGML_OP_VIEW)) {
+            t = t->src[0];
+        }
+        return t;
+    };
+
+    std::map<const ggml_tensor *, bool> lookups; // get_rows results of latent tables
+
+    for (int i = 0; i < ggml_graph_n_nodes(gf); ++i) {
+        const ggml_tensor * node = ggml_graph_node(gf, i);
+
+        if (node->op == GGML_OP_GET_ROWS && inverses.count(node->src[0])) {
+            lookups.emplace(node, false);
+            continue;
+        }
+
+        if (node->op != GGML_OP_MUL_MAT && node->op != GGML_OP_MUL_MAT_ID) {
+            continue;
+        }
+
+        if (node->op == GGML_OP_MUL_MAT && ((const int32_t *) node->op_params)[1] == GGML_HINT_SRC0_IS_HADAMARD) {
+            const auto lk = lookups.find(unwrap(node->src[1]));
+            if (lk != lookups.end()) {
+                lk->second = true;
+            }
+            continue;
+        }
+
+        const auto it = rotations.find(node->src[0]);
+        if (it == rotations.end()) {
+            continue;
+        }
+        const ggml_tensor * src = unwrap(node->src[1]);
+        const bool transformed = src && src->op == GGML_OP_MUL_MAT &&
+            ((const int32_t *) src->op_params)[1] == GGML_HINT_SRC0_IS_HADAMARD &&
+            src->src[0] == it->second.rot;
+        if (!transformed) {
+            throw std::runtime_error(format(
+                "Hadamard-folded weight '%s' is consumed without its activation transform; "
+                "this graph's matmul path does not support prism.hadamard folding",
+                node->src[0]->name));
+        }
+    }
+
+    for (const auto & [node, ok] : lookups) {
+        if (!ok) {
+            throw std::runtime_error(format(
+                "Hadamard-latent table '%s' is read without the inverse transform",
+                node->src[0]->name));
+        }
+    }
+}
+
 static llm_graph_type ctx_type_to_graph_type(llama_context_type ctx_type) {
     switch (ctx_type) {
         case LLAMA_CONTEXT_TYPE_DEFAULT: return LLM_GRAPH_TYPE_DEFAULT;
@@ -678,6 +740,12 @@ void llama_context::sched_reserve() {
 
         n_splits_pp = ggml_backend_sched_get_n_splits(sched.get());
         n_nodes_pp  = ggml_graph_n_nodes(gf);
+
+        // a folded weight consumed without its activation-side transform loads
+        // cleanly but computes wrong results; refuse to run such a graph
+        if (!model.hadamard_rotations.empty() || !model.hadamard_inverses.empty()) {
+            llama_verify_hadamard_graph(gf, model.hadamard_rotations, model.hadamard_inverses);
+        }
     }
 
     // reserve with tg (token generation) graph to get the number of splits and nodes

From ea2fa85649a4185407f5998bf18cf5cd5571a143 Mon Sep 17 00:00:00 2001
From: bri-prism <288398250+bri-prism@users.noreply.github.com>
Date: Sun, 23 Aug 2026 17:00:50 -0700
Subject: [PATCH 16/19] llama: run the Hadamard coverage check before graph
 scheduling

The scheduler splits cross-backend paths with copy tensors (a CPU-mapped
latent embedding feeding a GPU FWHT), which breaks the producer chain
the check follows and produced a false positive. Verify the pristine
graph right after build_graph, once per context, and log the consumers
of an unverified lookup before aborting.

(cherry picked from commit 8362c899672279557d786644f3267c32d51caad4)
---
 src/llama-context.cpp | 23 +++++++++++++++++------
 src/llama-context.h   |  3 +++
 2 files changed, 20 insertions(+), 6 deletions(-)

diff --git a/src/llama-context.cpp b/src/llama-context.cpp
index 060252eb089..2c8fbd9055c 100644
--- a/src/llama-context.cpp
+++ b/src/llama-context.cpp
@@ -83,6 +83,16 @@ static void llama_verify_hadamard_graph(
 
     for (const auto & [node, ok] : lookups) {
         if (!ok) {
+            for (int i = 0; i < ggml_graph_n_nodes(gf); ++i) {
+                const ggml_tensor * n2 = ggml_graph_node(gf, i);
+                for (int s = 0; s < GGML_MAX_SRC && n2->src[s]; ++s) {
+                    if (unwrap(n2->src[s]) == node) {
+                        LLAMA_LOG_WARN("%s: latent lookup '%s' consumed by op=%s name='%s' src%d hint=%d\n",
+                                __func__, node->name, ggml_op_name(n2->op), n2->name, s,
+                                ((const int32_t *) n2->op_params)[1]);
+                    }
+                }
+            }
             throw std::runtime_error(format(
                 "Hadamard-latent table '%s' is read without the inverse transform",
                 node->src[0]->name));
@@ -740,12 +750,6 @@ void llama_context::sched_reserve() {
 
         n_splits_pp = ggml_backend_sched_get_n_splits(sched.get());
         n_nodes_pp  = ggml_graph_n_nodes(gf);
-
-        // a folded weight consumed without its activation-side transform loads
-        // cleanly but computes wrong results; refuse to run such a graph
-        if (!model.hadamard_rotations.empty() || !model.hadamard_inverses.empty()) {
-            llama_verify_hadamard_graph(gf, model.hadamard_rotations, model.hadamard_inverses);
-        }
     }
 
     // reserve with tg (token generation) graph to get the number of splits and nodes
@@ -2530,6 +2534,13 @@ ggml_cgraph * llama_context::graph_reserve(
 
     auto * gf = model.build_graph(gparams);
 
+    // verify transform coverage on the pristine graph: after scheduling,
+    // cross-backend copies break the producer chain the check follows
+    if (!hadamard_verified && gf && (!model.hadamard_rotations.empty() || !model.hadamard_inverses.empty())) {
+        llama_verify_hadamard_graph(gf, model.hadamard_rotations, model.hadamard_inverses);
+        hadamard_verified = true;
+    }
+
     this->n_outputs = save_n_outputs;
 
     // initialize scheduler with the specified graph
diff --git a/src/llama-context.h b/src/llama-context.h
index bf91daa8b56..e2eb74de3db 100644
--- a/src/llama-context.h
+++ b/src/llama-context.h
@@ -367,6 +367,9 @@ struct llama_context {
     llm_graph_result_ptr gf_res_prev;
     llm_graph_result_ptr gf_res_reserve;
 
+    // one-time Hadamard transform-coverage check on the first built graph
+    bool hadamard_verified = false;
+
     // host buffer for the model output (logits and embeddings)
     ggml_backend_buffer_ptr buf_output;
 

From 5c5a4625dbc32c11420c9b6b1ae01d9746a5899b Mon Sep 17 00:00:00 2001
From: bri-prism <288398250+bri-prism@users.noreply.github.com>
Date: Sun, 23 Aug 2026 17:42:56 -0700
Subject: [PATCH 17/19] cuda: fuse the Hadamard sign flip into the FWHT load

The sign multiply was a separate full pass over the activation (4.3% of
traced GPU time, ~1250 launches per run). Detect the mul+reshape+FWHT-hint
subgraph in the CUDA graph evaluator and multiply the sign vector during
the transform kernel's load instead. Backend-internal: the graph and every
fallback path are unchanged. test-backend-ops gains fused-pattern cases
at widths 5120/6144/17408.

(cherry picked from commit deb4dfd92a22faf588a89606bf5d53c0a3dd932b)
---
 ggml/src/ggml-cuda/fwht.cu      | 132 ++++++++++++++++----------------
 ggml/src/ggml-cuda/fwht.cuh     |   2 +
 ggml/src/ggml-cuda/ggml-cuda.cu |  23 ++++++
 tests/test-backend-ops.cpp      |  69 +++++++++++++++++
 4 files changed, 161 insertions(+), 65 deletions(-)

diff --git a/ggml/src/ggml-cuda/fwht.cu b/ggml/src/ggml-cuda/fwht.cu
index 77421a71313..844e48cbc1c 100644
--- a/ggml/src/ggml-cuda/fwht.cu
+++ b/ggml/src/ggml-cuda/fwht.cu
@@ -11,9 +11,10 @@ __device__ __forceinline__ float fwht_load<half>(const half value) {
     return __half2float(value);
 }
 
-template <int N, typename T>
+template <int N, typename T, bool has_signs>
 __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) {
+__global__ void fwht_cuda(const T * src, float * dst, const int64_t n_rows, const float scale,
+                          const float * signs, const int n_blk) {
     constexpr int warp_size = ggml_cuda_get_physical_warp_size();
 
     const int64_t r = (int64_t) blockIdx.x * blockDim.y + threadIdx.y;
@@ -30,9 +31,13 @@ __global__ void fwht_cuda(const T * src, float * dst, const int64_t n_rows, cons
     const int lane = threadIdx.x;
 
     ggml_cuda_pdl_sync();
+    const float * signs_row = has_signs ? signs + (r % n_blk) * N : nullptr;
 #pragma unroll
     for (int i = 0; i < el_w; ++i) {
         reg[i] = fwht_load(src[i * warp_size + lane]) * scale;
+        if (has_signs) {
+            reg[i] *= signs_row[i * warp_size + lane];
+        }
     }
 
 #pragma unroll
@@ -68,81 +73,78 @@ __global__ void fwht_cuda(const T * src, float * dst, const int64_t n_rows, cons
     }
 }
 
-bool ggml_cuda_op_fwht(ggml_backend_cuda_context & ctx, const ggml_tensor * src, ggml_tensor * dst) {
-    GGML_ASSERT(ggml_are_same_shape(src, dst));
+template <typename T>
+static bool fwht_launch(ggml_backend_cuda_context & ctx, const T * src_d, float * dst_d,
+                        const int n, const int64_t rows, const float scale,
+                        const float * signs, const int n_blk) {
+    const int warp_size = ggml_cuda_info().devices[ggml_cuda_get_device()].warp_size;
+    const int rows_per_block = 4;
+    const int64_t num_blocks = (rows + rows_per_block - 1) / rows_per_block;
+    cudaStream_t stream = ctx.stream();
+    dim3 grid_dims(num_blocks, 1, 1);
+    dim3 block_dims(warp_size, rows_per_block, 1);
+    const ggml_cuda_kernel_launch_params launch_params =
+        ggml_cuda_kernel_launch_params(grid_dims, block_dims, 0, stream);
+
+    switch (n) {
+#define FWHT_CASE(NN) \
+        case NN: \
+            if (signs) { \
+                ggml_cuda_kernel_launch(fwht_cuda<NN, T, true>,  launch_params, src_d, dst_d, rows, scale, signs, n_blk); \
+            } else { \
+                ggml_cuda_kernel_launch(fwht_cuda<NN, T, false>, launch_params, src_d, dst_d, rows, scale, nullptr, 1); \
+            } \
+            return true;
+        FWHT_CASE(64)
+        FWHT_CASE(128)
+        FWHT_CASE(256)
+        FWHT_CASE(512)
+        FWHT_CASE(1024)
+        FWHT_CASE(2048)
+#undef FWHT_CASE
+        default:
+            return false;
+    }
+}
+
+static bool fwht_dispatch(ggml_backend_cuda_context & ctx, const ggml_tensor * src, ggml_tensor * dst,
+                          const ggml_tensor * signs_t) {
+    GGML_ASSERT(ggml_nelements(src) == ggml_nelements(dst));
     if (!ggml_is_contiguous(src) || !ggml_is_contiguous(dst)) {
         return false;
     }
-    const int     n    = src->ne[0];
-    const int64_t rows = ggml_nrows(src);
+    const int     n    = dst->ne[0];
+    const int64_t rows = ggml_nelements(dst) / n;
 
     if ((src->type != GGML_TYPE_F32 && src->type != GGML_TYPE_F16) || dst->type != GGML_TYPE_F32) {
         return false;
     }
 
-    const void * src_d = src->data;
-    float *       dst_d = (float *) dst->data;
-
-    const int warp_size = ggml_cuda_info().devices[ggml_cuda_get_device()].warp_size;
-    const int rows_per_block = 4;
-
-    const int64_t num_blocks = (rows + rows_per_block - 1) / rows_per_block;
-
-    cudaStream_t                         stream = ctx.stream();
-    dim3                                 grid_dims(num_blocks, 1, 1);
-    dim3                                 block_dims(warp_size, rows_per_block, 1);
-    const ggml_cuda_kernel_launch_params launch_params =
-        ggml_cuda_kernel_launch_params(grid_dims, block_dims, 0, stream);
+    const float * signs = nullptr;
+    int n_blk = 1;
+    if (signs_t) {
+        if (signs_t->type != GGML_TYPE_F32 || !ggml_is_contiguous(signs_t) || signs_t->ne[0] % n != 0) {
+            return false;
+        }
+        signs = (const float *) signs_t->data;
+        n_blk = signs_t->ne[0] / n;
+    }
 
+    float * dst_d = (float *) dst->data;
     const float scale = 1 / sqrtf(n);
 
     if (src->type == GGML_TYPE_F32) {
-        const float * src_f32 = (const float *) src_d;
-        switch (n) {
-            case 64:
-                ggml_cuda_kernel_launch(fwht_cuda<64, float>, launch_params, src_f32, dst_d, rows, scale);
-                return true;
-            case 128:
-                ggml_cuda_kernel_launch(fwht_cuda<128, float>, launch_params, src_f32, dst_d, rows, scale);
-                return true;
-            case 256:
-                ggml_cuda_kernel_launch(fwht_cuda<256, float>, launch_params, src_f32, dst_d, rows, scale);
-                return true;
-            case 512:
-                ggml_cuda_kernel_launch(fwht_cuda<512, float>, launch_params, src_f32, dst_d, rows, scale);
-                return true;
-            case 1024:
-                ggml_cuda_kernel_launch(fwht_cuda<1024, float>, launch_params, src_f32, dst_d, rows, scale);
-                return true;
-            case 2048:
-                ggml_cuda_kernel_launch(fwht_cuda<2048, float>, launch_params, src_f32, dst_d, rows, scale);
-                return true;
-            default:
-                return false;
-        }
+        return fwht_launch<float>(ctx, (const float *) src->data, dst_d, n, rows, scale, signs, n_blk);
     }
+    return fwht_launch<half>(ctx, (const half *) src->data, dst_d, n, rows, scale, signs, n_blk);
+}
 
-    const half * src_f16 = (const half *) src_d;
-    switch (n) {
-        case 64:
-            ggml_cuda_kernel_launch(fwht_cuda<64, half>, launch_params, src_f16, dst_d, rows, scale);
-            return true;
-        case 128:
-            ggml_cuda_kernel_launch(fwht_cuda<128, half>, launch_params, src_f16, dst_d, rows, scale);
-            return true;
-        case 256:
-            ggml_cuda_kernel_launch(fwht_cuda<256, half>, launch_params, src_f16, dst_d, rows, scale);
-            return true;
-        case 512:
-            ggml_cuda_kernel_launch(fwht_cuda<512, half>, launch_params, src_f16, dst_d, rows, scale);
-            return true;
-        case 1024:
-            ggml_cuda_kernel_launch(fwht_cuda<1024, half>, launch_params, src_f16, dst_d, rows, scale);
-            return true;
-        case 2048:
-            ggml_cuda_kernel_launch(fwht_cuda<2048, half>, launch_params, src_f16, dst_d, rows, scale);
-            return true;
-        default:
-            return false;
-    }
+bool ggml_cuda_op_fwht(ggml_backend_cuda_context & ctx, const ggml_tensor * src, ggml_tensor * dst) {
+    GGML_ASSERT(ggml_are_same_shape(src, dst));
+    return fwht_dispatch(ctx, src, dst, nullptr);
+}
+
+bool ggml_cuda_op_fwht_signed(ggml_backend_cuda_context & ctx, const ggml_tensor * src,
+                              const ggml_tensor * signs, ggml_tensor * dst) {
+    return fwht_dispatch(ctx, src, dst, signs);
 }
diff --git a/ggml/src/ggml-cuda/fwht.cuh b/ggml/src/ggml-cuda/fwht.cuh
index cf3df94cafa..62b2f288dab 100644
--- a/ggml/src/ggml-cuda/fwht.cuh
+++ b/ggml/src/ggml-cuda/fwht.cuh
@@ -2,3 +2,5 @@
 
 // Returns whether the Fast Walsh-Hadamard transform could be used.
 bool ggml_cuda_op_fwht(ggml_backend_cuda_context & ctx, const ggml_tensor * src, ggml_tensor * dst);
+bool ggml_cuda_op_fwht_signed(ggml_backend_cuda_context & ctx, const ggml_tensor * src,
+                              const ggml_tensor * signs, ggml_tensor * dst);
diff --git a/ggml/src/ggml-cuda/ggml-cuda.cu b/ggml/src/ggml-cuda/ggml-cuda.cu
index ad3acf8ab6f..9ef2a9fcfbf 100644
--- a/ggml/src/ggml-cuda/ggml-cuda.cu
+++ b/ggml/src/ggml-cuda/ggml-cuda.cu
@@ -3381,6 +3381,29 @@ static int ggml_cuda_try_fuse(ggml_backend_cuda_context * cuda_ctx, ggml_cgraph
         }
     }
 
+    // Hadamard sign flip + reshape + FWHT-hint matmul: multiply the sign
+    // vector during the transform's load instead of a separate full pass
+    if (ggml_can_fuse_subgraph(cgraph, i, { GGML_OP_MUL, GGML_OP_RESHAPE, GGML_OP_MUL_MAT }, { i + 2 })) {
+        const ggml_tensor * mul     = cgraph->nodes[i];
+        const ggml_tensor * reshape = cgraph->nodes[i + 1];
+        ggml_tensor *       mm      = cgraph->nodes[i + 2];
+
+        const ggml_tensor * x     = mul->src[0];
+        const ggml_tensor * signs = mul->src[1];
+
+        const bool pattern_ok = ggml_get_op_params_i32(mm, 1) == GGML_HINT_SRC0_IS_HADAMARD &&
+            mm->src[1] == reshape && reshape->src[0] == mul &&
+            signs->ne[1] == 1 && signs->ne[2] == 1 && signs->ne[3] == 1 &&
+            signs->type == GGML_TYPE_F32 && mul->type == GGML_TYPE_F32 &&
+            (x->type == GGML_TYPE_F32 || x->type == GGML_TYPE_F16) &&
+            ggml_is_contiguous(x) && ggml_is_contiguous(signs) &&
+            signs->ne[0] == x->ne[0] && signs->ne[0] % mm->src[0]->ne[0] == 0;
+
+        if (pattern_ok && ggml_cuda_op_fwht_signed(*cuda_ctx, x, signs, mm)) {
+            return 2;
+        }
+    }
+
     //RoPE + view + set-rows
     if (ggml_cuda_can_fuse(cgraph, i, { GGML_OP_ROPE, GGML_OP_VIEW, GGML_OP_SET_ROWS }, {})) {
         ggml_tensor * rope     = cgraph->nodes[i];
diff --git a/tests/test-backend-ops.cpp b/tests/test-backend-ops.cpp
index 265417eeefe..d3e301025b4 100644
--- a/tests/test-backend-ops.cpp
+++ b/tests/test-backend-ops.cpp
@@ -4673,6 +4673,71 @@ struct test_mul_mat_hadamard : public test_mul_mat {
     }
 };
 
+// sign flip + reshape + FWHT-hint matmul, the fusable Hadamard activation path
+struct test_fwht_signed : public test_case {
+    const int64_t blk;
+    const int64_t width;
+    const int64_t n_tokens;
+    const ggml_type type_x;
+
+    test_fwht_signed(int64_t blk = 1024, int64_t width = 5120, int64_t n_tokens = 7,
+                     ggml_type type_x = GGML_TYPE_F32)
+        : blk(blk), width(width), n_tokens(n_tokens), type_x(type_x) {}
+
+    std::string vars() override {
+        return VARS_TO_STR4(blk, width, n_tokens, type_x);
+    }
+
+    std::string op_desc(ggml_tensor * t) override {
+        GGML_UNUSED(t);
+        return "MUL_MAT_HADAMARD";
+    }
+
+    ggml_tensor * build_graph(ggml_context * ctx) override {
+        ggml_tensor * a = ggml_new_tensor_2d(ctx, GGML_TYPE_F32, blk, blk);
+        ggml_set_name(a, "a");
+        ggml_tensor * x = ggml_new_tensor_2d(ctx, type_x, width, n_tokens);
+        ggml_set_name(x, "x");
+        ggml_tensor * s = ggml_new_tensor_1d(ctx, GGML_TYPE_F32, width);
+        ggml_set_name(s, "s");
+
+        ggml_tensor * cur = ggml_mul(ctx, x, s);
+        cur = ggml_reshape_2d(ctx, cur, blk, width / blk * n_tokens);
+        ggml_tensor * out = ggml_mul_mat(ctx, a, cur);
+        ggml_mul_mat_set_hint(out, GGML_HINT_SRC0_IS_HADAMARD);
+        ggml_set_name(out, "out");
+        return out;
+    }
+
+    void initialize_tensors(ggml_context * ctx) override {
+        for (ggml_tensor * t = ggml_get_first_tensor(ctx); t != NULL; t = ggml_get_next_tensor(ctx, t)) {
+            if (strcmp(t->name, "a") == 0) {
+                const int64_t n_cols = t->ne[0];
+                const int64_t n_rows = ggml_nrows(t);
+                std::vector<float> data(n_cols * n_rows);
+                float scale = 1.0f / sqrtf((float)n_cols);
+                for (int64_t r = 0; r < n_rows; r++) {
+                    for (int64_t i = 0; i < n_cols; i++) {
+                        int pop = 0;
+                        int64_t val = r & i;
+                        while (val) { pop += (val & 1); val >>= 1; }
+                        data[r * n_cols + i] = (pop % 2 == 0) ? scale : -scale;
+                    }
+                }
+                ggml_backend_tensor_set(t, data.data(), 0, data.size() * sizeof(float));
+            } else if (strcmp(t->name, "s") == 0) {
+                std::vector<float> data(ggml_nelements(t));
+                for (size_t i = 0; i < data.size(); i++) {
+                    data[i] = (i % 3 == 0) ? -1.0f : 1.0f;
+                }
+                ggml_backend_tensor_set(t, data.data(), 0, data.size() * sizeof(float));
+            } else if (t->type == GGML_TYPE_F32 || t->type == GGML_TYPE_F16) {
+                init_tensor_uniform(t);
+            }
+        }
+    }
+};
+
 static void init_mul_mat_id_tensors(ggml_context * ctx, int n_mats) {
     std::random_device rd;
     std::default_random_engine rng(rd());
@@ -9169,6 +9234,10 @@ 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_F16, 64, 1, 64));
     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, 2048, 1, 2048));
+    test_cases.emplace_back(new test_fwht_signed(1024, 5120, 1));
+    test_cases.emplace_back(new test_fwht_signed(1024, 5120, 32));
+    test_cases.emplace_back(new test_fwht_signed(1024, 6144, 7, GGML_TYPE_F16));
+    test_cases.emplace_back(new test_fwht_signed(1024, 17408, 3));
     test_cases.emplace_back(new test_mul_mat_hadamard(GGML_TYPE_F32, GGML_TYPE_F32, 32, 1, 32)); // too small (N<64)
 
 #if 0

From ba54da65c5fe9241171f8a25ea2d58c3c934e427 Mon Sep 17 00:00:00 2001
From: bri-prism <288398250+bri-prism@users.noreply.github.com>
Date: Sun, 23 Aug 2026 19:02:35 -0700
Subject: [PATCH 18/19] llama, convert: allowlist gaps + device-side inverse
 transforms

Two fixes surfaced by a real folded checkpoint against the new
verified-path allowlist:

- The allowlist rejected two kinds that are on the build_lora_mm path and
  covered by the graph-walk check: output.weight (the lm head) and
  blk.N.attn_gate (the linear-attention z projection). A checkpoint
  folding either failed to load or convert.
- Inverse (lookup-side) transforms inherited the buffer type of a
  host-mapped embedding table, bouncing the per-token inverse across the
  PCIe boundary: measured 3x decode loss on an RTX 4090 (27.7 -> 80.9
  tok/s once device-side) while NVLink-class parts hid it. The inverse
  rotation and signs now use the buffer type the forward rotations chose.

(cherry picked from commit 52ad701b544f7bbf4a9b0b8cc43f8b353981e740)
---
 conversion/base.py  |  3 ++-
 src/llama-model.cpp | 18 ++++++++++++++++--
 2 files changed, 18 insertions(+), 3 deletions(-)

diff --git a/conversion/base.py b/conversion/base.py
index 2bf36156faf..97f6d77b314 100644
--- a/conversion/base.py
+++ b/conversion/base.py
@@ -693,8 +693,9 @@ def add_hadamard_metadata(self) -> None:
                 "the runtime would load the GGUF without applying the activation transform"
             )
         _HADAMARD_KINDS = re.compile(
+            r"output\.weight|"
             r"blk\.\d+\.("
-            r"attn_q|attn_k|attn_v|attn_qkv|attn_output"
+            r"attn_q|attn_k|attn_v|attn_qkv|attn_gate|attn_output"
             r"|ffn_gate|ffn_up|ffn_down"
             r"|ffn_gate_exps|ffn_up_exps|ffn_down_exps|ffn_gate_up_exps"
             r"|ffn_gate_shexp|ffn_up_shexp|ffn_down_shexp"
diff --git a/src/llama-model.cpp b/src/llama-model.cpp
index 6278cfcfd44..b95a1201cf6 100644
--- a/src/llama-model.cpp
+++ b/src/llama-model.cpp
@@ -1270,12 +1270,15 @@ void llama_model_base::load_hparams(llama_model_loader & ml) {
 
         const auto is_foldable_weight = [](const std::string & name) {
             static const char * kinds[] = {
-                "attn_q", "attn_k", "attn_v", "attn_qkv", "attn_output",
+                "attn_q", "attn_k", "attn_v", "attn_qkv", "attn_gate", "attn_output",
                 "ffn_gate", "ffn_up", "ffn_down",
                 "ffn_gate_exps", "ffn_up_exps", "ffn_down_exps", "ffn_gate_up_exps",
                 "ffn_gate_shexp", "ffn_up_shexp", "ffn_down_shexp",
                 "ssm_out",
             };
+            if (name == "output.weight") {
+                return true; // the output head is built through build_lora_mm in every arch
+            }
             if (name.compare(0, 4, "blk.") != 0) {
                 return false;
             }
@@ -1944,6 +1947,12 @@ bool llama_model_base::load_tensors(llama_model_loader & ml) {
             { &hadamard_weight_blocks,  &hadamard_rotations },
             { &hadamard_inverse_blocks, &hadamard_inverses  },
         };
+        // inverse (lookup-side) transforms must not inherit a host buffer
+        // type from a CPU-mapped table: the per-token transform would then
+        // ping-pong across the PCIe boundary. Prefer the buffer type the
+        // forward rotations live on (the GPU when layers are offloaded).
+        ggml_backend_buffer_type_t preferred_buft = nullptr;
+
         for (const auto & [blocks, target] : groups)
         for (const auto & entry : *blocks) {
             const std::string & weight_name = entry.first;
@@ -1961,7 +1970,12 @@ bool llama_model_base::load_tensors(llama_model_loader & ml) {
                 throw std::runtime_error(format("prism.hadamard weight has no buffer: %s", weight_name.c_str()));
             }
 
-            const ggml_backend_buffer_type_t buft = ggml_backend_buffer_get_type(weight->buffer);
+            ggml_backend_buffer_type_t buft = ggml_backend_buffer_get_type(weight->buffer);
+            if (target == &hadamard_rotations) {
+                preferred_buft = buft;
+            } else if (preferred_buft) {
+                buft = preferred_buft;
+            }
             auto it = std::find_if(rotations.begin(), rotations.end(),
                     [block_size, buft](const hadamard_rotation & rotation) {
                         return rotation.block_size == block_size && rotation.buft == buft;

From c0e4baf2b4bf41b824f49476ab19563c41e15e6f Mon Sep 17 00:00:00 2001
From: bri-prism <288398250+bri-prism@users.noreply.github.com>
Date: Wed, 26 Aug 2026 09:48:09 -0700
Subject: [PATCH 19/19] hadamard: fix dead F16 fusion gate, contract
 validation, and grouped-order scoping

Five review findings, all real.

The signed F16 fusion was unreachable. ggml_mul builds its result with
ggml_dup_tensor, so mul->type follows x->type, and the gate required
mul->type == F32 while allowing x to be F16. The condition collapsed to F32
only, so the templated signed half kernel never ran and the F16 signed test
covered the unfused path. The gate now requires mul->type == x->type with
either supported input type.

An explicit sign contract with no widths passed validation, left the sign table
empty, and was then read as identity mode, which silently changes the model
function. Explicit mode now requires at least one width.

The converter accepted zero-width and non-block-aligned sign vectors that the
runtime rejects, so a manifest could convert into a GGUF this same branch
cannot load. It now applies the runtime's width rule.

Grouped V order was selected by tensor-kind suffix, so one folded out_proj in
the manifest put every layer's out_proj in grouped order. A partially folded
checkpoint then computed the non-folded layers with mismatched ordering. The
match is now per-tensor, comparing full names modulo wrapper prefixes.

The prepacked NVFP4 repack path bypassed that guard entirely: it column-permuted
folded out-projections on the rotation axis, which cannot be refolded, and never
set the grouped flag. It now takes the same branch as modify_tensors.

test-backend-ops -o MUL_MAT_HADAMARD: 17/17 on Metal, 3/3 backends. The CUDA
gate change is not covered by that run; it needs an NVIDIA host.
---
 conversion/base.py              |  5 +++++
 conversion/qwen.py              | 30 ++++++++++++++++++++----------
 ggml/src/ggml-cuda/ggml-cuda.cu |  4 +++-
 src/llama-model.cpp             |  5 +++++
 4 files changed, 33 insertions(+), 11 deletions(-)

diff --git a/conversion/base.py b/conversion/base.py
index 97f6d77b314..3052e8a45ed 100644
--- a/conversion/base.py
+++ b/conversion/base.py
@@ -666,6 +666,11 @@ def add_hadamard_metadata(self) -> None:
                 raise ValueError("explicit sign mode requires a signs table")
             for width_str, vec in sorted(signs.items(), key=lambda kv: int(kv[0])):
                 width = int(width_str)
+                # same width rule the runtime enforces, so a manifest that converts also loads
+                if width <= 0 or width % block_size != 0:
+                    raise ValueError(
+                        f"sign width {width} must be positive and a multiple of block size {block_size}"
+                    )
                 if len(vec) != width or any(v not in (-1, 1) for v in vec):
                     raise ValueError(f"invalid sign vector for width {width}")
                 sign_widths.append(width)
diff --git a/conversion/qwen.py b/conversion/qwen.py
index bbb4672f860..ff72c61314f 100644
--- a/conversion/qwen.py
+++ b/conversion/qwen.py
@@ -547,11 +547,17 @@ def reorder_rows(qs: Tensor, scales: Tensor, head_dim: int) -> tuple[Tensor, Ten
         elif name.endswith((".linear_attn.in_proj_a.weight", ".linear_attn.in_proj_b.weight")):
             weight, scale = reorder_rows(weight, scale, 1)
         elif name.endswith(".linear_attn.out_proj.weight"):
-            col_perm = self._reorder_v_heads(
-                torch.arange(num_v_heads * head_v_dim, dtype=torch.long).unsqueeze(0),
-                1, num_k_heads, num_v_per_k, head_v_dim,
-            ).squeeze(0)
-            weight, scale = apply_col_perm(weight, scale, col_perm)
+            if self._hadamard_folds_tensor(name):
+                # folded latent: a column permutation on the rotation axis cannot be
+                # refolded, so keep the training (grouped) order and let the runtime
+                # permute the activation instead
+                self._hadamard_gdn_v_grouped = True
+            else:
+                col_perm = self._reorder_v_heads(
+                    torch.arange(num_v_heads * head_v_dim, dtype=torch.long).unsqueeze(0),
+                    1, num_k_heads, num_v_per_k, head_v_dim,
+                ).squeeze(0)
+                weight, scale = apply_col_perm(weight, scale, col_perm)
 
         return weight, scale
 
@@ -559,6 +565,10 @@ def _repack_nvfp4(self, name: str, weight: Tensor, scale: Tensor, scale2: Tensor
         weight, scale = self._transform_nvfp4_weight(name, weight, scale)
         super()._repack_nvfp4(name, weight, scale, scale2, input_scale)
 
+    def _hadamard_folds_tensor(self, name: str) -> bool:
+        # a manifest entry and `name` may differ only by leading wrapper prefixes
+        return any(name.endswith(n) or n.endswith(name) for n in self.hadamard_folded_names())
+
     def modify_tensors(self, data_torch: Tensor, name: str, bid: int | None) -> Iterable[tuple[str, Tensor]]:
         num_k_heads = self.hparams.get("linear_num_key_heads", 0)
         num_v_heads = self.hparams.get("linear_num_value_heads", 0)
@@ -605,11 +615,11 @@ def modify_tensors(self, data_torch: Tensor, name: str, bid: int | None) -> Iter
                 data_torch = torch.cat([qk_part, v_part], dim=0)
 
             elif ".out_proj." in name:
-                # the tensor name may have had wrapper prefixes stripped by the
-                # time it reaches modify_tensors, so match manifest entries by
-                # suffix rather than exact name
-                suffix = "linear_attn.out_proj.weight"
-                if name.endswith(suffix) and any(n.endswith(suffix) for n in self.hadamard_folded_names()):
+                # wrapper prefixes may already be stripped by the time the name reaches
+                # modify_tensors, so match on the full name modulo those prefixes. Matching
+                # on the tensor-kind suffix alone would treat every layer as folded as soon
+                # as one layer is.
+                if self._hadamard_folds_tensor(name):
                     # Hadamard-folded latent: the rotation axis must keep the
                     # training (grouped) V order; the runtime permutes the
                     # activation tiled->grouped before the transform instead.
diff --git a/ggml/src/ggml-cuda/ggml-cuda.cu b/ggml/src/ggml-cuda/ggml-cuda.cu
index 9ef2a9fcfbf..9178d32fe71 100644
--- a/ggml/src/ggml-cuda/ggml-cuda.cu
+++ b/ggml/src/ggml-cuda/ggml-cuda.cu
@@ -3394,8 +3394,10 @@ static int ggml_cuda_try_fuse(ggml_backend_cuda_context * cuda_ctx, ggml_cgraph
         const bool pattern_ok = ggml_get_op_params_i32(mm, 1) == GGML_HINT_SRC0_IS_HADAMARD &&
             mm->src[1] == reshape && reshape->src[0] == mul &&
             signs->ne[1] == 1 && signs->ne[2] == 1 && signs->ne[3] == 1 &&
-            signs->type == GGML_TYPE_F32 && mul->type == GGML_TYPE_F32 &&
+            signs->type == GGML_TYPE_F32 &&
             (x->type == GGML_TYPE_F32 || x->type == GGML_TYPE_F16) &&
+            // ggml_mul keeps src0's type, so an F16 x gives an F16 mul
+            mul->type == x->type &&
             ggml_is_contiguous(x) && ggml_is_contiguous(signs) &&
             signs->ne[0] == x->ne[0] && signs->ne[0] % mm->src[0]->ne[0] == 0;
 
diff --git a/src/llama-model.cpp b/src/llama-model.cpp
index b95a1201cf6..47a10db30f2 100644
--- a/src/llama-model.cpp
+++ b/src/llama-model.cpp
@@ -1230,6 +1230,11 @@ void llama_model_base::load_hparams(llama_model_loader & ml) {
             std::vector<int32_t> sign_values;
             ml.get_arr("prism.hadamard.sign_widths", sign_widths);
             ml.get_arr("prism.hadamard.sign_values", sign_values);
+            // explicit mode with no widths would leave the sign table empty, which reads
+            // as identity later and silently changes the model function
+            if (sign_widths.empty()) {
+                throw std::runtime_error("prism.hadamard.sign_mode is explicit but sign_widths is empty");
+            }
             size_t off = 0;
             for (const int32_t width : sign_widths) {
                 if (width <= 0 || (uint32_t) width % block_size != 0 || off + width > sign_values.size()) {