Add separate unit test project for the pinned host memory utils.

README.md

@@ -144,6 +144,18 @@ Running tests
     bin/test_grouped_gemm_splitk
 ```
 
+### Build and running unit tests
+
+At the root directory, run
+
+```bash
+cmake -S test/utility/ -B build_test -DCMAKE_VERBOSE_MAKEFILE=ON
+```
+
+```bash
+cmake --build build_test
+```
+
 ## Optional post-install steps
 
 * Build examples and tests:

include/ck/utility/host_memory_allocator.hpp

@@ -191,7 +191,7 @@ namespace memory {
             {
                 std::cerr << "[ StaticMemPool ] Memory pool exausted." << std::endl;
             }
-            throw std::runtime_error("Memory pool exausted");
+            throw std::runtime_error("[ StaticMemPool ] Memory pool exausted.");
         }
 
         void deallocate(void* p, std::size_t sizeInBytes) override

test/utility/CMakeLists.txt

+cmake_minimum_required(VERSION 3.27)
+
+project(utilities-tests LANGUAGES CXX HIP)
+
+enable_testing()
+include(../../cmake/gtest.cmake)
+
+add_definitions(-D__HIP_PLATFORM_AMD__)
+set(TEST_SOURCE_FILES test_utility.cpp)
+
+set(AMDDeviceLibs_DIR /opt/rocm/lib/cmake/AMDDeviceLibs)
+set(amd_comgr_DIR /opt/rocm/lib/cmake/amd_comgr)
+set(hsa-runtime64_DIR /opt/rocm/lib/cmake/hsa-runtime64)
+find_package(ROCM REQUIRED PATHS /opt/rocm)
+find_package(hip REQUIRED PATHS /opt/rocm)
+
+include_directories(
+    ${PROJECT_SOURCE_DIR}/../include/ck
+    ${PROJECT_SOURCE_DIR}/../include/ck/utility
+    ${PROJECT_SOURCE_DIR}/../include/ck/host_utility
+    /opt/rocm/include
+)
+
+set(CMAKE_CXX_STANDARD 17)
+set(CMAKE_CXX_COMPILER /opt/rocm/bin/hipcc)
+set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -g -O3")
+
+add_executable(utility_tests ${CSRC_FILES} ${TEST_SOURCE_FILES})
+target_link_libraries(utility_tests ${GTEST_LIBRARIES} GTest::gtest GTest::gtest_main hip::device pthread stdc++)
+
+set_target_properties(utility_tests PROPERTIES LINK_FLAGS "-Wl,-allow-shlib-undefined")
+
+add_test(NAME utility_tests COMMAND utility_tests)

test/utility/test_utility.cpp

+// SPDX-License-Identifier: MIT
+// Copyright (c) 2018-2024, Advanced Micro Devices, Inc. All rights reserved.
+#include <vector>
+#include <map>
+#include <queue>
+#include "gtest/gtest.h"
+#include "/workspaces/composable_kernel/include/ck/ck.hpp"
+#include "/workspaces/composable_kernel/include/ck/utility/host_memory_allocator.hpp"
+
+using namespace ck::memory;
+
+TEST(UtilityTests, PinnedHostMemoryAllocator_recycle_pinned_host_memory) 
+{
+    const size_t vSize = 10;
+    int* ptr1;
+    int* ptr2;
+    {
+        std::vector<int, PinnedHostMemoryAllocator<int>> v1(vSize);
+        std::vector<int, PinnedHostMemoryAllocator<int>> v2(2*vSize);
+        EXPECT_EQ(v1.size(), vSize);
+        EXPECT_EQ(v2.size(), 2*vSize);
+        EXPECT_TRUE(v1.data() != nullptr);
+        EXPECT_TRUE(v2.data() != nullptr);
+        EXPECT_TRUE(v1.data() != v2.data());
+
+        ptr1 = v1.data();
+        ptr2 = v2.data();
+    }
+    
+    {
+        // Check that for new vectors, the memory is reused.
+        std::vector<int, PinnedHostMemoryAllocator<int>> v3(vSize);
+        std::vector<int, PinnedHostMemoryAllocator<int>> v4(2*vSize);
+
+        EXPECT_EQ(v3.data(), ptr1);
+        EXPECT_EQ(v4.data(), ptr2);
+    }
+}
+
+TEST(UtilityTests, PinnedHostMemoryAllocator_access_elements) 
+{
+    const size_t vSize = 10;
+    {
+        std::vector<int, PinnedHostMemoryAllocator<int>> v(vSize);
+        for (size_t i = 0; i < vSize; ++i) {
+            v[i] = i;
+        }
+        for (size_t i = 0; i < vSize; ++i) {
+            EXPECT_EQ(v[i], i);
+        }
+    }
+
+    {
+        std::vector<int, PinnedHostMemoryAllocator<int>> v(vSize);
+        for (size_t i = 0; i < vSize; ++i) {
+            v[i] = 2*i;
+        }
+        for (size_t i = 0; i < vSize; ++i) {
+            EXPECT_EQ(v[i], 2*i);
+        }
+    }
+}
+
+TEST(UtilityTests, PinnedHostMemoryAllocator_complex_object) 
+{
+    struct ComplexObject {
+        int a;
+        float b;
+        double c;
+        std::string d;
+    };
+
+    const size_t vSize = 10;
+    {
+        std::vector<ComplexObject, PinnedHostMemoryAllocator<ComplexObject>> v(vSize);
+        for (int i = 0; i < vSize; ++i) {
+            v[i] = ComplexObject{i, 2.0f*i, 3.0*i, "hello" + std::to_string(i)};
+        }
+        for (size_t i = 0; i < vSize; ++i) {
+            EXPECT_EQ(v[i].a, i);
+            EXPECT_EQ(v[i].b, 2.0f*i);
+            EXPECT_EQ(v[i].c, 3.0*i);
+            EXPECT_EQ(v[i].d, "hello"  + std::to_string(i));
+        }
+    }
+}
+
+TEST(UtilityTests, PinnedHostMemoryAllocator_nested_vector) 
+{
+    const size_t vSize = 10;
+    using PinnedHostMemoryAllocatorInt = PinnedHostMemoryAllocator<int>;
+    using PinnedHostMemoryAllocatorVectorInt = PinnedHostMemoryAllocator<std::vector<int, PinnedHostMemoryAllocatorInt>>;
+    {
+        std::vector<std::vector<int, PinnedHostMemoryAllocatorInt>, PinnedHostMemoryAllocatorVectorInt> v(vSize);
+        for (size_t i = 0; i < vSize; ++i) {
+            v[i].resize(i+1);
+            for (size_t j = 0; j < i+1; ++j) {
+                v[i][j] = i*j;
+            }
+        }
+        for (size_t i = 0; i < vSize; ++i) {
+            for (size_t j = 0; j < i+1; ++j) {
+                EXPECT_EQ(v[i][j], i*j);
+            }
+        }
+    }
+}
+
+TEST(UtilityTests, PinnedHostMemoryAllocator_multiple_threads_create_vector_of_same_size)
+{
+    const size_t vSize = 10;
+    const size_t numThreads = 4;
+    std::vector<std::thread> threads;
+    for (size_t i = 0; i < numThreads; ++i) {
+        threads.push_back(std::thread([vSize, i](){
+            std::vector<int, PinnedHostMemoryAllocator<int>> v(vSize);
+            for (size_t j = 0; j < vSize; ++j) {
+                v[j] = i*j;
+            }
+        }));
+    }
+    for (size_t i = 0; i < numThreads; ++i) {
+        threads[i].join();
+    }
+}
+
+TEST(UtilityTests, PinnedHostMemoryAllocator_multiple_vectors_of_same_size_and_different_type)
+{
+    const size_t vSize = 10;
+    {
+        std::vector<int, PinnedHostMemoryAllocator<int>> v1(vSize);
+        std::vector<float, PinnedHostMemoryAllocator<float>> v2(vSize);
+        for (size_t i = 0; i < vSize; ++i) {
+            v1[i] = i;
+            v2[i] = 2.0f*i;
+        }
+        for (size_t i = 0; i < vSize; ++i) {
+            EXPECT_EQ(v1[i], i);
+            EXPECT_EQ(v2[i], 2.0f*i);
+        }
+    }
+}
\ No newline at end of file