mirror of
https://github.com/ROCm/composable_kernel.git
synced 2026-03-22 16:17:37 +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>
110 lines
3.7 KiB
C++
110 lines
3.7 KiB
C++
// Copyright (c) Advanced Micro Devices, Inc., or its affiliates.
|
|
// SPDX-License-Identifier: MIT
|
|
|
|
#include <ck/library/tensor_operation_instance/gpu/grouped_convolution_forward.hpp>
|
|
|
|
#include "testing_utils.hpp"
|
|
|
|
using ck_tile::test::HipError;
|
|
using ck_tile::test::HipSuccess;
|
|
using ck_tile::test::InstanceMatcher;
|
|
using ck_tile::test::InstanceSet;
|
|
using ck_tile::test::StringEqWithDiff;
|
|
|
|
TEST(InstanceSet, FromFactory)
|
|
{
|
|
using DeviceOp = ck::tensor_operation::device::DeviceGroupedConvFwdMultipleABD<
|
|
2, // NDimSpatial
|
|
ck::tensor_operation::device::instance::NHWGC, // InLayout
|
|
ck::tensor_operation::device::instance::GKYXC, // WeiLayout
|
|
ck::tensor_operation::device::instance::Empty_Tuple, // DsLayout
|
|
ck::tensor_operation::device::instance::NHWGK, // OutLayout
|
|
ck::half_t, // ADataType
|
|
ck::half_t, // BDataType
|
|
ck::Tuple<>, // DsDataType
|
|
ck::half_t, // EDataType
|
|
ck::tensor_operation::element_wise::PassThrough, // AElementwiseOperation
|
|
ck::tensor_operation::element_wise::PassThrough, // BElementwiseOperation
|
|
ck::tensor_operation::element_wise::PassThrough, // CDEElementwiseOperation
|
|
ck::half_t, // AComputeType
|
|
ck::half_t>; // BComputeType
|
|
|
|
const auto instances = InstanceSet::from_factory<DeviceOp>();
|
|
|
|
EXPECT_THAT(instances.instances.size(), testing::Gt(0));
|
|
|
|
const auto* el =
|
|
"DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle<2,NHWGC,GKYXC,EmptyTuple,NHWGK,fp16,fp16,"
|
|
"fp32,fp16,EmptyTuple,fp16,PassThrough,PassThrough,PassThrough,Default,MNKPadding,1,128,"
|
|
"128,128,32,8,8,32,32,4,2,Seq(4,32,1),Seq(1,0,2),Seq(1,0,2),2,8,8,true,Seq(4,32,1),"
|
|
"Seq(1,0,2),Seq(1,0,2),2,8,8,true,1,1,Seq(1,16,1,8),8,fp16,fp16,Default,1>";
|
|
EXPECT_THAT(instances.instances, testing::Contains(el));
|
|
}
|
|
|
|
TEST(InstanceMatcher, Basic)
|
|
{
|
|
auto a = InstanceSet{
|
|
"python",
|
|
"cobra",
|
|
"boa",
|
|
};
|
|
|
|
auto b = InstanceSet{
|
|
"cobra",
|
|
"boa",
|
|
"python",
|
|
};
|
|
|
|
auto c = InstanceSet{
|
|
"adder",
|
|
"boa",
|
|
"cobra",
|
|
};
|
|
|
|
auto d = InstanceSet{
|
|
"boa",
|
|
"python",
|
|
};
|
|
|
|
EXPECT_THAT(a, InstancesMatch(b));
|
|
EXPECT_THAT(c, Not(InstancesMatch(b)));
|
|
EXPECT_THAT(a, Not(InstancesMatch(d)));
|
|
EXPECT_THAT(d, Not(InstancesMatch(b)));
|
|
}
|
|
|
|
TEST(InstanceMatcher, ExplainMatchResult)
|
|
{
|
|
auto actual = InstanceSet{
|
|
"python",
|
|
"cobra",
|
|
"boa",
|
|
};
|
|
|
|
auto expected = InstanceSet{
|
|
"adder",
|
|
"boa",
|
|
"cobra",
|
|
"rattlesnake",
|
|
};
|
|
|
|
testing::StringMatchResultListener listener;
|
|
EXPECT_TRUE(!ExplainMatchResult(InstancesMatch(expected), actual, &listener));
|
|
|
|
EXPECT_THAT(listener.str(),
|
|
StringEqWithDiff("\n"
|
|
" Missing: 2\n"
|
|
"- adder\n"
|
|
"- rattlesnake\n"
|
|
"Unexpected: 1\n"
|
|
"- python\n"));
|
|
}
|
|
|
|
TEST(HipStatusMatcher, Basic)
|
|
{
|
|
EXPECT_THAT(hipSuccess, HipSuccess());
|
|
EXPECT_THAT(hipErrorInvalidValue, HipError(hipErrorInvalidValue));
|
|
EXPECT_THAT(hipErrorInvalidValue, Not(HipSuccess()));
|
|
EXPECT_THAT(hipSuccess, Not(HipError(hipErrorInvalidValue)));
|
|
EXPECT_THAT(hipErrorOutOfMemory, Not(HipError(hipErrorInvalidValue)));
|
|
}
|