From 0eab68951d3bb668e8f99beba2e2c07d7e1ca18a Mon Sep 17 00:00:00 2001 From: JiaLuo-CAN Date: Mon, 31 Aug 2026 10:52:21 -0400 Subject: [PATCH] fix(ck,ck_tile): Adapt to compiler change from Vector to Ext_Vector (#11476) Jira: ROCM-30233 LLVM PR: #4184 Fix CK build error caused by the compiler change from Vector to Ext_Vector. --- .../amd_buffer_addressing_builtins.hpp | 70 +++++----- .../include/ck/utility/amd_ck_fp8.hpp | 19 +-- .../include/ck/utility/amd_cluster_load.hpp | 4 +- .../include/ck/utility/amd_smfmac.hpp | 14 +- .../include/ck/utility/amd_transpose_load.hpp | 4 +- .../include/ck/utility/amd_wmma.hpp | 62 +++++++-- .../include/ck/utility/amd_xdlops.hpp | 128 ++++++++++++------ .../include/ck/utility/dtype_vector.hpp | 17 +++ .../include/ck/utility/mxf8_utils.hpp | 4 +- .../core/arch/amd_buffer_addressing.hpp | 12 +- .../arch/amd_buffer_addressing_builtins.hpp | 13 +- .../ck_tile/core/arch/mma/mfma/mfma_gfx9.hpp | 72 ++++++---- .../core/arch/mma/sparse/mfma/sparse_gfx9.hpp | 96 +++++++------ .../include/ck_tile/core/arch/utility.hpp | 4 +- .../ck_tile/core/numeric/vector_type.hpp | 4 + .../warp/warp_gemm_attribute_mfma_impl.hpp | 57 ++++---- ..._gemm_attribute_wmma_impl_16bit_traits.hpp | 36 +++-- 17 files changed, 389 insertions(+), 227 deletions(-) diff --git a/projects/composablekernel/include/ck/utility/amd_buffer_addressing_builtins.hpp b/projects/composablekernel/include/ck/utility/amd_buffer_addressing_builtins.hpp index 3f1681f4a61b..e9c2143e2ae0 100644 --- a/projects/composablekernel/include/ck/utility/amd_buffer_addressing_builtins.hpp +++ b/projects/composablekernel/include/ck/utility/amd_buffer_addressing_builtins.hpp @@ -248,32 +248,33 @@ amd_buffer_load_impl_raw(__amdgpu_buffer_rsrc_t src_wave_buffer_resource, } else if constexpr(N == 8) { - int32x2_t tmp = __builtin_amdgcn_raw_buffer_load_b64(src_wave_buffer_resource, - src_thread_addr_offset, - src_wave_addr_offset, - static_cast(coherence)); + auto tmp = __builtin_amdgcn_raw_buffer_load_b64(src_wave_buffer_resource, + src_thread_addr_offset, + src_wave_addr_offset, + static_cast(coherence)); return bit_cast(tmp); } else if constexpr(N == 16) { - int32x4_t tmp = __builtin_amdgcn_raw_buffer_load_b128(src_wave_buffer_resource, - src_thread_addr_offset, - src_wave_addr_offset, - static_cast(coherence)); + auto tmp = __builtin_amdgcn_raw_buffer_load_b128(src_wave_buffer_resource, + src_thread_addr_offset, + src_wave_addr_offset, + static_cast(coherence)); return bit_cast(tmp); } else if constexpr(N == 32) { - int32x4_t tmp0 = __builtin_amdgcn_raw_buffer_load_b128(src_wave_buffer_resource, - src_thread_addr_offset, - src_wave_addr_offset, - static_cast(coherence)); - int32x4_t tmp1 = + int32x4_t tmp0 = bit_cast( + __builtin_amdgcn_raw_buffer_load_b128(src_wave_buffer_resource, + src_thread_addr_offset, + src_wave_addr_offset, + static_cast(coherence))); + int32x4_t tmp1 = bit_cast( __builtin_amdgcn_raw_buffer_load_b128(src_wave_buffer_resource, src_thread_addr_offset, src_wave_addr_offset + 4 * sizeof(int32_t), - static_cast(coherence)); + static_cast(coherence))); vector_type tmp; tmp.AsType()(Number<0>{}) = tmp0; @@ -283,25 +284,26 @@ amd_buffer_load_impl_raw(__amdgpu_buffer_rsrc_t src_wave_buffer_resource, } else if constexpr(N == 64) { - int32x4_t tmp0 = __builtin_amdgcn_raw_buffer_load_b128(src_wave_buffer_resource, - src_thread_addr_offset, - src_wave_addr_offset, - static_cast(coherence)); - int32x4_t tmp1 = + int32x4_t tmp0 = bit_cast( + __builtin_amdgcn_raw_buffer_load_b128(src_wave_buffer_resource, + src_thread_addr_offset, + src_wave_addr_offset, + static_cast(coherence))); + int32x4_t tmp1 = bit_cast( __builtin_amdgcn_raw_buffer_load_b128(src_wave_buffer_resource, src_thread_addr_offset, src_wave_addr_offset + 4 * sizeof(int32_t), - static_cast(coherence)); - int32x4_t tmp2 = + static_cast(coherence))); + int32x4_t tmp2 = bit_cast( __builtin_amdgcn_raw_buffer_load_b128(src_wave_buffer_resource, src_thread_addr_offset, src_wave_addr_offset + 8 * sizeof(int32_t), - static_cast(coherence)); - int32x4_t tmp3 = + static_cast(coherence))); + int32x4_t tmp3 = bit_cast( __builtin_amdgcn_raw_buffer_load_b128(src_wave_buffer_resource, src_thread_addr_offset, src_wave_addr_offset + 12 * sizeof(int32_t), - static_cast(coherence)); + static_cast(coherence))); vector_type tmp; @@ -383,7 +385,7 @@ amd_buffer_store_impl_raw(const typename vector_type::type src_thread } else if constexpr(N == 8) { - __builtin_amdgcn_raw_buffer_store_b64(bit_cast(src_thread_data), + __builtin_amdgcn_raw_buffer_store_b64(bit_cast(src_thread_data), dst_wave_buffer_resource, dst_thread_addr_offset, dst_wave_addr_offset, @@ -391,7 +393,7 @@ amd_buffer_store_impl_raw(const typename vector_type::type src_thread } else if constexpr(N == 16) { - __builtin_amdgcn_raw_buffer_store_b128(bit_cast(src_thread_data), + __builtin_amdgcn_raw_buffer_store_b128(bit_cast(src_thread_data), dst_wave_buffer_resource, dst_thread_addr_offset, dst_wave_addr_offset, @@ -399,15 +401,15 @@ amd_buffer_store_impl_raw(const typename vector_type::type src_thread } else if constexpr(N == 32) { - vector_type tmp{bit_cast(src_thread_data)}; + vector_type tmp{bit_cast(src_thread_data)}; - __builtin_amdgcn_raw_buffer_store_b128(tmp.template AsType()[Number<0>{}], + __builtin_amdgcn_raw_buffer_store_b128(tmp.template AsType()[Number<0>{}], dst_wave_buffer_resource, dst_thread_addr_offset, dst_wave_addr_offset, static_cast(coherence)); - __builtin_amdgcn_raw_buffer_store_b128(tmp.template AsType()[Number<1>{}], + __builtin_amdgcn_raw_buffer_store_b128(tmp.template AsType()[Number<1>{}], dst_wave_buffer_resource, dst_thread_addr_offset, dst_wave_addr_offset + sizeof(int32_t) * 4, @@ -415,27 +417,27 @@ amd_buffer_store_impl_raw(const typename vector_type::type src_thread } else if constexpr(N == 64) { - vector_type tmp{bit_cast(src_thread_data)}; + vector_type tmp{bit_cast(src_thread_data)}; - __builtin_amdgcn_raw_buffer_store_b128(tmp.template AsType()[Number<0>{}], + __builtin_amdgcn_raw_buffer_store_b128(tmp.template AsType()[Number<0>{}], dst_wave_buffer_resource, dst_thread_addr_offset, dst_wave_addr_offset, static_cast(coherence)); - __builtin_amdgcn_raw_buffer_store_b128(tmp.template AsType()[Number<1>{}], + __builtin_amdgcn_raw_buffer_store_b128(tmp.template AsType()[Number<1>{}], dst_wave_buffer_resource, dst_thread_addr_offset, dst_wave_addr_offset + sizeof(int32_t) * 4, static_cast(coherence)); - __builtin_amdgcn_raw_buffer_store_b128(tmp.template AsType()[Number<2>{}], + __builtin_amdgcn_raw_buffer_store_b128(tmp.template AsType()[Number<2>{}], dst_wave_buffer_resource, dst_thread_addr_offset, dst_wave_addr_offset + sizeof(int32_t) * 8, static_cast(coherence)); - __builtin_amdgcn_raw_buffer_store_b128(tmp.template AsType()[Number<3>{}], + __builtin_amdgcn_raw_buffer_store_b128(tmp.template AsType()[Number<3>{}], dst_wave_buffer_resource, dst_thread_addr_offset, dst_wave_addr_offset + sizeof(int32_t) * 12, diff --git a/projects/composablekernel/include/ck/utility/amd_ck_fp8.hpp b/projects/composablekernel/include/ck/utility/amd_ck_fp8.hpp index dfd5b2a28021..c06a4cdf8f73 100644 --- a/projects/composablekernel/include/ck/utility/amd_ck_fp8.hpp +++ b/projects/composablekernel/include/ck/utility/amd_ck_fp8.hpp @@ -98,6 +98,7 @@ using bhalf_t = __bf16; using bhalf_t = ushort; #endif typedef bhalf_t bhalf2_t __attribute__((ext_vector_type(2))); +typedef __bf16 llvm_bf16x2_t __attribute__((ext_vector_type(2))); __host__ __device__ static inline constexpr bool fnuz_f8_is_nan(f8_fnuz_t a) { @@ -685,7 +686,7 @@ static __device__ fp8_storage_t cast_to_f8_from_f16(_Float16 v, unsigned int rng } } - val.half_vec = + val.i16_vec = __builtin_amdgcn_cvt_scalef32_pk_bf8_f16(i16x2val, val.half_vec, /* scale */ 1.f, 0); return val.i8val[0]; @@ -859,8 +860,8 @@ static __device__ fp8_storage_t cast_to_f8_from_bf16(ushort v, unsigned int rng } } - val.i16_vec = - __builtin_amdgcn_cvt_scalef32_pk_fp8_bf16(i16x2val, val.bhalf_vec, /* scale */ 1.f, 0); + val.i16_vec = __builtin_amdgcn_cvt_scalef32_pk_fp8_bf16( + i16x2val, bit_cast(val.bhalf_vec), /* scale */ 1.f, 0); return val.i8val[0]; } @@ -907,8 +908,8 @@ static __device__ fp8x2_storage_t cast_to_f8_from_bf16(ushortx2_t v, unsigned in } } - val.i16_vec = - __builtin_amdgcn_cvt_scalef32_pk_fp8_bf16(i16x2val, val.bhalf_vec, /* scale */ 1.f, 0); + val.i16_vec = __builtin_amdgcn_cvt_scalef32_pk_fp8_bf16( + i16x2val, bit_cast(val.bhalf_vec), /* scale */ 1.f, 0); return fp8x2_storage_t{val.i8val[0], val.i8val[1]}; #endif @@ -945,8 +946,8 @@ static __device__ fp8_storage_t cast_to_f8_from_bf16(ushort v, unsigned int rng } } - val.i16_vec = - __builtin_amdgcn_cvt_scalef32_pk_bf8_bf16(i16x2val, val.bhalf_vec, /* scale */ 1.f, 0); + val.i16_vec = __builtin_amdgcn_cvt_scalef32_pk_bf8_bf16( + i16x2val, bit_cast(val.bhalf_vec), /* scale */ 1.f, 0); return val.i8val[0]; } @@ -988,8 +989,8 @@ static __device__ fp8x2_storage_t cast_to_f8_from_bf16(ushortx2_t v, unsigned in } } - val.i16_vec = - __builtin_amdgcn_cvt_scalef32_pk_bf8_bf16(i16x2val, val.bhalf_vec, /* scale */ 1.f, 0); + val.i16_vec = __builtin_amdgcn_cvt_scalef32_pk_bf8_bf16( + i16x2val, bit_cast(val.bhalf_vec), /* scale */ 1.f, 0); return fp8x2_storage_t{val.i8val[0], val.i8val[1]}; } diff --git a/projects/composablekernel/include/ck/utility/amd_cluster_load.hpp b/projects/composablekernel/include/ck/utility/amd_cluster_load.hpp index 6a70281dbd3d..03c6aa10e89e 100644 --- a/projects/composablekernel/include/ck/utility/amd_cluster_load.hpp +++ b/projects/composablekernel/include/ck/utility/amd_cluster_load.hpp @@ -44,13 +44,13 @@ __device__ T cluster_multicast_load(const T* global_ptr, int mask) } else if constexpr(sizeof(T) == 8) { - using vec2i_t = __attribute__((vector_size(8))) int; + using vec2i_t = int __attribute__((ext_vector_type(2))); return bit_cast(__builtin_amdgcn_cluster_load_b64( detail::to_global(reinterpret_cast(global_ptr)), 0, mask)); } else if constexpr(sizeof(T) == 16) { - using vec4i_t = __attribute__((vector_size(16))) int; + using vec4i_t = int __attribute__((ext_vector_type(4))); return bit_cast(__builtin_amdgcn_cluster_load_b128( detail::to_global(reinterpret_cast(global_ptr)), 0, mask)); } diff --git a/projects/composablekernel/include/ck/utility/amd_smfmac.hpp b/projects/composablekernel/include/ck/utility/amd_smfmac.hpp index 4a09163038e7..15d1f0fdec8b 100644 --- a/projects/composablekernel/include/ck/utility/amd_smfmac.hpp +++ b/projects/composablekernel/include/ck/utility/amd_smfmac.hpp @@ -42,7 +42,12 @@ struct intrin_smfmac_f32_16x16x32bf16<16, 16> { #if defined(__gfx94__) reg_c.template AsType()(Number<0>{}) = __builtin_amdgcn_smfmac_f32_16x16x32_bf16( - reg_a, reg_b, reg_c.template AsType()[Number<0>{}], reg_idx, 0, abid); + bit_cast(reg_a), + bit_cast(reg_b), + reg_c.template AsType()[Number<0>{}], + reg_idx, + 0, + abid); #else ignore = reg_a; ignore = reg_b; @@ -86,7 +91,12 @@ struct intrin_smfmac_f32_32x32x16bf16<32, 32> { #if defined(__gfx94__) reg_c.template AsType()(Number<0>{}) = __builtin_amdgcn_smfmac_f32_32x32x16_bf16( - reg_a, reg_b, reg_c.template AsType()[Number<0>{}], reg_idx, 0, abid); + bit_cast(reg_a), + bit_cast(reg_b), + reg_c.template AsType()[Number<0>{}], + reg_idx, + 0, + abid); #else ignore = reg_a; ignore = reg_b; diff --git a/projects/composablekernel/include/ck/utility/amd_transpose_load.hpp b/projects/composablekernel/include/ck/utility/amd_transpose_load.hpp index 17de26a40112..f533296d50e2 100644 --- a/projects/composablekernel/include/ck/utility/amd_transpose_load.hpp +++ b/projects/composablekernel/include/ck/utility/amd_transpose_load.hpp @@ -13,7 +13,7 @@ __device__ auto amd_global_load_transpose_to_vgpr(const T* in_ptr) using vector_t = typename vector_type::type; if constexpr(sizeof(T) == 2) { - typedef __attribute__((__vector_size__(8 * sizeof(__fp16)))) __fp16 llvm_fp16x8_t; + typedef __fp16 llvm_fp16x8_t __attribute__((ext_vector_type(8))); __attribute__((address_space(1))) llvm_fp16x8_t* glb_ptr = reinterpret_cast<__attribute__((address_space(1))) llvm_fp16x8_t*>( reinterpret_cast(in_ptr)); @@ -21,7 +21,7 @@ __device__ auto amd_global_load_transpose_to_vgpr(const T* in_ptr) } else if constexpr(sizeof(T) == 1) { - typedef __attribute__((__vector_size__(2 * sizeof(int)))) int llvm_intx2_t; + typedef int llvm_intx2_t __attribute__((ext_vector_type(2))); __attribute__((address_space(1))) llvm_intx2_t* glb_ptr = reinterpret_cast<__attribute__((address_space(1))) llvm_intx2_t*>( reinterpret_cast(in_ptr)); diff --git a/projects/composablekernel/include/ck/utility/amd_wmma.hpp b/projects/composablekernel/include/ck/utility/amd_wmma.hpp index 72db5d144836..3a7f16c53616 100644 --- a/projects/composablekernel/include/ck/utility/amd_wmma.hpp +++ b/projects/composablekernel/include/ck/utility/amd_wmma.hpp @@ -64,7 +64,9 @@ struct intrin_wmma_f32_16x16x16_bf16_w32<16, 16> #if defined(__gfx11__) reg_c.template AsType()(Number<0>{}) = __builtin_amdgcn_wmma_f32_16x16x16_bf16_w32( - reg_a, reg_b, reg_c.template AsType()[Number<0>{}]); + bit_cast(reg_a), + bit_cast(reg_b), + reg_c.template AsType()[Number<0>{}]); #else ignore = reg_a; ignore = reg_b; @@ -112,8 +114,11 @@ struct intrin_wmma_bf16_16x16x16_bf16_w32<16, 16, Opsel> // true : D0.[16:31]= result #if defined(__gfx11__) reg_c.template AsType()(Number<0>{}) = - __builtin_amdgcn_wmma_bf16_16x16x16_bf16_w32( - reg_a, reg_b, reg_c.template AsType()[Number<0>{}], Opsel); + bit_cast(__builtin_amdgcn_wmma_bf16_16x16x16_bf16_w32( + bit_cast(reg_a), + bit_cast(reg_b), + bit_cast(reg_c.template AsType()[Number<0>{}]), + Opsel)); #else ignore = reg_a; ignore = reg_b; @@ -184,7 +189,9 @@ struct intrin_wmma_f32_16x16x16_bf16_w64<16, 16> #if defined(__gfx11__) reg_c.template AsType()(Number<0>{}) = __builtin_amdgcn_wmma_f32_16x16x16_bf16_w64( - reg_a, reg_b, reg_c.template AsType()[Number<0>{}]); + bit_cast(reg_a), + bit_cast(reg_b), + reg_c.template AsType()[Number<0>{}]); #else ignore = reg_a; ignore = reg_b; @@ -232,8 +239,11 @@ struct intrin_wmma_bf16_16x16x16_bf16_w64<16, 16, Opsel> // true : D0.[16:31]= result #if defined(__gfx11__) reg_c.template AsType()(Number<0>{}) = - __builtin_amdgcn_wmma_bf16_16x16x16_bf16_w64( - reg_a, reg_b, reg_c.template AsType()[Number<0>{}], Opsel); + bit_cast(__builtin_amdgcn_wmma_bf16_16x16x16_bf16_w64( + bit_cast(reg_a), + bit_cast(reg_b), + bit_cast(reg_c.template AsType()[Number<0>{}]), + Opsel)); #else ignore = reg_a; ignore = reg_b; @@ -311,7 +321,9 @@ struct intrin_wmma_f32_16x16x16_bf16_w32_gfx12<16, 16> #if defined(__gfx120__) reg_c.template AsType()(Number<0>{}) = __builtin_amdgcn_wmma_f32_16x16x16_bf16_w32_gfx12( - reg_a, reg_b, reg_c.template AsType()[Number<0>{}]); + bit_cast(reg_a), + bit_cast(reg_b), + reg_c.template AsType()[Number<0>{}]); #else ignore = reg_a; ignore = reg_b; @@ -483,8 +495,16 @@ struct intrin_wmma_bf16_16x16x32_bf16<16, 16> // false: D0.[0:15] = result // true : D0.[16:31]= result #if defined(__gfx125__) - reg_c.template AsType()(Number<0>{}) = __builtin_amdgcn_wmma_bf16_16x16x32_bf16( - 0, reg_a, 0, reg_b, 0, reg_c.template AsType()[Number<0>{}], false, false); + reg_c.template AsType()(Number<0>{}) = + bit_cast(__builtin_amdgcn_wmma_bf16_16x16x32_bf16( + 0, + bit_cast(reg_a), + 0, + bit_cast(reg_b), + 0, + bit_cast(reg_c.template AsType()[Number<0>{}]), + false, + false)); #else ignore = reg_a; ignore = reg_b; @@ -529,8 +549,15 @@ struct intrin_wmma_f32_16x16x32_bf16<16, 16> __device__ static void Run(const bhalf16_t& reg_a, const bhalf16_t& reg_b, FloatC& reg_c) { #if defined(__gfx125__) - reg_c.template AsType()(Number<0>{}) = __builtin_amdgcn_wmma_f32_16x16x32_bf16( - 0, reg_a, 0, reg_b, 0, reg_c.template AsType()[Number<0>{}], false, false); + reg_c.template AsType()(Number<0>{}) = + __builtin_amdgcn_wmma_f32_16x16x32_bf16(0, + bit_cast(reg_a), + 0, + bit_cast(reg_b), + 0, + reg_c.template AsType()[Number<0>{}], + false, + false); #else ignore = reg_a; ignore = reg_b; @@ -551,9 +578,16 @@ struct intrin_wmma_bf16f32_16x16x32_bf16<16, 16> Run(const bhalf16_t& reg_a, const bhalf16_t& reg_b, FloatC& reg_c, FloatD& reg_d) { #if defined(__gfx125__) - reg_d - .template AsType()(Number<0>{}) = __builtin_amdgcn_wmma_bf16f32_16x16x32_bf16( - 0, reg_a, 0, reg_b, 0, reg_c.template AsType()[Number<0>{}], false, false); + reg_d.template AsType()(Number<0>{}) = + bit_cast(__builtin_amdgcn_wmma_bf16f32_16x16x32_bf16( + 0, + bit_cast(reg_a), + 0, + bit_cast(reg_b), + 0, + reg_c.template AsType()[Number<0>{}], + false, + false)); #else ignore = reg_a; ignore = reg_b; diff --git a/projects/composablekernel/include/ck/utility/amd_xdlops.hpp b/projects/composablekernel/include/ck/utility/amd_xdlops.hpp index f8b1736801c0..5c7aa74ca197 100644 --- a/projects/composablekernel/include/ck/utility/amd_xdlops.hpp +++ b/projects/composablekernel/include/ck/utility/amd_xdlops.hpp @@ -274,8 +274,13 @@ struct intrin_mfma_f32_32x32x16bf16<32, 32> __device__ static void Run(const bhalf8_t& reg_a, const bhalf8_t& reg_b, FloatC& reg_c) { #if defined(__gfx950__) - reg_c.template AsType()(Number<0>{}) = __builtin_amdgcn_mfma_f32_32x32x16_bf16( - reg_a, reg_b, reg_c.template AsType()[Number<0>{}], 0, 0, 0); + reg_c.template AsType()(Number<0>{}) = + __builtin_amdgcn_mfma_f32_32x32x16_bf16(bit_cast(reg_a), + bit_cast(reg_b), + reg_c.template AsType()[Number<0>{}], + 0, + 0, + 0); #else ignore = reg_a; ignore = reg_b; @@ -294,8 +299,13 @@ struct intrin_mfma_f32_16x16x32bf16<16, 16> __device__ static void Run(const bhalf8_t& reg_a, const bhalf8_t& reg_b, FloatC& reg_c) { #if defined(__gfx950__) - reg_c.template AsType()(Number<0>{}) = __builtin_amdgcn_mfma_f32_16x16x32_bf16( - reg_a, reg_b, reg_c.template AsType()[Number<0>{}], 0, 0, 0); + reg_c.template AsType()(Number<0>{}) = + __builtin_amdgcn_mfma_f32_16x16x32_bf16(bit_cast(reg_a), + bit_cast(reg_b), + reg_c.template AsType()[Number<0>{}], + 0, + 0, + 0); #else ignore = reg_a; ignore = reg_b; @@ -314,7 +324,12 @@ struct intrin_mfma_f32_32x32x8bf16_1k<32, 32> __device__ static void Run(const bhalf4_t& reg_a, const bhalf4_t& reg_b, FloatC& reg_c) { reg_c.template AsType()(Number<0>{}) = __builtin_amdgcn_mfma_f32_32x32x8bf16_1k( - reg_a, reg_b, reg_c.template AsType()[Number<0>{}], 0, 0, 0); + bit_cast(reg_a), + bit_cast(reg_b), + reg_c.template AsType()[Number<0>{}], + 0, + 0, + 0); } }; @@ -328,7 +343,12 @@ struct intrin_mfma_f32_16x16x16bf16_1k<16, 16> __device__ static void Run(const bhalf4_t& reg_a, const bhalf4_t& reg_b, FloatC& reg_c) { reg_c.template AsType()(Number<0>{}) = __builtin_amdgcn_mfma_f32_16x16x16bf16_1k( - reg_a, reg_b, reg_c.template AsType()[Number<0>{}], 0, 0, 0); + bit_cast(reg_a), + bit_cast(reg_b), + reg_c.template AsType()[Number<0>{}], + 0, + 0, + 0); } }; @@ -341,8 +361,13 @@ struct intrin_mfma_f32_32x32x4bf16<32, 32> template __device__ static void Run(const bhalf2_t& reg_a, const bhalf2_t& reg_b, FloatC& reg_c) { - reg_c.template AsType()(Number<0>{}) = __builtin_amdgcn_mfma_f32_32x32x4bf16( - reg_a, reg_b, reg_c.template AsType()[Number<0>{}], 0, 0, 0); + reg_c.template AsType()(Number<0>{}) = + __builtin_amdgcn_mfma_f32_32x32x4bf16(bit_cast(reg_a), + bit_cast(reg_b), + reg_c.template AsType()[Number<0>{}], + 0, + 0, + 0); } }; @@ -355,8 +380,13 @@ struct intrin_mfma_f32_16x16x8bf16<16, 16> template __device__ static void Run(const bhalf2_t& reg_a, const bhalf2_t& reg_b, FloatC& reg_c) { - reg_c.template AsType()(Number<0>{}) = __builtin_amdgcn_mfma_f32_16x16x8bf16( - reg_a, reg_b, reg_c.template AsType()[Number<0>{}], 0, 0, 0); + reg_c.template AsType()(Number<0>{}) = + __builtin_amdgcn_mfma_f32_16x16x8bf16(bit_cast(reg_a), + bit_cast(reg_b), + reg_c.template AsType()[Number<0>{}], + 0, + 0, + 0); } }; @@ -408,8 +438,13 @@ struct intrin_mfma_i32_32x32x32i8<32, 32> __device__ static void Run(const int8x16_t& reg_a, const int8x16_t& reg_b, FloatC& reg_c) { #if defined(__gfx950__) - reg_c.template AsType()(Number<0>{}) = __builtin_amdgcn_mfma_i32_32x32x32_i8( - reg_a, reg_b, reg_c.template AsType()[Number<0>{}], 0, 0, 0); + reg_c.template AsType()(Number<0>{}) = + __builtin_amdgcn_mfma_i32_32x32x32_i8(bit_cast(reg_a), + bit_cast(reg_b), + reg_c.template AsType()[Number<0>{}], + 0, + 0, + 0); #else ignore = reg_a; ignore = reg_b; @@ -428,8 +463,13 @@ struct intrin_mfma_i32_16x16x64i8<16, 16> __device__ static void Run(const int8x16_t& reg_a, const int8x16_t& reg_b, FloatC& reg_c) { #if defined(__gfx950__) - reg_c.template AsType()(Number<0>{}) = __builtin_amdgcn_mfma_i32_16x16x64_i8( - reg_a, reg_b, reg_c.template AsType()[Number<0>{}], 0, 0, 0); + reg_c.template AsType()(Number<0>{}) = + __builtin_amdgcn_mfma_i32_16x16x64_i8(bit_cast(reg_a), + bit_cast(reg_b), + reg_c.template AsType()[Number<0>{}], + 0, + 0, + 0); #else ignore = reg_a; ignore = reg_b; @@ -514,8 +554,8 @@ struct intrin_mfma_f32_32x32x64f8f6f4<32, 32> #if defined(__gfx950__) reg_c.template AsType()(Number<0>{}) = __builtin_amdgcn_mfma_scale_f32_32x32x64_f8f6f4( - reg_a, - reg_b, + bit_cast(reg_a), + bit_cast(reg_b), reg_c.template AsType()[Number<0>{}], 0, // cbsz {0 FP8 E4M3; 1 FP8 E5M2; 2 FP6 E2M3; 3 FP6 E3M2; 4 FP4 E2M1} 0, // blgp @@ -536,8 +576,8 @@ struct intrin_mfma_f32_32x32x64f8f6f4<32, 32> #if defined(__gfx950__) reg_c.template AsType()(Number<0>{}) = __builtin_amdgcn_mfma_scale_f32_32x32x64_f8f6f4( - reg_a, - reg_b, + bit_cast(reg_a), + bit_cast(reg_b), reg_c.template AsType()[Number<0>{}], 1, // cbsz {0 FP8 E4M3; 1 FP8 E5M2; 2 FP6 E2M3; 3 FP6 E3M2; 4 FP4 E2M1} 1, // blgp @@ -558,8 +598,8 @@ struct intrin_mfma_f32_32x32x64f8f6f4<32, 32> #if defined(__gfx950__) reg_c.template AsType()(Number<0>{}) = __builtin_amdgcn_mfma_scale_f32_32x32x64_f8f6f4( - reg_a, - reg_b, + bit_cast(reg_a), + bit_cast(reg_b), reg_c.template AsType()[Number<0>{}], 1, // cbsz {0 FP8 E4M3; 1 FP8 E5M2; 2 FP6 E2M3; 3 FP6 E3M2; 4 FP4 E2M1} 0, // blgp @@ -580,8 +620,8 @@ struct intrin_mfma_f32_32x32x64f8f6f4<32, 32> #if defined(__gfx950__) reg_c.template AsType()(Number<0>{}) = __builtin_amdgcn_mfma_scale_f32_32x32x64_f8f6f4( - reg_a, - reg_b, + bit_cast(reg_a), + bit_cast(reg_b), reg_c.template AsType()[Number<0>{}], 0, // cbsz {0 FP8 E4M3; 1 FP8 E5M2; 2 FP6 E2M3; 3 FP6 E3M2; 4 FP4 E2M1} 1, // blgp @@ -698,8 +738,8 @@ struct intrin_mfma_scale_f32_32x32x64f8f6f4<32, 32, OpselA, OpselB> // https://github.com/ROCm/llvm-project/blob/656552edc693e2bb4abc9258399c39d190fce2b3/llvm/test/Verifier/AMDGPU/mfma-scale.ll#L10 reg_c.template AsType()(Number<0>{}) = __builtin_amdgcn_mfma_scale_f32_32x32x64_f8f6f4( - reg_a, - reg_b, + bit_cast(reg_a), + bit_cast(reg_b), reg_c.template AsType()[Number<0>{}], 0, // cbsz {0 FP8 E4M3; 1 FP8 E5M2; 2 FP6 E2M3; 3 FP6 E3M2; 4 FP4 E2M1} 0, // blgp @@ -735,8 +775,8 @@ struct intrin_mfma_scale_f32_32x32x64f8f6f4<32, 32, OpselA, OpselB> // https://github.com/ROCm/llvm-project/blob/656552edc693e2bb4abc9258399c39d190fce2b3/llvm/test/Verifier/AMDGPU/mfma-scale.ll#L10 reg_c.template AsType()(Number<0>{}) = __builtin_amdgcn_mfma_scale_f32_32x32x64_f8f6f4( - reg_a, - reg_b, + bit_cast(reg_a), + bit_cast(reg_b), reg_c.template AsType()[Number<0>{}], 1, // cbsz {0 FP8 E4M3; 1 FP8 E5M2; 2 FP6 E2M3; 3 FP6 E3M2; 4 FP4 E2M1} 1, // blgp @@ -772,8 +812,8 @@ struct intrin_mfma_scale_f32_32x32x64f8f6f4<32, 32, OpselA, OpselB> // https://github.com/ROCm/llvm-project/blob/656552edc693e2bb4abc9258399c39d190fce2b3/llvm/test/Verifier/AMDGPU/mfma-scale.ll#L10 reg_c.template AsType()(Number<0>{}) = __builtin_amdgcn_mfma_scale_f32_32x32x64_f8f6f4( - reg_a, - reg_b, + bit_cast(reg_a), + bit_cast(reg_b), reg_c.template AsType()[Number<0>{}], 1, // cbsz {0 FP8 E4M3; 1 FP8 E5M2; 2 FP6 E2M3; 3 FP6 E3M2; 4 FP4 E2M1} 0, // blgp @@ -918,8 +958,8 @@ struct intrin_mfma_scale_f32_16x16x128f8f6f4<16, 16, OpselA, OpselB> // https://github.com/ROCm/llvm-project/blob/656552edc693e2bb4abc9258399c39d190fce2b3/llvm/test/Verifier/AMDGPU/mfma-scale.ll#L10 reg_c.template AsType()(Number<0>{}) = __builtin_amdgcn_mfma_scale_f32_16x16x128_f8f6f4( - reg_a, - reg_b, + bit_cast(reg_a), + bit_cast(reg_b), reg_c.template AsType()[Number<0>{}], 0, // cbsz {0 FP8 E4M3; 1 FP8 E5M2; 2 FP6 E2M3; 3 FP6 E3M2; 4 FP4 E2M1} 0, // blgp @@ -947,8 +987,8 @@ struct intrin_mfma_scale_f32_16x16x128f8f6f4<16, 16, OpselA, OpselB> // https://github.com/ROCm/llvm-project/blob/656552edc693e2bb4abc9258399c39d190fce2b3/llvm/test/Verifier/AMDGPU/mfma-scale.ll#L10 reg_c.template AsType()(Number<0>{}) = __builtin_amdgcn_mfma_scale_f32_16x16x128_f8f6f4( - reg_a, - reg_b, + bit_cast(reg_a), + bit_cast(reg_b), reg_c.template AsType()[Number<0>{}], 1, // cbsz {0 FP8 E4M3; 1 FP8 E5M2; 2 FP6 E2M3; 3 FP6 E3M2; 4 FP4 E2M1} 1, // blgp @@ -976,8 +1016,8 @@ struct intrin_mfma_scale_f32_16x16x128f8f6f4<16, 16, OpselA, OpselB> // https://github.com/ROCm/llvm-project/blob/656552edc693e2bb4abc9258399c39d190fce2b3/llvm/test/Verifier/AMDGPU/mfma-scale.ll#L10 reg_c.template AsType()(Number<0>{}) = __builtin_amdgcn_mfma_scale_f32_16x16x128_f8f6f4( - reg_a, - reg_b, + bit_cast(reg_a), + bit_cast(reg_b), reg_c.template AsType()[Number<0>{}], 0, // cbsz {0 FP8 E4M3; 1 FP8 E5M2; 2 FP6 E2M3; 3 FP6 E3M2; 4 FP4 E2M1} 1, // blgp @@ -1005,8 +1045,8 @@ struct intrin_mfma_scale_f32_16x16x128f8f6f4<16, 16, OpselA, OpselB> // https://github.com/ROCm/llvm-project/blob/656552edc693e2bb4abc9258399c39d190fce2b3/llvm/test/Verifier/AMDGPU/mfma-scale.ll#L10 reg_c.template AsType()(Number<0>{}) = __builtin_amdgcn_mfma_scale_f32_16x16x128_f8f6f4( - reg_a, - reg_b, + bit_cast(reg_a), + bit_cast(reg_b), reg_c.template AsType()[Number<0>{}], 1, // cbsz {0 FP8 E4M3; 1 FP8 E5M2; 2 FP6 E2M3; 3 FP6 E3M2; 4 FP4 E2M1} 0, // blgp @@ -1236,8 +1276,8 @@ struct intrin_mfma_f32_16x16x128f8f6f4<16, 16> // https://github.com/ROCm/llvm-project/blob/656552edc693e2bb4abc9258399c39d190fce2b3/llvm/test/Verifier/AMDGPU/mfma-scale.ll#L10 reg_c.template AsType()(Number<0>{}) = __builtin_amdgcn_mfma_scale_f32_16x16x128_f8f6f4( - reg_a, - reg_b, + bit_cast(reg_a), + bit_cast(reg_b), reg_c.template AsType()[Number<0>{}], 0, // cbsz {0 FP8 E4M3; 1 FP8 E5M2; 2 FP6 E2M3; 3 FP6 E3M2; 4 FP4 E2M1} 0, // blgp @@ -1259,8 +1299,8 @@ struct intrin_mfma_f32_16x16x128f8f6f4<16, 16> // https://github.com/ROCm/llvm-project/blob/656552edc693e2bb4abc9258399c39d190fce2b3/llvm/test/Verifier/AMDGPU/mfma-scale.ll#L10 reg_c.template AsType()(Number<0>{}) = __builtin_amdgcn_mfma_scale_f32_16x16x128_f8f6f4( - reg_a, - reg_b, + bit_cast(reg_a), + bit_cast(reg_b), reg_c.template AsType()[Number<0>{}], 1, // cbsz {0 FP8 E4M3; 1 FP8 E5M2; 2 FP6 E2M3; 3 FP6 E3M2; 4 FP4 E2M1} 1, // blgp @@ -1282,8 +1322,8 @@ struct intrin_mfma_f32_16x16x128f8f6f4<16, 16> // https://github.com/ROCm/llvm-project/blob/656552edc693e2bb4abc9258399c39d190fce2b3/llvm/test/Verifier/AMDGPU/mfma-scale.ll#L10 reg_c.template AsType()(Number<0>{}) = __builtin_amdgcn_mfma_scale_f32_16x16x128_f8f6f4( - reg_a, - reg_b, + bit_cast(reg_a), + bit_cast(reg_b), reg_c.template AsType()[Number<0>{}], 1, // cbsz {0 FP8 E4M3; 1 FP8 E5M2; 2 FP6 E2M3; 3 FP6 E3M2; 4 FP4 E2M1} 0, // blgp @@ -1305,8 +1345,8 @@ struct intrin_mfma_f32_16x16x128f8f6f4<16, 16> // https://github.com/ROCm/llvm-project/blob/656552edc693e2bb4abc9258399c39d190fce2b3/llvm/test/Verifier/AMDGPU/mfma-scale.ll#L10 reg_c.template AsType()(Number<0>{}) = __builtin_amdgcn_mfma_scale_f32_16x16x128_f8f6f4( - reg_a, - reg_b, + bit_cast(reg_a), + bit_cast(reg_b), reg_c.template AsType()[Number<0>{}], 0, // cbsz {0 FP8 E4M3; 1 FP8 E5M2; 2 FP6 E2M3; 3 FP6 E3M2; 4 FP4 E2M1} 1, // blgp diff --git a/projects/composablekernel/include/ck/utility/dtype_vector.hpp b/projects/composablekernel/include/ck/utility/dtype_vector.hpp index f7d6701ccae8..13070ce7d117 100644 --- a/projects/composablekernel/include/ck/utility/dtype_vector.hpp +++ b/projects/composablekernel/include/ck/utility/dtype_vector.hpp @@ -733,6 +733,23 @@ using int32x16_t = typename vector_type::type; using int32x32_t = typename vector_type::type; using int32x64_t = typename vector_type::type; +// i16: the bf16 matrix-core builtins take their operands as vectors of short, regardless of how +// bhalf_t is spelled +using int16x2_t = NativeVectorT; +using int16x4_t = NativeVectorT; +using int16x8_t = NativeVectorT; +using int16x16_t = NativeVectorT; + +// take native __bf16 regardless of whether bfloat16_t is spelled __bf16 or ushort +using llvm_bf16x8_t = NativeVectorT<__bf16, 8>; +using llvm_bf16x16_t = NativeVectorT<__bf16, 16>; + +// u32 +using uint32x2_t = typename vector_type::type; +using uint32x4_t = typename vector_type::type; +using uint32x8_t = typename vector_type::type; +using uint32x16_t = typename vector_type::type; + // i8 using int8x2_t = typename vector_type::type; using int8x4_t = typename vector_type::type; diff --git a/projects/composablekernel/include/ck/utility/mxf8_utils.hpp b/projects/composablekernel/include/ck/utility/mxf8_utils.hpp index c73aeea4d007..03a501393467 100644 --- a/projects/composablekernel/include/ck/utility/mxf8_utils.hpp +++ b/projects/composablekernel/include/ck/utility/mxf8_utils.hpp @@ -174,7 +174,7 @@ static __device__ fp8_storage_t cast_to_f8_from_f32_scaled(float v, unsigned int union { uint32_t ival; - vector_type::type v2i16; + int16x2_t v2i16; fp8_storage_t v4i8[4]; } ret{}; @@ -227,7 +227,7 @@ static __device__ fp8x2_storage_t cast_to_f8_from_f32_scaled(float2_t v, union { uint32_t ival; - vector_type::type v2i16; + int16x2_t v2i16; StaticallyIndexedArray v2f8x2; } ret{}; diff --git a/projects/composablekernel/include/ck_tile/core/arch/amd_buffer_addressing.hpp b/projects/composablekernel/include/ck_tile/core/arch/amd_buffer_addressing.hpp index cb0529506f1b..1ba6492e69ea 100644 --- a/projects/composablekernel/include/ck_tile/core/arch/amd_buffer_addressing.hpp +++ b/projects/composablekernel/include/ck_tile/core/arch/amd_buffer_addressing.hpp @@ -157,8 +157,8 @@ struct buffer_load<16, pre_nop> using mbuf_t = typename impl::buffer_load_trait<16, T>::payload_t; #if HAS_RAW_BUFFER_BUILTINS index_t s_offset = i_offset; - reinterpret_cast(value) = __builtin_amdgcn_raw_buffer_load_b128( - cast_to_amdgpu_buffer_rsrc_t(res), v_offset, s_offset, 0); + reinterpret_cast(value) = bit_cast(__builtin_amdgcn_raw_buffer_load_b128( + cast_to_amdgpu_buffer_rsrc_t(res), v_offset, s_offset, 0)); #else if constexpr(pre_nop) asm volatile("s_nop 4\n" @@ -191,8 +191,8 @@ struct buffer_load<8, pre_nop> using mbuf_t = typename impl::buffer_load_trait<8, T>::payload_t; #if HAS_RAW_BUFFER_BUILTINS index_t s_offset = i_offset; - reinterpret_cast(value) = __builtin_amdgcn_raw_buffer_load_b64( - cast_to_amdgpu_buffer_rsrc_t(res), v_offset, s_offset, 0); + reinterpret_cast(value) = bit_cast(__builtin_amdgcn_raw_buffer_load_b64( + cast_to_amdgpu_buffer_rsrc_t(res), v_offset, s_offset, 0)); #else if constexpr(pre_nop) asm volatile("s_nop 4\n" @@ -3104,7 +3104,7 @@ amd_tdm_load(const TDMDescriptor& descriptor static constexpr auto I4 = number<4>{}; auto tdm_desc_grp = descriptor.getResourceDescriptorGroup(); - __builtin_amdgcn_tensor_load_to_lds(tdm_desc_grp.get(I0), + __builtin_amdgcn_tensor_load_to_lds(bit_cast(tdm_desc_grp.get(I0)), tdm_desc_grp.get(I1), tdm_desc_grp.get(I2), tdm_desc_grp.get(I3), @@ -3130,7 +3130,7 @@ amd_tdm_store(const TDMDescriptor& descripto static constexpr auto I4 = number<4>{}; auto tdm_desc_grp = descriptor.getResourceDescriptorGroup(); - __builtin_amdgcn_tensor_store_from_lds(tdm_desc_grp.get(I0), + __builtin_amdgcn_tensor_store_from_lds(bit_cast(tdm_desc_grp.get(I0)), tdm_desc_grp.get(I1), tdm_desc_grp.get(I2), tdm_desc_grp.get(I3), diff --git a/projects/composablekernel/include/ck_tile/core/arch/amd_buffer_addressing_builtins.hpp b/projects/composablekernel/include/ck_tile/core/arch/amd_buffer_addressing_builtins.hpp index 01e81f68f129..9a4d8a4f5d6b 100644 --- a/projects/composablekernel/include/ck_tile/core/arch/amd_buffer_addressing_builtins.hpp +++ b/projects/composablekernel/include/ck_tile/core/arch/amd_buffer_addressing_builtins.hpp @@ -178,8 +178,8 @@ struct buffer_load<16, pre_nop> using mbuf_t = typename impl::buffer_load_trait<16, T>::payload_t; #if HAS_RAW_BUFFER_BUILTINS index_t s_offset = i_offset; - reinterpret_cast(value) = __builtin_amdgcn_raw_buffer_load_b128( - cast_to_amdgpu_buffer_rsrc_t(res), v_offset, s_offset, 0); + reinterpret_cast(value) = bit_cast(__builtin_amdgcn_raw_buffer_load_b128( + cast_to_amdgpu_buffer_rsrc_t(res), v_offset, s_offset, 0)); #else if constexpr(pre_nop) asm volatile("s_nop 4\n" @@ -212,8 +212,8 @@ struct buffer_load<8, pre_nop> using mbuf_t = typename impl::buffer_load_trait<8, T>::payload_t; #if HAS_RAW_BUFFER_BUILTINS index_t s_offset = i_offset; - reinterpret_cast(value) = __builtin_amdgcn_raw_buffer_load_b64( - cast_to_amdgpu_buffer_rsrc_t(res), v_offset, s_offset, 0); + reinterpret_cast(value) = bit_cast(__builtin_amdgcn_raw_buffer_load_b64( + cast_to_amdgpu_buffer_rsrc_t(res), v_offset, s_offset, 0)); #else if constexpr(pre_nop) asm volatile("s_nop 4\n" @@ -3369,7 +3369,6 @@ __device__ auto amd_transpose_load_to_vgpr(const T* __restrict__ in_ptr) auto lds_ptr = reinterpret_cast<__LDS_ADDR llvm_bf16x4_t*>(in_ptr_); return bit_cast>(__builtin_amdgcn_ds_read_tr16_b64_v4bf16(lds_ptr)); #elif defined(__gfx125__) - typedef __attribute__((__vector_size__(8 * sizeof(__bf16)))) __bf16 llvm_bf16x8_t; auto lds_ptr = reinterpret_cast<__LDS_ADDR llvm_bf16x8_t*>(in_ptr_); return bit_cast>(__builtin_amdgcn_ds_load_tr16_b128_v8bf16(lds_ptr)); #else @@ -3427,7 +3426,7 @@ amd_tdm_load(const TDMDescriptor& descriptor static constexpr auto I4 = number<4>{}; auto tdm_desc_grp = descriptor.getResourceDescriptorGroup(); - __builtin_amdgcn_tensor_load_to_lds(tdm_desc_grp.get(I0), + __builtin_amdgcn_tensor_load_to_lds(bit_cast(tdm_desc_grp.get(I0)), tdm_desc_grp.get(I1), tdm_desc_grp.get(I2), tdm_desc_grp.get(I3), @@ -3453,7 +3452,7 @@ amd_tdm_store(const TDMDescriptor& descripto static constexpr auto I4 = number<4>{}; auto tdm_desc_grp = descriptor.getResourceDescriptorGroup(); - __builtin_amdgcn_tensor_store_from_lds(tdm_desc_grp.get(I0), + __builtin_amdgcn_tensor_store_from_lds(bit_cast(tdm_desc_grp.get(I0)), tdm_desc_grp.get(I1), tdm_desc_grp.get(I2), tdm_desc_grp.get(I3), diff --git a/projects/composablekernel/include/ck_tile/core/arch/mma/mfma/mfma_gfx9.hpp b/projects/composablekernel/include/ck_tile/core/arch/mma/mfma/mfma_gfx9.hpp index c86ab9771ae2..f741491d933d 100644 --- a/projects/composablekernel/include/ck_tile/core/arch/mma/mfma/mfma_gfx9.hpp +++ b/projects/composablekernel/include/ck_tile/core/arch/mma/mfma/mfma_gfx9.hpp @@ -536,7 +536,8 @@ struct amdgcn_mma; - return {__builtin_amdgcn_mfma_f32_32x32x2bf16(aVec, bVec, cVec, P::cbsz, P::abid, P::blgp)}; + return {__builtin_amdgcn_mfma_f32_32x32x2bf16( + bit_cast(aVec), bit_cast(bVec), cVec, P::cbsz, P::abid, P::blgp)}; } }; @@ -555,7 +556,8 @@ struct amdgcn_mma; - return {__builtin_amdgcn_mfma_f32_32x32x2bf16(aVec, bVec, cVec, P::cbsz, P::abid, P::blgp)}; + return {__builtin_amdgcn_mfma_f32_32x32x2bf16( + bit_cast(aVec), bit_cast(bVec), cVec, P::cbsz, P::abid, P::blgp)}; } }; @@ -574,7 +576,8 @@ struct amdgcn_mma; - return {__builtin_amdgcn_mfma_f32_16x16x2bf16(aVec, bVec, cVec, P::cbsz, P::abid, P::blgp)}; + return {__builtin_amdgcn_mfma_f32_16x16x2bf16( + bit_cast(aVec), bit_cast(bVec), cVec, P::cbsz, P::abid, P::blgp)}; } }; @@ -593,7 +596,8 @@ struct amdgcn_mma; - return {__builtin_amdgcn_mfma_f32_16x16x2bf16(aVec, bVec, cVec, P::cbsz, P::abid, P::blgp)}; + return {__builtin_amdgcn_mfma_f32_16x16x2bf16( + bit_cast(aVec), bit_cast(bVec), cVec, P::cbsz, P::abid, P::blgp)}; } }; @@ -612,7 +616,8 @@ struct amdgcn_mma; - return {__builtin_amdgcn_mfma_f32_4x4x2bf16(aVec, bVec, cVec, P::cbsz, P::abid, P::blgp)}; + return {__builtin_amdgcn_mfma_f32_4x4x2bf16( + bit_cast(aVec), bit_cast(bVec), cVec, P::cbsz, P::abid, P::blgp)}; } }; @@ -631,7 +636,8 @@ struct amdgcn_mma; - return {__builtin_amdgcn_mfma_f32_4x4x2bf16(aVec, bVec, cVec, P::cbsz, P::abid, P::blgp)}; + return {__builtin_amdgcn_mfma_f32_4x4x2bf16( + bit_cast(aVec), bit_cast(bVec), cVec, P::cbsz, P::abid, P::blgp)}; } }; @@ -650,7 +656,8 @@ struct amdgcn_mma; - return {__builtin_amdgcn_mfma_f32_32x32x4bf16(aVec, bVec, cVec, P::cbsz, P::abid, P::blgp)}; + return {__builtin_amdgcn_mfma_f32_32x32x4bf16( + bit_cast(aVec), bit_cast(bVec), cVec, P::cbsz, P::abid, P::blgp)}; } }; @@ -669,7 +676,8 @@ struct amdgcn_mma; - return {__builtin_amdgcn_mfma_f32_16x16x8bf16(aVec, bVec, cVec, P::cbsz, P::abid, P::blgp)}; + return {__builtin_amdgcn_mfma_f32_16x16x8bf16( + bit_cast(aVec), bit_cast(bVec), cVec, P::cbsz, P::abid, P::blgp)}; } }; @@ -688,8 +696,8 @@ struct amdgcn_mma; - return { - __builtin_amdgcn_mfma_f32_32x32x4bf16_1k(aVec, bVec, cVec, P::cbsz, P::abid, P::blgp)}; + return {__builtin_amdgcn_mfma_f32_32x32x4bf16_1k( + bit_cast(aVec), bit_cast(bVec), cVec, P::cbsz, P::abid, P::blgp)}; } }; @@ -708,8 +716,8 @@ struct amdgcn_mma; - return { - __builtin_amdgcn_mfma_f32_32x32x4bf16_1k(aVec, bVec, cVec, P::cbsz, P::abid, P::blgp)}; + return {__builtin_amdgcn_mfma_f32_32x32x4bf16_1k( + bit_cast(aVec), bit_cast(bVec), cVec, P::cbsz, P::abid, P::blgp)}; } }; @@ -728,8 +736,8 @@ struct amdgcn_mma; - return { - __builtin_amdgcn_mfma_f32_16x16x4bf16_1k(aVec, bVec, cVec, P::cbsz, P::abid, P::blgp)}; + return {__builtin_amdgcn_mfma_f32_16x16x4bf16_1k( + bit_cast(aVec), bit_cast(bVec), cVec, P::cbsz, P::abid, P::blgp)}; } }; @@ -748,8 +756,8 @@ struct amdgcn_mma; - return { - __builtin_amdgcn_mfma_f32_16x16x4bf16_1k(aVec, bVec, cVec, P::cbsz, P::abid, P::blgp)}; + return {__builtin_amdgcn_mfma_f32_16x16x4bf16_1k( + bit_cast(aVec), bit_cast(bVec), cVec, P::cbsz, P::abid, P::blgp)}; } }; @@ -768,8 +776,8 @@ struct amdgcn_mma; - return { - __builtin_amdgcn_mfma_f32_4x4x4bf16_1k(aVec, bVec, cVec, P::cbsz, P::abid, P::blgp)}; + return {__builtin_amdgcn_mfma_f32_4x4x4bf16_1k( + bit_cast(aVec), bit_cast(bVec), cVec, P::cbsz, P::abid, P::blgp)}; } }; @@ -788,8 +796,8 @@ struct amdgcn_mma; - return { - __builtin_amdgcn_mfma_f32_4x4x4bf16_1k(aVec, bVec, cVec, P::cbsz, P::abid, P::blgp)}; + return {__builtin_amdgcn_mfma_f32_4x4x4bf16_1k( + bit_cast(aVec), bit_cast(bVec), cVec, P::cbsz, P::abid, P::blgp)}; } }; @@ -808,8 +816,8 @@ struct amdgcn_mma; - return { - __builtin_amdgcn_mfma_f32_32x32x8bf16_1k(aVec, bVec, cVec, P::cbsz, P::abid, P::blgp)}; + return {__builtin_amdgcn_mfma_f32_32x32x8bf16_1k( + bit_cast(aVec), bit_cast(bVec), cVec, P::cbsz, P::abid, P::blgp)}; } }; @@ -828,8 +836,8 @@ struct amdgcn_mma; - return { - __builtin_amdgcn_mfma_f32_16x16x16bf16_1k(aVec, bVec, cVec, P::cbsz, P::abid, P::blgp)}; + return {__builtin_amdgcn_mfma_f32_16x16x16bf16_1k( + bit_cast(aVec), bit_cast(bVec), cVec, P::cbsz, P::abid, P::blgp)}; } }; @@ -1182,8 +1190,12 @@ struct amdgcn_mma; - return { - __builtin_amdgcn_mfma_f32_16x16x32_bf16(aVec, bVec, cVec, P::cbsz, P::abid, P::blgp)}; + return {__builtin_amdgcn_mfma_f32_16x16x32_bf16(bit_cast(aVec), + bit_cast(bVec), + cVec, + P::cbsz, + P::abid, + P::blgp)}; } }; @@ -1222,8 +1234,12 @@ struct amdgcn_mma; - return { - __builtin_amdgcn_mfma_f32_32x32x16_bf16(aVec, bVec, cVec, P::cbsz, P::abid, P::blgp)}; + return {__builtin_amdgcn_mfma_f32_32x32x16_bf16(bit_cast(aVec), + bit_cast(bVec), + cVec, + P::cbsz, + P::abid, + P::blgp)}; } }; diff --git a/projects/composablekernel/include/ck_tile/core/arch/mma/sparse/mfma/sparse_gfx9.hpp b/projects/composablekernel/include/ck_tile/core/arch/mma/sparse/mfma/sparse_gfx9.hpp index 1119ab8178e1..445c317ae934 100644 --- a/projects/composablekernel/include/ck_tile/core/arch/mma/sparse/mfma/sparse_gfx9.hpp +++ b/projects/composablekernel/include/ck_tile/core/arch/mma/sparse/mfma/sparse_gfx9.hpp @@ -94,7 +94,8 @@ struct amdgcn_mma; - return {__builtin_amdgcn_smfmac_f32_16x16x32_bf16(aVec, bVec, cVec, idx, P::cbsz, P::abid)}; + return {__builtin_amdgcn_smfmac_f32_16x16x32_bf16( + bit_cast(aVec), bit_cast(bVec), cVec, idx, P::cbsz, P::abid)}; } }; @@ -113,7 +114,8 @@ struct amdgcn_mma; - return {__builtin_amdgcn_smfmac_f32_32x32x16_bf16(aVec, bVec, cVec, idx, P::cbsz, P::abid)}; + return {__builtin_amdgcn_smfmac_f32_32x32x16_bf16( + bit_cast(aVec), bit_cast(bVec), cVec, idx, P::cbsz, P::abid)}; } }; @@ -132,7 +134,8 @@ struct amdgcn_mma; - return {__builtin_amdgcn_smfmac_i32_16x16x64_i8(aVec, bVec, cVec, idx, P::cbsz, P::abid)}; + return {__builtin_amdgcn_smfmac_i32_16x16x64_i8( + bit_cast(aVec), bit_cast(bVec), cVec, idx, P::cbsz, P::abid)}; } }; @@ -151,7 +154,8 @@ struct amdgcn_mma; - return {__builtin_amdgcn_smfmac_i32_32x32x32_i8(aVec, bVec, cVec, idx, P::cbsz, P::abid)}; + return {__builtin_amdgcn_smfmac_i32_32x32x32_i8( + bit_cast(aVec), bit_cast(bVec), cVec, idx, P::cbsz, P::abid)}; } }; @@ -170,8 +174,8 @@ struct amdgcn_mma; - return { - __builtin_amdgcn_smfmac_f32_16x16x64_bf8_bf8(aVec, bVec, cVec, idx, P::cbsz, P::abid)}; + return {__builtin_amdgcn_smfmac_f32_16x16x64_bf8_bf8( + bit_cast(aVec), bit_cast(bVec), cVec, idx, P::cbsz, P::abid)}; } }; @@ -190,8 +194,8 @@ struct amdgcn_mma; - return { - __builtin_amdgcn_smfmac_f32_16x16x64_bf8_fp8(aVec, bVec, cVec, idx, P::cbsz, P::abid)}; + return {__builtin_amdgcn_smfmac_f32_16x16x64_bf8_fp8( + bit_cast(aVec), bit_cast(bVec), cVec, idx, P::cbsz, P::abid)}; } }; @@ -210,8 +214,8 @@ struct amdgcn_mma; - return { - __builtin_amdgcn_smfmac_f32_16x16x64_fp8_bf8(aVec, bVec, cVec, idx, P::cbsz, P::abid)}; + return {__builtin_amdgcn_smfmac_f32_16x16x64_fp8_bf8( + bit_cast(aVec), bit_cast(bVec), cVec, idx, P::cbsz, P::abid)}; } }; @@ -230,8 +234,8 @@ struct amdgcn_mma; - return { - __builtin_amdgcn_smfmac_f32_16x16x64_fp8_fp8(aVec, bVec, cVec, idx, P::cbsz, P::abid)}; + return {__builtin_amdgcn_smfmac_f32_16x16x64_fp8_fp8( + bit_cast(aVec), bit_cast(bVec), cVec, idx, P::cbsz, P::abid)}; } }; @@ -250,8 +254,8 @@ struct amdgcn_mma; - return { - __builtin_amdgcn_smfmac_f32_32x32x32_bf8_bf8(aVec, bVec, cVec, idx, P::cbsz, P::abid)}; + return {__builtin_amdgcn_smfmac_f32_32x32x32_bf8_bf8( + bit_cast(aVec), bit_cast(bVec), cVec, idx, P::cbsz, P::abid)}; } }; @@ -270,8 +274,8 @@ struct amdgcn_mma; - return { - __builtin_amdgcn_smfmac_f32_32x32x32_bf8_fp8(aVec, bVec, cVec, idx, P::cbsz, P::abid)}; + return {__builtin_amdgcn_smfmac_f32_32x32x32_bf8_fp8( + bit_cast(aVec), bit_cast(bVec), cVec, idx, P::cbsz, P::abid)}; } }; @@ -290,8 +294,8 @@ struct amdgcn_mma; - return { - __builtin_amdgcn_smfmac_f32_32x32x32_fp8_bf8(aVec, bVec, cVec, idx, P::cbsz, P::abid)}; + return {__builtin_amdgcn_smfmac_f32_32x32x32_fp8_bf8( + bit_cast(aVec), bit_cast(bVec), cVec, idx, P::cbsz, P::abid)}; } }; @@ -310,8 +314,8 @@ struct amdgcn_mma; - return { - __builtin_amdgcn_smfmac_f32_32x32x32_fp8_fp8(aVec, bVec, cVec, idx, P::cbsz, P::abid)}; + return {__builtin_amdgcn_smfmac_f32_32x32x32_fp8_fp8( + bit_cast(aVec), bit_cast(bVec), cVec, idx, P::cbsz, P::abid)}; } }; @@ -368,7 +372,12 @@ struct amdgcn_mma; - return {__builtin_amdgcn_smfmac_f32_16x16x64_bf16(aVec, bVec, cVec, idx, P::cbsz, P::abid)}; + return {__builtin_amdgcn_smfmac_f32_16x16x64_bf16(bit_cast(aVec), + bit_cast(bVec), + cVec, + idx, + P::cbsz, + P::abid)}; } }; @@ -387,7 +396,12 @@ struct amdgcn_mma; - return {__builtin_amdgcn_smfmac_f32_32x32x32_bf16(aVec, bVec, cVec, idx, P::cbsz, P::abid)}; + return {__builtin_amdgcn_smfmac_f32_32x32x32_bf16(bit_cast(aVec), + bit_cast(bVec), + cVec, + idx, + P::cbsz, + P::abid)}; } }; @@ -406,7 +420,8 @@ struct amdgcn_mma; - return {__builtin_amdgcn_smfmac_i32_16x16x128_i8(aVec, bVec, cVec, idx, P::cbsz, P::abid)}; + return {__builtin_amdgcn_smfmac_i32_16x16x128_i8( + bit_cast(aVec), bit_cast(bVec), cVec, idx, P::cbsz, P::abid)}; } }; @@ -425,7 +440,8 @@ struct amdgcn_mma; - return {__builtin_amdgcn_smfmac_i32_32x32x64_i8(aVec, bVec, cVec, idx, P::cbsz, P::abid)}; + return {__builtin_amdgcn_smfmac_i32_32x32x64_i8( + bit_cast(aVec), bit_cast(bVec), cVec, idx, P::cbsz, P::abid)}; } }; @@ -444,8 +460,8 @@ struct amdgcn_mma; - return { - __builtin_amdgcn_smfmac_f32_16x16x128_bf8_bf8(aVec, bVec, cVec, idx, P::cbsz, P::abid)}; + return {__builtin_amdgcn_smfmac_f32_16x16x128_bf8_bf8( + bit_cast(aVec), bit_cast(bVec), cVec, idx, P::cbsz, P::abid)}; } }; @@ -464,8 +480,8 @@ struct amdgcn_mma; - return { - __builtin_amdgcn_smfmac_f32_16x16x128_bf8_fp8(aVec, bVec, cVec, idx, P::cbsz, P::abid)}; + return {__builtin_amdgcn_smfmac_f32_16x16x128_bf8_fp8( + bit_cast(aVec), bit_cast(bVec), cVec, idx, P::cbsz, P::abid)}; } }; @@ -484,8 +500,8 @@ struct amdgcn_mma; - return { - __builtin_amdgcn_smfmac_f32_16x16x128_fp8_bf8(aVec, bVec, cVec, idx, P::cbsz, P::abid)}; + return {__builtin_amdgcn_smfmac_f32_16x16x128_fp8_bf8( + bit_cast(aVec), bit_cast(bVec), cVec, idx, P::cbsz, P::abid)}; } }; @@ -504,8 +520,8 @@ struct amdgcn_mma; - return { - __builtin_amdgcn_smfmac_f32_16x16x128_fp8_fp8(aVec, bVec, cVec, idx, P::cbsz, P::abid)}; + return {__builtin_amdgcn_smfmac_f32_16x16x128_fp8_fp8( + bit_cast(aVec), bit_cast(bVec), cVec, idx, P::cbsz, P::abid)}; } }; @@ -524,8 +540,8 @@ struct amdgcn_mma; - return { - __builtin_amdgcn_smfmac_f32_32x32x64_bf8_bf8(aVec, bVec, cVec, idx, P::cbsz, P::abid)}; + return {__builtin_amdgcn_smfmac_f32_32x32x64_bf8_bf8( + bit_cast(aVec), bit_cast(bVec), cVec, idx, P::cbsz, P::abid)}; } }; @@ -544,8 +560,8 @@ struct amdgcn_mma; - return { - __builtin_amdgcn_smfmac_f32_32x32x64_bf8_fp8(aVec, bVec, cVec, idx, P::cbsz, P::abid)}; + return {__builtin_amdgcn_smfmac_f32_32x32x64_bf8_fp8( + bit_cast(aVec), bit_cast(bVec), cVec, idx, P::cbsz, P::abid)}; } }; @@ -564,8 +580,8 @@ struct amdgcn_mma; - return { - __builtin_amdgcn_smfmac_f32_32x32x64_fp8_bf8(aVec, bVec, cVec, idx, P::cbsz, P::abid)}; + return {__builtin_amdgcn_smfmac_f32_32x32x64_fp8_bf8( + bit_cast(aVec), bit_cast(bVec), cVec, idx, P::cbsz, P::abid)}; } }; @@ -584,8 +600,8 @@ struct amdgcn_mma; - return { - __builtin_amdgcn_smfmac_f32_32x32x64_fp8_fp8(aVec, bVec, cVec, idx, P::cbsz, P::abid)}; + return {__builtin_amdgcn_smfmac_f32_32x32x64_fp8_fp8( + bit_cast(aVec), bit_cast(bVec), cVec, idx, P::cbsz, P::abid)}; } }; diff --git a/projects/composablekernel/include/ck_tile/core/arch/utility.hpp b/projects/composablekernel/include/ck_tile/core/arch/utility.hpp index 647f5b4435c2..28368cc05ecb 100644 --- a/projects/composablekernel/include/ck_tile/core/arch/utility.hpp +++ b/projects/composablekernel/include/ck_tile/core/arch/utility.hpp @@ -64,8 +64,8 @@ CK_TILE_DEVICE auto warp_shuffle_down_pair(const T& v_local) { static_assert(sizeof(T) == sizeof(int32_t), "wrong!"); - const int32x2_t x = __builtin_amdgcn_permlane32_swap( - bit_cast(v_local), bit_cast(v_local), false, false); + const int32x2_t x = bit_cast(__builtin_amdgcn_permlane32_swap( + bit_cast(v_local), bit_cast(v_local), false, false)); thread_buffer v; v(0) = bit_cast(x[0]); diff --git a/projects/composablekernel/include/ck_tile/core/numeric/vector_type.hpp b/projects/composablekernel/include/ck_tile/core/numeric/vector_type.hpp index 9f5acf2ebb73..c92ce213670e 100644 --- a/projects/composablekernel/include/ck_tile/core/numeric/vector_type.hpp +++ b/projects/composablekernel/include/ck_tile/core/numeric/vector_type.hpp @@ -81,6 +81,10 @@ using bf16x16_t = bfloat16_t __attribute__((ext_vector_type(16))); using bf16x32_t = bfloat16_t __attribute__((ext_vector_type(32))); using bf16x64_t = bfloat16_t __attribute__((ext_vector_type(64))); +// take native __bf16 regardless of whether bfloat16_t is spelled __bf16 or ushort +using llvm_bf16x8_t = __bf16 __attribute__((ext_vector_type(8))); +using llvm_bf16x16_t = __bf16 __attribute__((ext_vector_type(16))); + // i32 // using int32_t = ... using int32x2_t = int32_t __attribute__((ext_vector_type(2))); diff --git a/projects/composablekernel/include/ck_tile/ops/gemm/warp/warp_gemm_attribute_mfma_impl.hpp b/projects/composablekernel/include/ck_tile/ops/gemm/warp/warp_gemm_attribute_mfma_impl.hpp index 615d16e9af4f..c62aa188fb39 100644 --- a/projects/composablekernel/include/ck_tile/ops/gemm/warp/warp_gemm_attribute_mfma_impl.hpp +++ b/projects/composablekernel/include/ck_tile/ops/gemm/warp/warp_gemm_attribute_mfma_impl.hpp @@ -750,9 +750,9 @@ struct WarpGemmAttributeMfmaImplBf16Bf16F32M32N32K8 static_for<0, 2, 1>{}([&](auto k) { c_vec = __builtin_amdgcn_mfma_f32_32x32x4bf16( reinterpret_cast&>(a_vec) - .template get_as>()[number{}], + .template get_as>()[number{}], reinterpret_cast&>(b_vec) - .template get_as>()[number{}], + .template get_as>()[number{}], c_vec, 0, 0, @@ -783,9 +783,9 @@ struct WarpGemmAttributeMfmaImplBf16Bf16F32M32N32K8 static_for<0, 2, 1>{}([&](auto k) { c_vec = __builtin_amdgcn_mfma_f32_32x32x4bf16( reinterpret_cast&>(a_vec) - .template get_as>()[number{}], + .template get_as>()[number{}], reinterpret_cast&>(b_vec) - .template get_as>()[number{}], + .template get_as>()[number{}], c_vec, 0, 0, @@ -848,9 +848,9 @@ struct WarpGemmAttributeMfmaImplBf16Bf16F32M16N16K16 static_for<0, 2, 1>{}([&](auto k) { c_vec = __builtin_amdgcn_mfma_f32_16x16x8bf16( reinterpret_cast&>(a_vec) - .template get_as>()[number{}], + .template get_as>()[number{}], reinterpret_cast&>(b_vec) - .template get_as>()[number{}], + .template get_as>()[number{}], c_vec, 0, 0, @@ -881,9 +881,9 @@ struct WarpGemmAttributeMfmaImplBf16Bf16F32M16N16K16 static_for<0, 2, 1>{}([&](auto k) { c_vec = __builtin_amdgcn_mfma_f32_16x16x8bf16( reinterpret_cast&>(a_vec) - .template get_as>()[number{}], + .template get_as>()[number{}], reinterpret_cast&>(b_vec) - .template get_as>()[number{}], + .template get_as>()[number{}], c_vec, 0, 0, @@ -948,9 +948,9 @@ struct WarpGemmAttributeMfmaImplBf16Bf16F32M4N64K4 static_for<0, 2, 1>{}([&](auto k) { c_vec = __builtin_amdgcn_mfma_f32_4x4x2bf16( reinterpret_cast&>(a_vec) - .template get_as>()[number{}], + .template get_as>()[number{}], reinterpret_cast&>(b_vec) - .template get_as>()[number{}], + .template get_as>()[number{}], c_vec, 0, 0, @@ -981,9 +981,9 @@ struct WarpGemmAttributeMfmaImplBf16Bf16F32M4N64K4 static_for<0, 2, 1>{}([&](auto k) { c_vec = __builtin_amdgcn_mfma_f32_4x4x2bf16( reinterpret_cast&>(a_vec) - .template get_as>()[number{}], + .template get_as>()[number{}], reinterpret_cast&>(b_vec) - .template get_as>()[number{}], + .template get_as>()[number{}], c_vec, 0, 0, @@ -1048,9 +1048,9 @@ struct WarpGemmAttributeMfmaImplBf16Bf16F32M64N4K4 static_for<0, 2, 1>{}([&](auto k) { c_vec = __builtin_amdgcn_mfma_f32_4x4x2bf16( reinterpret_cast&>(a_vec) - .template get_as>()[number{}], + .template get_as>()[number{}], reinterpret_cast&>(b_vec) - .template get_as>()[number{}], + .template get_as>()[number{}], c_vec, 0, 0, @@ -1081,9 +1081,9 @@ struct WarpGemmAttributeMfmaImplBf16Bf16F32M64N4K4 static_for<0, 2, 1>{}([&](auto k) { c_vec = __builtin_amdgcn_mfma_f32_4x4x2bf16( reinterpret_cast&>(a_vec) - .template get_as>()[number{}], + .template get_as>()[number{}], reinterpret_cast&>(b_vec) - .template get_as>()[number{}], + .template get_as>()[number{}], c_vec, 0, 0, @@ -1255,9 +1255,9 @@ struct WarpGemmAttributeMfmaImplBf16Bf16F32M32N32K16 static_for<0, 2, 1>{}([&](auto k) { c_vec = __builtin_amdgcn_mfma_f32_32x32x8bf16_1k( reinterpret_cast&>(a_vec) - .template get_as>()[number{}], + .template get_as>()[number{}], reinterpret_cast&>(b_vec) - .template get_as>()[number{}], + .template get_as>()[number{}], c_vec, 0, 0, @@ -1267,9 +1267,9 @@ struct WarpGemmAttributeMfmaImplBf16Bf16F32M32N32K16 static_for<0, 4, 1>{}([&](auto k) { c_vec = __builtin_amdgcn_mfma_f32_32x32x4bf16( reinterpret_cast&>(a_vec) - .template get_as>()[number{}], + .template get_as>()[number{}], reinterpret_cast&>(b_vec) - .template get_as>()[number{}], + .template get_as>()[number{}], c_vec, 0, 0, @@ -1294,9 +1294,9 @@ struct WarpGemmAttributeMfmaImplBf16Bf16F32M32N32K16 static_for<0, 2, 1>{}([&](auto k) { c_vec = __builtin_amdgcn_mfma_f32_32x32x8bf16_1k( reinterpret_cast&>(a_vec) - .template get_as>()[number{}], + .template get_as>()[number{}], reinterpret_cast&>(b_vec) - .template get_as>()[number{}], + .template get_as>()[number{}], c_vec, 0, 0, @@ -1308,9 +1308,9 @@ struct WarpGemmAttributeMfmaImplBf16Bf16F32M32N32K16 static_for<0, 4, 1>{}([&](auto k) { c_vec = __builtin_amdgcn_mfma_f32_32x32x4bf16( reinterpret_cast&>(a_vec) - .template get_as>()[number{}], + .template get_as>()[number{}], reinterpret_cast&>(b_vec) - .template get_as>()[number{}], + .template get_as>()[number{}], c_vec, 0, 0, @@ -1994,7 +1994,9 @@ struct WarpGemmAttributeMfmaImpl_i32_32x32x16_i8 type_convert(reinterpret_cast&>(b_vec) .template get_as()[number{}]); - c_vec = __builtin_amdgcn_mfma_f32_32x32x2f32(a_f32, b_f32, c_vec, 0, 0, 0); + // this fallback accumulates an i32 C tile through the f32 MFMA + c_vec = bit_cast(__builtin_amdgcn_mfma_f32_32x32x2f32( + a_f32, b_f32, bit_cast(c_vec), 0, 0, 0)); }); #else ck_tile::ignore = c_vec; @@ -2110,7 +2112,8 @@ struct WarpGemmAttributeMfmaImpl_i32_16x16x64_i8 else { #if defined(__gfx950__) - c_vec = __builtin_amdgcn_mfma_i32_16x16x64_i8(a_vec, b_vec, c_vec, 0, 0, 0); + c_vec = __builtin_amdgcn_mfma_i32_16x16x64_i8( + bit_cast(a_vec), bit_cast(b_vec), c_vec, 0, 0, 0); #else ck_tile::ignore = c_vec; ck_tile::ignore = a_vec; @@ -2168,7 +2171,7 @@ struct WarpGemmAttributeMfmaImpl_i32_32x32x32_i8 { #if defined(__gfx95__) c_vec = __builtin_amdgcn_mfma_i32_32x32x32_i8( - a_vec, bit_cast(b_vec), c_vec, 0, 0, 0); + bit_cast(a_vec), bit_cast(b_vec), c_vec, 0, 0, 0); #else ck_tile::ignore = c_vec; ck_tile::ignore = a_vec; diff --git a/projects/composablekernel/include/ck_tile/ops/gemm/warp/warp_gemm_attribute_wmma_impl_16bit_traits.hpp b/projects/composablekernel/include/ck_tile/ops/gemm/warp/warp_gemm_attribute_wmma_impl_16bit_traits.hpp index 3cde01ac08b2..d84ee96ea04e 100644 --- a/projects/composablekernel/include/ck_tile/ops/gemm/warp/warp_gemm_attribute_wmma_impl_16bit_traits.hpp +++ b/projects/composablekernel/include/ck_tile/ops/gemm/warp/warp_gemm_attribute_wmma_impl_16bit_traits.hpp @@ -41,7 +41,8 @@ struct WmmaTraits wmma_intrinsic(const AVecType& a_vec, const BVecType& b_vec, const CVecType& c_vec) { #ifdef __gfx11__ - return __builtin_amdgcn_wmma_f32_16x16x16_bf16_w32(a_vec, b_vec, c_vec); + return __builtin_amdgcn_wmma_f32_16x16x16_bf16_w32( + bit_cast(a_vec), bit_cast(b_vec), c_vec); #else ck_tile::ignore = a_vec; ck_tile::ignore = b_vec; @@ -85,7 +86,8 @@ struct WmmaTraits wmma_intrinsic(const AVecType& a_vec, const BVecType& b_vec, const CVecType& c_vec) { #ifdef __gfx120__ - return __builtin_amdgcn_wmma_f32_16x16x16_bf16_w32_gfx12(a_vec, b_vec, c_vec); + return __builtin_amdgcn_wmma_f32_16x16x16_bf16_w32_gfx12( + bit_cast(a_vec), bit_cast(b_vec), c_vec); #else ck_tile::ignore = a_vec; ck_tile::ignore = b_vec; @@ -132,8 +134,14 @@ struct WmmaTraits { #ifdef __gfx125__ using P = WarpGemmParamsParser; - return __builtin_amdgcn_wmma_f32_16x16x32_bf16( - 0, a_vec, 0, b_vec, 0, c_vec, P::reuse_a, P::reuse_b); + return __builtin_amdgcn_wmma_f32_16x16x32_bf16(0, + bit_cast(a_vec), + 0, + bit_cast(b_vec), + 0, + c_vec, + P::reuse_a, + P::reuse_b); #else ck_tile::ignore = a_vec; ck_tile::ignore = b_vec; @@ -153,8 +161,14 @@ struct WmmaTraits { #ifdef __gfx125__ using P = WarpGemmParamsParser; - return __builtin_amdgcn_wmma_bf16f32_16x16x32_bf16( - 0, a_vec, 0, b_vec, 0, c_vec, P::reuse_a, P::reuse_b); + return __builtin_amdgcn_wmma_bf16f32_16x16x32_bf16(0, + bit_cast(a_vec), + 0, + bit_cast(b_vec), + 0, + c_vec, + P::reuse_a, + P::reuse_b); #else ck_tile::ignore = a_vec; ck_tile::ignore = b_vec; @@ -177,8 +191,14 @@ struct WmmaTraits { #ifdef __gfx125__ using P = WarpGemmParamsParser; - return __builtin_amdgcn_wmma_bf16_16x16x32_bf16( - 0, a_vec, 0, b_vec, 0, c_vec, P::reuse_a, P::reuse_b); + return __builtin_amdgcn_wmma_bf16_16x16x32_bf16(0, + bit_cast(a_vec), + 0, + bit_cast(b_vec), + 0, + c_vec, + P::reuse_a, + P::reuse_b); #else ck_tile::ignore = a_vec; ck_tile::ignore = b_vec;