mirror of
https://github.com/ROCm/composable_kernel.git
synced 2026-04-19 22:39:03 +00:00
Clean up conv example, Instances, profiler and test (#324)
* convnd_fwd fp16 example * update example * update example * update instance * updating refernce conv * update reference conv * update conv fwd profiler * update conv 1d and 3d instance * update include path * clean * update profiler for conv bwd data and weight * update conv bwd weight * clean * update conv example * update profiler for conv bwd weight * update ckprofiler for conv bwd data * fix reference conv bwd data bug; update conv bwd data test * update examples * fix initialization issue * update test for conv fwd * clean * clean * remove test case too sensitive to error threshhold * fix test * clean * fix build * adding conv multiple d * adding conv multiple D * add matrix padder * add gemm padding to convnd * adding group conv * update gemm multi-d * refactor * refactor * refactor * clean * clean * refactor * refactor * reorg * add ds * add bias * clean * add G * adding group * adding group * adding group * update Tensor * clean * update example * update DeviceGemmMultipleD_Xdl_CShuffle * update conv bwd-data and bwd-weight * upate contraction example * update gemm and batch gemm with e permute * fix example build * instance for grouped conv1d * update example * adding group conv instance * update gemm bilinear instance * update gemm+add+add+fastgelu instance * update profiler * update profiler * update test * update test and client example * clean * add grouped conv into profiler * update profiler * clean * add test grouped conv, update all conv test to gtest * update test
This commit is contained in:
@@ -10,198 +10,147 @@
|
||||
#include "ck/tensor_operation/gpu/device/tensor_layout.hpp"
|
||||
|
||||
#include "ck/library/utility/check_err.hpp"
|
||||
#include "ck/library/utility/conv_util.hpp"
|
||||
#include "ck/library/utility/convolution_parameter.hpp"
|
||||
|
||||
namespace {
|
||||
|
||||
class TestConvUtil : public ::testing::Test
|
||||
{
|
||||
public:
|
||||
void SetNDParams(std::size_t ndims)
|
||||
void SetNDParams(std::size_t ndims, std::size_t s, std::size_t d, std::size_t p)
|
||||
{
|
||||
conv_params.num_dim_spatial_ = ndims;
|
||||
conv_params.filter_spatial_lengths_ = std::vector<ck::index_t>(ndims, 3);
|
||||
conv_params.input_spatial_lengths_ = std::vector<ck::index_t>(ndims, 71);
|
||||
conv_params.conv_filter_strides_ = std::vector<ck::index_t>(ndims, 2);
|
||||
conv_params.conv_filter_dilations_ = std::vector<ck::index_t>(ndims, 1);
|
||||
conv_params.input_left_pads_ = std::vector<ck::index_t>(ndims, 1);
|
||||
conv_params.input_right_pads_ = std::vector<ck::index_t>(ndims, 1);
|
||||
conv_params = ck::utils::conv::ConvParam(ndims,
|
||||
2,
|
||||
128,
|
||||
192,
|
||||
256,
|
||||
std::vector<ck::index_t>(ndims, 3),
|
||||
std::vector<ck::index_t>(ndims, 71),
|
||||
std::vector<ck::index_t>(ndims, s),
|
||||
std::vector<ck::index_t>(ndims, d),
|
||||
std::vector<ck::index_t>(ndims, p),
|
||||
std::vector<ck::index_t>(ndims, p));
|
||||
}
|
||||
|
||||
protected:
|
||||
// ------- default 2D -------
|
||||
// input NCHW {128,192,71,71},
|
||||
// weights KCYX {256,192,3,3},
|
||||
// stride {2,2},
|
||||
// dilations {1,1},
|
||||
// padding {{1,1}, {1,1}}
|
||||
ck::utils::conv::ConvParams conv_params;
|
||||
// input GNCHW {2, 128, 192, 71, 71},
|
||||
// weights GKCYX {2, 256, 192, 3, 3},
|
||||
// stride {s, s},
|
||||
// dilations {d, d},
|
||||
// padding {{p, p}, {p, p}
|
||||
ck::utils::conv::ConvParam conv_params;
|
||||
};
|
||||
|
||||
} // namespace
|
||||
|
||||
TEST_F(TestConvUtil, ConvParamsGetOutputSpatialLengths2D)
|
||||
{
|
||||
ck::utils::conv::ConvParams conv_params;
|
||||
std::vector<ck::index_t> out_spatial_len = conv_params.GetOutputSpatialLengths();
|
||||
EXPECT_TRUE(ck::utils::check_err(out_spatial_len,
|
||||
std::vector<ck::index_t>{36, 36},
|
||||
"Error: ConvParams 2D default constructor."));
|
||||
|
||||
conv_params.conv_filter_strides_ = std::vector<ck::index_t>{1, 1};
|
||||
out_spatial_len = conv_params.GetOutputSpatialLengths();
|
||||
EXPECT_TRUE(ck::utils::check_err(
|
||||
out_spatial_len, std::vector<ck::index_t>{71, 71}, "Error: ConvParams 2D stride {1,1}."));
|
||||
|
||||
conv_params.conv_filter_strides_ = std::vector<ck::index_t>{2, 2};
|
||||
conv_params.input_left_pads_ = std::vector<ck::index_t>{2, 2};
|
||||
conv_params.input_right_pads_ = std::vector<ck::index_t>{2, 2};
|
||||
out_spatial_len = conv_params.GetOutputSpatialLengths();
|
||||
EXPECT_TRUE(ck::utils::check_err(out_spatial_len,
|
||||
std::vector<ck::index_t>{37, 37},
|
||||
"Error: ConvParams 2D padding left/right {2,2}."));
|
||||
|
||||
conv_params.conv_filter_dilations_ = std::vector<ck::index_t>{2, 2};
|
||||
out_spatial_len = conv_params.GetOutputSpatialLengths();
|
||||
EXPECT_TRUE(ck::utils::check_err(
|
||||
out_spatial_len, std::vector<ck::index_t>{36, 36}, "Error: ConvParams 2D dilation {2,2}."));
|
||||
|
||||
conv_params.conv_filter_strides_ = std::vector<ck::index_t>{3, 3};
|
||||
conv_params.input_left_pads_ = std::vector<ck::index_t>{1, 1};
|
||||
conv_params.input_right_pads_ = std::vector<ck::index_t>{1, 1};
|
||||
conv_params.conv_filter_dilations_ = std::vector<ck::index_t>{2, 2};
|
||||
out_spatial_len = conv_params.GetOutputSpatialLengths();
|
||||
EXPECT_TRUE(
|
||||
ck::utils::check_err(out_spatial_len,
|
||||
std::vector<ck::index_t>{23, 23},
|
||||
"Error: ConvParams 2D strides{3,3}, padding {1,1}, dilations {2,2}."));
|
||||
}
|
||||
|
||||
TEST_F(TestConvUtil, ConvParamsGetOutputSpatialLengths1D)
|
||||
{
|
||||
SetNDParams(1);
|
||||
|
||||
// stride 2, dilation 1, pad 1
|
||||
SetNDParams(1, 2, 1, 1);
|
||||
std::vector<ck::index_t> out_spatial_len = conv_params.GetOutputSpatialLengths();
|
||||
EXPECT_TRUE(ck::utils::check_err(
|
||||
out_spatial_len, std::vector<ck::index_t>{36}, "Error: ConvParams 1D."));
|
||||
|
||||
conv_params.conv_filter_strides_ = std::vector<ck::index_t>{1};
|
||||
out_spatial_len = conv_params.GetOutputSpatialLengths();
|
||||
// stride 1, dilation 1, pad 1
|
||||
SetNDParams(1, 1, 1, 1);
|
||||
out_spatial_len = conv_params.GetOutputSpatialLengths();
|
||||
EXPECT_TRUE(ck::utils::check_err(
|
||||
out_spatial_len, std::vector<ck::index_t>{71}, "Error: ConvParams 1D stride {1}."));
|
||||
|
||||
conv_params.conv_filter_strides_ = std::vector<ck::index_t>{2};
|
||||
conv_params.input_left_pads_ = std::vector<ck::index_t>{2};
|
||||
conv_params.input_right_pads_ = std::vector<ck::index_t>{2};
|
||||
out_spatial_len = conv_params.GetOutputSpatialLengths();
|
||||
// stride 2, dilation 1, pad 2
|
||||
SetNDParams(1, 2, 1, 2);
|
||||
out_spatial_len = conv_params.GetOutputSpatialLengths();
|
||||
EXPECT_TRUE(ck::utils::check_err(out_spatial_len,
|
||||
std::vector<ck::index_t>{37},
|
||||
"Error: ConvParams 1D padding left/right {2}."));
|
||||
|
||||
conv_params.conv_filter_dilations_ = std::vector<ck::index_t>{2};
|
||||
out_spatial_len = conv_params.GetOutputSpatialLengths();
|
||||
// stride 2, dilation 2, pad 2
|
||||
SetNDParams(1, 2, 2, 2);
|
||||
out_spatial_len = conv_params.GetOutputSpatialLengths();
|
||||
EXPECT_TRUE(ck::utils::check_err(
|
||||
out_spatial_len, std::vector<ck::index_t>{36}, "Error: ConvParams 1D dilation {2}."));
|
||||
|
||||
conv_params.conv_filter_strides_ = std::vector<ck::index_t>{3};
|
||||
conv_params.input_left_pads_ = std::vector<ck::index_t>{1};
|
||||
conv_params.input_right_pads_ = std::vector<ck::index_t>{1};
|
||||
conv_params.conv_filter_dilations_ = std::vector<ck::index_t>{2};
|
||||
out_spatial_len = conv_params.GetOutputSpatialLengths();
|
||||
// stride 3, dilation 2, pad 1
|
||||
SetNDParams(1, 3, 2, 1);
|
||||
out_spatial_len = conv_params.GetOutputSpatialLengths();
|
||||
EXPECT_TRUE(
|
||||
ck::utils::check_err(out_spatial_len,
|
||||
std::vector<ck::index_t>{23},
|
||||
"Error: ConvParams 1D strides{3}, padding {1}, dilations {2}."));
|
||||
}
|
||||
|
||||
TEST_F(TestConvUtil, ConvParamsGetOutputSpatialLengths2D)
|
||||
{
|
||||
// stride 2, dilation 1, pad 1
|
||||
SetNDParams(2, 2, 1, 1);
|
||||
std::vector<ck::index_t> out_spatial_len = conv_params.GetOutputSpatialLengths();
|
||||
EXPECT_TRUE(ck::utils::check_err(out_spatial_len,
|
||||
std::vector<ck::index_t>{36, 36},
|
||||
"Error: ConvParams 2D default constructor."));
|
||||
|
||||
// stride 1, dilation 1, pad 1
|
||||
SetNDParams(2, 1, 1, 1);
|
||||
out_spatial_len = conv_params.GetOutputSpatialLengths();
|
||||
EXPECT_TRUE(ck::utils::check_err(
|
||||
out_spatial_len, std::vector<ck::index_t>{71, 71}, "Error: ConvParams 2D stride {1,1}."));
|
||||
|
||||
// stride 2, dilation 1, pad 2
|
||||
SetNDParams(2, 2, 1, 2);
|
||||
out_spatial_len = conv_params.GetOutputSpatialLengths();
|
||||
EXPECT_TRUE(ck::utils::check_err(out_spatial_len,
|
||||
std::vector<ck::index_t>{37, 37},
|
||||
"Error: ConvParams 2D padding left/right {2,2}."));
|
||||
|
||||
// stride 2, dilation 2, pad 2
|
||||
SetNDParams(2, 2, 2, 2);
|
||||
out_spatial_len = conv_params.GetOutputSpatialLengths();
|
||||
EXPECT_TRUE(ck::utils::check_err(
|
||||
out_spatial_len, std::vector<ck::index_t>{36, 36}, "Error: ConvParams 2D dilation {2,2}."));
|
||||
|
||||
// stride 3, dilation 2, pad 1
|
||||
SetNDParams(2, 3, 2, 1);
|
||||
out_spatial_len = conv_params.GetOutputSpatialLengths();
|
||||
EXPECT_TRUE(
|
||||
ck::utils::check_err(out_spatial_len,
|
||||
std::vector<ck::index_t>{23, 23},
|
||||
"Error: ConvParams 2D strides{3,3}, padding {1,1}, dilations {2,2}."));
|
||||
}
|
||||
|
||||
TEST_F(TestConvUtil, ConvParamsGetOutputSpatialLengths3D)
|
||||
{
|
||||
SetNDParams(3);
|
||||
|
||||
// stride 2, dilation 1, pad 1
|
||||
SetNDParams(3, 2, 1, 1);
|
||||
std::vector<ck::index_t> out_spatial_len = conv_params.GetOutputSpatialLengths();
|
||||
EXPECT_TRUE(ck::utils::check_err(
|
||||
out_spatial_len, std::vector<ck::index_t>{36, 36, 36}, "Error: ConvParams 3D."));
|
||||
|
||||
conv_params.conv_filter_strides_ = std::vector<ck::index_t>{1, 1, 1};
|
||||
out_spatial_len = conv_params.GetOutputSpatialLengths();
|
||||
// stride 1, dilation 1, pad 1
|
||||
SetNDParams(3, 1, 1, 1);
|
||||
out_spatial_len = conv_params.GetOutputSpatialLengths();
|
||||
EXPECT_TRUE(ck::utils::check_err(out_spatial_len,
|
||||
std::vector<ck::index_t>{71, 71, 71},
|
||||
"Error: ConvParams 3D stride {1, 1, 1}."));
|
||||
|
||||
conv_params.conv_filter_strides_ = std::vector<ck::index_t>{2, 2, 2};
|
||||
conv_params.input_left_pads_ = std::vector<ck::index_t>{2, 2, 2};
|
||||
conv_params.input_right_pads_ = std::vector<ck::index_t>{2, 2, 2};
|
||||
out_spatial_len = conv_params.GetOutputSpatialLengths();
|
||||
// stride 2, dilation 1, pad 2
|
||||
SetNDParams(3, 2, 1, 2);
|
||||
out_spatial_len = conv_params.GetOutputSpatialLengths();
|
||||
EXPECT_TRUE(ck::utils::check_err(out_spatial_len,
|
||||
std::vector<ck::index_t>{37, 37, 37},
|
||||
"Error: ConvParams 3D padding left/right {2, 2, 2}."));
|
||||
|
||||
conv_params.conv_filter_dilations_ = std::vector<ck::index_t>{2, 2, 2};
|
||||
out_spatial_len = conv_params.GetOutputSpatialLengths();
|
||||
// stride 2, dilation 2, pad 2
|
||||
SetNDParams(3, 2, 2, 2);
|
||||
out_spatial_len = conv_params.GetOutputSpatialLengths();
|
||||
EXPECT_TRUE(ck::utils::check_err(out_spatial_len,
|
||||
std::vector<ck::index_t>{36, 36, 36},
|
||||
"Error: ConvParams 3D dilation {2, 2, 2}."));
|
||||
|
||||
conv_params.conv_filter_strides_ = std::vector<ck::index_t>{3, 3, 3};
|
||||
conv_params.input_left_pads_ = std::vector<ck::index_t>{1, 1, 1};
|
||||
conv_params.input_right_pads_ = std::vector<ck::index_t>{1, 1, 1};
|
||||
conv_params.conv_filter_dilations_ = std::vector<ck::index_t>{2, 2, 2};
|
||||
out_spatial_len = conv_params.GetOutputSpatialLengths();
|
||||
// stride 3, dilation 2, pad 1
|
||||
SetNDParams(3, 3, 2, 1);
|
||||
out_spatial_len = conv_params.GetOutputSpatialLengths();
|
||||
EXPECT_TRUE(ck::utils::check_err(
|
||||
out_spatial_len,
|
||||
std::vector<ck::index_t>{23, 23, 23},
|
||||
"Error: ConvParams 3D strides{3, 3, 3}, padding {1, 1, 1}, dilations {2, 2, 2}."));
|
||||
}
|
||||
|
||||
TEST(ConvUtil, GetHostTensorDescriptor)
|
||||
{
|
||||
namespace tl = ck::tensor_layout::convolution;
|
||||
std::vector<std::size_t> dims{2, 3, 4, 5};
|
||||
HostTensorDescriptor h = ck::utils::conv::get_host_tensor_descriptor(dims, tl::NHWC{});
|
||||
EXPECT_TRUE(ck::utils::check_err(
|
||||
h.GetLengths(), {2, 3, 4, 5}, "Error: wrong NHWC dimensions lengths!"));
|
||||
EXPECT_TRUE(ck::utils::check_err(
|
||||
h.GetStrides(), {3 * 4 * 5, 1, 3 * 5, 3}, "Error: wrong NHWC dimensions strides!"));
|
||||
|
||||
h = ck::utils::conv::get_host_tensor_descriptor(dims, tl::NCHW{});
|
||||
EXPECT_TRUE(ck::utils::check_err(
|
||||
h.GetLengths(), {2, 3, 4, 5}, "Error: wrong NCHW dimensions lengths!"));
|
||||
EXPECT_TRUE(ck::utils::check_err(
|
||||
h.GetStrides(), {3 * 4 * 5, 4 * 5, 5, 1}, "Error: wrong NCHW dimensions strides!"));
|
||||
|
||||
dims = std::vector<std::size_t>{2, 3, 4};
|
||||
h = ck::utils::conv::get_host_tensor_descriptor(dims, tl::NWC{});
|
||||
EXPECT_TRUE(
|
||||
ck::utils::check_err(h.GetLengths(), {2, 3, 4}, "Error: wrong NWC dimensions lengths!"));
|
||||
EXPECT_TRUE(ck::utils::check_err(
|
||||
h.GetStrides(), {3 * 4, 1, 3}, "Error: wrong NWC dimensions strides!"));
|
||||
|
||||
h = ck::utils::conv::get_host_tensor_descriptor(dims, tl::NCW{});
|
||||
EXPECT_TRUE(
|
||||
ck::utils::check_err(h.GetLengths(), {2, 3, 4}, "Error: wrong NCW dimensions lengths!"));
|
||||
EXPECT_TRUE(ck::utils::check_err(
|
||||
h.GetStrides(), {3 * 4, 4, 1}, "Error: wrong NCW dimensions strides!"));
|
||||
|
||||
dims = std::vector<std::size_t>{2, 3, 4, 5, 6};
|
||||
h = ck::utils::conv::get_host_tensor_descriptor(dims, tl::NDHWC{});
|
||||
EXPECT_TRUE(
|
||||
ck::utils::check_err(h.GetLengths(), dims, "Error: wrong NDHWC dimensions lengths!"));
|
||||
EXPECT_TRUE(ck::utils::check_err(h.GetStrides(),
|
||||
{3 * 4 * 5 * 6, // N
|
||||
1, // C
|
||||
3 * 5 * 6, // D
|
||||
3 * 6, // H
|
||||
3}, // W
|
||||
"Error: wrong NDHWC dimensions strides!"));
|
||||
|
||||
h = ck::utils::conv::get_host_tensor_descriptor(dims, tl::NCDHW{});
|
||||
EXPECT_TRUE(
|
||||
ck::utils::check_err(h.GetLengths(), dims, "Error: wrong NCDHW dimensions lengths!"));
|
||||
EXPECT_TRUE(ck::utils::check_err(h.GetStrides(),
|
||||
{3 * 4 * 5 * 6, // N
|
||||
4 * 5 * 6, // C
|
||||
5 * 6, // D
|
||||
6, // H
|
||||
1}, // W
|
||||
"Error: wrong NCDHW dimensions strides!"));
|
||||
}
|
||||
|
||||
Reference in New Issue
Block a user