--- 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 <rocblas/rocblas.h>
+
+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