CUDA/HIP: add expicit conversion operator to support older versions of rocm

Author: IMbackK

ggml/src/ggml-cuda/convert.cu

@@ -31,8 +31,8 @@ static __global__ void dequantize_block(const void * __restrict__ vx, dst_t * __
     dequantize_kernel(vx, ib, iqs, v);
 
     const int64_t iy0 = ((i03*ne02 + i02)*ne01 + i01)*ne00 + iybs + iqs;
-    y[iy0 + 0]        = float(v.x);
-    y[iy0 + y_offset] = float(v.y);
+    y[iy0 + 0]        = ggml_cuda_convert_val<float, dst_t>(v.x);
+    y[iy0 + y_offset] = ggml_cuda_convert_val<float, dst_t>(v.y);
 }
 
 template <bool need_check>
@@ -630,7 +630,7 @@ static __global__ void convert_unary(
 
     const int64_t ix = i03*s03 + i02*s02 + i01*s01 + i00;
     const int64_t iy = ((i03*ne02 + i02)*ne01 + i01)*ne00 + i00;
-    y[iy] = float(x[ix]);
+    y[iy] = ggml_cuda_convert_val<src_t, dst_t>(x[ix]);
 }
 
 template <typename src_t, typename dst_t>

ggml/src/ggml-cuda/convert.cuh

@@ -29,3 +29,28 @@ typedef to_t_nc_cuda_t<nv_bfloat16> to_bf16_nc_cuda_t;
 to_fp32_nc_cuda_t ggml_get_to_fp32_nc_cuda(ggml_type type);
 to_fp16_nc_cuda_t ggml_get_to_fp16_nc_cuda(ggml_type type);
 to_bf16_nc_cuda_t ggml_get_to_bf16_nc_cuda(ggml_type type);
+
+
+template<typename src_t, typename dest_t>
+ __host__ __device__ inline dest_t ggml_cuda_convert_val(src_t x)
+{
+    return float(x);
+}
+
+template<>
+__host__ __device__ inline float ggml_cuda_convert_val<nv_bfloat16, float>(nv_bfloat16 x)
+{
+    return  __bfloat162float(x);
+}
+
+template<>
+__host__ __device__ inline nv_bfloat16 ggml_cuda_convert_val<nv_bfloat16, nv_bfloat16>(nv_bfloat16 x)
+{
+    return  x;
+}
+
+template<>
+__host__ __device__ inline nv_bfloat16 ggml_cuda_convert_val<float, nv_bfloat16>(float x)
+{
+    return __float2bfloat16(x);
+}

ggml/src/ggml-cuda/cpy-utils.cuh

@@ -1,13 +1,14 @@
 #pragma once
 
 #include "ggml-common.h"
+#include "convert.cuh"
 
 template<typename src_t, typename dst_t>
 static __device__ __forceinline__ void convert_flt(const src_t * src, dst_t * dst) {
     if constexpr (std::is_same_v<src_t, dst_t>) {
         *dst = *src;
     } else {
-        *dst = float(*src);
+        *dst = ggml_cuda_convert_val<src_t, dst_t>(*src);
     }
 }