From bd3d9f1bad163c9daec79df1f0e1b82ef756fa6b Mon Sep 17 00:00:00 2001 From: slaren Date: Fri, 29 Mar 2024 16:01:44 +0100 Subject: [PATCH] cuda : move GGML_CUDA_DMMV constants to dmmv.cuh --- 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 79e1d5424..b98d7cbd0 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 ea2950f3d..0b17e3cb9 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 3802678ff..4c5ebd475 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,