From 93ae47e1674c6383fc77abbff43ddb0786d278ca Mon Sep 17 00:00:00 2001
From: Kawrakow <iwankawrakow@gmail.com>
Date: Fri, 27 Mar 2026 10:37:35 +0100
Subject: [PATCH] Fix CUDA Hadamard transfrom bug (#1526)

---
 ggml/src/ggml-cuda.cu          | 3 ++-
 ggml/src/ggml-cuda/hadamard.cu | 2 +-
 2 files changed, 3 insertions(+), 2 deletions(-)

diff --git a/ggml/src/ggml-cuda.cu b/ggml/src/ggml-cuda.cu
index a3b0a418..8dd7e76d 100644
--- a/ggml/src/ggml-cuda.cu
+++ b/ggml/src/ggml-cuda.cu
@@ -4613,7 +4613,8 @@ GGML_CALL static bool ggml_backend_cuda_supports_op(ggml_backend_t backend, cons
         case GGML_OP_ARGMAX:
             return true;
         case GGML_OP_HADAMARD:
-            return (op->ne[0] == 64 || op->ne[0] == 128 || op->ne[0] == 256) && op->type == GGML_TYPE_F32 && op->src[0]->type == GGML_TYPE_F32;
+            return (op->op_params[0] == 64 || op->op_params[0] == 128 || op->op_params[0] == 256) && op->ne[0] % op->op_params[0] == 0 &&
+                op->type == GGML_TYPE_F32 && op->src[0]->type == GGML_TYPE_F32;
         case GGML_OP_DUP:
         case GGML_OP_REPEAT:
         case GGML_OP_CONCAT:
diff --git a/ggml/src/ggml-cuda/hadamard.cu b/ggml/src/ggml-cuda/hadamard.cu
index c1d6282c..e999417f 100644
--- a/ggml/src/ggml-cuda/hadamard.cu
+++ b/ggml/src/ggml-cuda/hadamard.cu
@@ -26,7 +26,7 @@ static __global__ void hadamard_f32(const char * src, char * dst, int ne0,
     float scale = ksqrt2;
 
 #pragma unroll
-    for (int h = 2; h < nh; h <<= 2) {
+    for (int h = 2; h < nh; h <<= 1) {
         __syncthreads();
         int ii = tid/h, jj = tid%h;
         int j = 2*h*ii+jj;