diff --git a/include/ck/utility/type_convert.hpp b/include/ck/utility/type_convert.hpp index 04ae046ac8..9b1321dea3 100644 --- a/include/ck/utility/type_convert.hpp +++ b/include/ck/utility/type_convert.hpp @@ -885,7 +885,14 @@ template <> inline __host__ __device__ float2_t type_convert(f8x2_ocp_t x) { #if CK_OCP_FP8_CVT_FAST_PATH +// __builtin_amdgcn_cvt_pk_f32_fp8 can produce incorrect results due to a compiler issue. +// TODO: Enable when SWDEV-532959 is fixed. +#if defined(__gfx1200__) || defined(__gfx1201__) + return float2_t{__builtin_amdgcn_cvt_f32_fp8(bit_cast(x), 0), + __builtin_amdgcn_cvt_f32_fp8(bit_cast(x), 1)}; +#else return __builtin_amdgcn_cvt_pk_f32_fp8(bit_cast(x), false); +#endif #else return float2_t{fp8_impl::cast_from_f8( x.AsType()[Number<0>{}]), @@ -1021,7 +1028,14 @@ template <> inline __host__ __device__ float2_t type_convert(bf8x2_ocp_t x) { #if CK_OCP_FP8_CVT_FAST_PATH +// __builtin_amdgcn_cvt_pk_f32_bf8 can produce incorrect results due to a compiler issue. +// TODO: Enable when SWDEV-532959 is fixed. +#if defined(__gfx1200__) || defined(__gfx1201__) + return float2_t{__builtin_amdgcn_cvt_f32_bf8(bit_cast(x), 0), + __builtin_amdgcn_cvt_f32_bf8(bit_cast(x), 1)}; +#else return __builtin_amdgcn_cvt_pk_f32_bf8(bit_cast(x), false); +#endif #else return float2_t{fp8_impl::cast_from_f8( x.AsType()[Number<0>{}]),