mirror of
https://github.com/ROCm/composable_kernel.git
synced 2026-05-14 02:02:46 +00:00
49965ba360
* Convolution ND
* Code unification across dimensions for generating tensor descriptors.
* Example
* Instances
* Move convnd f32 instance file to comply with repo structure.
* Conv 1D tensor layouts.
* Formatting and use ReferenceConv
* Reference ConvFwd supporting 1D and 2D convolution.
* Debug printing TensorLayout name.
* Conv fwd 1D instance f32
* Refactor conv ND example.
Needed to support various conv dimensio.
Needed to support various conv dimensions
* Rename conv nd example director to prevent conflicts.
* Refactor some common utility to single file.
Plus some tests.
* Refactor GetHostTensorDescriptor + UT.
* Add 1D test case.
* Test reference convolution 1d/2d
* Remove some leftovers.
* Fix convolution example error for 1D
* Refactor test check errors utility function.
* Test Conv2D Fwd XDL
* More UT for 1D case.
* Parameterize input & weight initializers.
* Rename example to prevent conflicts.
* Split convnd instance into separate files for 1d/2d
* Address review comments.
* Fix data type for flops/gbytes calculations.
* Assign example number 11.
Co-authored-by: Adam Osewski <aosewski@amd.com>
Co-authored-by: Chao Liu <chao.liu2@amd.com>
[ROCm/composable_kernel commit: 756a761727]
114 lines
1.9 KiB
C++
114 lines
1.9 KiB
C++
#ifndef TENSOR_LAYOUT_HPP
|
|
#define TENSOR_LAYOUT_HPP
|
|
|
|
namespace ck {
|
|
namespace tensor_layout {
|
|
|
|
struct BaseTensorLayout
|
|
{
|
|
};
|
|
|
|
namespace gemm {
|
|
|
|
struct RowMajor : public BaseTensorLayout
|
|
{
|
|
static constexpr const char* name = "RowMajor";
|
|
};
|
|
|
|
struct ColumnMajor : public BaseTensorLayout
|
|
{
|
|
static constexpr const char* name = "ColumnMajor";
|
|
};
|
|
} // namespace gemm
|
|
|
|
namespace convolution {
|
|
|
|
// 1D Conv
|
|
struct NWC : public BaseTensorLayout
|
|
{
|
|
static constexpr const char* name = "NWC";
|
|
};
|
|
|
|
struct KXC : public BaseTensorLayout
|
|
{
|
|
static constexpr const char* name = "KXC";
|
|
};
|
|
|
|
struct NWK : public BaseTensorLayout
|
|
{
|
|
static constexpr const char* name = "NWK";
|
|
};
|
|
|
|
struct NCW : public BaseTensorLayout
|
|
{
|
|
static constexpr const char* name = "NCW";
|
|
};
|
|
|
|
struct KCX : public BaseTensorLayout
|
|
{
|
|
static constexpr const char* name = "KCX";
|
|
};
|
|
|
|
struct NKW : public BaseTensorLayout
|
|
{
|
|
static constexpr const char* name = "NKW";
|
|
};
|
|
|
|
// 2D Conv
|
|
struct NHWC : public BaseTensorLayout
|
|
{
|
|
static constexpr const char* name = "NHWC";
|
|
};
|
|
|
|
struct KYXC : public BaseTensorLayout
|
|
{
|
|
static constexpr const char* name = "KYXC";
|
|
};
|
|
|
|
struct NHWK : public BaseTensorLayout
|
|
{
|
|
static constexpr const char* name = "NHWK";
|
|
};
|
|
|
|
struct NCHW : public BaseTensorLayout
|
|
{
|
|
static constexpr const char* name = "NCHW";
|
|
};
|
|
|
|
struct KCYX : public BaseTensorLayout
|
|
{
|
|
static constexpr const char* name = "KCYX";
|
|
};
|
|
|
|
struct NKHW : public BaseTensorLayout
|
|
{
|
|
static constexpr const char* name = "NKHW";
|
|
};
|
|
|
|
struct NDHWC : public BaseTensorLayout
|
|
{
|
|
};
|
|
|
|
struct KZYXC : public BaseTensorLayout
|
|
{
|
|
};
|
|
|
|
struct NDHWK : public BaseTensorLayout
|
|
{
|
|
};
|
|
|
|
} // namespace convolution
|
|
|
|
template <
|
|
typename Layout,
|
|
typename std::enable_if<std::is_base_of<BaseTensorLayout, Layout>::value, bool>::type = false>
|
|
std::ostream& operator<<(std::ostream& os, const Layout&)
|
|
{
|
|
os << Layout::name;
|
|
return os;
|
|
}
|
|
|
|
} // namespace tensor_layout
|
|
} // namespace ck
|
|
#endif
|