ROCm/composable_kernel
1
0
Fork 0
mirror of https://github.com/ROCm/composable_kernel.git synced 2026-06-30 11:47:48 +00:00
Code Issues Packages Projects Releases Wiki Activity

finish gelu,relu and silu

Browse Source
This commit is contained in:
letaoqin
2024-10-11 23:06:47 +08:00
parent 9857ca19c3
commit 52d313938a
4 changed files with 94 additions and 35 deletions
Download Patch File Download Diff File Expand all files Collapse all files

2
example/66_gemm_bias_activation/CMakeLists.txt
View File

@@ -1,6 +1,6 @@
set(GEMM_BIAS_ADD_SOURCES
gemm_bias_add_xdl_fp16.cpp
gemm_bias_add_fp16.cpp
gemm_bias_add_xdl_fp16.cpp
)
add_executable(example_gemm_bias_add_xdl_fp16 ${GEMM_BIAS_ADD_SOURCES})
target_link_libraries(example_gemm_bias_add_xdl_fp16 PRIVATE utility)

13
example/66_gemm_bias_activation/gemm_bias_add.hpp
View File

@@ -7,6 +7,8 @@
#include "ck/stream_config.hpp"
#include "ck/utility/data_type.hpp"
#include "ck/utility/type_convert.hpp"
#include "ck/tensor_operation/gpu/element/binary_element_wise_operation.hpp"
#include "ck/tensor_operation/gpu/element/unary_element_wise_operation.hpp"
namespace ck {
namespace impl {
@@ -63,6 +65,12 @@ struct AddActivation
} // namespace impl
} // namespace ck
using PassThrough = ck::tensor_operation::element_wise::PassThrough;
using Gelu = ck::tensor_operation::element_wise::Gelu;
using Relu = ck::tensor_operation::element_wise::Relu;
using Silu = ck::tensor_operation::element_wise::Silu;
using Sigmoid = ck::tensor_operation::element_wise::Sigmoid;
enum class ActivationType
{
Gelu = 0,
@@ -70,6 +78,7 @@ enum class ActivationType
Silu,
Swiglu,
Geglu,
Sigmoid,
Identity,
GeluNoneApproximate,
GeGluNoneApproximate,
@@ -86,4 +95,6 @@ struct GemmBiasAddArgs
ck::index_t K;
};
float gemm_bias_add_fp16(const GemmBiasAddArgs& args, const StreamConfig& config);
float gemm_bias_add_fp16(const GemmBiasAddArgs& args,
const StreamConfig& config,
ActivationType op_type);

56
example/66_gemm_bias_activation/gemm_bias_add_fp16.cpp
View File

@@ -23,16 +23,10 @@ using CShuffleDataType = F32;
using ALayout = Row;
using BLayout = Row;
using D0Layout = Row;
using DsLayout = ck::Tuple<D0Layout>;
using CLayout = Row;
using PassThrough = ck::tensor_operation::element_wise::PassThrough;
using Relu = ck::tensor_operation::element_wise::Relu;
using AElementOp = PassThrough;
using BElementOp = PassThrough;
using CDEElementOp = ck::impl::AddActivation<Relu>;
;
using AElementOp = PassThrough;
using BElementOp = PassThrough;
template <ck::index_t... Is>
using S = ck::Sequence<Is...>;
@@ -40,7 +34,7 @@ using S = ck::Sequence<Is...>;
static constexpr auto GemmSpec = ck::tensor_operation::device::GemmSpecialization::MNKPadding;
// clang-format off
template <typename ADataType, typename BDataType, typename DsDataType, typename CDataType>
template <typename DsLayout, typename ADataType, typename BDataType, typename DsDataType, typename CDataType, typename CDEElementOp>
using DeviceOpInstance_64_16_16_64 = ck::tensor_operation::device::DeviceGemmMultiD_Xdl_CShuffle_V3<
ALayout, BLayout, DsLayout, CLayout, ADataType, BDataType,
DsDataType, CDataType, AccDataType, CShuffleDataType,
@@ -57,7 +51,7 @@ using DeviceOpInstance_64_16_16_64 = ck::tensor_operation::device::DeviceGemmMul
1, 1,
S<1, 16, 1, 4>, S<4, 4>, ck::BlockGemmPipelineScheduler::Interwave, ck::BlockGemmPipelineVersion::v1, F16>;
template <typename ADataType, typename BDataType, typename DsDataType, typename CDataType>
template <typename DsLayout, typename ADataType, typename BDataType, typename DsDataType, typename CDataType, typename CDEElementOp>
using DeviceOpInstance_default = ck::tensor_operation::device::DeviceGemmMultiD_Xdl_CShuffle_V3<
ALayout, BLayout, DsLayout, CLayout, ADataType, BDataType,
DsDataType, CDataType, AccDataType, CShuffleDataType,
@@ -75,8 +69,8 @@ using DeviceOpInstance_default = ck::tensor_operation::device::DeviceGemmMultiD_
S<1, 16, 1, 4>, S<2, 2>, ck::BlockGemmPipelineScheduler::Interwave, ck::BlockGemmPipelineVersion::v1, F16>;
// clang-format on
float gemm_bias_add_fp16(const GemmBiasAddArgs& args, const StreamConfig& config)
template <typename DsLayout, typename CDEElementOp>
float run_impl(const GemmBiasAddArgs& args, const StreamConfig& config)
{
using ADataType = ck::half_t;
using BDataType = ck::half_t;
@@ -133,12 +127,46 @@ float gemm_bias_add_fp16(const GemmBiasAddArgs& args, const StreamConfig& config
return true;
};
auto gemm = DeviceOpInstance_64_16_16_64<ADataType, BDataType, DsDataType, CDataType>{};
auto gemm = DeviceOpInstance_64_16_16_64<DsLayout,
ADataType,
BDataType,
DsDataType,
CDataType,
CDEElementOp>{};
if(!Run(gemm))
{
auto gemm_def = DeviceOpInstance_default<ADataType, BDataType, DsDataType, CDataType>{};
auto gemm_def = DeviceOpInstance_default<DsLayout,
ADataType,
BDataType,
DsDataType,
CDataType,
CDEElementOp>{};
Run(gemm_def);
}
return ave_time;
}
float gemm_bias_add_fp16(const GemmBiasAddArgs& args,
const StreamConfig& config,
ActivationType op_type)
{
using DsLayout = ck::Tuple<D0Layout>;
switch(op_type)
{
case ActivationType::Gelu:
case ActivationType::Geglu:
case ActivationType::GeluNoneApproximate:
case ActivationType::GeGluNoneApproximate:
return run_impl<DsLayout, ck::impl::AddActivation<Gelu>>(args, config);
case ActivationType::Relu:
return run_impl<DsLayout, ck::impl::AddActivation<Relu>>(args, config);
case ActivationType::Silu:
case ActivationType::Swiglu:
return run_impl<DsLayout, ck::impl::AddActivation<Silu>>(args, config);
case ActivationType::Sigmoid:
return run_impl<DsLayout, ck::impl::AddActivation<Sigmoid>>(args, config);
case ActivationType::Identity:
case ActivationType::InvalidType:
default: return 0;
}
}

58
example/66_gemm_bias_activation/gemm_bias_add_xdl_fp16.cpp
View File

@@ -13,8 +13,6 @@
#include "ck/library/utility/host_tensor_generator.hpp"
#include "ck/library/utility/literals.hpp"
#include "ck/library/reference_tensor_operation/cpu/reference_gemm.hpp"
#include "ck/tensor_operation/gpu/element/binary_element_wise_operation.hpp"
#include "ck/tensor_operation/gpu/element/unary_element_wise_operation.hpp"
#include "ck/library/utility/check_err.hpp"
#include "ck/tensor_operation/gpu/device/tensor_layout.hpp"
@@ -39,12 +37,14 @@ using D0Layout = Row;
using DsLayout = ck::Tuple<D0Layout>;
using ELayout = Row;
using PassThrough = ck::tensor_operation::element_wise::PassThrough;
using Relu = ck::tensor_operation::element_wise::Relu;
// using PassThrough = ck::tensor_operation::element_wise::PassThrough;
// using Gelu = ck::tensor_operation::element_wise::Gelu;
// using Relu = ck::tensor_operation::element_wise::Relu;
// using Silu = ck::tensor_operation::element_wise::Silu;
// using Sigmoid = ck::tensor_operation::element_wise::Sigmoid;
using AElementOp = PassThrough;
using BElementOp = PassThrough;
using CElementOp = ck::impl::AddActivation<Relu>;
using ReferenceGemmInstance = ck::tensor_operation::host::ReferenceGemm<A0DataType,
B0DataType,
@@ -139,6 +139,7 @@ int main(int argc, char* argv[])
bool do_verification = true;
int init_method = 1;
bool time_kernel = true;
int op_type = 0;
// GEMM shape
ck::index_t M = 64;
@@ -155,22 +156,24 @@ int main(int argc, char* argv[])
init_method = std::stoi(argv[2]);
time_kernel = std::stoi(argv[3]);
}
else if(argc == 7)
else if(argc == 8)
{
do_verification = std::stoi(argv[1]);
init_method = std::stoi(argv[2]);
time_kernel = std::stoi(argv[3]);
M = std::stoi(argv[4]);
N = std::stoi(argv[5]);
K = std::stoi(argv[6]);
M = std::stoi(argv[4]);
N = std::stoi(argv[5]);
K = std::stoi(argv[6]);
op_type = std::stoi(argv[7]);
}
else
{
printf("arg1: verification (0=no, 1=yes)\n");
printf("arg2: initialization (0=no init, 1=integer value, 2=decimal value)\n");
printf("arg3: time kernel (0=no, 1=yes)\n");
printf("arg4 to 9: M (256x), N(128x), K(32x)m\n");
printf("arg4 to 9: M (256x), N(128x), K(32x)m, op_type(Gelu = 0, Relu, Silu, Swiglu, "
"Geglu, Identity, GeluNoneApproximate, GeGluNoneApproximate)\n");
exit(0);
}
@@ -235,7 +238,9 @@ int main(int argc, char* argv[])
N,
K};
float ave_time = gemm_bias_add_fp16(gemm_args, StreamConfig{nullptr, time_kernel, 20, 50});
float ave_time = gemm_bias_add_fp16(gemm_args,
StreamConfig{nullptr, time_kernel, 20, 50},
static_cast<ActivationType>(op_type));
// float ave_time = invoker.Run(argument, StreamConfig{nullptr, time_kernel, 20, 50});
std::size_t flop = std::size_t(2) * M * N * K;
@@ -253,9 +258,6 @@ int main(int argc, char* argv[])
if(do_verification)
{
// RunUnfusedTest(a0_m_k.mData, b0_k_n.mData, d0_m_n.mData, e_m_n_host_result.mData, K, M,
// N);
auto ref_gemm = ReferenceGemmInstance{};
auto ref_invoker = ref_gemm.MakeInvoker();
@@ -264,13 +266,31 @@ int main(int argc, char* argv[])
ref_invoker.Run(ref_argument);
CElementOp cde_element_op;
for(int m = 0; m < M; ++m)
{
for(int n = 0; n < N; ++n)
auto run_elementwise = [&](auto cde_element_op) {
for(int m = 0; m < M; ++m)
{
cde_element_op(e_m_n_host_result(m, n), e_m_n_host_result(m, n), d0_m_n(m, n));
for(int n = 0; n < N; ++n)
{
cde_element_op(e_m_n_host_result(m, n), e_m_n_host_result(m, n), d0_m_n(m, n));
}
}
};
ActivationType type = static_cast<ActivationType>(op_type);
switch(type)
{
case ActivationType::Gelu:
case ActivationType::Geglu:
case ActivationType::GeluNoneApproximate:
case ActivationType::GeGluNoneApproximate:
run_elementwise(ck::impl::AddActivation<Gelu>{});
break;
case ActivationType::Relu: run_elementwise(ck::impl::AddActivation<Relu>{}); break;
case ActivationType::Silu:
case ActivationType::Swiglu: run_elementwise(ck::impl::AddActivation<Silu>{}); break;
case ActivationType::Sigmoid: run_elementwise(ck::impl::AddActivation<Sigmoid>{}); break;
case ActivationType::Identity:
case ActivationType::InvalidType:
default: break;
}
e_device_buf.FromDevice(e_m_n_device_result.mData.data());