diff --git a/ggml/src/CMakeLists.txt b/ggml/src/CMakeLists.txt index cc42d8d0..536d7e75 100644 --- a/ggml/src/CMakeLists.txt +++ b/ggml/src/CMakeLists.txt @@ -455,8 +455,12 @@ if (GGML_HIPBLAS) list(APPEND GGML_SOURCES_ROCM "ggml-cuda.cu") file(GLOB SRCS "ggml-cuda/template-instances/fattn-wmma*.cu") list(APPEND GGML_SOURCES_ROCM ${SRCS}) + file(GLOB SRCS "ggml-cuda/template-instances/fattn-mma*.cu") + list(APPEND GGML_SOURCES_ROCM ${SRCS}) file(GLOB SRCS "ggml-cuda/template-instances/mmq*.cu") list(APPEND GGML_SOURCES_ROCM ${SRCS}) + file(GLOB SRCS "ggml-cuda/template-instances/mmvq-instance*.cu") + list(APPEND GGML_SOURCES_ROCM ${SRCS}) if (GGML_CUDA_FA_ALL_QUANTS) file(GLOB SRCS "ggml-cuda/template-instances/fattn-vec*.cu") @@ -469,14 +473,29 @@ if (GGML_HIPBLAS) list(APPEND GGML_SOURCES_ROCM ${SRCS}) file(GLOB SRCS "ggml-cuda/template-instances/fattn-vec*f16-f16.cu") list(APPEND GGML_SOURCES_ROCM ${SRCS}) + file(GLOB SRCS "ggml-cuda/template-instances/fattn-vec*q8_0-iq4_nl.cu") + list(APPEND GGML_SOURCES_ROCM ${SRCS}) + file(GLOB SRCS "ggml-cuda/template-instances/fattn-vec*iq4_nl-iq4_nl.cu") + list(APPEND GGML_SOURCES_ROCM ${SRCS}) + file(GLOB SRCS "ggml-cuda/template-instances/fattn-vec*q6_0-q5_0.cu") + list(APPEND GGML_SOURCES_ROCM ${SRCS}) + file(GLOB SRCS "ggml-cuda/template-instances/fattn-vec*q8_0-q6_0.cu") + list(APPEND GGML_SOURCES_ROCM ${SRCS}) endif() list(APPEND GGML_CDEF_PUBLIC GGML_USE_CUDA) add_compile_definitions(GGML_USE_HIPBLAS) + # the sources imported from upstream test for GGML_USE_HIP, this fork only had the older + # GGML_USE_HIPBLAS spelling. Without it mmq_id_common.cuh defines TURING_MMA_AVAILABLE, + # AMPERE_MMA_AVAILABLE, CP_ASYNC_AVAILABLE and FP16_MMA_AVAILABLE, i.e. the inline PTX paths + add_compile_definitions(GGML_USE_HIP) add_compile_definitions(GGML_CUDA_DMMV_X=${GGML_CUDA_DMMV_X}) add_compile_definitions(GGML_CUDA_MMV_Y=${GGML_CUDA_MMV_Y}) add_compile_definitions(K_QUANTS_PER_ITERATION=${GGML_CUDA_KQUANTS_ITER}) + add_compile_definitions(GGML_CUDA_PEER_MAX_BATCH_SIZE=${GGML_CUDA_PEER_MAX_BATCH_SIZE}) + add_compile_definitions(GGML_CUDA_MIN_BATCH_OFFLOAD=${GGML_CUDA_MIN_BATCH_OFFLOAD}) + add_compile_definitions(GGML_CUDA_FUSION=${GGML_CUDA_FUSION}) if (GGML_HIP_UMA) add_compile_definitions(GGML_HIP_UMA) @@ -494,6 +513,18 @@ if (GGML_HIPBLAS) add_compile_definitions(GGML_CUDA_FORCE_CUBLAS) endif() + if (GGML_CUDA_IQK_FORCE_BF16) + add_compile_definitions(GGML_CUDA_IQK_FORCE_BF16) + endif() + + if (DEFINED GGML_CUDA_DMMV_Y) + add_compile_definitions(GGML_CUDA_MMV_Y=${GGML_CUDA_DMMV_Y}) # for backwards compatibility + endif() + + if (GGML_CUDA_F16 OR GGML_CUDA_DMMV_F16) + add_compile_definitions(GGML_CUDA_F16) + endif() + if (GGML_CUDA_NO_PEER_COPY) add_compile_definitions(GGML_CUDA_NO_PEER_COPY) endif() diff --git a/ggml/src/ggml-cuda.cu b/ggml/src/ggml-cuda.cu index 78996594..9dd74fd5 100644 --- a/ggml/src/ggml-cuda.cu +++ b/ggml/src/ggml-cuda.cu @@ -307,6 +307,7 @@ const ggml_cuda_device_info & ggml_cuda_info() { /* ---------- hot-swap: invalidate all cached CUDA graphs ---------- */ extern "C" void ggml_backend_cuda_invalidate_graphs(const void * model) { +#ifdef USE_CUDA_GRAPH auto & info = const_cast(ggml_cuda_info()); if (auto it = info.all_ctx.find(model); it != info.all_ctx.end()) { for (auto ctx : it->second) { @@ -317,6 +318,9 @@ extern "C" void ggml_backend_cuda_invalidate_graphs(const void * model) { } else { fprintf(stderr, "================================= %s: did not find entry for model at %p\n", __func__, model); } +#else + GGML_UNUSED(model); +#endif // USE_CUDA_GRAPH //for (int i = 0; i < info.device_count; ++i) { // if (info.all_ctx[i]) { // info.all_ctx[i]->cuda_graphs.clear(); @@ -5122,11 +5126,7 @@ GGML_CALL static bool ggml_backend_cuda_supports_op(ggml_backend_t backend, cons case GGML_OP_LATENT_ATTN: return ggml_cuda_latent_attn_is_supported(op); case GGML_OP_FLASH_ATTN_EXT: -#if defined(GGML_USE_HIPBLAS) && defined(__HIP_PLATFORM_AMD__) - return (op->src[0]->ne[0] == 64 && op->src[1]->type == GGML_TYPE_F16) || op->src[0]->ne[0] == 128; -#else return ggml_cuda_fattn_is_supported(*cuda_ctx, op); -#endif // defined(GGML_USE_HIPBLAS) && defined(__HIP_PLATFORM_AMD__) default: return false; } diff --git a/ggml/src/ggml-cuda/argsort.cu b/ggml/src/ggml-cuda/argsort.cu index 0aaf5861..e641e972 100644 --- a/ggml/src/ggml-cuda/argsort.cu +++ b/ggml/src/ggml-cuda/argsort.cu @@ -577,6 +577,22 @@ void argsort_f32_i32_cuda_cub(ggml_cuda_pool & pool, } } } +#else + +// indexer_topk.cu calls this unconditionally, so it needs a definition when there is no CUB. +// The bitonic sort it falls back to is bounded by the shared memory size, which is the limit +// CUB is here to lift, so it asserts rather than sorts once ncols gets large +void argsort_f32_i32_cuda_cub(ggml_cuda_pool & pool, + const float * x, + int * dst, + const int ncols, + const int nrows, + ggml_sort_order order, + cudaStream_t stream) { + GGML_UNUSED(pool); + argsort_f32_T_cuda(x, dst, ncols, nrows, ncols, order, -1, 0.f, stream); +} + #endif // GGML_CUDA_USE_CUB void ggml_cuda_op_argsort(ggml_backend_cuda_context & ctx, ggml_tensor * dst) { @@ -592,13 +608,13 @@ void ggml_cuda_op_argsort(ggml_backend_cuda_context & ctx, ggml_tensor * dst) { const int64_t ncols = src0->ne[0]; const int64_t nrows = ggml_nrows(src0); + enum ggml_sort_order order = (enum ggml_sort_order) dst->op_params[0]; + #ifdef GGML_CUDA_USE_CUB const int ncols_pad = next_power_of_2(ncols); const size_t shared_mem = ncols_pad * sizeof(int); const size_t max_shared_mem = ggml_cuda_info().devices[ggml_cuda_get_device()].smpb; - enum ggml_sort_order order = (enum ggml_sort_order) dst->op_params[0]; - if (shared_mem > max_shared_mem || ncols > 1024) { ggml_cuda_pool & pool = ctx.pool(); argsort_f32_i32_cuda_cub(pool, src0_d, (int *) dst_d, ncols, nrows, order, stream); diff --git a/ggml/src/ggml-cuda/common.cuh b/ggml/src/ggml-cuda/common.cuh index 924100d1..f7e34e70 100644 --- a/ggml/src/ggml-cuda/common.cuh +++ b/ggml/src/ggml-cuda/common.cuh @@ -57,7 +57,9 @@ #define CC_RDNA1 (CC_OFFSET_AMD + 1010) #define CC_RDNA2 (CC_OFFSET_AMD + 1030) #define CC_RDNA3 (CC_OFFSET_AMD + 1100) -#define GGML_CUDA_CC_IS_NVIDIA(cc) (cc < CC_OFFSET_MTHREADS) +// CC_OFFSET_AMD is below CC_OFFSET_MTHREADS, so the AMD range has to be excluded explicitly. +// An NVIDIA cc is 100*major + 10*minor and never comes anywhere near either offset +#define GGML_CUDA_CC_IS_NVIDIA(cc) (cc < CC_OFFSET_AMD) #define GGML_CUDA_CC_IS_AMD(cc) (cc >= CC_OFFSET_AMD) #define MATRIX_ROW_PADDING 512 // last row of quant. matrices is a multiple of this to avoid out-of-bounds memory accesses @@ -393,6 +395,32 @@ static __device__ __forceinline__ int ggml_cuda_dp4a(const int a, const int b, i #endif // defined(GGML_USE_HIPBLAS) && defined(__HIP_PLATFORM_AMD__) } +#if defined(GGML_USE_HIPBLAS) && defined(__HIP_PLATFORM_AMD__) +// look up the 8 nibbles of q4 in a 16 entry table. HIP emulates __byte_perm() with an 8 byte union +// and four dynamically indexed byte loads, which end up in scratch, so the __byte_perm() version in +// get_int_from_table_16() is an order of magnitude slower here. v_perm_b32 does the same in one +// instruction, but takes one selector byte per output byte where prmt takes one nibble +static __device__ __forceinline__ int2 ggml_cuda_perm_table_16(const int & q4, const int8_t * values) { + const uint32_t * values32 = (const uint32_t *)values; + + const uint32_t q_even = q4; + const uint32_t q_odd = q4 >> 4; + + // indices 0-7 come from the lower half of the table, 8-15 from the upper half + const uint32_t v_even_low = __builtin_amdgcn_perm(values32[1], values32[0], q_even & 0x07070707); + const uint32_t v_odd_low = __builtin_amdgcn_perm(values32[1], values32[0], q_odd & 0x07070707); + const uint32_t v_even_high = __builtin_amdgcn_perm(values32[3], values32[2], q_even & 0x07070707); + const uint32_t v_odd_high = __builtin_amdgcn_perm(values32[3], values32[2], q_odd & 0x07070707); + + // bit 3 of each index picks the half, i.e. byte 0-3 or byte 4-7 of the pair above + const uint32_t mask_even = 0x03020100 | ((q_even & 0x08080808) >> 1); + const uint32_t mask_odd = 0x03020100 | ((q_odd & 0x08080808) >> 1); + + return make_int2(__builtin_amdgcn_perm(v_even_high, v_even_low, mask_even), + __builtin_amdgcn_perm(v_odd_high, v_odd_low, mask_odd)); +} +#endif // defined(GGML_USE_HIPBLAS) && defined(__HIP_PLATFORM_AMD__) + // TODO: move to ggml-common.h static constexpr __device__ int8_t kvalues_iq4nl[16] = {-127, -104, -83, -65, -49, -35, -22, -10, 1, 13, 25, 38, 53, 69, 89, 113}; diff --git a/ggml/src/ggml-cuda/dsa_attn.cu b/ggml/src/ggml-cuda/dsa_attn.cu index 57ad2ac1..0a0f7ec5 100644 --- a/ggml/src/ggml-cuda/dsa_attn.cu +++ b/ggml/src/ggml-cuda/dsa_attn.cu @@ -363,14 +363,14 @@ bool ggml_cuda_dsa_attn_ext(ggml_backend_cuda_context & ctx, ggml_tensor * dst) &alpha_32, k16.get() + v_offset, CUDA_R_16F, K->ne[0], K->ne[0]*indexer->ne[0], kq16.get(), CUDA_R_16F, indexer->ne[0], indexer->ne[0]*Q->ne[2], &beta_32, kqv32.get(), CUDA_R_32F, V->ne[0], V->ne[0]*Q->ne[2], nrows, - CUDA_R_32F, CUBLAS_GEMM_DEFAULT_TENSOR_OP)); + CUBLAS_COMPUTE_32F, CUBLAS_GEMM_DEFAULT_TENSOR_OP)); } else { CUBLAS_CHECK(cublasGemmStridedBatchedEx(ctx.cublas_handle(), CUBLAS_OP_N, CUBLAS_OP_N, V->ne[0], Q->ne[2], indexer->ne[0], &alpha_32, v16.get(), CUDA_R_16F, V->ne[0], V->ne[0]*indexer->ne[0], kq16.get(), CUDA_R_16F, indexer->ne[0], indexer->ne[0]*Q->ne[2], &beta_32, kqv32.get(), CUDA_R_32F, V->ne[0], V->ne[0]*Q->ne[2], nrows, - CUDA_R_32F, CUBLAS_GEMM_DEFAULT_TENSOR_OP)); + CUBLAS_COMPUTE_32F, CUBLAS_GEMM_DEFAULT_TENSOR_OP)); } { diff --git a/ggml/src/ggml-cuda/fattn-vec-f16.cu b/ggml/src/ggml-cuda/fattn-vec-f16.cu index d75a5601..c89a652c 100644 --- a/ggml/src/ggml-cuda/fattn-vec-f16.cu +++ b/ggml/src/ggml-cuda/fattn-vec-f16.cu @@ -106,7 +106,7 @@ bool ggml_cuda_fattn_vec_f16_is_supported([[maybe_unused]] ggml_backend_cuda_con auto K = dst->src[1]; auto V = dst->src[2]; if (K->ne[0] != V->ne[0]) { - if (K->ne[0] != 192 || V->ne[2] != 128) return false; + if (K->ne[0] != 192 || V->ne[0] != 128) return false; if (K->type != V->type) return false; return K->type == GGML_TYPE_F16 || K->type == GGML_TYPE_Q8_0; } diff --git a/ggml/src/ggml-cuda/fattn-vec-f32.cuh b/ggml/src/ggml-cuda/fattn-vec-f32.cuh index 6ed1e660..6ce30eb0 100644 --- a/ggml/src/ggml-cuda/fattn-vec-f32.cuh +++ b/ggml/src/ggml-cuda/fattn-vec-f32.cuh @@ -442,7 +442,14 @@ void ggml_cuda_flash_attn_ext_vec_f32_case(ggml_backend_cuda_context & ctx, ggml return; } +#if defined(GGML_USE_HIPBLAS) && defined(__HIP_PLATFORM_AMD__) + // both logit_softcap variants of the 8 column kernel in one module overflow the 16 bit branch + // offset of the AMDGPU backend, but only for this instance. Any block width covers any + // Q->ne[1], see the single column NVIDIA case above + constexpr int cols_per_block = (Dk == 128 && type_K == GGML_TYPE_F16 && type_V == GGML_TYPE_F16) ? 4 : 8; +#else constexpr int cols_per_block = 8; +#endif // defined(GGML_USE_HIPBLAS) && defined(__HIP_PLATFORM_AMD__) if (logit_softcap == 0.0f) { constexpr bool use_logit_softcap = false; ggml_cuda_flash_attn_ext_vec_f32_case_impl(ctx, dst); diff --git a/ggml/src/ggml-cuda/iqk_mmvq_templates.cuh b/ggml/src/ggml-cuda/iqk_mmvq_templates.cuh index 842e4065..a642bd56 100644 --- a/ggml/src/ggml-cuda/iqk_mmvq_templates.cuh +++ b/ggml/src/ggml-cuda/iqk_mmvq_templates.cuh @@ -449,7 +449,9 @@ static __device__ __forceinline__ void get_int_from_table_16_shift(const uint32_ } static __device__ __forceinline__ int2 get_int_from_table_16(const int & q4, const int8_t * values) { -#if defined(__CUDA_ARCH__) +#if defined(GGML_USE_HIPBLAS) && defined(__HIP_PLATFORM_AMD__) + return ggml_cuda_perm_table_16(q4, values); +#elif defined(__CUDA_ARCH__) uint32_t v1, v2, v3, v4, mask; const uint32_t * values32 = (const uint32_t *)values; diff --git a/ggml/src/ggml-cuda/solve_tri.cu b/ggml/src/ggml-cuda/solve_tri.cu index 60cefa08..f14e67ac 100644 --- a/ggml/src/ggml-cuda/solve_tri.cu +++ b/ggml/src/ggml-cuda/solve_tri.cu @@ -2,7 +2,6 @@ #include "ggml.h" #include "solve_tri.cuh" #include "ggml-cuda.h" -#include #include #define MAX_N_FAST 64 diff --git a/ggml/src/ggml-cuda/template-instances/mmvq-instance-iq1_bn.cu b/ggml/src/ggml-cuda/template-instances/mmvq-instance-iq1_bn.cu index 8c8184fe..07d89461 100644 --- a/ggml/src/ggml-cuda/template-instances/mmvq-instance-iq1_bn.cu +++ b/ggml/src/ggml-cuda/template-instances/mmvq-instance-iq1_bn.cu @@ -32,7 +32,7 @@ __device__ __forceinline__ void vec_dot_iq1_bn_q8_1( uint16_t v = (mult[l]*bq1->extra) & 0xff; v += v << 1; *a = v >> 8; - sumi = __dp4a(val[0], q8[4*l+0], __dp4a(val[1], q8[4*l+1], __dp4a(val[2], q8[4*l+2], __dp4a(val[3], q8[4*l+3], sumi)))); + sumi = ggml_cuda_dp4a(val[0], q8[4*l+0], ggml_cuda_dp4a(val[1], q8[4*l+1], ggml_cuda_dp4a(val[2], q8[4*l+2], ggml_cuda_dp4a(val[3], q8[4*l+3], sumi)))); } float2 d8 = __half22float2(bq8_1[iqs].ds); *result += scale * (d8.x * sumi - d8.y); diff --git a/ggml/src/ggml-cuda/template-instances/mmvq-instance-iq2_bn.cu b/ggml/src/ggml-cuda/template-instances/mmvq-instance-iq2_bn.cu index 93bc5621..16703f27 100644 --- a/ggml/src/ggml-cuda/template-instances/mmvq-instance-iq2_bn.cu +++ b/ggml/src/ggml-cuda/template-instances/mmvq-instance-iq2_bn.cu @@ -16,10 +16,10 @@ __device__ __forceinline__ void vec_dot_iq2_bn_q8_1( for (int j = 0; j < 2; ++j) { int vl = qs[j]; int vh = qs[j] >> 4; - sumi1 = __dp4a(vl & 0x03030303, q8l[j+0], sumi1); - sumi2 = __dp4a(vl & 0x0c0c0c0c, q8l[j+4], sumi2); - sumi3 = __dp4a(vh & 0x03030303, q8h[j+0], sumi3); - sumi4 = __dp4a(vh & 0x0c0c0c0c, q8h[j+4], sumi4); + sumi1 = ggml_cuda_dp4a(vl & 0x03030303, q8l[j+0], sumi1); + sumi2 = ggml_cuda_dp4a(vl & 0x0c0c0c0c, q8l[j+4], sumi2); + sumi3 = ggml_cuda_dp4a(vh & 0x03030303, q8h[j+0], sumi3); + sumi4 = ggml_cuda_dp4a(vh & 0x0c0c0c0c, q8h[j+4], sumi4); } auto d8l = __half22float2(bq8_1[0].ds); auto d8h = __half22float2(bq8_1[1].ds); diff --git a/ggml/src/ggml-cuda/vecdotq.cuh b/ggml/src/ggml-cuda/vecdotq.cuh index a565f139..491f545f 100644 --- a/ggml/src/ggml-cuda/vecdotq.cuh +++ b/ggml/src/ggml-cuda/vecdotq.cuh @@ -1127,7 +1127,9 @@ static __device__ __forceinline__ float vec_dot_iq1_m_q8_1( } static __device__ __forceinline__ int2 get_int_from_table_16(const int & q4, const int8_t * values) { -#if defined(__CUDA_ARCH__) +#if defined(GGML_USE_HIPBLAS) && defined(__HIP_PLATFORM_AMD__) + return ggml_cuda_perm_table_16(q4, values); +#elif defined(__CUDA_ARCH__) uint32_t v1, v2, v3, v4, mask; const uint32_t * values32 = (const uint32_t *)values; diff --git a/ggml/src/ggml-cuda/vendors/hip.h b/ggml/src/ggml-cuda/vendors/hip.h index ec24fcd3..7d746f65 100644 --- a/ggml/src/ggml-cuda/vendors/hip.h +++ b/ggml/src/ggml-cuda/vendors/hip.h @@ -1,5 +1,8 @@ #pragma once +// ROCm 6 and later provide __shfl_sync() and friends as templates that static_assert on the +// width of the mask, so the shims below have to replace them rather than shadow them +#define HIP_DISABLE_WARP_SYNC_BUILTINS 1 #include #include #include @@ -8,9 +11,6 @@ // for rocblas_initialize() #include "rocblas/rocblas.h" #endif // __HIP_PLATFORM_AMD__ -#define CUBLAS_COMPUTE_16F HIPBLAS_R_16F -#define CUBLAS_COMPUTE_32F HIPBLAS_R_32F -#define CUBLAS_COMPUTE_32F_FAST_16F HIPBLAS_R_32F #define CUBLAS_GEMM_DEFAULT HIPBLAS_GEMM_DEFAULT #define CUBLAS_GEMM_DEFAULT_TENSOR_OP HIPBLAS_GEMM_DEFAULT #define CUBLAS_OP_N HIPBLAS_OP_N @@ -18,21 +18,33 @@ #define CUBLAS_STATUS_SUCCESS HIPBLAS_STATUS_SUCCESS #define CUBLAS_TF32_TENSOR_OP_MATH 0 #define CUDA_R_16F HIPBLAS_R_16F +#define CUDA_R_16BF HIPBLAS_R_16B #define CUDA_R_32F HIPBLAS_R_32F +#define CUBLAS_SIDE_RIGHT HIPBLAS_SIDE_RIGHT +#define CUBLAS_FILL_MODE_UPPER HIPBLAS_FILL_MODE_UPPER +#define CUBLAS_DIAG_NON_UNIT HIPBLAS_DIAG_NON_UNIT +#define __shfl_sync(mask, var, laneMask, width) __shfl(var, laneMask, width) +#define __shfl_up_sync(mask, var, laneMask, width) __shfl_up(var, laneMask, width) #define __shfl_xor_sync(mask, var, laneMask, width) __shfl_xor(var, laneMask, width) -#define cublasComputeType_t hipblasDatatype_t //deprecated, new hipblasComputeType_t not in 5.6 +#define __all_sync(mask, var) __all(var) +#define __any_sync(mask, var) __any(var) #define cublasCreate hipblasCreate #define cublasDestroy hipblasDestroy #define cublasGemmEx hipblasGemmEx #define cublasGemmBatchedEx hipblasGemmBatchedEx #define cublasGemmStridedBatchedEx hipblasGemmStridedBatchedEx #define cublasHandle_t hipblasHandle_t +#define cublasOperation_t hipblasOperation_t +#define cublasStrsmBatched hipblasStrsmBatched +// hipBLAS has no equivalent of the cuBLAS math mode, TF32 is not a thing here +#define cublasMath_t int +#define CUBLAS_DEFAULT_MATH 0 +#define cublasGetMathMode(handle, mode) (*(mode) = CUBLAS_DEFAULT_MATH, CUBLAS_STATUS_SUCCESS) #define cublasSetMathMode(handle, mode) CUBLAS_STATUS_SUCCESS #define cublasSetStream hipblasSetStream #define cublasSgemm hipblasSgemm #define cublasSgemmStridedBatched hipblasSgemmStridedBatched #define cublasStatus_t hipblasStatus_t -#define cudaDataType_t hipblasDatatype_t //deprecated, new hipblasDatatype not in 5.6 #define cudaDeviceCanAccessPeer hipDeviceCanAccessPeer #define cudaDeviceDisablePeerAccess hipDeviceDisablePeerAccess #define cudaDeviceEnablePeerAccess hipDeviceEnablePeerAccess @@ -49,6 +61,9 @@ #define cudaEventDestroy hipEventDestroy #define cudaFree hipFree #define cudaFreeHost hipHostFree +// cudaFuncSetAttribute() is a template over the kernel type, hipFuncSetAttribute() takes a const void * +#define cudaFuncSetAttribute(func, attr, value) hipFuncSetAttribute((const void *)(func), attr, value) +#define cudaFuncAttributeMaxDynamicSharedMemorySize hipFuncAttributeMaxDynamicSharedMemorySize #define cudaGetDevice hipGetDevice #define cudaGetDeviceCount hipGetDeviceCount #define cudaGetDeviceProperties hipGetDeviceProperties @@ -72,6 +87,7 @@ #define cudaMemset hipMemset #define cudaMemsetAsync hipMemsetAsync #define cudaMemGetInfo hipMemGetInfo +#define cudaOccupancyMaxActiveBlocksPerMultiprocessor hipOccupancyMaxActiveBlocksPerMultiprocessor #define cudaOccupancyMaxPotentialBlockSize hipOccupancyMaxPotentialBlockSize #define cudaSetDevice hipSetDevice #define cudaStreamCreateWithFlags hipStreamCreateWithFlags @@ -80,7 +96,7 @@ #define cudaStreamNonBlocking hipStreamNonBlocking #define cudaStreamPerThread hipStreamPerThread #define cudaStreamSynchronize hipStreamSynchronize -#define cudaStreamWaitEvent(stream, event, flags) hipStreamWaitEvent(stream, event, flags) +#define cudaStreamWaitEvent hipStreamWaitEvent #define cudaStream_t hipStream_t #define cudaSuccess hipSuccess #define __trap() do { abort(); __builtin_unreachable(); } while(0) @@ -94,6 +110,22 @@ #define CUBLAS_STATUS_INTERNAL_ERROR HIPBLAS_STATUS_INTERNAL_ERROR #define CUBLAS_STATUS_NOT_SUPPORTED HIPBLAS_STATUS_NOT_SUPPORTED +// hipblasDatatype_t was deprecated in ROCm 6.5 and the GEMM entry points now take the +// dedicated compute and data types instead +#if HIP_VERSION >= 60500000 +#define CUBLAS_COMPUTE_16F HIPBLAS_COMPUTE_16F +#define CUBLAS_COMPUTE_32F HIPBLAS_COMPUTE_32F +#define CUBLAS_COMPUTE_32F_FAST_16F HIPBLAS_COMPUTE_32F_FAST_16F +#define cublasComputeType_t hipblasComputeType_t +#define cudaDataType_t hipDataType +#else +#define CUBLAS_COMPUTE_16F HIPBLAS_R_16F +#define CUBLAS_COMPUTE_32F HIPBLAS_R_32F +#define CUBLAS_COMPUTE_32F_FAST_16F HIPBLAS_R_32F +#define cublasComputeType_t hipblasDatatype_t +#define cudaDataType_t hipblasDatatype_t +#endif // HIP_VERSION >= 60500000 + #define __CUDA_ARCH__ 1300 #if defined(__gfx1100__) || defined(__gfx1101__) || defined(__gfx1102__) || defined(__gfx1103__) || \ @@ -114,6 +146,22 @@ #define __has_builtin(x) 0 #endif +typedef __hip_bfloat16 nv_bfloat16; +typedef __hip_bfloat162 nv_bfloat162; + +// hipBLAS spells a half as an unsigned short, so the pointers need a cast the cuBLAS call does not +static inline hipblasStatus_t ggml_cuda_hgemm_strided_batched( + hipblasHandle_t handle, hipblasOperation_t transA, hipblasOperation_t transB, int m, int n, int k, + const half * alpha, const half * AP, int lda, long long strideA, + const half * BP, int ldb, long long strideB, + const half * beta, half * CP, int ldc, long long strideC, int batchCount) { + return hipblasHgemmStridedBatched(handle, transA, transB, m, n, k, + (const hipblasHalf *)alpha, (const hipblasHalf *)AP, lda, strideA, + (const hipblasHalf *)BP, ldb, strideB, + (const hipblasHalf *)beta, (hipblasHalf *)CP, ldc, strideC, batchCount); +} +#define cublasHgemmStridedBatched ggml_cuda_hgemm_strided_batched + typedef int8_t int8x4_t __attribute__((ext_vector_type(4))); typedef uint8_t uint8x4_t __attribute__((ext_vector_type(4))); static __device__ __forceinline__ int __vsubss4(const int a, const int b) {