CUDA/HIP: spechalize ggml_cuda_convert_val for half <-> bf16 conversions

Author: IMbackK

ggml/src/ggml-cuda/convert.cuh

@@ -34,23 +34,39 @@ 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);
+    if constexpr (std::is_same_v<src_t, dest_t>) {
+        return x;
+    } else {
+        return float(x);
+    }
 }
 
 template<>
 __host__ __device__ inline float ggml_cuda_convert_val<nv_bfloat16, float>(nv_bfloat16 x)
 {
-    return  __bfloat162float(x);
+    return __bfloat162float(x);
 }
 
 template<>
 __host__ __device__ inline nv_bfloat16 ggml_cuda_convert_val<nv_bfloat16, nv_bfloat16>(nv_bfloat16 x)
 {
-    return  x;
+    return x;
 }
 
 template<>
 __host__ __device__ inline nv_bfloat16 ggml_cuda_convert_val<float, nv_bfloat16>(float x)
 {
     return __float2bfloat16(x);
 }
+
+template<>
+__host__ __device__ inline half ggml_cuda_convert_val<nv_bfloat16, half>(nv_bfloat16 x)
+{
+    return half(__bfloat162float(x));
+}
+
+template<>
+__host__ __device__ inline nv_bfloat16 ggml_cuda_convert_val<half, nv_bfloat16>(half x)
+{
+    return __float2bfloat16(float(x));
+}