From 037f474361cc6168a2e145b57197c06bbf25328f Mon Sep 17 00:00:00 2001 From: qin letao Date: Fri, 25 Apr 2025 06:18:05 +0000 Subject: [PATCH] change kpack value --- ...multiply_multiply_xdl_fp16_bpreshuffle.cpp | 65 ++----------------- ...profile_gemm_multiply_multiply_wp_impl.hpp | 20 +++--- 2 files changed, 16 insertions(+), 69 deletions(-) diff --git a/example/65_gemm_multiply_multiply/gemm_multiply_multiply_xdl_fp16_bpreshuffle.cpp b/example/65_gemm_multiply_multiply/gemm_multiply_multiply_xdl_fp16_bpreshuffle.cpp index 35c3fa41d8..82060a4da3 100644 --- a/example/65_gemm_multiply_multiply/gemm_multiply_multiply_xdl_fp16_bpreshuffle.cpp +++ b/example/65_gemm_multiply_multiply/gemm_multiply_multiply_xdl_fp16_bpreshuffle.cpp @@ -49,58 +49,10 @@ using D1Layout = Col; using DsLayout = ck::Tuple; using ELayout = Row; -struct MultiplyMultiply -{ - template - __host__ __device__ constexpr void - operator()(E& e, const C& c, const D0& d0, const D1& d1) const; - - template <> - __host__ __device__ constexpr void operator()(F16& e, - const float& c, - const float& d0, - const float& d1) const - { - const float x0_f = c * d0 * d1; - - e = ck::type_convert(x0_f); - } - - template <> - __host__ __device__ constexpr void operator()(BF16& e, - const float& c, - const float& d0, - const float& d1) const - { - const float x0_f = c * d0 * d1; - - e = ck::type_convert(x0_f); - } - - template <> - __host__ __device__ constexpr void operator()( - ck::half_t& e, const int& c, const float& d0, const float& d1) const - { - const float x0_f = - ck::type_convert(c) * ck::type_convert(d0) * ck::type_convert(d1); - - e = ck::type_convert(x0_f); - } - - template <> - __host__ __device__ constexpr void operator()( - ck::bhalf_t& e, const int& c, const float& d0, const float& d1) const - { - const float x0_f = - ck::type_convert(c) * ck::type_convert(d0) * ck::type_convert(d1); - - e = ck::type_convert(x0_f); - } -}; template void preShuffleBuffer(const DataType* src, DataType* dst, int N, int K, int NXdl) { - int KPack = 16; + int KPack = 16 / sizeof(DataType); int NLane = NXdl; int KLane = 64 / NLane; @@ -132,7 +84,7 @@ using PassThrough = ck::tensor_operation::element_wise::PassThrough; using AElementOp = PassThrough; using BElementOp = PassThrough; -using CDEElementOp = MultiplyMultiply; +using CDEElementOp = ck::tensor_operation::element_wise::MultiplyMultiply; static constexpr auto GemmSpec = ck::tensor_operation::device::GemmSpecialization::Default; @@ -262,10 +214,10 @@ int main(int argc, char* argv[]) { case 0: break; case 1: - a0_m_k.GenerateTensorValue(GeneratorTensor_2{-2, 2}); - b0_k_n.GenerateTensorValue(GeneratorTensor_2{0, 2}); - d0_m_n.GenerateTensorValue(GeneratorTensor_2{-2, 2}); - d1_m_n.GenerateTensorValue(GeneratorTensor_2{-2, 2}); + a0_m_k.GenerateTensorValue(GeneratorTensor_3{-2, 2}); + b0_k_n.GenerateTensorValue(GeneratorTensor_3{-2, 2}); + d0_m_n.GenerateTensorValue(GeneratorTensor_3{-2, 2}); + d1_m_n.GenerateTensorValue(GeneratorTensor_3{-2, 2}); break; case 2: a0_m_k.GenerateTensorValue(GeneratorTensor_1{}); @@ -386,10 +338,7 @@ int main(int argc, char* argv[]) e_device_buf.FromDevice(e_m_n_device_result.mData.data()); - return ck::utils::check_err( - e_m_n_device_result, e_m_n_host_result, "Error: Incorrect results!", 1e-3, 5e-2) - ? 0 - : 1; + return ck::utils::check_err(e_m_n_device_result, e_m_n_host_result) ? 0 : 1; } return 0; diff --git a/profiler/include/profiler/profile_gemm_multiply_multiply_wp_impl.hpp b/profiler/include/profiler/profile_gemm_multiply_multiply_wp_impl.hpp index 16cc2df3ec..8724a51a4c 100644 --- a/profiler/include/profiler/profile_gemm_multiply_multiply_wp_impl.hpp +++ b/profiler/include/profiler/profile_gemm_multiply_multiply_wp_impl.hpp @@ -28,7 +28,7 @@ template void preShuffleBuffer( const InOutDataType* src, InOutDataType* dst, int N, int K, int NXdl, int Knew) { - int KPack = 16; + int KPack = 16 / sizeof(InOutDataType); int NLane = NXdl; int KLane = 64 / NLane; @@ -154,10 +154,10 @@ bool profile_gemm_multiply_multiply_weight_preshuffle_impl(int do_verification, { case 0: break; case 1: - a_m_k.GenerateTensorValue(GeneratorTensor_2{-1, 2}); - b_k_n.GenerateTensorValue(GeneratorTensor_2{-1, 2}); - d0_m_n.GenerateTensorValue(GeneratorTensor_2{-5, 5}); - d1_m_n.GenerateTensorValue(GeneratorTensor_2{-1, 1}); + a_m_k.GenerateTensorValue(GeneratorTensor_3{-2, 2}); + b_k_n.GenerateTensorValue(GeneratorTensor_3{-2, 2}); + d0_m_n.GenerateTensorValue(GeneratorTensor_3{-2, 2}); + d1_m_n.GenerateTensorValue(GeneratorTensor_3{-2, 2}); break; default: a_m_k.GenerateTensorValue(GeneratorTensor_3{0.0, 1.0}); @@ -228,8 +228,7 @@ bool profile_gemm_multiply_multiply_weight_preshuffle_impl(int do_verification, AccDataType, AElementOp, BElementOp, - PassThrough, - ComputeDataType>; + PassThrough>; auto ref_gemm = ReferenceGemmInstance{}; auto ref_invoker = ref_gemm.MakeInvoker(); @@ -364,11 +363,10 @@ bool profile_gemm_multiply_multiply_weight_preshuffle_impl(int do_verification, << std::endl; } } - if(!pass) + if(pass) { - continue; + pass_count++; } - pass_count++; std::string op_name = op_ptr->GetTypeString(); float ave_time = invoker_ptr->Run(argument_ptr.get(), @@ -410,6 +408,7 @@ bool profile_gemm_multiply_multiply_weight_preshuffle_impl(int do_verification, } } } + std::cout << "\nPass instance: " << pass_count << " in: " << op_ptrs.size() << std::endl; if constexpr(is_same::value) { @@ -446,7 +445,6 @@ bool profile_gemm_multiply_multiply_weight_preshuffle_impl(int do_verification, std::cout << " BLayout = ColumnMajor"; } - std::cout << "\nPass instance: " << pass_count << " in: " << op_ptrs.size() << std::endl; std::cout << " M = " << M << " N = " << N << " K = " << K << " StrideA = " << StrideA << " StrideB = " << StrideB << " StrideE = " << StrideE << " KBatch = " << best_kbatch << " : " << best_ave_time << " ms, " << best_tflops << " TFlops, " << best_gb_per_sec