diff --git a/ggml/src/ggml-cuda/common.cuh b/ggml/src/ggml-cuda/common.cuh index 6b36ca90b..a46c4d3f0 100644 --- a/ggml/src/ggml-cuda/common.cuh +++ b/ggml/src/ggml-cuda/common.cuh @@ -82,6 +82,7 @@ #define GGML_CUDA_CC_RDNA2 (GGML_CUDA_CC_OFFSET_AMD + 0x1030) // RX 6000, minimum for dp4a #define GGML_CUDA_CC_RDNA3 (GGML_CUDA_CC_OFFSET_AMD + 0x1100) // RX 7000, minimum for WMMA #define GGML_CUDA_CC_RDNA3_5 (GGML_CUDA_CC_OFFSET_AMD + 0x1150) // AI 370, AI Max 395 laptops. +#define GGML_CUDA_CC_RDNA4_M (GGML_CUDA_CC_OFFSET_AMD + 0x1170) // gfx117X #define GGML_CUDA_CC_RDNA4 (GGML_CUDA_CC_OFFSET_AMD + 0x1200) // RX 9000 #define GGML_CUDA_CC_IS_AMD(cc) (cc >= GGML_CUDA_CC_OFFSET_AMD) @@ -89,8 +90,9 @@ #define GGML_CUDA_CC_IS_RDNA1(cc) (cc >= GGML_CUDA_CC_RDNA1 && cc < GGML_CUDA_CC_RDNA2) #define GGML_CUDA_CC_IS_RDNA2(cc) (cc >= GGML_CUDA_CC_RDNA2 && cc < GGML_CUDA_CC_RDNA3) #define GGML_CUDA_CC_IS_RDNA3_0(cc) (cc >= GGML_CUDA_CC_RDNA3 && cc < GGML_CUDA_CC_RDNA3_5) -#define GGML_CUDA_CC_IS_RDNA3_5(cc) (cc >= GGML_CUDA_CC_RDNA3_5 && cc < GGML_CUDA_CC_RDNA4) +#define GGML_CUDA_CC_IS_RDNA3_5(cc) (cc >= GGML_CUDA_CC_RDNA3_5 && cc < GGML_CUDA_CC_RDNA4_M) #define GGML_CUDA_CC_IS_RDNA3(cc) (GGML_CUDA_CC_IS_RDNA3_0(cc) || GGML_CUDA_CC_IS_RDNA3_5(cc)) +#define GGML_CUDA_CC_IS_RDNA4_M(cc) (cc >= GGML_CUDA_CC_RDNA4_M && cc < GGML_CUDA_CC_RDNA4) #define GGML_CUDA_CC_IS_RDNA4(cc) (cc >= GGML_CUDA_CC_RDNA4) #define GGML_CUDA_CC_IS_GCN_APU(cc) ((cc) == GGML_CUDA_CC_GFX909 || (cc) == GGML_CUDA_CC_GFX90C) #define GGML_CUDA_CC_IS_GCN(cc) ((cc > GGML_CUDA_CC_OFFSET_AMD && cc < GGML_CUDA_CC_CDNA1) || GGML_CUDA_CC_IS_GCN_APU(cc)) @@ -277,9 +279,17 @@ static const char * cu_get_error_str(CUresult err) { #define AMD_MFMA_AVAILABLE #endif // defined(GGML_USE_HIP) && defined(CDNA) && !defined(GGML_HIP_NO_MMQ_MFMA) -#if defined(GGML_USE_HIP) && (defined(RDNA4) || defined(RDNA3)) +#if defined(GGML_USE_HIP) && (defined(RDNA4) || defined(RDNA4_M)) +#define AMD_WMMA_128B_AVAILABLE +#endif // defined(GGML_USE_HIP) && (defined(RDNA4) || defined(RDNA4_M)) + +#if defined(GGML_USE_HIP) && defined(RDNA3) +#define AMD_WMMA_256B_AVAILABLE +#endif // defined(GGML_USE_HIP) && defined(RDNA3) + +#if defined(AMD_WMMA_128B_AVAILABLE) || defined(AMD_WMMA_256B_AVAILABLE) #define AMD_WMMA_AVAILABLE -#endif // defined(GGML_USE_HIP) && defined(RDNA4) +#endif // defined(AMD_WMMA_128B_AVAILABLE) || defined(AMD_WMMA_256B_AVAILABLE) // The Volta instructions are in principle available on Turing or newer but they are effectively unusable: #if !defined(GGML_USE_HIP) && __CUDA_ARCH__ == GGML_CUDA_CC_VOLTA @@ -326,7 +336,7 @@ static bool fast_fp16_hardware_available(const int cc) { // To be used for feature selection of external libraries, e.g. cuBLAS. static bool fp16_mma_hardware_available(const int cc) { return (GGML_CUDA_CC_IS_NVIDIA(cc) && cc >= GGML_CUDA_CC_VOLTA) || - GGML_CUDA_CC_IS_CDNA(cc) || GGML_CUDA_CC_IS_RDNA3(cc) || GGML_CUDA_CC_IS_RDNA4(cc) || + GGML_CUDA_CC_IS_CDNA(cc) || GGML_CUDA_CC_IS_RDNA3(cc) || GGML_CUDA_CC_IS_RDNA4_M(cc) || GGML_CUDA_CC_IS_RDNA4(cc) || (GGML_CUDA_CC_IS_MTHREADS(cc) && cc >= GGML_CUDA_CC_QY2); } @@ -355,7 +365,7 @@ static bool amd_mfma_available(const int cc) { } static bool amd_wmma_available(const int cc) { - return (GGML_CUDA_CC_IS_RDNA4(cc) || GGML_CUDA_CC_IS_RDNA3(cc)); + return GGML_CUDA_CC_IS_RDNA3(cc) || GGML_CUDA_CC_IS_RDNA4_M(cc) || GGML_CUDA_CC_IS_RDNA4(cc); } static bool volta_mma_available(const int cc) { @@ -713,7 +723,7 @@ static __device__ __forceinline__ int ggml_cuda_dp4a(const int a, const int b, i #if defined(GGML_USE_HIP) #if defined(CDNA) || defined(RDNA2) || defined(__gfx906__) c = __builtin_amdgcn_sdot4(a, b, c, false); -#elif defined(RDNA3) || defined(RDNA4) +#elif defined(RDNA3) || defined(RDNA4_M) || defined(RDNA4) c = __builtin_amdgcn_sudot4( true, a, true, b, c, false); #elif defined(RDNA1) || defined(__gfx900__) int tmp1; @@ -758,9 +768,9 @@ static __device__ __forceinline__ void ggml_cuda_mad(float & acc, const float2 v acc += v.y*u.y; } -#if defined(GGML_USE_HIP) && (defined(RDNA2) || defined(RDNA3) || defined(RDNA4) || defined(__gfx906__) || defined(CDNA)) +#if defined(GGML_USE_HIP) && (defined(RDNA2) || defined(RDNA3) || defined(RDNA4_M) || defined(RDNA4) || defined(__gfx906__) || defined(CDNA)) #define V_DOT2_F32_F16_AVAILABLE -#endif // defined(GGML_USE_HIP) && (defined(RDNA2) || defined(RDNA3) || defined(RDNA4) || defined(__gfx906__) || defined(CDNA)) +#endif // defined(GGML_USE_HIP) && (defined(RDNA2) || defined(RDNA3) || defined(RDNA4_M) || defined(RDNA4) || defined(__gfx906__) || defined(CDNA)) static __device__ __forceinline__ void ggml_cuda_mad(float & acc, const half2 v, const half2 u) { #ifdef V_DOT2_F32_F16_AVAILABLE diff --git a/ggml/src/ggml-cuda/fattn-mma-f16.cuh b/ggml/src/ggml-cuda/fattn-mma-f16.cuh index 31a273919..56decfa2b 100644 --- a/ggml/src/ggml-cuda/fattn-mma-f16.cuh +++ b/ggml/src/ggml-cuda/fattn-mma-f16.cuh @@ -1255,7 +1255,7 @@ static __device__ __forceinline__ void flash_attn_ext_f16_process_tile( T_B_KQ Q_B[(Q_in_reg ? DKQ/(2*T_B_KQ::J) : 1)]; #if defined(TURING_MMA_AVAILABLE) T_C_VKQ VKQ_C[cols_per_warp == 8 ? DV/T_C_VKQ::I : DV/(2*T_C_VKQ::J)]; -#elif defined(AMD_WMMA_AVAILABLE) && defined(RDNA3) +#elif defined(AMD_WMMA_256B_AVAILABLE) T_C_VKQ VKQ_C[DV % 32 != 0 ? DV/T_C_VKQ::J : DV/(2*T_C_VKQ::J)]; #elif defined(AMD_MFMA_AVAILABLE) T_C_VKQ VKQ_C[ DV/T_C_VKQ::J]; diff --git a/ggml/src/ggml-cuda/ggml-cuda.cu b/ggml/src/ggml-cuda/ggml-cuda.cu index 72e961d5c..fa740915d 100644 --- a/ggml/src/ggml-cuda/ggml-cuda.cu +++ b/ggml/src/ggml-cuda/ggml-cuda.cu @@ -1565,7 +1565,7 @@ static void ggml_cuda_mul_mat_cublas_impl(ggml_backend_cuda_context & ctx, const if (compute_type == GGML_TYPE_F16) { prefer_f32_output = cc == GGML_CUDA_CC_VOLTA || GGML_CUDA_CC_IS_RDNA4(cc) || GGML_CUDA_CC_IS_CDNA(cc); } else if (compute_type == GGML_TYPE_BF16) { - prefer_f32_output = !GGML_CUDA_CC_IS_RDNA3(cc) && !GGML_CUDA_CC_IS_CDNA(cc); + prefer_f32_output = !GGML_CUDA_CC_IS_RDNA3(cc) && !GGML_CUDA_CC_IS_RDNA4_M(cc) && !GGML_CUDA_CC_IS_CDNA(cc); } if (prefer_f32_output) { diff --git a/ggml/src/ggml-cuda/mma.cuh b/ggml/src/ggml-cuda/mma.cuh index c72f12131..524db9d85 100644 --- a/ggml/src/ggml-cuda/mma.cuh +++ b/ggml/src/ggml-cuda/mma.cuh @@ -76,11 +76,11 @@ namespace ggml_cuda_mma { // For the A/C matrices this means I major == row major, J major == column major. // For the B matrix this means I major == column major, J major == row major. // MIRRORED == Each data value is held exactly once per thread subgroup. - DATA_LAYOUT_I_MAJOR = 0, // Always used for Turing, Ampere, Ada Lovelace, consumer Blackwell, matrix A&B for RDNA4 and CDNA. - DATA_LAYOUT_J_MAJOR = 10, // Matrix C for CDNA and RDNA4, int and float matrix C for RDNA3. + DATA_LAYOUT_I_MAJOR = 0, // Always used for Turing, Ampere, Ada Lovelace, consumer Blackwell, matrix A&B for WMMA128 and CDNA. + DATA_LAYOUT_J_MAJOR = 10, // Matrix C for CDNA and WMMA128, int and float matrix C for WMMA256. DATA_LAYOUT_I_MAJOR_MIRRORED = 20, // Volta, matrix A&B for RDNA3. DATA_LAYOUT_J_MAJOR_MIRRORED = 30, - DATA_LAYOUT_I_MAJOR_SCRAMBLED = 40, // Scrambled matrix C for faster transposition (RDNA4/CDNA), convert to float to unscramble. + DATA_LAYOUT_I_MAJOR_SCRAMBLED = 40, // Scrambled matrix C for faster transposition (WMMA128/CDNA), convert to float to unscramble. }; // Implemented mma combinations are: // - (I_MAJOR, I_MAJOR) -> I_MAJOR @@ -88,11 +88,11 @@ namespace ggml_cuda_mma { // - (I_MAJOR, J_MAJOR_MIRRORED) -> I_MAJOR static constexpr __device__ data_layout get_input_data_layout() { -#if defined(RDNA3) || defined(VOLTA_MMA_AVAILABLE) +#if defined(AMD_WMMA_256B_AVAILABLE) || defined(VOLTA_MMA_AVAILABLE) return DATA_LAYOUT_I_MAJOR_MIRRORED; #else return DATA_LAYOUT_I_MAJOR; -#endif // defined(RDNA3) || defined(VOLTA_MMA_AVAILABLE) +#endif // defined(AMD_WMMA_256B_AVAILABLE) || defined(VOLTA_MMA_AVAILABLE) } template @@ -201,7 +201,7 @@ namespace ggml_cuda_mma { static __device__ __forceinline__ int get_j(const int l) { if constexpr (I == 16 && J == 16) { -#if defined(RDNA3) +#if defined(AMD_WMMA_256B_AVAILABLE) if constexpr (std::is_same_v || std::is_same_v) { // matrix C return 2 * l + (threadIdx.x / 16); @@ -210,11 +210,11 @@ namespace ggml_cuda_mma { return l; } #else - // matrix C is the transposed matrix A&B on RDNA4 + // matrix C is the transposed matrix A&B on WMMA128 return ne * (threadIdx.x / 16) + l; -#endif // defined(RDNA3) +#endif // defined(AMD_WMMA_256B_AVAILABLE) } else if constexpr (I == 16 && J == 8) { - // mmq input for RDNA4 + // mmq input for WMMA128 return ne * (threadIdx.x / 16) + l; } else if constexpr (I == 16 && J == 4) { return ne * (threadIdx.x / 16) + l; @@ -573,7 +573,7 @@ namespace ggml_cuda_mma { static constexpr int I = I_; static constexpr int J = J_; static constexpr data_layout dl = DATA_LAYOUT_I_MAJOR_MIRRORED; -#if defined(RDNA3) +#if defined(AMD_WMMA_256B_AVAILABLE) static constexpr int ne = tile::ne; half2 x[ne] = {{0.0f, 0.0f}}; @@ -616,7 +616,7 @@ namespace ggml_cuda_mma { return -1; } } -#endif // defined(RDNA3) +#endif // defined(AMD_WMMA_256B_AVAILABLE) }; template @@ -694,7 +694,7 @@ namespace ggml_cuda_mma { }; static __device__ __forceinline__ tile<16, 16, half2, DATA_LAYOUT_I_MAJOR> unscramble(const tile<16, 16, half2, DATA_LAYOUT_I_MAJOR_SCRAMBLED> & t) { -#if defined(AMD_MFMA_AVAILABLE) || (defined(AMD_WMMA_AVAILABLE) && defined(RDNA4)) +#if defined(AMD_MFMA_AVAILABLE) || defined(AMD_WMMA_128B_AVAILABLE) tile<16, 16, half2, DATA_LAYOUT_I_MAJOR> ret; #pragma unroll for (int l0 = 0; l0 < t.ne/2; ++l0) { @@ -705,7 +705,7 @@ namespace ggml_cuda_mma { #else NO_DEVICE_CODE; GGML_UNUSED(t); -#endif // defined(AMD_MFMA_AVAILABLE) || (defined(AMD_WMMA_AVAILABLE) && defined(RDNA4)) +#endif // defined(AMD_MFMA_AVAILABLE) || defined(AMD_WMMA_128B_AVAILABLE) } #if defined(TURING_MMA_AVAILABLE) @@ -726,7 +726,7 @@ namespace ggml_cuda_mma { return ret; } -#elif defined(AMD_WMMA_AVAILABLE) && defined(RDNA3) +#elif defined(AMD_WMMA_256B_AVAILABLE) static __device__ __forceinline__ tile<16, 8, half2, DATA_LAYOUT_I_MAJOR_MIRRORED> get_half2( const tile<16, 16, float, DATA_LAYOUT_I_MAJOR> & tile_float) { tile<16, 8, half2, DATA_LAYOUT_I_MAJOR_MIRRORED> ret; @@ -1171,7 +1171,7 @@ namespace ggml_cuda_mma { : "r"(Axi[2]), "r"(Axi[3]), "r"(Bxi[3])); #endif // __CUDA_ARCH__ >= GGML_CUDA_CC_AMPERE #elif defined(AMD_WMMA_AVAILABLE) -#if defined(RDNA4) +#if defined(AMD_WMMA_128B_AVAILABLE) using halfx8_t = __attribute__((ext_vector_type(8))) _Float16; halfx8_t& acc_frag = reinterpret_cast(D.x[0]); const halfx8_t& a_frag = reinterpret_cast(A.x[0]); @@ -1180,7 +1180,7 @@ namespace ggml_cuda_mma { #else GGML_UNUSED_VARS(D, A, B); NO_DEVICE_CODE; -#endif // defined(RDNA4) +#endif // defined(AMD_WMMA_128B_AVAILABLE) #elif defined(AMD_MFMA_AVAILABLE) // MFMA: FP16 input, FP32 accumulate, convert back to half2. using halfx4_t = __attribute__((ext_vector_type(4))) _Float16; @@ -1218,7 +1218,7 @@ namespace ggml_cuda_mma { static __device__ __forceinline__ void mma( tile<16, 16, half2, DATA_LAYOUT_I_MAJOR_SCRAMBLED> & D, const tile<32, 8, half2, DATA_LAYOUT_I_MAJOR> & A, const tile<16, 8, half2, DATA_LAYOUT_I_MAJOR> & B) { -#if defined(AMD_MFMA_AVAILABLE) || (defined(AMD_WMMA_AVAILABLE) && defined(RDNA4)) +#if defined(AMD_MFMA_AVAILABLE) || defined(AMD_WMMA_128B_AVAILABLE) tile<16, 8, half2> * D16 = (tile<16, 8, half2> *) &D; const tile<16, 8, half2> * A16 = (const tile<16, 8, half2> *) &A; mma(D16[0], A16[0], B); @@ -1226,7 +1226,7 @@ namespace ggml_cuda_mma { #else GGML_UNUSED_VARS(D, A, B); NO_DEVICE_CODE; -#endif // defined(AMD_MFMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) && defined(RDNA4) +#endif // defined(AMD_MFMA_AVAILABLE) || defined(AMD_WMMA_128B_AVAILABLE) } template @@ -1373,14 +1373,14 @@ namespace ggml_cuda_mma { : "r"(Axi[2]), "r"(Axi[3]), "r"(Bxi[3])); #endif // __CUDA_ARCH__ >= GGML_CUDA_CC_AMPERE #elif defined(AMD_WMMA_AVAILABLE) -#if defined(RDNA4) +#if defined(AMD_WMMA_128B_AVAILABLE) using halfx8_t = __attribute__((ext_vector_type(8))) _Float16; using floatx8_t = __attribute__((ext_vector_type(8))) float; floatx8_t& acc_frag = reinterpret_cast(D.x[0]); const halfx8_t& a_frag = reinterpret_cast(A.x[0]); const halfx8_t& b_frag = reinterpret_cast(B.x[0]); acc_frag = __builtin_amdgcn_wmma_f32_16x16x16_f16_w32_gfx12(a_frag, b_frag, acc_frag); -#elif defined(RDNA3) +#elif defined(AMD_WMMA_256B_AVAILABLE) using halfx16_t = __attribute__((ext_vector_type(16))) _Float16; using floatx8_t = __attribute__((ext_vector_type(8))) float; floatx8_t& acc_frag = reinterpret_cast(D.x[0]); @@ -1390,7 +1390,7 @@ namespace ggml_cuda_mma { #else GGML_UNUSED_VARS(D, A, B); NO_DEVICE_CODE; -#endif // RDNA4 +#endif // AMD_WMMA_128B_AVAILABLE #elif defined(AMD_MFMA_AVAILABLE) using halfx4_t = __attribute__((ext_vector_type(4))) _Float16; using floatx4_t = __attribute__((ext_vector_type(4))) float; @@ -1408,14 +1408,14 @@ namespace ggml_cuda_mma { static __device__ __forceinline__ void mma( tile<16, 16, float, dl_d> & D, const tile<16, 8, nv_bfloat162, dl_ab> & A, const tile<16, 8, nv_bfloat162, dl_ab> & B) { #if defined(AMD_WMMA_AVAILABLE) -#if defined(RDNA4) +#if defined(AMD_WMMA_128B_AVAILABLE) using bf16x8_t = __attribute__((ext_vector_type(8))) short; using floatx8_t = __attribute__((ext_vector_type(8))) float; floatx8_t& acc_frag = reinterpret_cast(D.x[0]); const bf16x8_t& a_frag = reinterpret_cast(A.x[0]); const bf16x8_t& b_frag = reinterpret_cast(B.x[0]); acc_frag = __builtin_amdgcn_wmma_f32_16x16x16_bf16_w32_gfx12(a_frag, b_frag, acc_frag); -#elif defined(RDNA3) +#elif defined(AMD_WMMA_256B_AVAILABLE) using bf16x16_t = __attribute__((ext_vector_type(16))) short; using floatx8_t = __attribute__((ext_vector_type(8))) float; floatx8_t& acc_frag = reinterpret_cast(D.x[0]); @@ -1425,7 +1425,7 @@ namespace ggml_cuda_mma { #else GGML_UNUSED_VARS(D, A, B); NO_DEVICE_CODE; -#endif // defined(RDNA4) +#endif // defined(AMD_WMMA_128B_AVAILABLE) #elif defined(AMD_MFMA_AVAILABLE) using floatx4_t = __attribute__((ext_vector_type(4))) float; floatx4_t& acc_frag = reinterpret_cast(D.x[0]); @@ -1467,19 +1467,19 @@ namespace ggml_cuda_mma { #elif defined(AMD_WMMA_AVAILABLE) using int32x8_t = __attribute__((__vector_size__(8 * sizeof(int)))) int; int32x8_t * acc = (int32x8_t *) D.x; -#if defined(RDNA4) +#if defined(AMD_WMMA_128B_AVAILABLE) using int32x2_t = __attribute__((__vector_size__(2 * sizeof(int)))) int; int32x2_t * a_vec = (int32x2_t *) A.x; int32x2_t * b_vec = (int32x2_t *) B.x; acc[0] = __builtin_amdgcn_wmma_i32_16x16x16_iu8_w32_gfx12(true, a_vec[0], true, b_vec[0], acc[0], true); acc[0] = __builtin_amdgcn_wmma_i32_16x16x16_iu8_w32_gfx12(true, a_vec[1], true, b_vec[1], acc[0], true); -#elif defined(RDNA3) +#elif defined(AMD_WMMA_256B_AVAILABLE) using int32x4_t = __attribute__((__vector_size__(4 * sizeof(int)))) int; int32x4_t * a_vec = (int32x4_t *) A.x; int32x4_t * b_vec = (int32x4_t *) B.x; acc[0] = __builtin_amdgcn_wmma_i32_16x16x16_iu8_w32(true, a_vec[0], true, b_vec[0], acc[0], true); acc[0] = __builtin_amdgcn_wmma_i32_16x16x16_iu8_w32(true, a_vec[1], true, b_vec[1], acc[0], true); -#endif // RDNA4 +#endif // AMD_WMMA_128B_AVAILABLE #else GGML_UNUSED_VARS(D, A, B); NO_DEVICE_CODE; @@ -1556,7 +1556,7 @@ namespace ggml_cuda_mma { static __device__ __forceinline__ void mma( tile<16, 16, half2, DATA_LAYOUT_I_MAJOR> & D, const tile<32, 8, half2, DATA_LAYOUT_I_MAJOR_MIRRORED> & A, const tile<16, 8, half2, DATA_LAYOUT_I_MAJOR_MIRRORED> & B) { -#if defined(AMD_WMMA_AVAILABLE) && defined(RDNA3) +#if defined(AMD_WMMA_256B_AVAILABLE) using halfx16_t = __attribute__((ext_vector_type(16))) _Float16; halfx16_t * xD = (halfx16_t *) D.x; const halfx16_t * xA = (const halfx16_t *) A.x; @@ -1585,17 +1585,17 @@ namespace ggml_cuda_mma { #elif defined(AMD_WMMA_AVAILABLE) using int32x8_t = __attribute__((__vector_size__(8 * sizeof(int)))) int; int32x8_t * acc = (int32x8_t *) D.x; -#if defined(RDNA4) +#if defined(AMD_WMMA_128B_AVAILABLE) using int32x2_t = __attribute__((__vector_size__(2 * sizeof(int)))) int; int32x2_t * a_vec = (int32x2_t *) A.x; int32x2_t * b_vec = (int32x2_t *) B.x; acc[0] = __builtin_amdgcn_wmma_i32_16x16x16_iu8_w32_gfx12(true, a_vec[0], true, b_vec[0], acc[0], false); -#elif defined(RDNA3) +#elif defined(AMD_WMMA_256B_AVAILABLE) using int32x4_t = __attribute__((__vector_size__(4 * sizeof(int)))) int; int32x4_t * a_vec = (int32x4_t *) A.x; int32x4_t * b_vec = (int32x4_t *) B.x; acc[0] = __builtin_amdgcn_wmma_i32_16x16x16_iu8_w32(true, a_vec[0], true, b_vec[0], acc[0], false); -#endif // RDNA4 +#endif // AMD_WMMA_128B_AVAILABLE #else GGML_UNUSED(D); GGML_UNUSED(A); diff --git a/ggml/src/ggml-cuda/mmq.cu b/ggml/src/ggml-cuda/mmq.cu index 4af4f36e8..49e0ce905 100644 --- a/ggml/src/ggml-cuda/mmq.cu +++ b/ggml/src/ggml-cuda/mmq.cu @@ -242,7 +242,7 @@ void ggml_cuda_mul_mat_q( // Each expert only sees ne12*n_expert_used/ne02 tokens on average. // On RDNA3 and RDNA4 it is faster to pick the tile size against this value instead of ne12. - if (GGML_CUDA_CC_IS_RDNA3(cc) || GGML_CUDA_CC_IS_RDNA4(cc)) { + if (GGML_CUDA_CC_IS_RDNA3(cc) || GGML_CUDA_CC_IS_RDNA4_M(cc) || GGML_CUDA_CC_IS_RDNA4(cc)) { ncols_opt = (ne12*n_expert_used + ne02 - 1) / ne02; } } @@ -470,7 +470,7 @@ bool ggml_cuda_should_use_mmq(enum ggml_type type, int cc, int64_t ne11, int64_t } if (amd_wmma_available(cc)) { - if (GGML_CUDA_CC_IS_RDNA3(cc)) { + if (GGML_CUDA_CC_IS_RDNA3(cc) || GGML_CUDA_CC_IS_RDNA4_M(cc)) { // High expert counts are almost always better on MMQ due to // the synchronization overhead in the cuBLAS/hipBLAS path: // https://github.com/ggml-org/llama.cpp/pull/18202 @@ -487,7 +487,7 @@ bool ggml_cuda_should_use_mmq(enum ggml_type type, int cc, int64_t ne11, int64_t return ne11 <= (GGML_CUDA_CC_IS_RDNA3_0(cc) ? 128 : 256); case GGML_TYPE_IQ2_XS: case GGML_TYPE_IQ2_S: - return GGML_CUDA_CC_IS_RDNA3_5(cc) || ne11 <= 128; + return GGML_CUDA_CC_IS_RDNA3_5(cc) || GGML_CUDA_CC_IS_RDNA4_M(cc) || ne11 <= 128; default: return true; } diff --git a/ggml/src/ggml-cuda/mmq.cuh b/ggml/src/ggml-cuda/mmq.cuh index 565ed1481..308968be9 100644 --- a/ggml/src/ggml-cuda/mmq.cuh +++ b/ggml/src/ggml-cuda/mmq.cuh @@ -240,6 +240,9 @@ static __host__ ggml_cuda_mmq_config ggml_cuda_mmq_get_config(const ggml_type ty if (GGML_CUDA_CC_IS_RDNA4(cc)) { return ggml_cuda_mmq_get_config_rdna4(type, J, fallback); } + if (GGML_CUDA_CC_IS_RDNA4_M(cc)) { + return ggml_cuda_mmq_get_config_rdna3_5(type, J, fallback); + } if (GGML_CUDA_CC_IS_RDNA3_5(cc)) { return ggml_cuda_mmq_get_config_rdna3_5(type, J, fallback); } @@ -272,6 +275,8 @@ static constexpr __device__ ggml_cuda_mmq_config ggml_cuda_mmq_get_config(ggml_t return ggml_cuda_mmq_get_config_cdna(type, J, fallback); #elif defined(RDNA4) return ggml_cuda_mmq_get_config_rdna4(type, J, fallback); +#elif defined(RDNA4_M) + return ggml_cuda_mmq_get_config_rdna3_5(type, J, fallback); #elif defined(RDNA3_5) return ggml_cuda_mmq_get_config_rdna3_5(type, J, fallback); #elif defined(RDNA3) diff --git a/ggml/src/ggml-cuda/mmvf.cu b/ggml/src/ggml-cuda/mmvf.cu index 752d28c4f..a62bf53e8 100644 --- a/ggml/src/ggml-cuda/mmvf.cu +++ b/ggml/src/ggml-cuda/mmvf.cu @@ -844,7 +844,7 @@ bool ggml_cuda_should_use_mmvf(enum ggml_type type, int cc, int warp_size, const return ne11 <= MMVF_MAX_BATCH_SIZE; } if (fp16_mma_hardware_available(cc)) { - if (GGML_CUDA_CC_IS_RDNA3(cc)) { + if (GGML_CUDA_CC_IS_RDNA3(cc) || GGML_CUDA_CC_IS_RDNA4_M(cc)) { return ne11 <= 3; } if (GGML_CUDA_CC_IS_RDNA4(cc)) { diff --git a/ggml/src/ggml-cuda/mmvq.cu b/ggml/src/ggml-cuda/mmvq.cu index 930f1adc4..1a9cf55cb 100644 --- a/ggml/src/ggml-cuda/mmvq.cu +++ b/ggml/src/ggml-cuda/mmvq.cu @@ -108,7 +108,7 @@ static constexpr __device__ mmvq_parameter_table_id get_device_table_id() { return MMVQ_PARAMETERS_RDNA4; #elif defined(RDNA3_0) return MMVQ_PARAMETERS_RDNA3_0; -#elif defined(RDNA2) || defined(RDNA3_5) +#elif defined(RDNA2) || defined(RDNA3_5) || defined(RDNA4_M) return MMVQ_PARAMETERS_RDNA2; #elif defined(GCN) || defined(CDNA) return MMVQ_PARAMETERS_GCN; @@ -128,7 +128,7 @@ static __host__ mmvq_parameter_table_id get_device_table_id(int cc) { if (GGML_CUDA_CC_IS_RDNA3_0(cc)) { return MMVQ_PARAMETERS_RDNA3_0; } - if (GGML_CUDA_CC_IS_RDNA2(cc) || GGML_CUDA_CC_IS_RDNA3_5(cc)) { + if (GGML_CUDA_CC_IS_RDNA2(cc) || GGML_CUDA_CC_IS_RDNA3_5(cc) || GGML_CUDA_CC_IS_RDNA4_M(cc)) { return MMVQ_PARAMETERS_RDNA2; } if (GGML_CUDA_CC_IS_GCN(cc) || GGML_CUDA_CC_IS_CDNA(cc)) { @@ -299,7 +299,7 @@ int get_mmvq_mmid_max_batch(ggml_type type, int cc) { if (GGML_CUDA_CC_IS_RDNA4(cc)) { return get_mmvq_mmid_max_batch_rdna4(type); } - if (GGML_CUDA_CC_IS_RDNA3(cc)) { + if (GGML_CUDA_CC_IS_RDNA3(cc) || GGML_CUDA_CC_IS_RDNA4_M(cc)) { return get_mmvq_mmid_max_batch_rdna3(type); } if (GGML_CUDA_CC_IS_RDNA1(cc) || GGML_CUDA_CC_IS_RDNA2(cc)) { @@ -432,7 +432,7 @@ template static constexpr __device__ int get_mmvq_mmid_max_batch_for_device() { #if defined(RDNA4) return get_mmvq_mmid_max_batch_rdna4(type); -#elif defined(RDNA3) +#elif defined(RDNA3) || defined(RDNA4_M) return get_mmvq_mmid_max_batch_rdna3(type); #elif defined(RDNA2) || defined(RDNA1) return get_mmvq_mmid_max_batch_rdna1_rdna2(type); diff --git a/ggml/src/ggml-cuda/vendors/hip.h b/ggml/src/ggml-cuda/vendors/hip.h index 61769f400..0db8e97c9 100644 --- a/ggml/src/ggml-cuda/vendors/hip.h +++ b/ggml/src/ggml-cuda/vendors/hip.h @@ -216,9 +216,13 @@ #define RDNA4 #endif // defined(__GFX12__) -#if defined(__GFX11__) +#if defined(__gfx1170__) || defined(__gfx1171__) +#define RDNA4_M +#endif // defined(__gfx1170__) || defined(__gfx1171__) + +#if defined(__GFX11__) && !defined(RDNA4_M) #define RDNA3 -#endif // defined(__GFX11__) +#endif // defined(__GFX11__) && !defined(RDNA4_M) #if defined(__gfx1150__) || defined(__gfx1151__) || defined(__gfx1152__) || defined(__gfx1153__) #define RDNA3_5 @@ -237,9 +241,9 @@ #define RDNA1 #endif // defined(__gfx1010__) || defined(__gfx1012__) -#if defined(RDNA4) || defined(RDNA3) || defined(RDNA2) || defined(RDNA1) +#if defined(RDNA4) || defined(RDNA4_M) || defined(RDNA3) || defined(RDNA2) || defined(RDNA1) #define RDNA // For the entire family -#endif // defined(RDNA4) || defined(RDNA3) || defined(RDNA2) || defined(RDNA1) +#endif // defined(RDNA4) || defined(RDNA4_M) || defined(RDNA3) || defined(RDNA2) || defined(RDNA1) #ifndef __has_builtin #define __has_builtin(x) 0