ROCm/composable_kernel
1
0
Fork 0
mirror of https://github.com/ROCm/composable_kernel.git synced 2026-07-15 11:34:54 +00:00
Code Issues Packages Projects Releases Wiki Activity

change kpack value

Browse Source
This commit is contained in:
qin letao
2025-04-25 06:18:05 +00:00
parent 114b403b6b
commit 037f474361
2 changed files with 16 additions and 69 deletions
Download Patch File Download Diff File Expand all files Collapse all files

65
example/65_gemm_multiply_multiply/gemm_multiply_multiply_xdl_fp16_bpreshuffle.cpp
View File

@@ -49,58 +49,10 @@ using D1Layout = Col;
using DsLayout = ck::Tuple<D0Layout, D1Layout>;
using ELayout = Row;
struct MultiplyMultiply
{
template <typename E, typename C, typename D0, typename D1>
__host__ __device__ constexpr void
operator()(E& e, const C& c, const D0& d0, const D1& d1) const;
template <>
__host__ __device__ constexpr void operator()<F16, float, float, float>(F16& e,
const float& c,
const float& d0,
const float& d1) const
{
const float x0_f = c * d0 * d1;
e = ck::type_convert<F16>(x0_f);
}
template <>
__host__ __device__ constexpr void operator()<BF16, float, float, float>(BF16& e,
const float& c,
const float& d0,
const float& d1) const
{
const float x0_f = c * d0 * d1;
e = ck::type_convert<BF16>(x0_f);
}
template <>
__host__ __device__ constexpr void operator()<ck::half_t, int, float, float>(
ck::half_t& e, const int& c, const float& d0, const float& d1) const
{
const float x0_f =
ck::type_convert<float>(c) * ck::type_convert<float>(d0) * ck::type_convert<float>(d1);
e = ck::type_convert<ck::half_t>(x0_f);
}
template <>
__host__ __device__ constexpr void operator()<ck::bhalf_t, int, float, float>(
ck::bhalf_t& e, const int& c, const float& d0, const float& d1) const
{
const float x0_f =
ck::type_convert<float>(c) * ck::type_convert<float>(d0) * ck::type_convert<float>(d1);
e = ck::type_convert<ck::bhalf_t>(x0_f);
}
};
template <typename DataType>
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<A0DataType>{-2, 2});
b0_k_n.GenerateTensorValue(GeneratorTensor_2<B0DataType>{0, 2});
d0_m_n.GenerateTensorValue(GeneratorTensor_2<D0DataType>{-2, 2});
d1_m_n.GenerateTensorValue(GeneratorTensor_2<D1DataType>{-2, 2});
a0_m_k.GenerateTensorValue(GeneratorTensor_3<A0DataType>{-2, 2});
b0_k_n.GenerateTensorValue(GeneratorTensor_3<B0DataType>{-2, 2});
d0_m_n.GenerateTensorValue(GeneratorTensor_3<D0DataType>{-2, 2});
d1_m_n.GenerateTensorValue(GeneratorTensor_3<D1DataType>{-2, 2});
break;
case 2:
a0_m_k.GenerateTensorValue(GeneratorTensor_1<A0DataType>{});
@@ -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;

20
profiler/include/profiler/profile_gemm_multiply_multiply_wp_impl.hpp
View File

@@ -28,7 +28,7 @@ template <typename InOutDataType>
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<ADataType>{-1, 2});
b_k_n.GenerateTensorValue(GeneratorTensor_2<BDataType>{-1, 2});
d0_m_n.GenerateTensorValue(GeneratorTensor_2<D0DataType>{-5, 5});
d1_m_n.GenerateTensorValue(GeneratorTensor_2<D1DataType>{-1, 1});
a_m_k.GenerateTensorValue(GeneratorTensor_3<ADataType>{-2, 2});
b_k_n.GenerateTensorValue(GeneratorTensor_3<BDataType>{-2, 2});
d0_m_n.GenerateTensorValue(GeneratorTensor_3<D0DataType>{-2, 2});
d1_m_n.GenerateTensorValue(GeneratorTensor_3<D1DataType>{-2, 2});
break;
default:
a_m_k.GenerateTensorValue(GeneratorTensor_3<ADataType>{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<EDataType, float>::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