From 3301ef501fbac2d25918bb432b6a0cc62a14dcfb Mon Sep 17 00:00:00 2001 From: apoorva Date: Thu, 19 Jun 2025 09:32:02 +0000 Subject: [PATCH] removed the changes as per the MR review comment. --- .../tensor_operation/gpu/warp/wmma_gemm.hpp | 12 -- include/ck/utility/amd_wmma.hpp | 151 ++++++++++-------- include/ck/utility/dtype_vector.hpp | 10 -- 3 files changed, 83 insertions(+), 90 deletions(-) diff --git a/include/ck/tensor_operation/gpu/warp/wmma_gemm.hpp b/include/ck/tensor_operation/gpu/warp/wmma_gemm.hpp index f1baf223a1..429df2413f 100644 --- a/include/ck/tensor_operation/gpu/warp/wmma_gemm.hpp +++ b/include/ck/tensor_operation/gpu/warp/wmma_gemm.hpp @@ -16,13 +16,11 @@ enum struct WmmaInstr wmma_f32_16x16x16_bf16, wmma_f16_16x16x16_f16, wmma_bf16_16x16x16_bf16, - wmma_i32_16x16x16_iu16, wmma_i32_16x16x16_iu8, wmma_i32_16x16x16_iu4, // gfx12 wmma_f32_16x16x16_f16_gfx12, wmma_f32_16x16x16_bf16_gfx12, - wmma_i32_16x16x16_iu16_gfx12, wmma_i32_16x16x16_iu8_gfx12, wmma_f32_16x16x16_f8f8_gfx12, wmma_f32_16x16x16_f8bf8_gfx12, @@ -592,16 +590,6 @@ struct WmmaSelector return WmmaInstr::wmma_bf16_16x16x16_bf16; } - template <> - constexpr auto GetWmma() - { -#ifdef __gfx12__ - return WmmaInstr::wmma_i32_16x16x16_iu16_gfx12; -#else - return WmmaInstr::wmma_i32_16x16x16_iu16; -#endif - } - template <> constexpr auto GetWmma() { diff --git a/include/ck/utility/amd_wmma.hpp b/include/ck/utility/amd_wmma.hpp index f6f6712b9a..e14c0d62a8 100644 --- a/include/ck/utility/amd_wmma.hpp +++ b/include/ck/utility/amd_wmma.hpp @@ -143,33 +143,6 @@ struct intrin_wmma_i32_16x16x16_iu8_w32<16, 16, neg_a, neg_b, clamp> } }; -// src: iu8, dst: i32 -template -struct intrin_wmma_i32_16x16x16_iu16_w32; - -template -struct intrin_wmma_i32_16x16x16_iu16_w32<16, 16, neg_a, neg_b, clamp> -{ - template - __device__ static void Run(const int16x16_t& reg_a, const int16x16_t& reg_b, FloatC& reg_c) - { -#if defined(__gfx11__) - reg_c.template AsType()(Number<0>{}) = - __builtin_amdgcn_wmma_i32_16x16x16_iu16_w32( - neg_a, - bit_cast(reg_a), - neg_b, - bit_cast(reg_b), - reg_c.template AsType()[Number<0>{}], - clamp); -#else - ignore = reg_a; - ignore = reg_b; - ignore = reg_c; -#endif - } -}; - /********************************WAVE64 MODE***********************************************/ template @@ -290,33 +263,6 @@ struct intrin_wmma_i32_16x16x16_iu8_w64<16, 16, neg_a, neg_b, clamp> } }; -// src: iu16, dst: i32 -template -struct intrin_wmma_i32_16x16x16_iu16_w64; - -template -struct intrin_wmma_i32_16x16x16_iu16_w64<16, 16, neg_a, neg_b, clamp> -{ - template - __device__ static void Run(const int16x16_t& reg_a, const int16x16_t& reg_b, FloatC& reg_c) - { -#if defined(__gfx11__) - reg_c.template AsType()(Number<0>{}) = - __builtin_amdgcn_wmma_i32_16x16x16_iu8_w64( - neg_a, - bit_cast(reg_a), - neg_b, - bit_cast(reg_b), - reg_c.template AsType()[Number<0>{}], - clamp); -#else - ignore = reg_a; - ignore = reg_b; - ignore = reg_c; -#endif - } -}; - // gfx12 /********************************WAVE32 MODE***********************************************/ @@ -395,25 +341,94 @@ struct intrin_wmma_i32_16x16x16_iu8_w32_gfx12<16, 16, neg_a, neg_b, clamp> } }; -// src: iu16, dst: i32 -template -struct intrin_wmma_i32_16x16x16_iu16_w32_gfx12; +// src: f8, f8, dst: fp32 +template +struct intrin_wmma_f32_16x16x16_f8f8_w32_gfx12; -template -struct intrin_wmma_i32_16x16x16_iu16_w32_gfx12<16, 16, neg_a, neg_b, clamp> +template <> +struct intrin_wmma_f32_16x16x16_f8f8_w32_gfx12<16, 16> { template - __device__ static void Run(const int16x8_t& reg_a, const int16x8_t& reg_b, FloatC& reg_c) + __device__ static void Run(const f8x8_t& reg_a, const f8x8_t& reg_b, FloatC& reg_c) { #if defined(__gfx12__) - reg_c.template AsType()(Number<0>{}) = - __builtin_amdgcn_wmma_i32_16x16x16_iu8_w32_gfx12( - neg_a, - bit_cast(reg_a), - neg_b, - bit_cast(reg_b), - reg_c.template AsType()[Number<0>{}], - clamp); + reg_c.template AsType()(Number<0>{}) = + __builtin_amdgcn_wmma_f32_16x16x16_fp8_fp8_w32_gfx12( + bit_cast(reg_a), + bit_cast(reg_b), + reg_c.template AsType()[Number<0>{}]); +#else + ignore = reg_a; + ignore = reg_b; + ignore = reg_c; +#endif + } +}; + +// src: f8, bf8, dst: fp32 +template +struct intrin_wmma_f32_16x16x16_f8bf8_w32_gfx12; + +template <> +struct intrin_wmma_f32_16x16x16_f8bf8_w32_gfx12<16, 16> +{ + template + __device__ static void Run(const f8x8_t& reg_a, const bf8x8_t& reg_b, FloatC& reg_c) + { +#if defined(__gfx12__) + reg_c.template AsType()(Number<0>{}) = + __builtin_amdgcn_wmma_f32_16x16x16_fp8_bf8_w32_gfx12( + bit_cast(reg_a), + bit_cast(reg_b), + reg_c.template AsType()[Number<0>{}]); +#else + ignore = reg_a; + ignore = reg_b; + ignore = reg_c; +#endif + } +}; + +// src: bf8, f8, dst: fp32 +template +struct intrin_wmma_f32_16x16x16_bf8f8_w32_gfx12; + +template <> +struct intrin_wmma_f32_16x16x16_bf8f8_w32_gfx12<16, 16> +{ + template + __device__ static void Run(const bf8x8_t& reg_a, const f8x8_t& reg_b, FloatC& reg_c) + { +#if defined(__gfx12__) + reg_c.template AsType()(Number<0>{}) = + __builtin_amdgcn_wmma_f32_16x16x16_bf8_fp8_w32_gfx12( + bit_cast(reg_a), + bit_cast(reg_b), + reg_c.template AsType()[Number<0>{}]); +#else + ignore = reg_a; + ignore = reg_b; + ignore = reg_c; +#endif + } +}; + +// src: bf8, bf8, dst: fp32 +template +struct intrin_wmma_f32_16x16x16_bf8bf8_w32_gfx12; + +template <> +struct intrin_wmma_f32_16x16x16_bf8bf8_w32_gfx12<16, 16> +{ + template + __device__ static void Run(const bf8x8_t& reg_a, const bf8x8_t& reg_b, FloatC& reg_c) + { +#if defined(__gfx12__) + reg_c.template AsType()(Number<0>{}) = + __builtin_amdgcn_wmma_f32_16x16x16_bf8_bf8_w32_gfx12( + bit_cast(reg_a), + bit_cast(reg_b), + reg_c.template AsType()[Number<0>{}]); #else ignore = reg_a; ignore = reg_b; diff --git a/include/ck/utility/dtype_vector.hpp b/include/ck/utility/dtype_vector.hpp index 3342e43a48..65eed0624c 100644 --- a/include/ck/utility/dtype_vector.hpp +++ b/include/ck/utility/dtype_vector.hpp @@ -2093,8 +2093,6 @@ struct vector_type()>> } }; -using int64_t = long; - // fp32 using float2_t = typename vector_type::type; using float4_t = typename vector_type::type; @@ -2126,14 +2124,6 @@ using int32x16_t = typename vector_type::type; using int32x32_t = typename vector_type::type; using int32x64_t = typename vector_type::type; -// i16 -using int16x2_t = typename vector_type::type; -using int16x4_t = typename vector_type::type; -using int16x8_t = typename vector_type::type; -using int16x16_t = typename vector_type::type; -using int16x32_t = typename vector_type::type; -using int16x64_t = typename vector_type::type; - // i8 using int8x2_t = typename vector_type::type; using int8x4_t = typename vector_type::type;