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

bwd weight dl v4 - v0

Browse Source
This commit is contained in:
Qun Lin
2025-06-02 10:39:13 +08:00
parent 044df51b56
commit b27dddfb59
3 changed files with 728 additions and 408 deletions
Download Patch File Download Diff File Expand all files Collapse all files

4
example/20_grouped_conv_bwd_weight/CMakeLists.txt
View File

@@ -20,3 +20,7 @@ add_example_dependencies(example_grouped_conv_bwd_weight example_grouped_conv_bw
add_example_executable(example_grouped_conv_bwd_weight_v3_xdl_fp16 grouped_conv_bwd_weight_v3_xdl_fp16.cpp)
add_example_dependencies(example_grouped_conv_bwd_weight example_grouped_conv_bwd_weight_v3_xdl_fp16)
add_example_executable(example_grouped_conv_bwd_weight_dl_v4_fp16 grouped_conv_bwd_weight_dl_v4_fp16.cpp)
target_compile_options(example_grouped_conv_bwd_weight_dl_v4_fp16 PRIVATE -save-temps=obj -Wno-gnu-line-marker)

1120
example/20_grouped_conv_bwd_weight/device_grouped_conv_bwd_weight_dl_v4.hpp
View File

File diff suppressed because it is too large Load Diff

12
example/20_grouped_conv_bwd_weight/grouped_conv_bwd_weight_dl_v4_fp16.cpp
View File

@@ -23,15 +23,25 @@ using ELayout = ck::tensor_layout::convolution::GNHWK;
template <ck::index_t NDimSpatial>
using DeviceConvBwdWeightInstance =
ck::tensor_operation::device::DeviceGroupedConvBwdWeightDlV4<NDimSpatial,
256,
ALayout,
BLayout,
ELayout,
InDataType,
WeiDataType,
OutDataType,
S<28, 28>,
5,
ck::Tuple<S<1,1>, S<1,1>, S<2,2>>,
InElementOp,
WeiElementOp,
OutElementOp>;
OutElementOp,
2, // N batch
1, // NumWavePerTile
4, // InScalarPerVector
4, // OutScalarPerVector
2, // DstScalarPerVector
false>;
template <ck::index_t NDimSpatial>
using HostConvBwdWeightInstance = ck::tensor_operation::host::ReferenceConvBwdWeight<NDimSpatial,