tqcq/llama.cpp
0
0
Fork 0
mirror of https://github.com/ggml-org/llama.cpp.git synced 2025-06-30 04:45:17 +00:00
Code Issues Packages Projects Releases Wiki Activity

add CUDA_GLU_BLOCK_SIZE [no ci]

Browse Source
This commit is contained in:
Sigbjørn Skjæret
2025-06-13 16:10:03 +02:00
committed by Akarshan
parent d9ddeb9dfd
commit cfa9c7a47a
2 changed files with 3 additions and 2 deletions
Download Patch File Download Diff File Expand all files Collapse all files

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

@ -213,8 +213,8 @@ static __global__ void unary_gated_op_kernel(const T * x, T * dst, const int k,
template <float (*op)(float), typename T> template <float (*op)(float), typename T>
static void unary_gated_cuda(const T * x, T * dst, const int k, const int n, const int o, cudaStream_t stream) { static void unary_gated_cuda(const T * x, T * dst, const int k, const int n, const int o, cudaStream_t stream) {
const int num_blocks = (k + CUDA_NEG_BLOCK_SIZE - 1) / CUDA_NEG_BLOCK_SIZE; const int num_blocks = (k + CUDA_GLU_BLOCK_SIZE - 1) / CUDA_GLU_BLOCK_SIZE;
unary_gated_op_kernel<op><<<num_blocks, CUDA_NEG_BLOCK_SIZE, 0, stream>>>(x, dst, k, n, o); unary_gated_op_kernel<op><<<num_blocks, CUDA_GLU_BLOCK_SIZE, 0, stream>>>(x, dst, k, n, o);
} }
template <float (*op)(float)> template <float (*op)(float)>

1
ggml/src/ggml-cuda/unary.cuh
View File

@ -15,6 +15,7 @@
#define CUDA_SQRT_BLOCK_SIZE 256 #define CUDA_SQRT_BLOCK_SIZE 256
#define CUDA_SIN_BLOCK_SIZE 256 #define CUDA_SIN_BLOCK_SIZE 256
#define CUDA_COS_BLOCK_SIZE 256 #define CUDA_COS_BLOCK_SIZE 256
#define CUDA_GLU_BLOCK_SIZE 256
void ggml_cuda_op_abs(ggml_backend_cuda_context & ctx, ggml_tensor * dst); void ggml_cuda_op_abs(ggml_backend_cuda_context & ctx, ggml_tensor * dst);