Use asm instruction instead of a builtin

deda3f92 · Rostyslav Geyyer · 306fd506 · deda3f92
Commit deda3f92 authored Nov 02, 2023 by Rostyslav Geyyer
Hide whitespace changes
Inline Side-by-side

Showing with 2 additions and 2 deletions

include/ck/utility/type_convert.hpp include/ck/utility/type_convert.hpp +2 -2

No files found.
--- a/include/ck/utility/type_convert.hpp
+++ b/include/ck/utility/type_convert.hpp
@@ -131,8 +131,8 @@ inline __host__ __device__ float type_convert<float, f8_t>(f8_t x)
 #if defined(__gfx940__) || defined(__gfx941__) || defined(__gfx942__)
    float fval;
    uint32_t i32val = static_cast<uint32_t>(x);
-    fval            = __builtin_amdgcn_cvt_f32_fp8(i32val, 0);
+    // fval            = __builtin_amdgcn_cvt_f32_fp8(i32val, 0);
-    // asm volatile("v_cvt_f32_fp8 %0, %1 src0_sel:BYTE_0" : "=v"(fval) : "v"(i32val));
+    asm volatile("v_cvt_f32_fp8 %0, %1 src0_sel:BYTE_0" : "=v"(fval) : "v"(i32val));
    return fval;
 #else
    constexpr bool negative_zero_nan = true;