nicolas
/
ik_llama_opt
1
0
Fork 0
Code Issues Pull Requests Packages Projects Releases Wiki Activity

Fix CUDA Hadamard transfrom bug (#1526)

This commit is contained in:
Kawrakow 2026-03-27 10:37:35 +01:00 committed by GitHub
parent 78977c0c0d
commit 93ae47e167
No known key found for this signature in database
GPG Key ID: B5690EEEBB952194
2 changed files with 3 additions and 2 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

3
ggml/src/ggml-cuda.cu
View File

@ -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:

2
ggml/src/ggml-cuda/hadamard.cu
View File

@ -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;