From 465569dff8b49a195450a0eb1974fd72a32fcebc Mon Sep 17 00:00:00 2001 From: Kawrakow Date: Mon, 12 May 2025 07:49:00 +0300 Subject: Faster DeepSeek FA on CUDA (#408) * New DeepSeek FlashMLA Does not work because the RoPE portion is stored at the end in our case, while in mainline it is stored at the beginning, and the FA kernel assumes that. * Rearrange MLA K cache so it first new CUDA FA implementation * constexpr and minor changes --------- Co-authored-by: Iwan Kawrakow --- ggml/src/ggml-cuda/cp-async.cuh | 10 ++++++++++ 1 file changed, 10 insertions(+) (limited to 'ggml/src/ggml-cuda/cp-async.cuh') diff --git a/ggml/src/ggml-cuda/cp-async.cuh b/ggml/src/ggml-cuda/cp-async.cuh index ecb65999..a87dc247 100644 --- a/ggml/src/ggml-cuda/cp-async.cuh +++ b/ggml/src/ggml-cuda/cp-async.cuh @@ -2,6 +2,16 @@ #include "common.cuh" +static __device__ __forceinline__ unsigned int ggml_cuda_cvta_generic_to_shared(void * generic_ptr) { +#ifdef CP_ASYNC_AVAILABLE + return __cvta_generic_to_shared(generic_ptr); +#else + GGML_UNUSED(generic_ptr); + NO_DEVICE_CODE; + return 0; +#endif // CP_ASYNC_AVAILABLE +} + // Copies data from global to shared memory, cg == cache global. // Both the src and dst pointers must be aligned to 16 bit. // Shared memory uses 32 bit addressing, the pointer is passed as unsigned int. -- cgit v1.2.3