ROCm/composable_kernel
1
0
Fork 0
mirror of https://github.com/ROCm/composable_kernel.git synced 2026-07-14 19:18:35 +00:00
Code Issues Packages Projects Releases Wiki Activity

fwd v1

Browse Source
This commit is contained in:
Qun Lin
2025-06-10 17:00:26 +08:00
parent 01c1138bfe
commit 169e4277b5
7 changed files with 538 additions and 497 deletions
Download Patch File Download Diff File Expand all files Collapse all files

8
example/20_grouped_conv_bwd_weight/device_grouped_conv_bwd_weight_dl_v4.hpp
View File

@@ -4,7 +4,7 @@
#include "ck/utility/common_header.hpp"
#include "ck/tensor_operation/gpu/device/device_grouped_conv_bwd_weight.hpp"
#include "ck/tensor_operation/gpu/element/unary_element_wise_operation.hpp"
//#include "ck/tensor_operation/gpu/element/unary_element_wise_operation.hpp"
#include "ck/host_utility/kernel_launch.hpp"
#define ENABLE_PIPELINE_V2 1
@@ -1057,9 +1057,9 @@ struct DeviceGroupedConvBwdWeightDlV4 : public DeviceGroupedConvBwdWeight<NDimSp
static_assert(FilterSize * FilterSize < 64);
static_assert(RequirePadding == false);
static_assert(NBatch % DstScalarPerVector == 0);
static_assert(is_same_v<InElementwiseOperation, element_wise::PassThrough>);
static_assert(is_same_v<WeiElementwiseOperation, element_wise::PassThrough>);
static_assert(is_same_v<OutElementwiseOperation, element_wise::PassThrough>);
//static_assert(is_same_v<InElementwiseOperation, element_wise::PassThrough>);
//static_assert(is_same_v<WeiElementwiseOperation, element_wise::PassThrough>);
//static_assert(is_same_v<OutElementwiseOperation, element_wise::PassThrough>);
using GridwiseConvBwdWeight = GridwiseGroupedConv2DBwdWeightDlV4<BlockSize,
InDataType,

3
example/30_grouped_conv_fwd_multiple_d/CMakeLists.txt
View File

@@ -21,3 +21,6 @@ endif() # USE_BITINT_EXTENSION_INT4
add_example_executable(example_grouped_conv_fwd_bias_relu_add_wmma_fp16 grouped_conv_fwd_bias_relu_add_wmma_fp16.cpp)
add_example_executable(example_grouped_conv_fwd_bias_relu_add_wmma_int8 grouped_conv_fwd_bias_relu_add_wmma_int8.cpp)
add_example_executable(example_grouped_conv_fwd_dl_fp16 grouped_conv_fwd_dl_fp16.cpp)
target_compile_options(example_grouped_conv_fwd_dl_fp16 PRIVATE -save-temps=obj -Wno-gnu-line-marker)

2
example/30_grouped_conv_fwd_multiple_d/common.hpp
View File

@@ -64,7 +64,7 @@ struct CommonLayoutSettingSelector<1> final
template <>
struct CommonLayoutSettingSelector<2> final
: CommonLayoutSetting<ctl::G_NHW_C, ctl::G_K_YX_C, ctl::G_NHW_K>
: CommonLayoutSetting<ctl::GNHWC, ctl::GKYXC, ctl::GNHWK>
{
};

924
example/30_grouped_conv_fwd_multiple_d/device_grouped_conv_fwd_dl_v4.hpp
View File

File diff suppressed because it is too large Load Diff

25
example/30_grouped_conv_fwd_multiple_d/grouped_conv_fwd_dl_fp16.cpp
View File

@@ -1,6 +1,7 @@
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2023, Advanced Micro Devices, Inc. All rights reserved.
#include "device_grouped_conv_fwd_dl_v4.hpp"
#include "common.hpp"
// kernel data types
@@ -19,6 +20,30 @@ using InElementOp = PassThrough;
using WeiElementOp = PassThrough;
using OutElementOp = PassThrough;
template <ck::index_t NDimSpatial>
using DeviceConvFwdInstance =
ck::tensor_operation::device::DeviceGroupedConvFwdDlV4<
NDimSpatial,
64,
InKernelDataType,
WeiKernelDataType,
AccDataType,
OutKernelDataType,
S<28, 28>,
5,
ck::Tuple<S<1,1>, S<1,1>, S<2,2>>,
InElementOp,
WeiElementOp,
OutElementOp,
2,
4,4,
4,
4,
false>;
#include "run_grouped_conv_fwd_example.inc"
int main(int argc, char* argv[]) { return !run_grouped_conv_fwd_example(argc, argv); }

3
example/30_grouped_conv_fwd_multiple_d/grouped_conv_fwd_xdl_fp16.cpp
View File

@@ -1,6 +1,6 @@
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2023, Advanced Micro Devices, Inc. All rights reserved.
#include "device_grouped_conv_fwd_dl_v4.hpp"
#include "common.hpp"
// kernel data types
@@ -19,6 +19,7 @@ using InElementOp = PassThrough;
using WeiElementOp = PassThrough;
using OutElementOp = PassThrough;
template <ck::index_t NDimSpatial>
using DeviceConvFwdInstance =
ck::tensor_operation::device::DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle<

70
example/30_grouped_conv_fwd_multiple_d/run_grouped_conv_fwd_example.inc
View File

@@ -1,54 +1,6 @@
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2023, Advanced Micro Devices, Inc. All rights reserved.
template <ck::index_t NDimSpatial>
using DeviceConvFwdInstance =
ck::tensor_operation::device::DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle<
NDimSpatial,
InputLayout<NDimSpatial>,
WeightLayout<NDimSpatial>,
ck::Tuple<>,
OutputLayout<NDimSpatial>,
InKernelDataType,
WeiKernelDataType,
AccDataType,
CShuffleDataType,
ck::Tuple<>,
OutKernelDataType,
InElementOp,
WeiElementOp,
OutElementOp,
ConvSpec, // ConvForwardSpecialization
GemmSpec, // GemmSpecialization
1, //
256, // BlockSize
128, // MPerBlock
256, // NPerBlock
16, // KPerBlock
4, // AK1
4, // BK1
32, // MPerXdl
32, // NPerXdl
2, // MXdlPerWave
4, // NXdlPerWave
S<4, 64, 1>, // ABlockTransferThreadClusterLengths_AK0_M_AK1
S<1, 0, 2>, // ABlockTransferThreadClusterArrangeOrder
S<1, 0, 2>, // ABlockTransferSrcAccessOrder
2, // ABlockTransferSrcVectorDim
4, // ABlockTransferSrcScalarPerVector
4, // ABlockTransferDstScalarPerVector_AK1
1, // ABlockLdsExtraM
S<4, 64, 1>, // BBlockTransferThreadClusterLengths_BK0_N_BK1
S<1, 0, 2>, // BBlockTransferThreadClusterArrangeOrder
S<1, 0, 2>, // BBlockTransferSrcAccessOrder
2, // BBlockTransferSrcVectorDim
4, // BBlockTransferSrcScalarPerVector
4, // BBlockTransferDstScalarPerVector_BK1
1, // BBlockLdsExtraN
1,
1,
S<1, 16, 1, 16>,
4>;
template <ck::index_t NDimSpatial>
using HostConvFwdInstance = ck::tensor_operation::host::ReferenceConvFwd<NDimSpatial,
@@ -65,9 +17,21 @@ bool run_grouped_conv_fwd(const ExecutionConfig& config,
{
static_assert(1 <= NDimSpatial && NDimSpatial <= 3, "Unsupported NDimSpatial");
const auto in_g_n_c_wis_desc = make_input_descriptor(conv_param);
const auto wei_g_k_c_xs_desc = make_weight_descriptor(conv_param);
const auto out_g_n_k_wos_desc = make_output_descriptor(conv_param);
const auto in_g_n_c_wis_desc =
ck::utils::conv::make_input_host_tensor_descriptor_g_n_c_wis_packed<
InputLayout<NDimSpatial>>(conv_param);
const auto wei_g_k_c_xs_desc =
ck::utils::conv::make_weight_host_tensor_descriptor_g_k_c_xs_packed<
WeightLayout<NDimSpatial>>(conv_param);
const auto out_g_n_k_wos_desc =
ck::utils::conv::make_output_host_tensor_descriptor_g_n_k_wos_packed<
OutputLayout<NDimSpatial>>(conv_param);
//const auto in_g_n_c_wis_desc = make_input_descriptor(conv_param);
//const auto wei_g_k_c_xs_desc = make_weight_descriptor(conv_param);
//const auto out_g_n_k_wos_desc = make_output_descriptor(conv_param);
Tensor<InUserDataType> in(in_g_n_c_wis_desc);
Tensor<WeiUserDataType> wei(wei_g_k_c_xs_desc);
@@ -214,9 +178,9 @@ bool run_grouped_conv_fwd_example(int argc, char* argv[])
switch(conv_param.num_dim_spatial_)
{
case 1: return run_grouped_conv_fwd<1>(config, conv_param);
//case 1: return run_grouped_conv_fwd<1>(config, conv_param);
case 2: return run_grouped_conv_fwd<2>(config, conv_param);
case 3: return run_grouped_conv_fwd<3>(config, conv_param);
//case 3: return run_grouped_conv_fwd<3>(config, conv_param);
}
return false;