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

refine test example_grouped_conv_bwd_weight_xdl_fp16

Browse Source
This commit is contained in:
Lin, Qun
2025-05-24 02:01:52 -05:00
committed by Qun Lin
parent 94ae9ae73f
commit f5377edfe5
7 changed files with 102 additions and 14 deletions
Download Patch File Download Diff File Expand all files Collapse all files

1
example/20_grouped_conv_bwd_weight/CMakeLists.txt
View File

@@ -1,6 +1,7 @@
add_custom_target(example_grouped_conv_bwd_weight)
add_example_executable(example_grouped_conv_bwd_weight_xdl_fp16 grouped_conv_bwd_weight_xdl_fp16.cpp)
add_example_dependencies(example_grouped_conv_bwd_weight example_grouped_conv_bwd_weight_xdl_fp16)
target_compile_options(example_grouped_conv_bwd_weight_xdl_fp16 PRIVATE -save-temps=obj -Wno-gnu-line-marker)
add_example_executable(example_grouped_conv_bwd_weight_xdl_bf16 grouped_conv_bwd_weight_xdl_bf16.cpp)
add_example_dependencies(example_grouped_conv_bwd_weight example_grouped_conv_bwd_weight_xdl_bf16)

10
example/20_grouped_conv_bwd_weight/common.hpp
View File

@@ -41,13 +41,15 @@ struct CommonLayoutSetting
using WeightLayout = WeightLay;
using OutputLayout = OutputLay;
};
using ALayout = ck::tensor_layout::convolution::NHWGC;
using BLayout = ck::tensor_layout::convolution::GKYXC;
using ELayout = ck::tensor_layout::convolution::NHWGK;
namespace ctl = ck::tensor_layout::convolution;
template <ck::index_t NDimSpatial>
struct CommonLayoutSettingSelector
: CommonLayoutSetting<ck::tuple_element_t<NDimSpatial - 1,
ck::Tuple<ck::tensor_layout::convolution::GNWC,
ck::tensor_layout::convolution::GNHWC,
ck::tensor_layout::convolution::NHWGC,
ck::tensor_layout::convolution::GNDHWC>>,
ck::tuple_element_t<NDimSpatial - 1,
ck::Tuple<ck::tensor_layout::convolution::GKXC,
@@ -55,7 +57,7 @@ struct CommonLayoutSettingSelector
ck::tensor_layout::convolution::GKZYXC>>,
ck::tuple_element_t<NDimSpatial - 1,
ck::Tuple<ck::tensor_layout::convolution::GNWK,
ck::tensor_layout::convolution::GNHWK,
ck::tensor_layout::convolution::NHWGK,
ck::tensor_layout::convolution::GNDHWK>>>
{
};
@@ -123,7 +125,7 @@ inline bool parse_cmd_args(int argc,
const ck::index_t num_dim_spatial = std::stoi(argv[4]);
conv_param = ck::utils::conv::parse_conv_param(
num_dim_spatial, threshold_to_catch_partial_args, argv);
num_dim_spatial, threshold_to_catch_partial_args + 1, argv);
}
else
{

85
example/20_grouped_conv_bwd_weight/grouped_conv_bwd_weight_xdl_fp16.cpp
View File

@@ -3,7 +3,9 @@
#include "common.hpp"
#include "ck/tensor_operation/gpu/device/impl/device_grouped_conv_bwd_weight_xdl_cshuffle.hpp"
//#include "ck/tensor_operation/gpu/device/impl/device_grouped_conv_bwd_weight_xdl_cshuffle.hpp"
#include "ck/utility/blkgemmpipe_scheduler.hpp"
#include "ck/tensor_operation/gpu/device/impl/device_grouped_conv_bwd_weight_two_stage_xdl_cshuffle.hpp"
using InDataType = F16;
using WeiDataType = F16;
@@ -14,6 +16,7 @@ using InElementOp = PassThrough;
using WeiElementOp = PassThrough;
using OutElementOp = PassThrough;
#if 0
template <ck::index_t NDimSpatial>
using DeviceConvBwdWeightInstance =
ck::tensor_operation::device::DeviceGroupedConvBwdWeight_Xdl_CShuffle<
@@ -51,20 +54,86 @@ using DeviceConvBwdWeightInstance =
S<0, 3, 1, 2>, // ABlockTransferThreadClusterArrangeOrder
S<0, 2, 1, 3>, // ABlockTransferSrcAccessOrder
2, // ABlockTransferSrcVectorDim
8, // ABlockTransferSrcScalarPerVector
2, // ABlockTransferDstScalarPerVector_K1
1, // ABlockTransferSrcScalarPerVector
1, // ABlockTransferDstScalarPerVector_K1
true, // ABlockLdsAddExtraM
S<1, 4, 16, 4>, // BBlockTransferThreadClusterLengths_K0_N_K1
S<0, 3, 1, 2>, // BBlockTransferThreadClusterArrangeOrder
S<0, 2, 1, 3>, // BBlockTransferSrcAccessOrder
2, // BBlockTransferSrcVectorDim
8, // BBlockTransferSrcScalarPerVector
2, // BBlockTransferDstScalarPerVector_K1
1, // BBlockTransferSrcScalarPerVector
1, // BBlockTransferDstScalarPerVector_K1
true, // BBlockLdsAddExtraN
1, // CShuffleMXdlPerWavePerShuffle
1, // CShuffleNXdlPerWavePerShuffle
S<1, 32, 1, 4>, // CBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock
128 / (sizeof(WeiDataType) * CHAR_BIT)>; // CBlockTransferScalarPerVector_NWaveNPerXdl
1>; // CBlockTransferScalarPerVector_NWaveNPerXdl
#endif
// DeviceGroupedConvBwdWeightTwoStage_Xdl_CShuffle< NDimSpatial, ALayout, BLayout, ELayout, F16, F16, F16, F32, PassThrough, PassThrough, PassThrough, ConvSpec, 64, 32, 32, 32, 8, 32, 32, 1, 1, S<4, 8, 1>, S<2, 0, 1>, S<1, 0, 2>, 1, 2, 2, false, S<4, 16, 1>, S<2, 0, 1>, S<1, 0, 2>, 1, 2, 2, false, 1, 1, S<1, 8, 1, 8>, 1, Scheduler, PipelineVersion, 2>,
//ConvBwdWeightDefault,
//is_NHWGC_GKYXC_NHWGK
using ALayout = ck::tensor_layout::convolution::NHWGC;
using BLayout = ck::tensor_layout::convolution::GKYXC;
using ELayout = ck::tensor_layout::convolution::NHWGK;
//using Scheduler =ck::BlockGemmPipelineScheduler::Intrawave;
//using PipelineVersion =ck::BlockGemmPipelineVersion::v1;
template <ck::index_t NDimSpatial>
using DeviceConvBwdWeightInstance =
ck::tensor_operation::device::DeviceGroupedConvBwdWeightTwoStage_Xdl_CShuffle<
NDimSpatial,
ALayout,
BLayout,
ELayout,
F16,
F16,
F16,
F32,
PassThrough,
PassThrough,
PassThrough,
ConvBwdWeightDefault,
64,
32,
64,
32,
8,
32,
32,
1,
2,
S<4, 8, 1>,
S<2, 0, 1>,
S<1, 0, 2>,
1,
2,
2,
false,
S<4, 16, 1>,
S<2, 0, 1>,
S<1, 0, 2>,
1,
2,
2,
false,
1,
1,
S<1, 8, 1, 8>,
1,
ck::BlockGemmPipelineScheduler::Intrawave,
ck::BlockGemmPipelineVersion::v2,
2>;
#if 0
DeviceGroupedConvBwdWeightTwoStage_Xdl_CShuffle< NDimSpatial, ALayout, BLayout, ELayout, F16, F16, F16, F32, PassThrough, PassThrough, PassThrough, ConvSpec, 64, 32, 64, 32, 8, 32, 32, 1, 2, S<4, 8, 1>, S<2, 0, 1>, S<1, 0, 2>, 1, 2, 2, false, S<4, 16, 1>, S<2, 0, 1>, S<1, 0, 2>, 1, 2, 2, false, 1, 1, S<1, 8, 1, 8>, 1, Scheduler, PipelineVersion, 2>,
64, 16, 16, 32, 8, 16, 16, 1, 1, S<4, 8, 1>, S<2, 0, 1>, S<1, 0, 2>, 1, 1, 4, false, S<4, 8, 1>,
S<2, 0, 1>, S<1, 0, 2>, 1, 1, 4, false, 1, 1, S<1, 8, 1, 8>, 1,
ck::BlockGemmPipelineScheduler::Intrawave, ck::BlockGemmPipelineVersion::v1, 8 > ;
#endif
template <ck::index_t NDimSpatial>
using HostConvBwdWeightInstance = ck::tensor_operation::host::ReferenceConvBwdWeight<NDimSpatial,
@@ -89,9 +158,9 @@ int main(int argc, char* argv[])
switch(conv_param.num_dim_spatial_)
{
case 1: return !run_grouped_conv_bwd_weight<1>(config, conv_param);
case 1: break;//return !run_grouped_conv_bwd_weight<1>(config, conv_param);
case 2: return !run_grouped_conv_bwd_weight<2>(config, conv_param);
case 3: return !run_grouped_conv_bwd_weight<3>(config, conv_param);
case 3: break;//return !run_grouped_conv_bwd_weight<3>(config, conv_param);
default: break;
}

5
example/20_grouped_conv_bwd_weight/run_grouped_conv_bwd_weight_example.inc
View File

@@ -95,6 +95,9 @@ bool run_grouped_conv_bwd_weight(const ExecutionConfig& config,
WeiElementOp{},
OutElementOp{},
split_k);
invoker.ShowInfo(argument);
DeviceMem gemm_workspace_dev(conv.GetWorkSpaceSize(&argument));
conv.SetWorkSpacePointer(&argument, gemm_workspace_dev.GetDeviceBuffer());
if(!conv.IsSupportedArgument(argument))
{
@@ -104,7 +107,7 @@ bool run_grouped_conv_bwd_weight(const ExecutionConfig& config,
return true;
}
invoker.Run(argument, StreamConfig{nullptr, false});
invoker.Run(argument, StreamConfig{nullptr, config.time_kernel});
if(config.do_verification)
{

1
include/ck/host_utility/device_prop.hpp
View File

@@ -54,6 +54,7 @@ inline std::string get_device_name()
inline bool is_xdl_supported()
{
return true;
return ck::get_device_name() == "gfx908" || ck::get_device_name() == "gfx90a" ||
ck::get_device_name() == "gfx942" || ck::get_device_name() == "gfx950";
}

12
include/ck/tensor_operation/gpu/device/impl/device_grouped_conv_bwd_weight_two_stage_xdl_cshuffle.hpp
View File

@@ -1422,6 +1422,8 @@ struct DeviceGroupedConvBwdWeightTwoStage_Xdl_CShuffle
float Run(const Argument& arg, const StreamConfig& stream_config = StreamConfig{})
{
#if 1
//#ifdef NDEBUG
float avg_time = 0.f;
auto launch_elementwise_kernel = [&]() {
const AccDataType* p_c_grid = type_convert<const AccDataType*>(arg.p_workspace_);
@@ -1543,6 +1545,11 @@ struct DeviceGroupedConvBwdWeightTwoStage_Xdl_CShuffle
avg_time += RunGemmV3(arg, stream_config);
avg_time += launch_elementwise_kernel();
return avg_time;
#else
ignore = arg;
ignore = stream_config;
return 0;
#endif
}
float Run(const BaseArgument* p_arg,
@@ -1596,6 +1603,8 @@ struct DeviceGroupedConvBwdWeightTwoStage_Xdl_CShuffle
}
if constexpr(NDimSpatial == 2)
{
static_assert(is_NHWGC_GKYXC_NHWGK<InLayout, WeiLayout, OutLayout>() ||
is_NGCHW_NGKHW<InLayout, WeiLayout, OutLayout>());
if constexpr(!(is_NHWGC_GKYXC_NHWGK<InLayout, WeiLayout, OutLayout>() ||
is_NGCHW_NGKHW<InLayout, WeiLayout, OutLayout>()))
{
@@ -1632,8 +1641,10 @@ struct DeviceGroupedConvBwdWeightTwoStage_Xdl_CShuffle
if constexpr(NumGroupsToMerge > 1)
{
// support only if whole M and N can be proccessed on one block
//static_assert(GemmM <= MPerBlock && GemmN <= NPerBlock);
if(!(GemmM <= MPerBlock && GemmN <= NPerBlock))
{
printf("%d, %d, %d, %d\n",GemmM,MPerBlock, GemmN,NPerBlock);
return false;
}
if(!(arg.Conv_C_ == 1 && arg.Conv_K_ == 1))
@@ -1669,6 +1680,7 @@ struct DeviceGroupedConvBwdWeightTwoStage_Xdl_CShuffle
}
// vector load A/B matrix from global memory
static_assert(ABlockTransferSrcVectorDim == 1 && BBlockTransferSrcVectorDim == 1);
if(!(ABlockTransferSrcVectorDim == 1 && BBlockTransferSrcVectorDim == 1))
{
return false;

2
include/ck/utility/generic_memory_space_atomic.hpp
View File

@@ -12,7 +12,7 @@ namespace ck {
// instantiate this template. The purpose is to make the implementation of atomic_add explicit for
// each datatype.
template <typename X>
__device__ X atomic_add(X* p_dst, const X& x);
__device__ X atomic_add(X* p_dst, const X& x) { ignore = p_dst; ignore = x; return 0;};
template <>
__device__ int32_t atomic_add<int32_t>(int32_t* p_dst, const int32_t& x)