KBaba7/llama.cpp
0
1#pragma once2 3#include "ggml.h"4#include "ggml-cuda.h"5 6#include <cstdint>7#include <memory>8 9#if defined(GGML_USE_HIP)10#define GGML_COMMON_DECL_HIP11#define GGML_COMMON_IMPL_HIP12#else13#define GGML_COMMON_DECL_CUDA14#define GGML_COMMON_IMPL_CUDA15#if defined(GGML_USE_MUSA)16#define GGML_COMMON_DECL_MUSA17#define GGML_COMMON_IMPL_MUSA18#endif19#endif20#include "ggml-common.h"21 22#include <cstdio>23#include <array>24#include <cassert>25#include <cfloat>26#include <string>27#include <vector>28 29#if defined(GGML_USE_HIP)30#include "vendors/hip.h"31#elif defined(GGML_USE_MUSA)32#include "vendors/musa.h"33#else34#include "vendors/cuda.h"35#endif // defined(GGML_USE_HIP)36 37#define STRINGIZE_IMPL(...) #__VA_ARGS__38#define STRINGIZE(...) STRINGIZE_IMPL(__VA_ARGS__)39 40#define WARP_SIZE 3241#define CUDART_HMAX 11070 // CUDA 11.7, min. ver. for which __hmax and __hmax2 are known to work (may be higher than needed)42#define CUDART_HMASK 12000 // CUDA 12.0, min. ver. for half2 -> uint mask comparisons43 44#define GGML_CUDA_CC_PASCAL 60045#define GGML_CUDA_CC_DP4A 610 // minimum compute capability for __dp4a, an intrinsic for byte-wise dot products46#define GGML_CUDA_CC_VOLTA 70047#define GGML_CUDA_CC_TURING 75048#define GGML_CUDA_CC_AMPERE 80049#define GGML_CUDA_CC_OFFSET_AMD 0x100000050 51// GCN/CNDA, wave size is 6452#define GGML_CUDA_CC_GCN4 (GGML_CUDA_CC_OFFSET_AMD + 0x803) // Tonga, Fiji, Polaris, minimum for fast fp1653#define GGML_CUDA_CC_VEGA (GGML_CUDA_CC_OFFSET_AMD + 0x900) // Vega56/64, minimum for fp16 dual issue54#define GGML_CUDA_CC_VEGA20 (GGML_CUDA_CC_OFFSET_AMD + 0x906) // MI50/Radeon VII, minimum for dp4a55#define GGML_CUDA_CC_CDNA (GGML_CUDA_CC_OFFSET_AMD + 0x908) // MI100, minimum for MFMA, acc registers56#define GGML_CUDA_CC_CDNA2 (GGML_CUDA_CC_OFFSET_AMD + 0x910) // MI210, minimum acc register renameing57#define GGML_CUDA_CC_CDNA3 (GGML_CUDA_CC_OFFSET_AMD + 0x942) // MI30058 59// RNDA removes MFMA, dp4a, xnack, acc registers, wave size is 3260#define GGML_CUDA_CC_RDNA1 (GGML_CUDA_CC_OFFSET_AMD + 0x1010) // RX 500061#define GGML_CUDA_CC_RDNA2 (GGML_CUDA_CC_OFFSET_AMD + 0x1030) // RX 6000, minimum for dp4a62#define GGML_CUDA_CC_RDNA3 (GGML_CUDA_CC_OFFSET_AMD + 0x1100) // RX 7000, minimum for WMMA63 64#define GGML_CUDA_CC_IS_RDNA(cc) (cc >= GGML_CUDA_CC_RDNA1)65#define GGML_CUDA_CC_IS_RDNA1(cc) (cc >= GGML_CUDA_CC_RDNA1 && cc < GGML_CUDA_CC_RDNA2)66#define GGML_CUDA_CC_IS_RDNA2(cc) (cc >= GGML_CUDA_CC_RDNA2 && cc < GGML_CUDA_CC_RDNA3)67#define GGML_CUDA_CC_IS_RDNA3(cc) (cc >= GGML_CUDA_CC_RDNA3)68#define GGML_CUDA_CC_IS_GCN(cc) (cc > GGML_CUDA_CC_OFFSET_AMD && cc < GGML_CUDA_CC_CDNA)69#define GGML_CUDA_CC_IS_CDNA(cc) (cc >= GGML_CUDA_CC_CDNA && cc < GGML_CUDA_CC_RDNA1)70 71#define GGML_CUDA_CC_QY1 21072#define GGML_CUDA_CC_QY2 22073 74#define MATRIX_ROW_PADDING 512 // last row of quant. matrices is a multiple of this to avoid out-of-bounds memory accesses75 76#if defined(_MSC_VER)77#pragma warning(disable: 4244 4267) // possible loss of data78#endif79 80#define GGML_CUDA_MAX_STREAMS 881 82[[noreturn]]83void ggml_cuda_error(const char * stmt, const char * func, const char * file, int line, const char * msg);84 85#define CUDA_CHECK_GEN(err, success, error_fn) \86 do { \87 auto err_ = (err); \88 if (err_ != (success)) { \89 ggml_cuda_error(#err, __func__, __FILE__, __LINE__, error_fn(err_)); \90 } \91 } while (0)92 93#define CUDA_CHECK(err) CUDA_CHECK_GEN(err, cudaSuccess, cudaGetErrorString)94 95#if CUDART_VERSION >= 12000 || defined(GGML_USE_MUSA)96 static const char * cublas_get_error_str(const cublasStatus_t err) {97 return cublasGetStatusString(err);98 }99#else100 static const char * cublas_get_error_str(const cublasStatus_t err) {101 switch (err) {102 case CUBLAS_STATUS_SUCCESS: return "CUBLAS_STATUS_SUCCESS";103 case CUBLAS_STATUS_NOT_INITIALIZED: return "CUBLAS_STATUS_NOT_INITIALIZED";104 case CUBLAS_STATUS_ALLOC_FAILED: return "CUBLAS_STATUS_ALLOC_FAILED";105 case CUBLAS_STATUS_INVALID_VALUE: return "CUBLAS_STATUS_INVALID_VALUE";106 case CUBLAS_STATUS_ARCH_MISMATCH: return "CUBLAS_STATUS_ARCH_MISMATCH";107 case CUBLAS_STATUS_MAPPING_ERROR: return "CUBLAS_STATUS_MAPPING_ERROR";108 case CUBLAS_STATUS_EXECUTION_FAILED: return "CUBLAS_STATUS_EXECUTION_FAILED";109 case CUBLAS_STATUS_INTERNAL_ERROR: return "CUBLAS_STATUS_INTERNAL_ERROR";110 case CUBLAS_STATUS_NOT_SUPPORTED: return "CUBLAS_STATUS_NOT_SUPPORTED";111 default: return "unknown error";112 }113 }114#endif // CUDART_VERSION >= 12000115 116#define CUBLAS_CHECK(err) CUDA_CHECK_GEN(err, CUBLAS_STATUS_SUCCESS, cublas_get_error_str)117 118#if !defined(GGML_USE_HIP)119static const char * cu_get_error_str(CUresult err) {120 const char * err_str;121 cuGetErrorString(err, &err_str);122 return err_str;123}124#define CU_CHECK(err) CUDA_CHECK_GEN(err, CUDA_SUCCESS, cu_get_error_str)125#endif126 127#if CUDART_VERSION >= 11100 || defined(GGML_USE_MUSA)128#define GGML_CUDA_ASSUME(x) __builtin_assume(x)129#else130#define GGML_CUDA_ASSUME(x)131#endif // CUDART_VERSION >= 11100132 133#ifdef GGML_CUDA_F16134typedef half dfloat; // dequantize float135typedef half2 dfloat2;136#else137typedef float dfloat; // dequantize float138typedef float2 dfloat2;139#endif // GGML_CUDA_F16140 141#if (!defined(GGML_USE_HIP) && !defined(GGML_CUDA_NO_VMM)) || (defined(GGML_USE_HIP) && !defined(GGML_HIP_NO_VMM))142#define GGML_USE_VMM143#endif // (!defined(GGML_USE_HIP) && !defined(GGML_CUDA_NO_VMM)) || (defined(GGML_USE_HIP) && !defined(GGML_HIP_NO_VMM))144 145#if (defined(GGML_USE_HIP) && defined(__HIP_PLATFORM_AMD__)) || __CUDA_ARCH__ >= GGML_CUDA_CC_PASCAL146#define FP16_AVAILABLE147#endif // (defined(GGML_USE_HIP) && defined(__HIP_PLATFORM_AMD__)) || __CUDA_ARCH__ >= GGML_CUDA_CC_PASCAL148 149#if defined(FP16_AVAILABLE) && __CUDA_ARCH__ != 610150#define FAST_FP16_AVAILABLE151#endif // defined(FP16_AVAILABLE) && __CUDA_ARCH__ != 610152 153#if !(defined(GGML_USE_HIP) && defined(__HIP_PLATFORM_AMD__)) && __CUDA_ARCH__ >= GGML_CUDA_CC_VOLTA154#define FP16_MMA_AVAILABLE155#endif // !(defined(GGML_USE_HIP) && defined(__HIP_PLATFORM_AMD__)) && __CUDA_ARCH__ >= GGML_CUDA_CC_VOLTA156 157#if !(defined(GGML_USE_HIP) && defined(__HIP_PLATFORM_AMD__)) && __CUDA_ARCH__ >= GGML_CUDA_CC_TURING158#define NEW_MMA_AVAILABLE159#endif // !(defined(GGML_USE_HIP) && defined(__HIP_PLATFORM_AMD__)) && __CUDA_ARCH__ >= GGML_CUDA_CC_TURING160 161#if !(defined(GGML_USE_MUSA) && __MUSA_ARCH__ <= GGML_CUDA_CC_QY1)162#define FLASH_ATTN_AVAILABLE163#endif // !(defined(GGML_USE_MUSA) && __MUSA_ARCH__ <= GGML_CUDA_CC_QY1)164 165static constexpr bool fast_fp16_available(const int cc) {166 return cc >= GGML_CUDA_CC_PASCAL && cc != 610;167}168 169// Any FP16 tensor cores are available.170static constexpr bool fp16_mma_available(const int cc) {171 return cc < GGML_CUDA_CC_OFFSET_AMD && cc >= GGML_CUDA_CC_VOLTA;172}173 174// Volta technically had FP16 tensor cores but they work very differently compared to Turing and later.175static constexpr bool new_mma_available(const int cc) {176 return cc < GGML_CUDA_CC_OFFSET_AMD && cc >= GGML_CUDA_CC_TURING;177}178 179static constexpr __device__ int ggml_cuda_get_physical_warp_size() {180#if defined(GGML_USE_HIP) && defined(__HIP_PLATFORM_AMD__)181 return __AMDGCN_WAVEFRONT_SIZE;182#else183 return 32;184#endif // defined(GGML_USE_HIP) && defined(__HIP_PLATFORM_AMD__)185}186 187[[noreturn]]188static __device__ void no_device_code(189 const char * file_name, const int line, const char * function_name, const int arch, const char * arch_list) {190 191#if defined(GGML_USE_HIP) && defined(__HIP_PLATFORM_AMD__)192 printf("%s:%d: ERROR: HIP kernel %s has no device code compatible with HIP arch %d.\n",193 file_name, line, function_name, arch);194 GGML_UNUSED(arch_list);195#else196 printf("%s:%d: ERROR: CUDA kernel %s has no device code compatible with CUDA arch %d. ggml-cuda.cu was compiled for: %s\n",197 file_name, line, function_name, arch, arch_list);198#endif // defined(GGML_USE_HIP) && defined(__HIP_PLATFORM_AMD__)199 __trap();200 201 GGML_UNUSED(no_device_code); // suppress unused function warning202}203 204#ifdef __CUDA_ARCH__205#define NO_DEVICE_CODE no_device_code(__FILE__, __LINE__, __FUNCTION__, __CUDA_ARCH__, STRINGIZE(__CUDA_ARCH_LIST__))206#else207#define NO_DEVICE_CODE //GGML_ABORT("NO_DEVICE_CODE not valid in host code.")208#endif // __CUDA_ARCH__209 210template<int width = WARP_SIZE>211static __device__ __forceinline__ int warp_reduce_sum(int x) {212#if !(defined(GGML_USE_HIP) && defined(__HIP_PLATFORM_AMD__)) && __CUDA_ARCH__ >= GGML_CUDA_CC_AMPERE213 return __reduce_add_sync(0xffffffff, x);214#else215#pragma unroll216 for (int offset = width/2; offset > 0; offset >>= 1) {217 x += __shfl_xor_sync(0xffffffff, x, offset, width);218 }219 return x;220#endif // !(defined(GGML_USE_HIP) && defined(__HIP_PLATFORM_AMD__)) && __CUDA_ARCH__ >= GGML_CUDA_CC_AMPERE221}222 223template<int width = WARP_SIZE>224static __device__ __forceinline__ float warp_reduce_sum(float x) {225#pragma unroll226 for (int offset = width/2; offset > 0; offset >>= 1) {227 x += __shfl_xor_sync(0xffffffff, x, offset, width);228 }229 return x;230}231 232template<int width = WARP_SIZE>233static __device__ __forceinline__ float2 warp_reduce_sum(float2 a) {234#pragma unroll235 for (int offset = width/2; offset > 0; offset >>= 1) {236 a.x += __shfl_xor_sync(0xffffffff, a.x, offset, width);237 a.y += __shfl_xor_sync(0xffffffff, a.y, offset, width);238 }239 return a;240}241 242template<int width = WARP_SIZE>243static __device__ __forceinline__ half2 warp_reduce_sum(half2 a) {244#ifdef FP16_AVAILABLE245#pragma unroll246 for (int offset = width/2; offset > 0; offset >>= 1) {247 a = __hadd2(a, __shfl_xor_sync(0xffffffff, a, offset, width));248 }249 return a;250 251#else252 NO_DEVICE_CODE;253 return a;254#endif // FP16_AVAILABLE255}256 257template<int width = WARP_SIZE>258static __device__ __forceinline__ float warp_reduce_max(float x) {259#pragma unroll260 for (int offset = width/2; offset > 0; offset >>= 1) {261 x = fmaxf(x, __shfl_xor_sync(0xffffffff, x, offset, width));262 }263 return x;264}265 266static __device__ __forceinline__ half ggml_cuda_hmax(const half a, const half b) {267#ifdef FP16_AVAILABLE268 269#if !(defined(GGML_USE_HIP) && defined(__HIP_PLATFORM_AMD__)) && CUDART_VERSION < CUDART_HMAX270 return __float2half(fmaxf(__half2float(a), __half2float(b)));271#else272 return __hmax(a, b);273#endif // !(defined(GGML_USE_HIP) && defined(__HIP_PLATFORM_AMD__)) && CUDART_VERSION < CUDART_HMAX274 275#else276 NO_DEVICE_CODE;277 GGML_UNUSED(b);278 return a;279#endif // FP16_AVAILABLE280}281 282static __device__ __forceinline__ half2 ggml_cuda_hmax2(const half2 a, const half2 b) {283#if defined(GGML_USE_HIP) && HIP_VERSION >= 50700000284 return half2(__hmax(a.x, b.x), __hmax(a.y, b.y));285#elif !defined(GGML_USE_HIP) && CUDART_VERSION >= CUDART_HMAX286 return __hmax2(a, b);287#elif !defined(GGML_USE_HIP)288 half2 ret;289 reinterpret_cast<half&>(ret.x) = __float2half(fmaxf( __low2float(a), __low2float(b)));290 reinterpret_cast<half&>(ret.y) = __float2half(fmaxf(__high2float(a), __high2float(b)));291 return ret;292#else293 GGML_UNUSED(a);294 GGML_UNUSED(b);295 NO_DEVICE_CODE;296#endif297}298 299template<int width = WARP_SIZE>300static __device__ __forceinline__ half2 warp_reduce_max(half2 x) {301#if !(defined(GGML_USE_HIP) && defined(__HIP_PLATFORM_AMD__)) && __CUDA_ARCH__ >= GGML_CUDA_CC_PASCAL || (defined(GGML_USE_HIP) && HIP_VERSION >= 50700000)302#pragma unroll303 for (int offset = width/2; offset > 0; offset >>= 1) {304 x = ggml_cuda_hmax2(x, __shfl_xor_sync(0xffffffff, x, offset, width));305 }306 return x;307#else308 GGML_UNUSED(x);309 NO_DEVICE_CODE;310#endif // !(defined(GGML_USE_HIP) && defined(__HIP_PLATFORM_AMD__)) && __CUDA_ARCH__ >= GGML_CUDA_CC_PASCAL || (defined(GGML_USE_HIP) && HIP_VERSION >= 50700000)311}312 313#if CUDART_VERSION < CUDART_HMASK314static __device__ __forceinline__ uint32_t __hgt2_mask(const half2 a, const half2 b) {315 const uint32_t mask_low = 0x0000FFFF * (float( __low2half(a)) > float( __low2half(b)));316 const uint32_t mask_high = 0xFFFF0000 * (float(__high2half(a)) > float(__high2half(b)));317 return mask_low | mask_high;318}319#endif // CUDART_VERSION < CUDART_HMASK320 321static __device__ __forceinline__ int ggml_cuda_dp4a(const int a, const int b, int c) {322#if defined(GGML_USE_HIP) && defined(__HIP_PLATFORM_AMD__)323#if defined(__gfx906__) || defined(__gfx908__) || defined(__gfx90a__) || defined(RDNA2)324 c = __builtin_amdgcn_sdot4(a, b, c, false);325#elif defined(RDNA3)326 c = __builtin_amdgcn_sudot4( true, a, true, b, c, false);327#elif defined(__gfx1010__) || defined(__gfx900__)328 int tmp1;329 int tmp2;330 asm("\n \331 v_mul_i32_i24 %1, sext(%3), sext(%4) dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:BYTE_0 src1_sel:BYTE_0 \n \332 v_mul_i32_i24 %2, sext(%3), sext(%4) dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:BYTE_1 src1_sel:BYTE_1 \n \333 v_add3_u32 %0, %1, %2, %0 \n \334 v_mul_i32_i24 %1, sext(%3), sext(%4) dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:BYTE_2 src1_sel:BYTE_2 \n \335 v_mul_i32_i24 %2, sext(%3), sext(%4) dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:BYTE_3 src1_sel:BYTE_3 \n \336 v_add3_u32 %0, %1, %2, %0 \n \337 "338 : "+v"(c), "=&v"(tmp1), "=&v"(tmp2)339 : "v"(a), "v"(b)340 );341#else342 const int8x4_t va = reinterpret_cast<const int8x4_t&>(a);343 const int8x4_t vb = reinterpret_cast<const int8x4_t&>(b);344 c += va[0] * vb[0] + va[1] * vb[1] + va[2] * vb[2] + va[3] * vb[3];345#endif346 return c;347 348#else // defined(GGML_USE_HIP) && defined(__HIP_PLATFORM_AMD__)349 350#if __CUDA_ARCH__ >= GGML_CUDA_CC_DP4A351 return __dp4a(a, b, c);352#else // __CUDA_ARCH__ >= GGML_CUDA_CC_DP4A353 const int8_t * a8 = (const int8_t *) &a;354 const int8_t * b8 = (const int8_t *) &b;355 return c + a8[0]*b8[0] + a8[1]*b8[1] + a8[2]*b8[2] + a8[3]*b8[3];356#endif // __CUDA_ARCH__ >= GGML_CUDA_CC_DP4A357 358#endif // defined(GGML_USE_HIP) && defined(__HIP_PLATFORM_AMD__)359}360 361// TODO: move to ggml-common.h362static constexpr __device__ int8_t kvalues_iq4nl[16] = {-127, -104, -83, -65, -49, -35, -22, -10, 1, 13, 25, 38, 53, 69, 89, 113};363 364typedef void (*dequantize_kernel_t)(const void * vx, const int64_t ib, const int iqs, dfloat2 & v);365 366static __device__ __forceinline__ float get_alibi_slope(367 const float max_bias, const uint32_t h, const uint32_t n_head_log2, const float m0, const float m1368) {369 if (max_bias <= 0.0f) {370 return 1.0f;371 }372 const float base = h < n_head_log2 ? m0 : m1;373 const int exph = h < n_head_log2 ? h + 1 : 2*(h - n_head_log2) + 1;374 375 return powf(base, exph);376}377 378template <ggml_type type>379struct ggml_cuda_type_traits;380 381template<>382struct ggml_cuda_type_traits<GGML_TYPE_F16> {383 static constexpr int qk = 1;384 static constexpr int qr = 1;385};386 387template<>388struct ggml_cuda_type_traits<GGML_TYPE_Q4_0> {389 static constexpr int qk = QK4_0;390 static constexpr int qr = QR4_0;391 static constexpr int qi = QI4_0;392};393 394template<>395struct ggml_cuda_type_traits<GGML_TYPE_Q4_1> {396 static constexpr int qk = QK4_1;397 static constexpr int qr = QR4_1;398 static constexpr int qi = QI4_1;399};400 401template<>402struct ggml_cuda_type_traits<GGML_TYPE_Q5_0> {403 static constexpr int qk = QK5_0;404 static constexpr int qr = QR5_0;405 static constexpr int qi = QI5_0;406};407 408template<>409struct ggml_cuda_type_traits<GGML_TYPE_Q5_1> {410 static constexpr int qk = QK5_1;411 static constexpr int qr = QR5_1;412 static constexpr int qi = QI5_1;413};414 415template<>416struct ggml_cuda_type_traits<GGML_TYPE_Q8_0> {417 static constexpr int qk = QK8_0;418 static constexpr int qr = QR8_0;419 static constexpr int qi = QI8_0;420};421 422template<>423struct ggml_cuda_type_traits<GGML_TYPE_Q2_K> {424 static constexpr int qk = QK_K;425 static constexpr int qr = QR2_K;426 static constexpr int qi = QI2_K;427};428 429template<>430struct ggml_cuda_type_traits<GGML_TYPE_Q3_K> {431 static constexpr int qk = QK_K;432 static constexpr int qr = QR3_K;433 static constexpr int qi = QI3_K;434};435 436template<>437struct ggml_cuda_type_traits<GGML_TYPE_Q4_K> {438 static constexpr int qk = QK_K;439 static constexpr int qr = QR4_K;440 static constexpr int qi = QI4_K;441};442 443template<>444struct ggml_cuda_type_traits<GGML_TYPE_Q5_K> {445 static constexpr int qk = QK_K;446 static constexpr int qr = QR5_K;447 static constexpr int qi = QI5_K;448};449 450template<>451struct ggml_cuda_type_traits<GGML_TYPE_Q6_K> {452 static constexpr int qk = QK_K;453 static constexpr int qr = QR6_K;454 static constexpr int qi = QI6_K;455};456 457template<>458struct ggml_cuda_type_traits<GGML_TYPE_IQ2_XXS> {459 static constexpr int qk = QK_K;460 static constexpr int qr = QR2_XXS;461 static constexpr int qi = QI2_XXS;462};463 464template<>465struct ggml_cuda_type_traits<GGML_TYPE_IQ2_XS> {466 static constexpr int qk = QK_K;467 static constexpr int qr = QR2_XS;468 static constexpr int qi = QI2_XS;469};470 471template<>472struct ggml_cuda_type_traits<GGML_TYPE_IQ2_S> {473 static constexpr int qk = QK_K;474 static constexpr int qr = QR2_S;475 static constexpr int qi = QI2_S;476};477 478template<>479struct ggml_cuda_type_traits<GGML_TYPE_IQ3_XXS> {480 static constexpr int qk = QK_K;481 static constexpr int qr = QR3_XXS;482 static constexpr int qi = QI3_XXS;483};484 485template<>486struct ggml_cuda_type_traits<GGML_TYPE_IQ1_S> {487 static constexpr int qk = QK_K;488 static constexpr int qr = QR1_S;489 static constexpr int qi = QI1_S;490};491 492template<>493struct ggml_cuda_type_traits<GGML_TYPE_IQ1_M> {494 static constexpr int qk = QK_K;495 static constexpr int qr = QR1_M;496 static constexpr int qi = QI1_M;497};498 499template<>500struct ggml_cuda_type_traits<GGML_TYPE_IQ4_NL> {501 static constexpr int qk = QK4_NL;502 static constexpr int qr = QR4_NL;503 static constexpr int qi = QI4_NL;504};505 506template<>507struct ggml_cuda_type_traits<GGML_TYPE_IQ4_XS> {508 static constexpr int qk = QK_K;509 static constexpr int qr = QR4_XS;510 static constexpr int qi = QI4_XS;511};512 513template<>514struct ggml_cuda_type_traits<GGML_TYPE_IQ3_S> {515 static constexpr int qk = QK_K;516 static constexpr int qr = QR3_S;517 static constexpr int qi = QI3_S;518};519 520//////////////////////521 522struct ggml_cuda_device_info {523 int device_count;524 525 struct cuda_device_info {526 int cc; // compute capability527 int nsm; // number of streaming multiprocessors528 size_t smpb; // max. shared memory per block529 size_t smpbo; // max. shared memory per block (with opt-in)530 bool vmm; // virtual memory support531 size_t vmm_granularity; // granularity of virtual memory532 size_t total_vram;533 int warp_size; // Number of threads in a dispatch534 };535 536 cuda_device_info devices[GGML_CUDA_MAX_DEVICES] = {};537 538 std::array<float, GGML_CUDA_MAX_DEVICES> default_tensor_split = {};539};540 541const ggml_cuda_device_info & ggml_cuda_info();542 543void ggml_cuda_set_device(int device);544int ggml_cuda_get_device();545 546struct ggml_cuda_pool {547 virtual ~ggml_cuda_pool() = default;548 549 virtual void * alloc(size_t size, size_t * actual_size) = 0;550 virtual void free(void * ptr, size_t size) = 0;551};552 553template<typename T>554struct ggml_cuda_pool_alloc {555 ggml_cuda_pool * pool = nullptr;556 T * ptr = nullptr;557 size_t actual_size = 0;558 559 ggml_cuda_pool_alloc() = default;560 561 explicit ggml_cuda_pool_alloc(ggml_cuda_pool & pool) : pool(&pool) {562 }563 564 ggml_cuda_pool_alloc(ggml_cuda_pool & pool, size_t size) : pool(&pool) {565 alloc(size);566 }567 568 ~ggml_cuda_pool_alloc() {569 if (ptr != nullptr) {570 pool->free(ptr, actual_size);571 }572 }573 574 // size is in number of elements575 T * alloc(size_t size) {576 GGML_ASSERT(pool != nullptr);577 GGML_ASSERT(ptr == nullptr);578 ptr = (T *) pool->alloc(size * sizeof(T), &this->actual_size);579 return ptr;580 }581 582 T * alloc(ggml_cuda_pool & pool, size_t size) {583 this->pool = &pool;584 return alloc(size);585 }586 587 T * get() {588 return ptr;589 }590 591 ggml_cuda_pool_alloc(const ggml_cuda_pool_alloc &) = delete;592 ggml_cuda_pool_alloc(ggml_cuda_pool_alloc &&) = delete;593 ggml_cuda_pool_alloc& operator=(const ggml_cuda_pool_alloc &) = delete;594 ggml_cuda_pool_alloc& operator=(ggml_cuda_pool_alloc &&) = delete;595};596 597 598// backend interface599 600struct ggml_tensor_extra_gpu {601 void * data_device[GGML_CUDA_MAX_DEVICES]; // 1 pointer for each device for split tensors602 cudaEvent_t events[GGML_CUDA_MAX_DEVICES][GGML_CUDA_MAX_STREAMS]; // events for synchronizing multiple GPUs603};604 605 606#if ((CUDART_VERSION >= 12000) && defined(GGML_CUDA_USE_GRAPHS)) || defined(GGML_HIP_GRAPHS)607#define USE_CUDA_GRAPH608#endif609 610struct ggml_graph_node_properties {611 void * node_address;612 ggml_op node_op;613 int64_t ne[GGML_MAX_DIMS];614 size_t nb[GGML_MAX_DIMS];615 void * src_address[GGML_MAX_SRC];616 int32_t op_params[GGML_MAX_OP_PARAMS / sizeof(int32_t)];617};618 619struct ggml_cuda_graph {620#ifdef USE_CUDA_GRAPH621 ~ggml_cuda_graph() {622 if (instance != nullptr) {623 CUDA_CHECK(cudaGraphExecDestroy(instance));624 }625 if (graph != nullptr) {626 CUDA_CHECK(cudaGraphDestroy(graph));627 }628 }629 cudaGraph_t graph = nullptr;630 cudaGraphExec_t instance = nullptr;631 size_t num_nodes = 0;632 std::vector<cudaGraphNode_t> nodes;633 std::vector<cudaKernelNodeParams> params;634 bool disable_due_to_gpu_arch = false;635 bool disable_due_to_too_many_updates = false;636 bool disable_due_to_failed_graph_capture = false;637 int number_consecutive_updates = 0;638 std::vector<ggml_graph_node_properties> ggml_graph_properties;639 std::vector<char **> updated_kernel_arg;640#endif641};642 643struct ggml_backend_cuda_context {644 int device;645 std::string name;646 cudaEvent_t copy_event = nullptr;647 648 cudaStream_t streams[GGML_CUDA_MAX_DEVICES][GGML_CUDA_MAX_STREAMS] = { { nullptr } };649 cublasHandle_t cublas_handles[GGML_CUDA_MAX_DEVICES] = {nullptr};650 651 std::unique_ptr<ggml_cuda_graph> cuda_graph;652 653 explicit ggml_backend_cuda_context(int device) :654 device(device),655 name(GGML_CUDA_NAME + std::to_string(device)) {656 }657 658 ~ggml_backend_cuda_context() {659 if (copy_event != nullptr) {660 CUDA_CHECK(cudaEventDestroy(copy_event));661 }662 for (int i = 0; i < GGML_CUDA_MAX_DEVICES; ++i) {663 for (int j = 0; j < GGML_CUDA_MAX_STREAMS; ++j) {664 if (streams[i][j] != nullptr) {665 CUDA_CHECK(cudaStreamDestroy(streams[i][j]));666 }667 }668 if (cublas_handles[i] != nullptr) {669 CUBLAS_CHECK(cublasDestroy(cublas_handles[i]));670 }671 }672 }673 674 cudaStream_t stream(int device, int stream) {675 if (streams[device][stream] == nullptr) {676 ggml_cuda_set_device(device);677 CUDA_CHECK(cudaStreamCreateWithFlags(&streams[device][stream], cudaStreamNonBlocking));678 }679 return streams[device][stream];680 }681 682 cudaStream_t stream() {683 return stream(device, 0);684 }685 686 cublasHandle_t cublas_handle(int device) {687 if (cublas_handles[device] == nullptr) {688 ggml_cuda_set_device(device);689 CUBLAS_CHECK(cublasCreate(&cublas_handles[device]));690 CUBLAS_CHECK(cublasSetMathMode(cublas_handles[device], CUBLAS_TF32_TENSOR_OP_MATH));691 }692 return cublas_handles[device];693 }694 695 cublasHandle_t cublas_handle() {696 return cublas_handle(device);697 }698 699 // pool700 std::unique_ptr<ggml_cuda_pool> pools[GGML_CUDA_MAX_DEVICES];701 702 static std::unique_ptr<ggml_cuda_pool> new_pool_for_device(int device);703 704 ggml_cuda_pool & pool(int device) {705 if (pools[device] == nullptr) {706 pools[device] = new_pool_for_device(device);707 }708 return *pools[device];709 }710 711 ggml_cuda_pool & pool() {712 return pool(device);713 }714};715