From 7daf21081effcd034debbf1a51e0afad4597c9ff Mon Sep 17 00:00:00 2001 From: Rostyslav Geyyer Date: Tue, 18 Feb 2025 19:56:20 +0000 Subject: [PATCH] Fix vector sr conversion --- include/ck/utility/type_convert.hpp | 1 + test/data_type/test_mx_fp4.cpp | 4 ++-- 2 files changed, 3 insertions(+), 2 deletions(-) diff --git a/include/ck/utility/type_convert.hpp b/include/ck/utility/type_convert.hpp index 8411f2db06..907e85804d 100644 --- a/include/ck/utility/type_convert.hpp +++ b/include/ck/utility/type_convert.hpp @@ -1022,6 +1022,7 @@ inline __host__ __device__ f4x32_t f4_convert_sr(float32_t x, float scale = 1.0f float2_t floatx2_array[16]; float32_t floatx32_array; } float_values{{0}}; + float_values.floatx32_array = x; // TODO: pack in a loop tmp_values.bitwise = __builtin_amdgcn_cvt_scalef32_sr_pk_fp4_f32( tmp_values.bitwise, float_values.floatx2_array[0], rng, scale, 0); diff --git a/test/data_type/test_mx_fp4.cpp b/test/data_type/test_mx_fp4.cpp index e12fe2de47..1509778533 100644 --- a/test/data_type/test_mx_fp4.cpp +++ b/test/data_type/test_mx_fp4.cpp @@ -438,9 +438,9 @@ __global__ void test_mx_fp4x32_device_scaled_convert_sr(float* p_test, uint64_t* ck::static_for<0, N / 2, 1>{}([&](auto ii) { p_test[i++] = type_convert( - f4x32.AsType()(ck::Number{}).template unpack<>(ck::Number<0>{})); + f4_t(f4x32.AsType()(ck::Number{}).template unpack<>(ck::Number<0>{}))); p_test[i++] = type_convert( - f4x32.AsType()(ck::Number{}).template unpack<>(ck::Number<1>{})); + f4_t(f4x32.AsType()(ck::Number{}).template unpack<>(ck::Number<1>{}))); }); }