mirror of
https://github.com/ROCm/composable_kernel.git
synced 2026-05-14 02:02:46 +00:00
a7b7eae2a1
* ck-builder: restructure testing conv
In order to prepare for bwd of conv testing, this commit moves some
files and types around so that we can reuse ckt::Args for both forward
and backwards convolution.
* ck-builder: decouple fwd_ck.hpp and fwd_reference.hpp from fwd.hpp
This will allow us to more easily include fwd.hpp from backwards
definitions, which is required for initializing bwd values.
* ck-builder: fix layout of test_ckb_conv_bwd_weight_xdl_cshuffle_v3
Turns out that the supplied layout isn't actually supported...
* ck-builder: ck and reference conv integration for bwd weight
* ck-builder: ck bwd weight execution test
* ck-builder: ckt::run support for ck-tile bwd weight
* ck-builder: ck tile bwd weight execution test
* ck-builder: extra debug printing in MatchesReference
* ck-builder: make ckt::run return RunResult
This type is more convenient than std::tuple, as it will allow us to
use google test matchers with this in the future.
* ck-builder: RunResult matcher
Using EXPECT_THAT(..., SuccessfulRun()) will generate a check and a nice error
message about how and why running an algorithm failed.
* ck-builder: doc fixes
* ck-builder: add missing headers
[ROCm/composable_kernel commit: cc75948d1c]
66 lines
4.0 KiB
C++
66 lines
4.0 KiB
C++
#include "../../builder/test/utils/ckb_conv_tile_test_configs.hpp"
|
|
#include "ck_tile/builder/testing/conv/fwd.hpp"
|
|
#include "ck_tile/builder/testing/conv/ck_tile.hpp"
|
|
|
|
namespace ckb = ck_tile::builder;
|
|
namespace ckt = ck_tile::builder::test;
|
|
namespace cku = ck_tile::builder::test_utils;
|
|
|
|
namespace ck_tile::builder::profiling {
|
|
|
|
constexpr auto SIGNATURE_NHWGC_FP32_FWD =
|
|
ckt::ConvSignature{.spatial_dim = 2,
|
|
.direction = ckb::ConvDirection::FORWARD,
|
|
.data_type = ckb::DataType::FP32,
|
|
.accumulation_data_type = ckb::DataType::FP32,
|
|
.input = {.config = {.layout = ckb::TensorLayout::NHWGC}},
|
|
.weight = {.config = {.layout = ckb::TensorLayout::GKYXC}},
|
|
.output = {.config = {.layout = ckb::TensorLayout::NHWGK}}};
|
|
|
|
constexpr auto SIGNATURE_NHWGC_BF16_FWD =
|
|
ckt::ConvSignature{.spatial_dim = 2,
|
|
.direction = ckb::ConvDirection::FORWARD,
|
|
.data_type = ckb::DataType::BF16,
|
|
.accumulation_data_type = ckb::DataType::FP32,
|
|
.input = {.config = {.layout = ckb::TensorLayout::NHWGC}},
|
|
.weight = {.config = {.layout = ckb::TensorLayout::GKYXC}},
|
|
.output = {.config = {.layout = ckb::TensorLayout::NHWGK}}};
|
|
|
|
constexpr auto SIGNATURE_NHWGC_FP16_FWD =
|
|
ckt::ConvSignature{.spatial_dim = 2,
|
|
.direction = ckb::ConvDirection::FORWARD,
|
|
.data_type = ckb::DataType::FP16,
|
|
.accumulation_data_type = ckb::DataType::FP32,
|
|
.input = {.config = {.layout = ckb::TensorLayout::NHWGC}},
|
|
.weight = {.config = {.layout = ckb::TensorLayout::GKYXC}},
|
|
.output = {.config = {.layout = ckb::TensorLayout::NHWGK}}};
|
|
|
|
constexpr auto SIGNATURE_NDHWGC_FP32_FWD =
|
|
ckt::ConvSignature{.spatial_dim = 3,
|
|
.direction = ckb::ConvDirection::FORWARD,
|
|
.data_type = ckb::DataType::FP32,
|
|
.accumulation_data_type = ckb::DataType::FP32,
|
|
.input = {.config = {.layout = ckb::TensorLayout::NDHWGC}},
|
|
.weight = {.config = {.layout = ckb::TensorLayout::GKZYXC}},
|
|
.output = {.config = {.layout = ckb::TensorLayout::NDHWGK}}};
|
|
|
|
constexpr auto SIGNATURE_NDHWGC_BF16_FWD =
|
|
ckt::ConvSignature{.spatial_dim = 3,
|
|
.direction = ckb::ConvDirection::FORWARD,
|
|
.data_type = ckb::DataType::BF16,
|
|
.accumulation_data_type = ckb::DataType::FP32,
|
|
.input = {.config = {.layout = ckb::TensorLayout::NDHWGC}},
|
|
.weight = {.config = {.layout = ckb::TensorLayout::GKZYXC}},
|
|
.output = {.config = {.layout = ckb::TensorLayout::NDHWGK}}};
|
|
|
|
constexpr auto SIGNATURE_NDHWGC_FP16_FWD =
|
|
ckt::ConvSignature{.spatial_dim = 3,
|
|
.direction = ckb::ConvDirection::FORWARD,
|
|
.data_type = ckb::DataType::FP16,
|
|
.accumulation_data_type = ckb::DataType::FP32,
|
|
.input = {.config = {.layout = ckb::TensorLayout::NDHWGC}},
|
|
.weight = {.config = {.layout = ckb::TensorLayout::GKZYXC}},
|
|
.output = {.config = {.layout = ckb::TensorLayout::NDHWGK}}};
|
|
|
|
} // namespace ck_tile::builder::profiling
|