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

Enable Hadamard tranform for head size of 512 (#1598)

This commit is contained in:
Kawrakow 2026-04-08 12:04:38 +02:00 committed by GitHub
parent 3de81530c5
commit fac404509c
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

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

@ -4616,8 +4616,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->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;
return (op->op_params[0] == 64 || op->op_params[0] == 128 || op->op_params[0] == 256 || op->op_params[0] == 512)
&& 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:

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

@ -50,6 +50,7 @@ static void hadamard_f32_cuda(int nh, const char * x, char * y, int ne0, int ne1
case 64: hadamard_f32< 64><<<num_blocks, 32, 0, stream>>>(x, y, ne0, nb01, nb02, nb03, nb1, nb2, nb3); break;
case 128: hadamard_f32<128><<<num_blocks, 64, 0, stream>>>(x, y, ne0, nb01, nb02, nb03, nb1, nb2, nb3); break;
case 256: hadamard_f32<256><<<num_blocks, 128, 0, stream>>>(x, y, ne0, nb01, nb02, nb03, nb1, nb2, nb3); break;
case 512: hadamard_f32<512><<<num_blocks, 256, 0, stream>>>(x, y, ne0, nb01, nb02, nb03, nb1, nb2, nb3); break;
default: GGML_ABORT("Unsupported Hadamard block size");
}
}