From 97878e53fd30580bbc3e96479691de5b78527bb0 Mon Sep 17 00:00:00 2001 From: Georgi Gerganov Date: Sat, 6 Apr 2024 17:50:21 +0300 Subject: [PATCH] sync : llama.cpp (skip) ggml-ci --- ggml-cuda/common.cuh | 6 ------ ggml-cuda/dmmv.cu | 4 ---- ggml-cuda/dmmv.cuh | 11 +++++++++++ 3 files changed, 11 insertions(+), 10 deletions(-) diff --git a/ggml-cuda/common.cuh b/ggml-cuda/common.cuh index 79e1d54..b98d7cb 100644 --- a/ggml-cuda/common.cuh +++ b/ggml-cuda/common.cuh @@ -231,12 +231,6 @@ typedef float dfloat; // dequantize float typedef float2 dfloat2; #endif //GGML_CUDA_F16 -// dmmv = dequantize_mul_mat_vec -// TODO: remove this? -#ifndef GGML_CUDA_DMMV_X -#define GGML_CUDA_DMMV_X 32 -#endif - [[noreturn]] static __device__ void no_device_code( const char * file_name, const int line, const char * function_name, const int arch, const char * arch_list) { diff --git a/ggml-cuda/dmmv.cu b/ggml-cuda/dmmv.cu index ea2950f..0b17e3c 100644 --- a/ggml-cuda/dmmv.cu +++ b/ggml-cuda/dmmv.cu @@ -2,10 +2,6 @@ #include "dequantize.cuh" #include "convert.cuh" -#ifndef GGML_CUDA_MMV_Y -#define GGML_CUDA_MMV_Y 1 -#endif - #ifndef K_QUANTS_PER_ITERATION #define K_QUANTS_PER_ITERATION 2 #else diff --git a/ggml-cuda/dmmv.cuh b/ggml-cuda/dmmv.cuh index 3802678..4c5ebd4 100644 --- a/ggml-cuda/dmmv.cuh +++ b/ggml-cuda/dmmv.cuh @@ -1,5 +1,16 @@ #include "common.cuh" +// dmmv = dequantize_mul_mat_vec + +// TODO: remove this? +#ifndef GGML_CUDA_DMMV_X +#define GGML_CUDA_DMMV_X 32 +#endif + +#ifndef GGML_CUDA_MMV_Y +#define GGML_CUDA_MMV_Y 1 +#endif + void ggml_cuda_op_dequantize_mul_mat_vec( ggml_backend_cuda_context & ctx, const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst, const char * src0_dd_i, const float * src1_ddf_i,