mirror of
https://github.com/ROCm/composable_kernel.git
synced 2026-03-21 23:57:39 +00:00
6219b12730
* Add README.md for testing * Add tensor_memory_manager. * ck-builder: tensor memory manager rebase fixes This fixes some issues caused by the API being changed recently. Also, this streamlines the ckt namespace to always be ck_tile::builder::test, as this is already being used by other tests Really, this commit should be squashed into the previous, but I'm keeping it separate for brevity. * ck-builder: test arguments initial prototype * ck-builder: test system initial prototype * ck-builder: fix non-standardized copyright comments * ck-builder: new prototype * ck-builder: group testing inputs/outputs into a separate structure This is basically the return of the tensor memory manager after all, except that the design is more closely tied to the actual operation. Using a struct allows us to add additional input/output tensors without breaking code (by defaulting those new parameters). Note that the tensors are split into a separate inputs/outputs because we usually want to allocate the output _twice_: once for the real computation and once for the reference computation. * ck-builder: simplify prototype naming; start docs * ck-builder: update testing readme * ck-builder: testing documentation * ck-builder: HipStatusMatcher This matcher can be used to check HIP status codes and provide nice and readable error messages. * ck-builder: tensor_buffer.hpp tests * ck-builder: conv_fwd.hpp tests * ck-builder: add example end-to-end test in conv fwd 2d fp16 * ck-builder: simplify extent usage * ck-builder: update testing doc * ck-builder: skip end to end test on non-gfx9 * fix check_copyright_year interpreter /bin/bash is not guaranteed to exist on Linux. Signed, a NixOS user * ck-builder: fix copyrights * ck-builder: reduce conv fwd testing size This test allocated 24GB of memory, too much for 16GB cards. --------- Co-authored-by: John Shumway <jshumway@amd.com>
84 lines
2.9 KiB
C++
84 lines
2.9 KiB
C++
// Copyright (c) Advanced Micro Devices, Inc., or its affiliates.
|
|
// SPDX-License-Identifier: MIT
|
|
|
|
#include "impl/conv_signature_types.hpp"
|
|
#include "testing_utils.hpp"
|
|
#include "ck_tile/builder/testing/conv_fwd.hpp"
|
|
#include <gtest/gtest.h>
|
|
#include <gmock/gmock.h>
|
|
#include <vector>
|
|
|
|
namespace ckb = ck_tile::builder;
|
|
namespace ckt = ck_tile::builder::test;
|
|
|
|
using ::testing::ElementsAreArray;
|
|
using ::testing::NotNull;
|
|
|
|
constexpr auto SIGNATURE =
|
|
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 ckt::Args<SIGNATURE> ARGS = {
|
|
.lengths =
|
|
{
|
|
.batch_size = 17,
|
|
.groups = 5,
|
|
.input_channels = 13,
|
|
.output_channels = 44,
|
|
.image =
|
|
{
|
|
.width = 99,
|
|
.height = 125,
|
|
},
|
|
.filter =
|
|
{
|
|
.width = 9,
|
|
.height = 4,
|
|
},
|
|
},
|
|
.filter_strides = {.width = 1, .height = 1},
|
|
.filter_dilation = {.width = 1, .height = 1},
|
|
.input_left_pad = {.width = 0, .height = 0},
|
|
.input_right_pad = {.width = 0, .height = 0},
|
|
.a_elementwise_op = {},
|
|
.b_elementwise_op = {},
|
|
.cde_elementwise_op = {},
|
|
};
|
|
|
|
using Inputs = ckt::Inputs<SIGNATURE>;
|
|
using Outputs = ckt::Outputs<SIGNATURE>;
|
|
using UniqueInputs = ckt::UniqueInputs<SIGNATURE>;
|
|
using UniqueOutputs = ckt::UniqueOutputs<SIGNATURE>;
|
|
|
|
static_assert(ckt::ValidUniqueInputs<SIGNATURE>);
|
|
static_assert(ckt::ValidUniqueOutputs<SIGNATURE>);
|
|
|
|
TEST(ConvFwdTesting, MakeDescriptors)
|
|
{
|
|
const auto get_lengths = [](const auto& descriptor) {
|
|
const auto lengths = descriptor.get_lengths();
|
|
// Google Test cannot print std::span, so turn it into a vector for
|
|
// legibility.
|
|
return std::vector(lengths.begin(), lengths.end());
|
|
};
|
|
|
|
EXPECT_THAT(get_lengths(ARGS.make_input_descriptor()), ElementsAreArray({5, 17, 13, 125, 99}));
|
|
EXPECT_THAT(get_lengths(ARGS.make_weight_descriptor()), ElementsAreArray({5, 44, 13, 4, 9}));
|
|
EXPECT_THAT(get_lengths(ARGS.make_output_descriptor()), ElementsAreArray({5, 17, 44, 122, 91}));
|
|
}
|
|
|
|
TEST(ConvFwdTesting, Alloc)
|
|
{
|
|
auto inputs = alloc_inputs(ARGS);
|
|
auto outputs = alloc_outputs(ARGS);
|
|
|
|
EXPECT_THAT(inputs.get().input, NotNull());
|
|
EXPECT_THAT(inputs.get().weight, NotNull());
|
|
EXPECT_THAT(outputs.get().output, NotNull());
|
|
}
|