mirror of
https://github.com/ROCm/composable_kernel.git
synced 2026-05-20 21:09:08 +00:00
Unified implementation of 1d/2d/3d conv bwd-data. fp32/fp16/bfp16/int8 (#134)
* start convnd bwd data
* add 3d laoyout name
* add conv1d reference
* add con3d reference
* finished example client code
* conv1d kernel finished
* fix input error
* add conv3d
* add 3d layout in conv_utils.hpp
* fix sepecial check
* addconvnd lib
* add test for bwd data
* finished test
* add check slice length
* convnd bwd data start
* profiler can be compiled
* fix some bug
* set input to zero
* modify readme for example
* fix test_convnd_bwd_data bug
* test_convnd_bwd_data parameter desc
* workaround for 1d
* workaroud for 2d
* change init value
* workaround for 3d int8
* fix init value bug
* remove workaround
* fix acc data type
* add int32
* change select function to template
* tilda to tilde
* remove int32 instance
* fix commit for device hpp
* fix comments for profiler
* using profile imp to test
* add pass verification
* fix conv2d reference
* fix conflict
* remove double batched_gemm
* fix exampel conv2d data and test convnd
* format
* change conv2d_bwd_data return value
* remove repeat = 1
* remove conv bwd data
Co-authored-by: ltqin <letaoqin@amd.com>
Co-authored-by: Chao Liu <chao.liu2@amd.com>
[ROCm/composable_kernel commit: 0536f2b312]
This commit is contained in:
ltqin
@@ -41,5 +41,4 @@ add_subdirectory(gemm_reduce)
|
||||
add_subdirectory(batched_gemm)
|
||||
add_subdirectory(grouped_gemm)
|
||||
add_subdirectory(convnd_fwd)
|
||||
add_subdirectory(conv2d_bwd_data)
|
||||
add_subdirectory(reduce)
|
||||
|
||||
@@ -121,15 +121,17 @@ int main(int argc, char* argv[])
|
||||
exit(1);
|
||||
}
|
||||
|
||||
auto Run = [&](auto input_type, auto wei_type, auto out_type) {
|
||||
auto Run = [&](auto input_type, auto wei_type, auto out_type, auto acc_type) {
|
||||
using InDataType = decltype(input_type);
|
||||
using WeiDataType = decltype(wei_type);
|
||||
using OutDataType = decltype(out_type);
|
||||
using AccDataType = decltype(acc_type);
|
||||
|
||||
using ReferenceConvBwdInstance =
|
||||
ck::tensor_operation::host::ReferenceConvBwdData<InDataType,
|
||||
WeiDataType,
|
||||
OutDataType,
|
||||
AccDataType,
|
||||
InElementOp,
|
||||
WeiElementOp,
|
||||
OutElementOp>;
|
||||
@@ -293,33 +295,33 @@ int main(int argc, char* argv[])
|
||||
if(success)
|
||||
{
|
||||
std::cout << "test conv2d bwd : Pass" << std::endl;
|
||||
return 0;
|
||||
}
|
||||
else
|
||||
{
|
||||
std::cout << "test conv2d bwd: Fail " << std::endl;
|
||||
return -1;
|
||||
}
|
||||
};
|
||||
|
||||
if(data_type == 0)
|
||||
{
|
||||
Run(F32(), F32(), F32());
|
||||
return Run(F32(), F32(), F32(), F32());
|
||||
}
|
||||
else if(data_type == 1)
|
||||
{
|
||||
Run(F16(), F16(), F16());
|
||||
return Run(F16(), F16(), F16(), F32());
|
||||
}
|
||||
else if(data_type == 2)
|
||||
{
|
||||
Run(BF16(), BF16(), BF16());
|
||||
return Run(BF16(), BF16(), BF16(), F32());
|
||||
}
|
||||
else if(data_type == 3)
|
||||
{
|
||||
Run(INT8(), INT8(), INT8());
|
||||
return Run(INT8(), INT8(), INT8(), int());
|
||||
}
|
||||
else
|
||||
{
|
||||
return 1;
|
||||
}
|
||||
|
||||
return 0;
|
||||
}
|
||||
|
||||
8
test/convnd_bwd_data/CMakeLists.txt
Normal file
8
test/convnd_bwd_data/CMakeLists.txt
Normal file
@@ -0,0 +1,8 @@
|
||||
include_directories(BEFORE
|
||||
${PROJECT_SOURCE_DIR}/profiler/include
|
||||
${PROJECT_SOURCE_DIR}/external/include/half
|
||||
)
|
||||
|
||||
add_test_executable(test_convnd_bwd_data convnd_bwd_data.cpp)
|
||||
target_link_libraries(test_convnd_bwd_data PRIVATE host_tensor)
|
||||
target_link_libraries(test_convnd_bwd_data PRIVATE device_convnd_bwd_data_instance)
|
||||
330
test/convnd_bwd_data/convnd_bwd_data.cpp
Normal file
330
test/convnd_bwd_data/convnd_bwd_data.cpp
Normal file
@@ -0,0 +1,330 @@
|
||||
#include <iostream>
|
||||
#include <numeric>
|
||||
#include <initializer_list>
|
||||
#include <cstdlib>
|
||||
#include <stdlib.h>
|
||||
#include <half.hpp>
|
||||
#include <vector>
|
||||
|
||||
#include "profile_convnd_bwd_data_impl.hpp"
|
||||
|
||||
int main()
|
||||
{
|
||||
bool pass = true;
|
||||
// check 1d
|
||||
std::vector<ck::conv_util::ConvParams> params;
|
||||
params.push_back({1, 128, 128, 256, {1}, {14}, {2}, {1}, {0}, {0}});
|
||||
params.push_back({1, 128, 128, 256, {3}, {28}, {1}, {1}, {1}, {1}});
|
||||
params.push_back({1, 128, 128, 256, {1}, {3}, {1}, {1}, {0}, {0}});
|
||||
|
||||
for(auto& param : params)
|
||||
{
|
||||
pass &= ck::profiler::profile_convnd_bwd_data_impl<1,
|
||||
float,
|
||||
float,
|
||||
float,
|
||||
float,
|
||||
ck::tensor_layout::convolution::NWC,
|
||||
ck::tensor_layout::convolution::KXC,
|
||||
ck::tensor_layout::convolution::NWK>(
|
||||
1, // do_verification,
|
||||
1, // init_method,
|
||||
0, // do_log,
|
||||
1, // nrepeat,
|
||||
param.N,
|
||||
param.K,
|
||||
param.C,
|
||||
param.input_spatial_lengths,
|
||||
param.filter_spatial_lengths,
|
||||
param.GetOutputSpatialLengths(),
|
||||
param.conv_filter_strides,
|
||||
param.conv_filter_dilations,
|
||||
param.input_left_pads,
|
||||
param.input_right_pads);
|
||||
|
||||
pass &= ck::profiler::profile_convnd_bwd_data_impl<1,
|
||||
ck::half_t,
|
||||
ck::half_t,
|
||||
ck::half_t,
|
||||
float,
|
||||
ck::tensor_layout::convolution::NWC,
|
||||
ck::tensor_layout::convolution::KXC,
|
||||
ck::tensor_layout::convolution::NWK>(
|
||||
1, // do_verification,
|
||||
1, // init_method,
|
||||
0, // do_log,
|
||||
1, // nrepeat,
|
||||
param.N,
|
||||
param.K,
|
||||
param.C,
|
||||
param.input_spatial_lengths,
|
||||
param.filter_spatial_lengths,
|
||||
param.GetOutputSpatialLengths(),
|
||||
param.conv_filter_strides,
|
||||
param.conv_filter_dilations,
|
||||
param.input_left_pads,
|
||||
param.input_right_pads);
|
||||
|
||||
pass &= ck::profiler::profile_convnd_bwd_data_impl<1,
|
||||
ck::bhalf_t,
|
||||
ck::bhalf_t,
|
||||
ck::bhalf_t,
|
||||
float,
|
||||
ck::tensor_layout::convolution::NWC,
|
||||
ck::tensor_layout::convolution::KXC,
|
||||
ck::tensor_layout::convolution::NWK>(
|
||||
1, // do_verification,
|
||||
1, // init_method,
|
||||
0, // do_log,
|
||||
1, // nrepeat,
|
||||
param.N,
|
||||
param.K,
|
||||
param.C,
|
||||
param.input_spatial_lengths,
|
||||
param.filter_spatial_lengths,
|
||||
param.GetOutputSpatialLengths(),
|
||||
param.conv_filter_strides,
|
||||
param.conv_filter_dilations,
|
||||
param.input_left_pads,
|
||||
param.input_right_pads);
|
||||
|
||||
pass &= ck::profiler::profile_convnd_bwd_data_impl<1,
|
||||
int8_t,
|
||||
int8_t,
|
||||
int8_t,
|
||||
int,
|
||||
ck::tensor_layout::convolution::NWC,
|
||||
ck::tensor_layout::convolution::KXC,
|
||||
ck::tensor_layout::convolution::NWK>(
|
||||
1, // do_verification,
|
||||
1, // init_method,
|
||||
0, // do_log,
|
||||
1, // nrepeat,
|
||||
param.N,
|
||||
param.K,
|
||||
param.C,
|
||||
param.input_spatial_lengths,
|
||||
param.filter_spatial_lengths,
|
||||
param.GetOutputSpatialLengths(),
|
||||
param.conv_filter_strides,
|
||||
param.conv_filter_dilations,
|
||||
param.input_left_pads,
|
||||
param.input_right_pads);
|
||||
}
|
||||
|
||||
// check 2d
|
||||
params.clear();
|
||||
params.push_back({2, 128, 128, 256, {1, 1}, {7, 7}, {2, 2}, {1, 1}, {0, 0}, {0, 0}});
|
||||
params.push_back({2, 128, 128, 256, {3, 3}, {14, 14}, {1, 1}, {1, 1}, {1, 1}, {1, 1}});
|
||||
params.push_back({2, 128, 128, 256, {1, 1}, {3, 3}, {1, 1}, {1, 1}, {0, 0}, {0, 0}});
|
||||
|
||||
for(auto& param : params)
|
||||
{
|
||||
pass &= ck::profiler::profile_convnd_bwd_data_impl<2,
|
||||
float,
|
||||
float,
|
||||
float,
|
||||
float,
|
||||
ck::tensor_layout::convolution::NHWC,
|
||||
ck::tensor_layout::convolution::KYXC,
|
||||
ck::tensor_layout::convolution::NHWK>(
|
||||
1, // do_verification,
|
||||
1, // init_method,
|
||||
0, // do_log,
|
||||
1, // nrepeat,
|
||||
param.N,
|
||||
param.K,
|
||||
param.C,
|
||||
param.input_spatial_lengths,
|
||||
param.filter_spatial_lengths,
|
||||
param.GetOutputSpatialLengths(),
|
||||
param.conv_filter_strides,
|
||||
param.conv_filter_dilations,
|
||||
param.input_left_pads,
|
||||
param.input_right_pads);
|
||||
|
||||
pass &= ck::profiler::profile_convnd_bwd_data_impl<2,
|
||||
ck::half_t,
|
||||
ck::half_t,
|
||||
ck::half_t,
|
||||
float,
|
||||
ck::tensor_layout::convolution::NHWC,
|
||||
ck::tensor_layout::convolution::KYXC,
|
||||
ck::tensor_layout::convolution::NHWK>(
|
||||
1, // do_verification,
|
||||
1, // init_method,
|
||||
0, // do_log,
|
||||
1, // nrepeat,
|
||||
param.N,
|
||||
param.K,
|
||||
param.C,
|
||||
param.input_spatial_lengths,
|
||||
param.filter_spatial_lengths,
|
||||
param.GetOutputSpatialLengths(),
|
||||
param.conv_filter_strides,
|
||||
param.conv_filter_dilations,
|
||||
param.input_left_pads,
|
||||
param.input_right_pads);
|
||||
|
||||
pass &= ck::profiler::profile_convnd_bwd_data_impl<2,
|
||||
ck::bhalf_t,
|
||||
ck::bhalf_t,
|
||||
ck::bhalf_t,
|
||||
float,
|
||||
ck::tensor_layout::convolution::NHWC,
|
||||
ck::tensor_layout::convolution::KYXC,
|
||||
ck::tensor_layout::convolution::NHWK>(
|
||||
1, // do_verification,
|
||||
1, // init_method,
|
||||
0, // do_log,
|
||||
1, // nrepeat,
|
||||
param.N,
|
||||
param.K,
|
||||
param.C,
|
||||
param.input_spatial_lengths,
|
||||
param.filter_spatial_lengths,
|
||||
param.GetOutputSpatialLengths(),
|
||||
param.conv_filter_strides,
|
||||
param.conv_filter_dilations,
|
||||
param.input_left_pads,
|
||||
param.input_right_pads);
|
||||
|
||||
pass &= ck::profiler::profile_convnd_bwd_data_impl<2,
|
||||
int8_t,
|
||||
int8_t,
|
||||
int8_t,
|
||||
int,
|
||||
ck::tensor_layout::convolution::NHWC,
|
||||
ck::tensor_layout::convolution::KYXC,
|
||||
ck::tensor_layout::convolution::NHWK>(
|
||||
1, // do_verification,
|
||||
1, // init_method,
|
||||
0, // do_log,
|
||||
1, // nrepeat,
|
||||
param.N,
|
||||
param.K,
|
||||
param.C,
|
||||
param.input_spatial_lengths,
|
||||
param.filter_spatial_lengths,
|
||||
param.GetOutputSpatialLengths(),
|
||||
param.conv_filter_strides,
|
||||
param.conv_filter_dilations,
|
||||
param.input_left_pads,
|
||||
param.input_right_pads);
|
||||
}
|
||||
|
||||
// check 3d
|
||||
params.clear();
|
||||
params.push_back(
|
||||
{3, 128, 128, 256, {1, 1, 1}, {7, 7, 7}, {2, 2, 2}, {1, 1, 1}, {0, 0, 0}, {0, 0, 0}});
|
||||
params.push_back(
|
||||
{3, 128, 128, 256, {3, 3, 3}, {14, 14, 14}, {1, 1, 1}, {1, 1, 1}, {1, 1, 1}, {1, 1, 1}});
|
||||
params.push_back(
|
||||
{3, 128, 128, 256, {1, 1, 1}, {3, 3, 3}, {1, 1, 1}, {1, 1, 1}, {0, 0, 0}, {0, 0, 0}});
|
||||
|
||||
for(auto& param : params)
|
||||
{
|
||||
pass &= ck::profiler::profile_convnd_bwd_data_impl<3,
|
||||
float,
|
||||
float,
|
||||
float,
|
||||
float,
|
||||
ck::tensor_layout::convolution::NDHWC,
|
||||
ck::tensor_layout::convolution::KZYXC,
|
||||
ck::tensor_layout::convolution::NDHWK>(
|
||||
1, // do_verification,
|
||||
1, // init_method,
|
||||
0, // do_log,
|
||||
1, // nrepeat,
|
||||
param.N,
|
||||
param.K,
|
||||
param.C,
|
||||
param.input_spatial_lengths,
|
||||
param.filter_spatial_lengths,
|
||||
param.GetOutputSpatialLengths(),
|
||||
param.conv_filter_strides,
|
||||
param.conv_filter_dilations,
|
||||
param.input_left_pads,
|
||||
param.input_right_pads);
|
||||
|
||||
pass &= ck::profiler::profile_convnd_bwd_data_impl<3,
|
||||
ck::half_t,
|
||||
ck::half_t,
|
||||
ck::half_t,
|
||||
float,
|
||||
ck::tensor_layout::convolution::NDHWC,
|
||||
ck::tensor_layout::convolution::KZYXC,
|
||||
ck::tensor_layout::convolution::NDHWK>(
|
||||
1, // do_verification,
|
||||
1, // init_method,
|
||||
0, // do_log,
|
||||
1, // nrepeat,
|
||||
param.N,
|
||||
param.K,
|
||||
param.C,
|
||||
param.input_spatial_lengths,
|
||||
param.filter_spatial_lengths,
|
||||
param.GetOutputSpatialLengths(),
|
||||
param.conv_filter_strides,
|
||||
param.conv_filter_dilations,
|
||||
param.input_left_pads,
|
||||
param.input_right_pads);
|
||||
|
||||
pass &= ck::profiler::profile_convnd_bwd_data_impl<3,
|
||||
ck::bhalf_t,
|
||||
ck::bhalf_t,
|
||||
ck::bhalf_t,
|
||||
float,
|
||||
ck::tensor_layout::convolution::NDHWC,
|
||||
ck::tensor_layout::convolution::KZYXC,
|
||||
ck::tensor_layout::convolution::NDHWK>(
|
||||
1, // do_verification,
|
||||
1, // init_method,
|
||||
0, // do_log,
|
||||
1, // nrepeat,
|
||||
param.N,
|
||||
param.K,
|
||||
param.C,
|
||||
param.input_spatial_lengths,
|
||||
param.filter_spatial_lengths,
|
||||
param.GetOutputSpatialLengths(),
|
||||
param.conv_filter_strides,
|
||||
param.conv_filter_dilations,
|
||||
param.input_left_pads,
|
||||
param.input_right_pads);
|
||||
|
||||
pass &= ck::profiler::profile_convnd_bwd_data_impl<3,
|
||||
int8_t,
|
||||
int8_t,
|
||||
int8_t,
|
||||
int,
|
||||
ck::tensor_layout::convolution::NDHWC,
|
||||
ck::tensor_layout::convolution::KZYXC,
|
||||
ck::tensor_layout::convolution::NDHWK>(
|
||||
1, // do_verification,
|
||||
1, // init_method,
|
||||
0, // do_log,
|
||||
1, // nrepeat,
|
||||
param.N,
|
||||
param.K,
|
||||
param.C,
|
||||
param.input_spatial_lengths,
|
||||
param.filter_spatial_lengths,
|
||||
param.GetOutputSpatialLengths(),
|
||||
param.conv_filter_strides,
|
||||
param.conv_filter_dilations,
|
||||
param.input_left_pads,
|
||||
param.input_right_pads);
|
||||
}
|
||||
|
||||
if(pass)
|
||||
{
|
||||
std::cout << "test convnd bwd : Pass" << std::endl;
|
||||
return 0;
|
||||
}
|
||||
else
|
||||
{
|
||||
std::cout << "test convnd bwd: Fail " << std::endl;
|
||||
return -1;
|
||||
}
|
||||
}
|
||||
Reference in New Issue
Block a user