Add exception for HIP code

Author: am17an

ggml/src/ggml-cuda/common.cuh

@@ -532,7 +532,6 @@ static __device__ __forceinline__ int ggml_cuda_dp4a(const int a, const int b, i
 #endif // defined(GGML_USE_HIP)
 }
 
-
 static __device__ __forceinline__ void ggml_cuda_mad(float & acc, const float v, const float u) {
     acc += v*u;
 }
@@ -571,16 +570,6 @@ static __device__ __forceinline__ void ggml_cuda_mad(half2 & acc, const half2 v,
 #endif // FAST_FP16_AVAILABLE
 }
 
-
-#if defined(GGML_USE_HIP)
-static __device__ __forceinline__ void ggml_cuda_mad(float & acc, const __hip_bfloat162 v, const __hip_bfloat162 u) {
-    const float2 tmpv = __bfloat162float2(v);
-    const float2 tmpu = __bfloat162float2(u);
-    acc += tmpv.x * tmpu.x;
-    acc += tmpv.y * tmpu.y;
-}
-#endif
-
 // Aligned memory transfers of 8/16 bytes can be faster than 2 transfers with 4 bytes, especially on AMD.
 template <int nbytes, int alignment = 0>
 static __device__ __forceinline__ void ggml_cuda_memcpy_1(void * __restrict__ dst, const void * __restrict__ src) {

ggml/src/ggml-cuda/mmvf.cu

@@ -88,6 +88,21 @@ static __global__ void mul_mat_vec_f(
 #endif // FP16_AVAILABLE
         }
     } else if constexpr (std::is_same_v<T, nv_bfloat16>) {
+//TODO: add support for ggml_cuda_mad for hip_bfloat162
+#if defined(GGML_USE_HIP)
+        const int * x2 = (const int *) x;
+        for (int col2 = tid; col2 < ncols2; col2 += block_size) {
+            const int tmpx = x2[col2];
+#pragma unroll
+            for (int j = 0; j < ncols_dst; ++j) {
+                const float2 tmpy = y2[j*stride_col_y2 + col2];
+                const float tmpx0 = ggml_cuda_cast<float>(reinterpret_cast<const nv_bfloat16 *>(&tmpx)[0]);
+                const float tmpx1 = ggml_cuda_cast<float>(reinterpret_cast<const nv_bfloat16 *>(&tmpx)[1]);
+                ggml_cuda_mad(sumf[j], tmpx0, tmpy.x);
+                ggml_cuda_mad(sumf[j], tmpx1, tmpy.y);
+            }
+        }
+#else
         const nv_bfloat162 * x2 = (const nv_bfloat162 *) x;
         for (int col2 = tid; col2 < ncols2; col2 += block_size) {
             const nv_bfloat162 tmpx = x2[col2];
@@ -98,6 +113,7 @@ static __global__ void mul_mat_vec_f(
                 ggml_cuda_mad(sumf[j], tmpx.y, tmpy.y);
             }
         }
+#endif
     } else {
         static_assert(std::is_same_v<T, void>, "unsupported type");
     }