mirror of
https://github.com/ROCm/composable_kernel.git
synced 2026-05-14 10:09:41 +00:00
ccbeebe9ea
* 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>
[ROCm/composable_kernel commit: 6219b12730]
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());
|
|
}
|