--- a/ggml/src/ggml-cuda/vendors/hip.h 2026-06-29 10:58:51.000000000 +0200 +++ b/ggml/src/ggml-cuda/vendors/hip.h 2026-07-08 11:10:05.428269296 +0200 @@ -41,13 +41,80 @@ #define cublasStrsmBatched hipblasStrsmBatched #define cublasCreate hipblasCreate #define cublasDestroy hipblasDestroy -#define cublasGemmEx hipblasGemmEx +#define cublasGemmEx ggml_hip_gemm_ex #define cublasGemmBatchedEx hipblasGemmBatchedEx #define cublasGemmStridedBatchedEx hipblasGemmStridedBatchedEx #define cublasHandle_t hipblasHandle_t #define cublasSetMathMode(handle, mode) CUBLAS_STATUS_SUCCESS #define cublasSetStream hipblasSetStream -#define cublasSgemm hipblasSgemm + + + +// Workaround: hipBLAS GEMM functions crash on gfx1201 (RDNA4) with non-default +// CUDA streams. Route sgemm and GemmEx through rocBLAS directly which works. +// Uses a thread-local cached rocBLAS handle to avoid per-call create/destroy. +#include + +static inline rocblas_handle ggml_hip_rocblas_handle(hipblasHandle_t hb) { + thread_local rocblas_handle cached = nullptr; + if (!cached) { + rocblas_create_handle(&cached); + } + // Sync stream from hipBLAS handle to rocBLAS handle + hipStream_t stream = nullptr; + hipblasGetStream(hb, &stream); + rocblas_set_stream(cached, stream); + return cached; +} + +static inline hipblasStatus_t ggml_hip_sgemm( + hipblasHandle_t handle, hipblasOperation_t ta, hipblasOperation_t tb, + int m, int n, int k, const float* alpha, + const float* A, int lda, const float* B, int ldb, + const float* beta, float* C, int ldc) { + rocblas_handle rh = ggml_hip_rocblas_handle(handle); + auto ra = (ta == HIPBLAS_OP_T) ? rocblas_operation_transpose : rocblas_operation_none; + auto rb = (tb == HIPBLAS_OP_T) ? rocblas_operation_transpose : rocblas_operation_none; + auto rs = rocblas_sgemm(rh, ra, rb, m, n, k, alpha, A, lda, B, ldb, beta, C, ldc); + return (rs == rocblas_status_success) ? HIPBLAS_STATUS_SUCCESS : HIPBLAS_STATUS_EXECUTION_FAILED; +} + +static inline rocblas_datatype ggml_hip_to_rocblas_type(hipDataType t) { + switch (t) { + case HIP_R_32F: return rocblas_datatype_f32_r; + case HIP_R_16F: return rocblas_datatype_f16_r; + case HIP_R_16BF: return rocblas_datatype_bf16_r; + default: return rocblas_datatype_f32_r; + } +} +static inline rocblas_datatype ggml_hip_compute_to_rocblas(hipblasComputeType_t c) { + // HIPBLAS_COMPUTE_16F -> f16, HIPBLAS_COMPUTE_32F -> f32 + if (c == HIPBLAS_COMPUTE_16F) return rocblas_datatype_f16_r; + return rocblas_datatype_f32_r; +} + +static inline hipblasStatus_t ggml_hip_gemm_ex( + hipblasHandle_t handle, hipblasOperation_t ta, hipblasOperation_t tb, + int m, int n, int k, const void* alpha, + const void* A, hipDataType At, int lda, + const void* B, hipDataType Bt, int ldb, + const void* beta, void* C, hipDataType Ct, int ldc, + hipblasComputeType_t comp, hipblasGemmAlgo_t algo) { + rocblas_handle rh = ggml_hip_rocblas_handle(handle); + auto ra = (ta == HIPBLAS_OP_T) ? rocblas_operation_transpose : rocblas_operation_none; + auto rb = (tb == HIPBLAS_OP_T) ? rocblas_operation_transpose : rocblas_operation_none; + auto rAt = ggml_hip_to_rocblas_type(At); + auto rBt = ggml_hip_to_rocblas_type(Bt); + auto rCt = ggml_hip_to_rocblas_type(Ct); + auto rComp = ggml_hip_compute_to_rocblas(comp); + auto rs = rocblas_gemm_ex(rh, ra, rb, m, n, k, + alpha, A, rAt, lda, B, rBt, ldb, + beta, C, rCt, ldc, C, rCt, ldc, + rComp, rocblas_gemm_algo_standard, 0, 0); + return (rs == rocblas_status_success) ? HIPBLAS_STATUS_SUCCESS : HIPBLAS_STATUS_EXECUTION_FAILED; +} + +#define cublasSgemm ggml_hip_sgemm #define cublasSgemmBatched hipblasSgemmBatched #define cublasSgemmStridedBatched hipblasSgemmStridedBatched #define cublasStatus_t hipblasStatus_t