mirror of
https://github.com/ROCm/composable_kernel.git
synced 2026-03-22 08:07:38 +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>
82 lines
2.2 KiB
C++
82 lines
2.2 KiB
C++
// Copyright (c) Advanced Micro Devices, Inc., or its affiliates.
|
|
// SPDX-License-Identifier: MIT
|
|
|
|
#include "ck_tile/builder/testing/tensor_buffer.hpp"
|
|
#include "testing_utils.hpp"
|
|
#include <gtest/gtest.h>
|
|
#include <gmock/gmock.h>
|
|
#include <vector>
|
|
|
|
namespace ckb = ck_tile::builder;
|
|
namespace ckt = ck_tile::builder::test;
|
|
|
|
using ck_tile::test::HipError;
|
|
using ck_tile::test::HipSuccess;
|
|
using ::testing::Eq;
|
|
using ::testing::IsNull;
|
|
using ::testing::NotNull;
|
|
using ::testing::Throws;
|
|
|
|
TEST(DeviceBuffer, DefaultToNull)
|
|
{
|
|
ckt::DeviceBuffer buffer;
|
|
EXPECT_THAT(buffer.get(), IsNull());
|
|
}
|
|
|
|
TEST(DeviceBuffer, AllocBuffer)
|
|
{
|
|
const auto size = 12345;
|
|
auto buffer = ckt::alloc_buffer(size);
|
|
|
|
// Pointer should be non-null
|
|
EXPECT_THAT(buffer.get(), NotNull());
|
|
|
|
// Actually, the pointer should be a device pointer
|
|
hipPointerAttribute_t attr;
|
|
EXPECT_THAT(hipPointerGetAttributes(&attr, buffer.get()), HipSuccess());
|
|
|
|
EXPECT_THAT(attr.devicePointer, NotNull());
|
|
EXPECT_THAT(attr.type, Eq(hipMemoryTypeDevice));
|
|
|
|
// Memory should be writable without error
|
|
EXPECT_THAT(hipMemset(buffer.get(), 0xFF, size), HipSuccess());
|
|
}
|
|
|
|
TEST(DeviceBuffer, AutoFree)
|
|
{
|
|
const auto size = 12345;
|
|
std::byte* ptr = nullptr;
|
|
|
|
{
|
|
auto buffer = ckt::alloc_buffer(size);
|
|
ptr = buffer.get();
|
|
}
|
|
|
|
// Trying to use a pointer after freeing should return en error in HIP.
|
|
EXPECT_THAT(hipMemset(ptr, 0xFF, size), HipError(hipErrorInvalidValue));
|
|
}
|
|
|
|
TEST(DeviceBuffer, ThrowsOnOom)
|
|
{
|
|
const auto size = size_t{1} << 60; // 1 exabyte
|
|
|
|
auto check = [] { auto buffer = ckt::alloc_buffer(size); };
|
|
EXPECT_THAT(check, Throws<ckt::OutOfDeviceMemoryError>());
|
|
}
|
|
|
|
TEST(DeviceBuffer, AllocTensorBuffer)
|
|
{
|
|
std::vector<size_t> lengths = {128, 128, 128};
|
|
std::vector<size_t> strides = {128 * 128, 128, 1};
|
|
ckt::TensorDescriptor<ckb::DataType::FP32> descriptor(lengths, strides);
|
|
|
|
auto buffer = ckt::alloc_tensor_buffer(descriptor);
|
|
|
|
// Pointer should be non-null
|
|
EXPECT_THAT(buffer.get(), NotNull());
|
|
|
|
// Memory should be writable without error
|
|
EXPECT_THAT(hipMemset(buffer.get(), 0xFF, descriptor.get_element_space_size_in_bytes()),
|
|
HipSuccess());
|
|
}
|