ROCm/composable_kernel
1
0
Fork 0
mirror of https://github.com/ROCm/composable_kernel.git synced 2026-06-29 19:28:33 +00:00
Code Issues Packages Projects Releases Wiki Activity

add a tuning instance

Browse Source
This commit is contained in:
letaoqin
2024-10-15 16:19:30 +08:00
parent 12c3c3c0d4
commit 3b22355a7e
1 changed files with 41 additions and 5 deletions
Download Patch File Download Diff File Expand all files Collapse all files

46
example/66_gemm_bias_activation/gemm_bias_add_fp16.cpp
View File

@@ -36,7 +36,7 @@ static constexpr auto GemmSpec = ck::tensor_operation::device::GemmSpecializatio
// clang-format off
template <typename ADataType, typename BDataType, typename DsDataType, typename CDataType, typename CDEElementOp>
using DeviceOpInstance_64_16_16_64 = ck::tensor_operation::device::DeviceGemmMultiD_Xdl_CShuffle_V3<
using DeviceOpInstance_128_32_64_64 = ck::tensor_operation::device::DeviceGemmMultiD_Xdl_CShuffle_V3<
ALayout, BLayout, DsLayout, CLayout, ADataType, BDataType,
DsDataType, CDataType, AccDataType, CShuffleDataType,
AElementOp, BElementOp, CDEElementOp, GemmSpec,
@@ -52,6 +52,24 @@ using DeviceOpInstance_64_16_16_64 = ck::tensor_operation::device::DeviceGemmMul
1, 1,
S<1, 16, 1, 8>, S<8, 8>, ck::BlockGemmPipelineScheduler::Interwave, ck::BlockGemmPipelineVersion::v2, F16>;
template <typename ADataType, typename BDataType, typename DsDataType, typename CDataType, typename CDEElementOp>
using DeviceOpInstance_256_128_128_64 = ck::tensor_operation::device::DeviceGemmMultiD_Xdl_CShuffle_V3<
ALayout, BLayout, DsLayout, CLayout, ADataType, BDataType,
DsDataType, CDataType, AccDataType, CShuffleDataType,
AElementOp, BElementOp, CDEElementOp, GemmSpec,
256,
128, 128, 64,
8, 4,
32, 32,
2, 2,
S<8, 32, 1>, S<1, 0, 2>, S<1, 0, 2>,
2, 8, 8, 0,
S<16, 16, 1>, S<0, 2, 1>, S<0, 2, 1>,
1, 8, 4, 0,
1, 1,
S<1, 32, 1, 8>, S<8, 8>,
ck::BlockGemmPipelineScheduler::Intrawave, ck::BlockGemmPipelineVersion::v3, F16>;
template <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,
@@ -128,14 +146,32 @@ float run_impl(const GemmBiasAddArgs& args, const StreamConfig& config)
return true;
};
auto gemm =
DeviceOpInstance_64_16_16_64<ADataType, BDataType, DsDataType, CDataType, CDEElementOp>{};
if(!Run(gemm))
do
{
if(args.M <= 512)
{
auto gemm = DeviceOpInstance_128_32_64_64<ADataType,
BDataType,
DsDataType,
CDataType,
CDEElementOp>{};
if(Run(gemm))
break;
}
else
{
auto gemm = DeviceOpInstance_256_128_128_64<ADataType,
BDataType,
DsDataType,
CDataType,
CDEElementOp>{};
if(Run(gemm))
break;
}
auto gemm_def =
DeviceOpInstance_default<ADataType, BDataType, DsDataType, CDataType, CDEElementOp>{};
Run(gemm_def);
}
} while(0);
return ave_time;
}