From 843c092b632d956262e977d02a1ed29fea7d039b Mon Sep 17 00:00:00 2001 From: yadaish Date: Mon, 24 Nov 2025 10:59:09 +0000 Subject: [PATCH 1/2] update --- include/ck_tile/core/numeric/pk_int4.hpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/include/ck_tile/core/numeric/pk_int4.hpp b/include/ck_tile/core/numeric/pk_int4.hpp index f0a681cb17..47297b1ef2 100644 --- a/include/ck_tile/core/numeric/pk_int4.hpp +++ b/include/ck_tile/core/numeric/pk_int4.hpp @@ -180,7 +180,7 @@ CK_TILE_HOST_DEVICE bf16x2_t pk_int4_t_to_bfloat16x2_t(const pk_int4_t& x) CK_TILE_HOST_DEVICE bf16x2_t pk_int4_t_to_bfloat16x2_t(const pk_int4_t& x, float scale) { - auto float_vec2 = pk_int4_t_to_fp32x2_t(x); + auto float_vec2 = pk_int4_t_to_fp32x2_t_signed_conversion(x); float_vec2.x = float_vec2.x * scale; float_vec2.y = float_vec2.y * scale; return bf16x2_t{type_convert(float_vec2.x), type_convert(float_vec2.y)}; From 1575d849bd0af24478c12c01b29acbf72861f22b Mon Sep 17 00:00:00 2001 From: yadaish Date: Tue, 25 Nov 2025 05:23:47 +0000 Subject: [PATCH 2/2] change endian seems working --- include/ck_tile/core/numeric/pk_int4.hpp | 3 +++ 1 file changed, 3 insertions(+) diff --git a/include/ck_tile/core/numeric/pk_int4.hpp b/include/ck_tile/core/numeric/pk_int4.hpp index 47297b1ef2..3e9df1806f 100644 --- a/include/ck_tile/core/numeric/pk_int4.hpp +++ b/include/ck_tile/core/numeric/pk_int4.hpp @@ -127,11 +127,14 @@ CK_TILE_HOST_DEVICE fp32x2_t pk_int4_t_to_fp32x2_t_signed_conversion(const pk_in x_l = x_l > 7 ? x_l - 16 : x_l; x_h = x_h > 7 ? x_h - 16 : x_h; + /* #ifdef CK_TILE_USE_PK4_LAYOUT_SHUFFLE fp32x2_t res = {x_h, x_l}; #elif fp32x2_t res = {x_l, x_h}; #endif + */ + fp32x2_t res = {x_l, x_h}; return res; }