ROCm/composable_kernel
1
0
Fork 0
mirror of https://github.com/ROCm/composable_kernel.git synced 2026-05-20 21:09:08 +00:00
Code Issues Packages Projects Releases Wiki Activity

Add Grouped Conv Fwd Large Tensor kernel (#1432)

Browse Source 
* Support 64 bit indexing

* Add new grouped conv fwd kernel for large tensors

* Add instances large tensor

* Fixes for transform conv to gemm

* Fixes

* fixes

* Remove not needed instances

* examples fixes

* Remove not need ds arrays

* Fix tests

* Add 2GB check in gridwise dl

* Fixes

[ROCm/composable_kernel commit: 4ec5c52a0c]
This commit is contained in:
Bartłomiej Kocot
2024-08-06 10:06:10 +02:00
committed by GitHub
parent 1152542dab
commit 458d8bef26
42 changed files with 3220 additions and 451 deletions
Download Patch File Download Diff File Expand all files Collapse all files

56
test/conv_util/conv_util.cpp
View File

@@ -1,5 +1,5 @@
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2023, Advanced Micro Devices, Inc. All rights reserved.
// Copyright (c) 2018-2024, Advanced Micro Devices, Inc. All rights reserved.
#include <iostream>
#include <string>
@@ -24,12 +24,12 @@ class TestConvUtil : public ::testing::Test
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));
std::vector<ck::long_index_t>(ndims, 3),
std::vector<ck::long_index_t>(ndims, 71),
std::vector<ck::long_index_t>(ndims, s),
std::vector<ck::long_index_t>(ndims, d),
std::vector<ck::long_index_t>(ndims, p),
std::vector<ck::long_index_t>(ndims, p));
}
protected:
@@ -48,35 +48,35 @@ TEST_F(TestConvUtil, ConvParamsGetOutputSpatialLengths1D)
{
// stride 2, dilation 1, pad 1
SetNDParams(1, 2, 1, 1);
std::vector<ck::index_t> out_spatial_len = conv_params.GetOutputSpatialLengths();
std::vector<ck::long_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."));
out_spatial_len, std::vector<ck::long_index_t>{36}, "Error: ConvParams 1D."));
// 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}."));
out_spatial_len, std::vector<ck::long_index_t>{71}, "Error: ConvParams 1D stride {1}."));
// 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},
std::vector<ck::long_index_t>{37},
"Error: ConvParams 1D padding left/right {2}."));
// 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}."));
out_spatial_len, std::vector<ck::long_index_t>{36}, "Error: ConvParams 1D dilation {2}."));
// 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},
std::vector<ck::long_index_t>{23},
"Error: ConvParams 1D strides{3}, padding {1}, dilations {2}."));
}
@@ -84,36 +84,38 @@ 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();
std::vector<ck::long_index_t> out_spatial_len = conv_params.GetOutputSpatialLengths();
EXPECT_TRUE(ck::utils::check_err(out_spatial_len,
std::vector<ck::index_t>{36, 36},
std::vector<ck::long_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}."));
EXPECT_TRUE(ck::utils::check_err(out_spatial_len,
std::vector<ck::long_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},
std::vector<ck::long_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}."));
EXPECT_TRUE(ck::utils::check_err(out_spatial_len,
std::vector<ck::long_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},
std::vector<ck::long_index_t>{23, 23},
"Error: ConvParams 2D strides{3,3}, padding {1,1}, dilations {2,2}."));
}
@@ -121,29 +123,29 @@ TEST_F(TestConvUtil, ConvParamsGetOutputSpatialLengths3D)
{
// stride 2, dilation 1, pad 1
SetNDParams(3, 2, 1, 1);
std::vector<ck::index_t> out_spatial_len = conv_params.GetOutputSpatialLengths();
std::vector<ck::long_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."));
out_spatial_len, std::vector<ck::long_index_t>{36, 36, 36}, "Error: ConvParams 3D."));
// 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},
std::vector<ck::long_index_t>{71, 71, 71},
"Error: ConvParams 3D stride {1, 1, 1}."));
// 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},
std::vector<ck::long_index_t>{37, 37, 37},
"Error: ConvParams 3D padding left/right {2, 2, 2}."));
// 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},
std::vector<ck::long_index_t>{36, 36, 36},
"Error: ConvParams 3D dilation {2, 2, 2}."));
// stride 3, dilation 2, pad 1
@@ -151,6 +153,6 @@ TEST_F(TestConvUtil, ConvParamsGetOutputSpatialLengths3D)
out_spatial_len = conv_params.GetOutputSpatialLengths();
EXPECT_TRUE(ck::utils::check_err(
out_spatial_len,
std::vector<ck::index_t>{23, 23, 23},
std::vector<ck::long_index_t>{23, 23, 23},
"Error: ConvParams 3D strides{3, 3, 3}, padding {1, 1, 1}, dilations {2, 2, 2}."));
}

52
test/grouped_convnd_fwd/test_grouped_convnd_fwd.cpp
View File

@@ -17,6 +17,7 @@ class TestGroupedConvndFwd : public ::testing::Test
using InLayout = std::tuple_element_t<1, Tuple>;
using WeiLayout = std::tuple_element_t<2, Tuple>;
using OutLayout = std::tuple_element_t<3, Tuple>;
using IndexType = std::tuple_element_t<4, Tuple>;
std::vector<ck::utils::conv::ConvParam> conv_params;
@@ -33,7 +34,10 @@ class TestGroupedConvndFwd : public ::testing::Test
OutLayout,
DataType,
DataType,
DataType>(
DataType,
DataType,
DataType,
IndexType>(
true, // do_verification
1, // init_method: integer value
false, // do_log
@@ -46,30 +50,31 @@ class TestGroupedConvndFwd : public ::testing::Test
using namespace ck::tensor_layout::convolution;
using KernelTypes1d = ::testing::Types<std::tuple<float, GNWC, GKXC, GNWK>,
std::tuple<ck::half_t, GNWC, GKXC, GNWK>,
std::tuple<ck::bhalf_t, GNWC, GKXC, GNWK>,
std::tuple<int8_t, GNWC, GKXC, GNWK>>;
using KernelTypes1d = ::testing::Types<std::tuple<float, GNWC, GKXC, GNWK, ck::index_t>,
std::tuple<ck::half_t, GNWC, GKXC, GNWK, ck::index_t>,
std::tuple<ck::bhalf_t, GNWC, GKXC, GNWK, ck::index_t>,
std::tuple<int8_t, GNWC, GKXC, GNWK, ck::index_t>>;
using KernelTypes2d = ::testing::Types<std::tuple<float, GNHWC, GKYXC, GNHWK>,
std::tuple<ck::half_t, GNHWC, GKYXC, GNHWK>,
std::tuple<ck::bhalf_t, GNHWC, GKYXC, GNHWK>,
std::tuple<int8_t, GNHWC, GKYXC, GNHWK>,
std::tuple<float, NHWGC, GKYXC, NHWGK>,
std::tuple<ck::half_t, NHWGC, GKYXC, NHWGK>,
std::tuple<ck::bhalf_t, NHWGC, GKYXC, NHWGK>,
std::tuple<int8_t, NHWGC, GKYXC, NHWGK>>;
using KernelTypes2d = ::testing::Types<std::tuple<float, GNHWC, GKYXC, GNHWK, ck::index_t>,
std::tuple<ck::half_t, GNHWC, GKYXC, GNHWK, ck::index_t>,
std::tuple<ck::bhalf_t, GNHWC, GKYXC, GNHWK, ck::index_t>,
std::tuple<int8_t, GNHWC, GKYXC, GNHWK, ck::index_t>,
std::tuple<float, NHWGC, GKYXC, NHWGK, ck::index_t>,
std::tuple<ck::half_t, NHWGC, GKYXC, NHWGK, ck::index_t>,
std::tuple<ck::bhalf_t, NHWGC, GKYXC, NHWGK, ck::index_t>,
std::tuple<int8_t, NHWGC, GKYXC, NHWGK, ck::index_t>>;
using KernelTypes3d = ::testing::Types<std::tuple<float, GNDHWC, GKZYXC, GNDHWK>,
std::tuple<ck::half_t, GNDHWC, GKZYXC, GNDHWK>,
std::tuple<ck::bhalf_t, GNDHWC, GKZYXC, GNDHWK>,
std::tuple<int8_t, GNDHWC, GKZYXC, GNDHWK>,
std::tuple<float, NDHWGC, GKZYXC, NDHWGK>,
std::tuple<ck::half_t, NDHWGC, GKZYXC, NDHWGK>,
std::tuple<ck::bhalf_t, NDHWGC, GKZYXC, NDHWGK>,
std::tuple<int8_t, NDHWGC, GKZYXC, NDHWGK>>;
using KernelTypes3d = ::testing::Types<std::tuple<float, GNDHWC, GKZYXC, GNDHWK, ck::index_t>,
std::tuple<ck::half_t, GNDHWC, GKZYXC, GNDHWK, ck::index_t>,
std::tuple<ck::bhalf_t, GNDHWC, GKZYXC, GNDHWK, ck::index_t>,
std::tuple<int8_t, GNDHWC, GKZYXC, GNDHWK, ck::index_t>,
std::tuple<float, NDHWGC, GKZYXC, NDHWGK, ck::index_t>,
std::tuple<ck::half_t, NDHWGC, GKZYXC, NDHWGK, ck::index_t>,
std::tuple<ck::bhalf_t, NDHWGC, GKZYXC, NDHWGK, ck::index_t>,
std::tuple<int8_t, NDHWGC, GKZYXC, NDHWGK, ck::index_t>>;
using KernelTypes2dLargeCases = ::testing::Types<std::tuple<float, NHWGC, GKYXC, NHWGK>>;
using KernelTypes2dLargeCases =
::testing::Types<std::tuple<float, NHWGC, GKYXC, NHWGK, ck::long_index_t>>;
template <typename Tuple>
class TestGroupedConvndFwd1d : public TestGroupedConvndFwd<Tuple>
@@ -153,5 +158,8 @@ TYPED_TEST(TestGroupedConvndFwd2dLargeCases, Test2DLargeCases)
// With supported NumGroupsToMerge > 1
this->conv_params.push_back(
{2, 32, 64, 1, 1, {2, 2}, {672, 672}, {672, 672}, {1, 1}, {0, 0}, {0, 0}});
// When image is larger than 2GB
this->conv_params.push_back(
{2, 1, 1, 256, 256, {3, 3}, {4096, 2048}, {1024, 1024}, {3, 3}, {1, 1}, {1, 1}});
this->template Run<2>();
}