Merge commit '6219b12730e29c357a02177dbee6e565987fcc56' into develop

experimental/builder/test/unit_device_buffer.cpp

@@ -0,0 +1,81 @@
+// 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());
+}