EXSWHTEC-251 - Implement Performance Tests for Stream APIs (#142)

Change-Id: I790d8cadb3a94a220a52c8bdc0a23b1a63931232

catch/include/hip_test_defgroups.hh

@@ -115,7 +115,14 @@ THE SOFTWARE.
  * @}
  */
 
- /**
+/**
+ * @defgroup PerformanceTest Performance tests
+ * @{
+ * This section describes performance tests for the target API groups and use-cases.
+ * @}
+ */
+
+/**
  * @defgroup ShflTest warp shuffle function Management
  * @{
  * This section describes the warp shuffle types & functions of HIP runtime API.

catch/include/performance_common.hh

@@ -34,6 +34,7 @@ THE SOFTWARE.
 #include <resource_guards.hh>
 
 #pragma clang diagnostic ignored "-Wunused-but-set-variable"
+#pragma clang diagnostic ignored "-Wunused-parameter"
 #pragma clang diagnostic ignored "-Wunused-function"
 
 #if defined(_WIN32)

catch/performance/CMakeLists.txt

@@ -18,5 +18,6 @@
 # OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
 # THE SOFTWARE.
 
+add_subdirectory(stream)
 add_subdirectory(event)
 add_subdirectory(example)

catch/performance/stream/CMakeLists.txt

@@ -0,0 +1,63 @@
+# Copyright (c) 2022 Advanced Micro Devices, Inc. All Rights Reserved.
+#
+# Permission is hereby granted, free of charge, to any person obtaining a copy
+# of this software and associated documentation files (the "Software"), to deal
+# in the Software without restriction, including without limitation the rights
+# to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+# copies of the Software, and to permit persons to whom the Software is
+# furnished to do so, subject to the following conditions:
+#
+# The above copyright notice and this permission notice shall be included in
+# all copies or substantial portions of the Software.
+#
+# THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+# IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+# FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+# AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+# LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+# OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+# THE SOFTWARE.
+
+if(HIP_PLATFORM MATCHES "amd")
+set(TEST_SRC
+    hipStreamWaitEvent.cc
+    hipStreamGetFlags.cc
+    hipStreamGetPriority.cc
+    hipExtStreamCreateWithCUMask.cc
+    hipExtStreamGetCUMask.cc
+    hipStreamAddCallback.cc
+    hipStreamWaitValue.cc
+    hipStreamWriteValue.cc
+    hipMallocAsync.cc
+    hipFreeAsync.cc
+    hipMemPoolCreate.cc
+    hipMemPoolDestroy.cc
+    hipMemPoolTrimTo.cc
+    hipMemPoolSetAttribute.cc
+    hipMemPoolGetAttribute.cc
+    hipMemPoolSetAccess.cc
+    hipMallocFromPoolAsync.cc
+    hipMemPoolExportToShareableHandle.cc
+    hipMemPoolImportFromShareableHandle.cc
+    hipMemPoolExportPointer.cc
+    hipMemPoolImportPointer.cc
+    hipStreamBasic.cc
+)
+else()
+set(TEST_SRC
+    hipStreamWaitEvent.cc
+    hipStreamGetFlags.cc
+    hipStreamGetPriority.cc
+    hipStreamAddCallback.cc
+    hipStreamWaitValue.cc
+    hipStreamWriteValue.cc
+    hipMallocAsync.cc
+    hipFreeAsync.cc
+    hipStreamBasic.cc
+)
+endif()
+
+hip_add_exe_to_target(NAME StreamPerformance
+                      TEST_SRC ${TEST_SRC}
+                      TEST_TARGET_NAME build_tests
+                      COMPILE_OPTIONS -std=c++17)

catch/performance/stream/hipExtStreamCreateWithCUMask.cc

@@ -0,0 +1,65 @@
+/*
+Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved.
+Permission is hereby granted, free of charge, to any person obtaining a copy
+of this software and associated documentation files (the "Software"), to deal
+in the Software without restriction, including without limitation the rights
+to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+copies of the Software, and to permit persons to whom the Software is
+furnished to do so, subject to the following conditions:
+The above copyright notice and this permission notice shall be included in
+all copies or substantial portions of the Software.
+THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+THE SOFTWARE.
+*/
+
+#include <hip_test_common.hh>
+#include <performance_common.hh>
+
+/**
+ * @addtogroup stream stream
+ * @{
+ * @ingroup PerformanceTest
+ * Contains performance tests for all stream management HIP APIs.
+ */
+
+class ExtStreamCreateWithCUMaskBenchmark : public Benchmark<ExtStreamCreateWithCUMaskBenchmark> {
+ public:
+  void operator()() {
+    hipDeviceProp_t props;
+    HIP_CHECK(hipGetDeviceProperties(&props, 0));
+    std::vector<uint32_t> cu_mask(props.multiProcessorCount, 0);
+    hipStream_t stream{};
+
+    TIMED_SECTION(kTimerTypeCpu) {
+      HIP_CHECK(hipExtStreamCreateWithCUMask(&stream, cu_mask.size(), cu_mask.data()));
+    }
+
+    HIP_CHECK(hipStreamDestroy(stream));
+  }
+};
+
+static void RunBenchmark() {
+  ExtStreamCreateWithCUMaskBenchmark benchmark;
+  benchmark.Run();
+}
+
+/**
+ * Test Description
+ * ------------------------
+ *  - Executes `hipExtStreamCreateWithCUMask`.
+ * Test source
+ * ------------------------
+ *  - performance/stream/hipExtStreamCreateWithCUMask.cc
+ * Test requirements
+ * ------------------------
+ *  - Platform specific (AMD)
+ *  - HIP_VERSION >= 5.2
+ */
+TEST_CASE("Performance_hipExtStreamCreateWithCUMask") {
+  RunBenchmark();
+}

catch/performance/stream/hipExtStreamGetCUMask.cc

@@ -0,0 +1,67 @@
+/*
+Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved.
+Permission is hereby granted, free of charge, to any person obtaining a copy
+of this software and associated documentation files (the "Software"), to deal
+in the Software without restriction, including without limitation the rights
+to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+copies of the Software, and to permit persons to whom the Software is
+furnished to do so, subject to the following conditions:
+The above copyright notice and this permission notice shall be included in
+all copies or substantial portions of the Software.
+THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+THE SOFTWARE.
+*/
+
+#include <hip_test_common.hh>
+#include <performance_common.hh>
+
+/**
+ * @addtogroup stream stream
+ * @{
+ * @ingroup PerformanceTest
+ */
+
+class ExtStreamGetCUMaskBenchmark : public Benchmark<ExtStreamGetCUMaskBenchmark> {
+ public:
+  void operator()() {
+    hipDeviceProp_t props;
+    HIP_CHECK(hipGetDeviceProperties(&props, 0));
+    std::vector<uint32_t> cu_mask(props.multiProcessorCount, 0);
+    hipStream_t stream{};
+    HIP_CHECK(hipExtStreamCreateWithCUMask(&stream, cu_mask.size(), cu_mask.data()));
+    std::vector<uint32_t> new_cu_mask(cu_mask.size(), 0);
+
+    TIMED_SECTION(kTimerTypeCpu) {
+      HIP_CHECK(hipExtStreamGetCUMask(stream, new_cu_mask.size(), new_cu_mask.data()));
+    }
+
+    HIP_CHECK(hipStreamDestroy(stream));
+  }
+};
+
+static void RunBenchmark() {
+  ExtStreamGetCUMaskBenchmark benchmark;
+  benchmark.Run();
+}
+
+/**
+ * Test Description
+ * ------------------------
+ *  - Executes `hipExtStreamGetCUMask`.
+ *  - Creates basic mask and gets it into the new one.
+ * Test source
+ * ------------------------
+ *  - performance/stream/hipExtStreamGetCUMask.cc
+ * Test requirements
+ * ------------------------
+ *  - Platform specific (AMD)
+ *  - HIP_VERSION >= 5.2
+ */
+TEST_CASE("Performance_hipExtStreamGetCUMask") {
+  RunBenchmark();
+}

catch/performance/stream/hipFreeAsync.cc

@@ -0,0 +1,69 @@
+/*
+Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved.
+Permission is hereby granted, free of charge, to any person obtaining a copy
+of this software and associated documentation files (the "Software"), to deal
+in the Software without restriction, including without limitation the rights
+to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+copies of the Software, and to permit persons to whom the Software is
+furnished to do so, subject to the following conditions:
+The above copyright notice and this permission notice shall be included in
+all copies or substantial portions of the Software.
+THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+THE SOFTWARE.
+*/
+
+#include <hip_test_common.hh>
+#include <performance_common.hh>
+
+/**
+ * @addtogroup stream stream
+ * @{
+ * @ingroup PerformanceTest
+ */
+
+class FreeAsyncBenchmark : public Benchmark<FreeAsyncBenchmark> {
+ public:
+  void operator()(const size_t array_size) {
+    const StreamGuard stream_guard{Streams::created};
+    const hipStream_t stream = stream_guard.stream();
+    float* dev_ptr{nullptr};
+    HIP_CHECK(hipMallocAsync(reinterpret_cast<void**>(&dev_ptr), array_size * sizeof(float), stream));
+
+    TIMED_SECTION_STREAM(kTimerTypeEvent, stream) {
+      HIP_CHECK(hipFreeAsync(dev_ptr, stream));
+    }
+
+    HIP_CHECK(hipStreamSynchronize(stream));
+  }
+};
+
+static void RunBenchmark(const size_t array_size) {
+  FreeAsyncBenchmark benchmark;
+  benchmark.AddSectionName(std::to_string(array_size));
+  benchmark.Run(array_size);
+}
+
+/**
+ * Test Description
+ * ------------------------
+ *  - Executes `hipFreeAsync` with created stream:
+ *    -# Allocation size:
+ *      - 4 KB
+ *      - 4 MB
+ *      - 16 MB
+ * Test source
+ * ------------------------
+ *  - performance/stream/hipFreeAsync.cc
+ * Test requirements
+ * ------------------------
+ *  - HIP_VERSION >= 5.2
+ */
+TEST_CASE("Performance_hipFreeAsync") {
+  size_t array_size = GENERATE(4_KB, 4_MB, 16_MB);
+  RunBenchmark(array_size);
+}

catch/performance/stream/hipMallocAsync.cc

@@ -0,0 +1,68 @@
+/*
+Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved.
+Permission is hereby granted, free of charge, to any person obtaining a copy
+of this software and associated documentation files (the "Software"), to deal
+in the Software without restriction, including without limitation the rights
+to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+copies of the Software, and to permit persons to whom the Software is
+furnished to do so, subject to the following conditions:
+The above copyright notice and this permission notice shall be included in
+all copies or substantial portions of the Software.
+THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+THE SOFTWARE.
+*/
+
+#include <hip_test_common.hh>
+#include <performance_common.hh>
+
+/**
+ * @addtogroup stream stream
+ * @{
+ * @ingroup PerformanceTest
+ */
+
+class MallocAsyncBenchmark : public Benchmark<MallocAsyncBenchmark> {
+ public:
+  void operator()(const size_t array_size) {
+    const StreamGuard stream_guard{Streams::created};
+    const hipStream_t stream = stream_guard.stream();
+    float* dev_ptr{nullptr};
+
+    TIMED_SECTION_STREAM(kTimerTypeEvent, stream) {
+      HIP_CHECK(hipMallocAsync(reinterpret_cast<void**>(&dev_ptr), array_size * sizeof(float), stream));
+    }
+    HIP_CHECK(hipStreamSynchronize(stream));
+    HIP_CHECK(hipFree(dev_ptr));
+  }
+};
+
+static void RunBenchmark(const size_t array_size) {
+  MallocAsyncBenchmark benchmark;
+  benchmark.AddSectionName(std::to_string(array_size));
+  benchmark.Run(array_size);
+}
+
+/**
+ * Test Description
+ * ------------------------
+ *  - Executes `hipMallocAsync` with created stream:
+ *    -# Allocation size:
+ *      - 4 KB
+ *      - 4 MB
+ *      - 16 MB
+ * Test source
+ * ------------------------
+ *  - performance/stream/hipMallocAsync.cc
+ * Test requirements
+ * ------------------------
+ *  - HIP_VERSION >= 5.2
+ */
+TEST_CASE("Performance_hipMallocAsync") {
+  size_t array_size = GENERATE(4_KB, 4_MB, 16_MB);
+  RunBenchmark(array_size);
+}

catch/performance/stream/hipMallocFromPoolAsync.cc

@@ -0,0 +1,82 @@
+/*
+Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved.
+Permission is hereby granted, free of charge, to any person obtaining a copy
+of this software and associated documentation files (the "Software"), to deal
+in the Software without restriction, including without limitation the rights
+to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+copies of the Software, and to permit persons to whom the Software is
+furnished to do so, subject to the following conditions:
+The above copyright notice and this permission notice shall be included in
+all copies or substantial portions of the Software.
+THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+THE SOFTWARE.
+*/
+
+#include "mem_pools_performance_common.hh"
+
+/**
+ * @addtogroup stream stream
+ * @{
+ * @ingroup PerformanceTest
+ */
+
+class MallocFromPoolAsyncBenchmark : public Benchmark<MallocFromPoolAsyncBenchmark> {
+ public:
+  void operator()(const size_t array_size) {
+    const StreamGuard stream_guard{Streams::created};
+    const hipStream_t stream = stream_guard.stream();
+
+    hipMemPool_t mem_pool{nullptr};
+    hipMemPoolProps pool_props = CreateMemPoolProps(0, hipMemHandleTypeNone);
+    HIP_CHECK(hipMemPoolCreate(&mem_pool, &pool_props));
+
+    float* array_ptr{nullptr};
+
+    TIMED_SECTION_STREAM(kTimerTypeEvent, stream) {
+      HIP_CHECK(hipMallocFromPoolAsync(&array_ptr, array_size * sizeof(float), mem_pool, stream));
+    }
+
+    REQUIRE(array_ptr != nullptr);
+
+    HIP_CHECK(hipFreeAsync(array_ptr, stream));
+    HIP_CHECK(hipStreamSynchronize(stream));
+    HIP_CHECK(hipMemPoolDestroy(mem_pool));
+  }
+};
+
+static void RunBenchmark(const size_t array_size) {
+  MallocFromPoolAsyncBenchmark benchmark;
+  benchmark.AddSectionName(std::to_string(array_size));
+  benchmark.Run(array_size);
+}
+
+/**
+ * Test Description
+ * ------------------------
+ *  - Executes `hipMallocFromPoolAsync`:
+ *    -# Allocation size:
+ *      - 4 KB
+ *      - 4 MB
+ *      - 16 MB
+ * Test source
+ * ------------------------
+ *  - performance/stream/hipMallocFromPoolAsync.cc
+ * Test requirements
+ * ------------------------
+ *  - Device supports memory pools
+ *  - HIP_VERSION >= 5.2
+ */
+TEST_CASE("Performance_hipMallocFromPoolAsync") {
+  if (!AreMemPoolsSupported(0)) {
+    HipTest::HIP_SKIP_TEST("GPU 0 doesn't support hipDeviceAttributeMemoryPoolsSupported "
+                           "attribute. Hence skipping the testing with Pass result.\n");
+    return;
+  }
+  size_t array_size = GENERATE(4_KB, 4_MB, 16_MB);
+  RunBenchmark(array_size);
+}

catch/performance/stream/hipMemPoolCreate.cc

@@ -0,0 +1,71 @@
+/*
+Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved.
+Permission is hereby granted, free of charge, to any person obtaining a copy
+of this software and associated documentation files (the "Software"), to deal
+in the Software without restriction, including without limitation the rights
+to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+copies of the Software, and to permit persons to whom the Software is
+furnished to do so, subject to the following conditions:
+The above copyright notice and this permission notice shall be included in
+all copies or substantial portions of the Software.
+THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+THE SOFTWARE.
+*/
+
+#include "mem_pools_performance_common.hh"
+
+/**
+ * @addtogroup stream stream
+ * @{
+ * @ingroup PerformanceTest
+ */
+
+class MemPoolCreateBenchmark : public Benchmark<MemPoolCreateBenchmark> {
+ public:
+  void operator()() {
+    hipMemPool_t mem_pool{nullptr};
+    hipMemPoolProps pool_props = CreateMemPoolProps(0, hipMemHandleTypeNone);
+
+    TIMED_SECTION(kTimerTypeCpu) {
+      HIP_CHECK(hipMemPoolCreate(&mem_pool, &pool_props));
+    }
+
+    REQUIRE(mem_pool != nullptr);
+    HIP_CHECK(hipMemPoolDestroy(mem_pool));
+  }
+};
+
+static void RunBenchmark() {
+  MemPoolCreateBenchmark benchmark;
+  benchmark.Run();
+}
+
+/**
+ * @warning **MemPool APIs are not fully implemented within current version
+ *          or HIP and therefore they cannot be appropriately executed on AMD and NVIDIA platforms.
+ *          Therefore, all tests related to MemPool APIs are implemented without formal
+ *          verification and will be verified once HIP fully supports MemPool APIs.**
+ * Test Description
+ * ------------------------
+ *  - Executes `hipMemPoolCreate`.
+ * Test source
+ * ------------------------
+ *  - performance/stream/hipMemPoolCreate.cc
+ * Test requirements
+ * ------------------------
+ *  - Device supports memory pools
+ *  - HIP_VERSION >= 5.2
+ */
+TEST_CASE("Performance_hipMemPoolCreate") {
+  if (!AreMemPoolsSupported(0)) {
+    HipTest::HIP_SKIP_TEST("GPU 0 doesn't support hipDeviceAttributeMemoryPoolsSupported "
+                           "attribute. Hence skipping the testing with Pass result.\n");
+    return;
+  }
+  RunBenchmark();
+}

catch/performance/stream/hipMemPoolDestroy.cc

@@ -0,0 +1,70 @@
+/*
+Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved.
+Permission is hereby granted, free of charge, to any person obtaining a copy
+of this software and associated documentation files (the "Software"), to deal
+in the Software without restriction, including without limitation the rights
+to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+copies of the Software, and to permit persons to whom the Software is
+furnished to do so, subject to the following conditions:
+The above copyright notice and this permission notice shall be included in
+all copies or substantial portions of the Software.
+THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+THE SOFTWARE.
+*/
+
+#include "mem_pools_performance_common.hh"
+
+/**
+ * @addtogroup stream stream
+ * @{
+ * @ingroup PerformanceTest
+ */
+
+class MemPoolDestroyBenchmark : public Benchmark<MemPoolDestroyBenchmark> {
+ public:
+  void operator()() {
+    hipMemPool_t mem_pool{nullptr};
+    hipMemPoolProps pool_props = CreateMemPoolProps(0, hipMemHandleTypeNone);
+    HIP_CHECK(hipMemPoolCreate(&mem_pool, &pool_props));
+
+    TIMED_SECTION(kTimerTypeCpu) {
+      HIP_CHECK(hipMemPoolDestroy(mem_pool));
+    }
+  }
+};
+
+static void RunBenchmark() {
+  MemPoolDestroyBenchmark benchmark;
+  benchmark.Run();
+}
+
+/**
+ * @warning **MemPool APIs are not fully implemented within current version
+ *          or HIP and therefore they cannot be appropriately executed on AMD and NVIDIA platforms.
+ *          Therefore, all tests related to MemPool APIs are implemented without formal
+ *          verification and will be verified once HIP fully supports MemPool APIs.**
+ * Test Description
+ * ------------------------
+ *  - Creates new mem pool.
+ *  - Executes `hipMemPoolDestroy`.
+ * Test source
+ * ------------------------
+ *  - performance/stream/hipMemPoolDestroy.cc
+ * Test requirements
+ * ------------------------
+ *  - Device supports memory pools
+ *  - HIP_VERSION >= 5.2
+ */
+TEST_CASE("Performance_hipMemPoolDestroy") {
+  if (!AreMemPoolsSupported(0)) {
+    HipTest::HIP_SKIP_TEST("GPU 0 doesn't support hipDeviceAttributeMemoryPoolsSupported "
+                           "attribute. Hence skipping the testing with Pass result.\n");
+    return;
+  }
+  RunBenchmark();
+}

catch/performance/stream/hipMemPoolExportPointer.cc

@@ -0,0 +1,84 @@
+/*
+Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved.
+Permission is hereby granted, free of charge, to any person obtaining a copy
+of this software and associated documentation files (the "Software"), to deal
+in the Software without restriction, including without limitation the rights
+to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+copies of the Software, and to permit persons to whom the Software is
+furnished to do so, subject to the following conditions:
+The above copyright notice and this permission notice shall be included in
+all copies or substantial portions of the Software.
+THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+THE SOFTWARE.
+*/
+
+#include "mem_pools_performance_common.hh"
+
+/**
+ * @addtogroup stream stream
+ * @{
+ * @ingroup PerformanceTest
+ */
+
+class MemPoolExportPointerBenchmark : public Benchmark<MemPoolExportPointerBenchmark> {
+ public:
+  void operator()(const size_t array_size) {
+    float* device_ptr{nullptr};
+    hipMemPool_t mem_pool{nullptr};
+    hipMemPoolPtrExportData exp_data;
+
+    hipMemPoolProps props = CreateMemPoolProps(0, kHandleType);
+    HIP_CHECK(hipMemPoolCreate(&mem_pool, &props));
+    HIP_CHECK(hipMallocFromPoolAsync(&device_ptr, array_size * sizeof(float), mem_pool, nullptr));
+    HIP_CHECK(hipStreamSynchronize(nullptr));
+
+    TIMED_SECTION(kTimerTypeCpu) {
+      HIP_CHECK(hipMemPoolExportPointer(&exp_data, device_ptr));
+    }
+
+    HIP_CHECK(hipFreeAsync(device_ptr, nullptr));
+    HIP_CHECK(hipMemPoolDestroy(mem_pool));
+  }
+};
+
+static void RunBenchmark(const size_t array_size) {
+  MemPoolExportPointerBenchmark benchmark;
+  benchmark.AddSectionName(std::to_string(array_size));
+  benchmark.Run(array_size);
+}
+
+/**
+ * @warning **MemPool APIs are not fully implemented within current version
+ *          or HIP and therefore they cannot be appropriately executed on AMD and NVIDIA platforms.
+ *          Therefore, all tests related to MemPool APIs are implemented without formal
+ *          verification and will be verified once HIP fully supports MemPool APIs.**
+ * Test Description
+ * ------------------------
+ *  - Executes `hipMemPoolExportPointer`:
+ *    -# Allocation size:
+ *      - 4 KB
+ *      - 4 MB
+ *      - 16 MB
+ *  - Uses the same process for import and export operations.
+ * Test source
+ * ------------------------
+ *  - performance/stream/hipMemPoolExportPointer.cc
+ * Test requirements
+ * ------------------------
+ *  - Device supports memory pools
+ *  - HIP_VERSION >= 5.2
+ */
+TEST_CASE("Performance_hipMemPoolExportPointer") {
+  if (!AreMemPoolsSupported(0)) {
+    HipTest::HIP_SKIP_TEST("GPU 0 doesn't support hipDeviceAttributeMemoryPoolsSupported "
+                           "attribute. Hence skipping the testing with Pass result.\n");
+    return;
+  }
+  size_t array_size = GENERATE(4_KB, 4_MB, 16_MB);
+  RunBenchmark(array_size);
+}

catch/performance/stream/hipMemPoolExportToShareableHandle.cc

@@ -0,0 +1,74 @@
+/*
+Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved.
+Permission is hereby granted, free of charge, to any person obtaining a copy
+of this software and associated documentation files (the "Software"), to deal
+in the Software without restriction, including without limitation the rights
+to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+copies of the Software, and to permit persons to whom the Software is
+furnished to do so, subject to the following conditions:
+The above copyright notice and this permission notice shall be included in
+all copies or substantial portions of the Software.
+THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+THE SOFTWARE.
+*/
+
+#include "mem_pools_performance_common.hh"
+
+/**
+ * @addtogroup stream stream
+ * @{
+ * @ingroup PerformanceTest
+ */
+
+class MemPoolExportToShareableHandleBenchmark : public Benchmark<MemPoolExportToShareableHandleBenchmark> {
+ public:
+  void operator()() {
+    hipMemPool_t mem_pool{nullptr};
+    int share_handle;
+
+    hipMemPoolProps props = CreateMemPoolProps(0, kHandleType);
+    HIP_CHECK(hipMemPoolCreate(&mem_pool, &props));
+
+    TIMED_SECTION(kTimerTypeCpu) {
+      HIP_CHECK(hipMemPoolExportToShareableHandle(&share_handle, mem_pool, kHandleType, 0));
+    }
+
+    HIP_CHECK(hipMemPoolDestroy(mem_pool));
+  }
+};
+
+static void RunBenchmark() {
+  MemPoolExportToShareableHandleBenchmark benchmark;
+  benchmark.Run();
+}
+
+/**
+ * @warning **MemPool APIs are not fully implemented within current version
+ *          or HIP and therefore they cannot be appropriately executed on AMD and NVIDIA platforms.
+ *          Therefore, all tests related to MemPool APIs are implemented without formal
+ *          verification and will be verified once HIP fully supports MemPool APIs.**
+ * Test Description
+ * ------------------------
+ *  - Executes `hipMemPoolExportToShareableHandle`.
+ *  - Uses the same process for import and export operations.
+ * Test source
+ * ------------------------
+ *  - performance/stream/hipMemPoolExportToShareableHandle.cc
+ * Test requirements
+ * ------------------------
+ *  - Device supports memory pools
+ *  - HIP_VERSION >= 5.2
+ */
+TEST_CASE("Performance_hipMemPoolExportToShareableHandle") {
+  if (!AreMemPoolsSupported(0)) {
+    HipTest::HIP_SKIP_TEST("GPU 0 doesn't support hipDeviceAttributeMemoryPoolsSupported "
+                           "attribute. Hence skipping the testing with Pass result.\n");
+    return;
+  }
+  RunBenchmark();
+}

catch/performance/stream/hipMemPoolGetAccess.cc

@@ -0,0 +1,76 @@
+/*
+Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved.
+Permission is hereby granted, free of charge, to any person obtaining a copy
+of this software and associated documentation files (the "Software"), to deal
+in the Software without restriction, including without limitation the rights
+to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+copies of the Software, and to permit persons to whom the Software is
+furnished to do so, subject to the following conditions:
+The above copyright notice and this permission notice shall be included in
+all copies or substantial portions of the Software.
+THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+THE SOFTWARE.
+*/
+
+#include "mem_pools_performance_common.hh"
+
+/**
+ * @addtogroup stream stream
+ * @{
+ * @ingroup PerformanceTest
+ */
+
+class MemPoolGetAccessBenchmark : public Benchmark<MemPoolGetAccessBenchmark> {
+ public:
+  void operator()() {
+    hipMemPool_t mem_pool{nullptr};
+    hipMemPoolProps pool_props = CreateMemPoolProps(0, hipMemHandleTypeNone);
+    HIP_CHECK(hipMemPoolCreate(&mem_pool, &pool_props));
+
+    hipMemAccessFlags flags = hipMemAccessFlagsProtNone;
+    hipMemLocation location = {
+      hipMemLocationTypeDevice,
+      0
+    };
+    TIMED_SECTION(kTimerTypeCpu) {
+      HIP_CHECK(hipMemPoolGetAccess(&flags, mem_pool, location));
+    }
+
+    HIP_CHECK(hipMemPoolDestroy(mem_pool));
+  }
+};
+
+static void RunBenchmark() {
+  MemPoolGetAccessBenchmark benchmark;
+  benchmark.Run();
+}
+
+/**
+ * @warning **MemPool APIs are not fully implemented within current version
+ *          or HIP and therefore they cannot be appropriately executed on AMD and NVIDIA platforms.
+ *          Therefore, all tests related to MemPool APIs are implemented without formal
+ *          verification and will be verified once HIP fully supports MemPool APIs.**
+ * Test Description
+ * ------------------------
+ *  - Executes `hipMemPoolGetAccess`.
+ * Test source
+ * ------------------------
+ *  - performance/stream/hipMemPoolGetAccess.cc
+ * Test requirements
+ * ------------------------
+ *  - Device supports memory pools
+ *  - HIP_VERSION >= 5.2
+ */
+TEST_CASE("Performance_hipMemPoolGetAccess") {
+  if (!AreMemPoolsSupported(0)) {
+    HipTest::HIP_SKIP_TEST("GPU 0 doesn't support hipDeviceAttributeMemoryPoolsSupported "
+                           "attribute. Hence skipping the testing with Pass result.\n");
+    return;
+  }
+  RunBenchmark();
+}

catch/performance/stream/hipMemPoolGetAttribute.cc

@@ -0,0 +1,83 @@
+/*
+Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved.
+Permission is hereby granted, free of charge, to any person obtaining a copy
+of this software and associated documentation files (the "Software"), to deal
+in the Software without restriction, including without limitation the rights
+to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+copies of the Software, and to permit persons to whom the Software is
+furnished to do so, subject to the following conditions:
+The above copyright notice and this permission notice shall be included in
+all copies or substantial portions of the Software.
+THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+THE SOFTWARE.
+*/
+
+#include "mem_pools_performance_common.hh"
+
+/**
+ * @addtogroup stream stream
+ * @{
+ * @ingroup PerformanceTest
+ */
+
+class MemPoolGetAttributeBenchmark : public Benchmark<MemPoolGetAttributeBenchmark> {
+ public:
+  void operator()(const hipMemPoolAttr attribute) {
+    hipMemPool_t mem_pool{nullptr};
+    hipMemPoolProps pool_props = CreateMemPoolProps(0, hipMemHandleTypeNone);
+    HIP_CHECK(hipMemPoolCreate(&mem_pool, &pool_props));
+
+    uint64_t value{0};
+
+    TIMED_SECTION(kTimerTypeCpu) {
+      HIP_CHECK(hipMemPoolGetAttribute(mem_pool, attribute, &value));
+    }
+
+    HIP_CHECK(hipMemPoolDestroy(mem_pool));
+  }
+};
+
+static void RunBenchmark(const hipMemPoolAttr attribute) {
+  MemPoolGetAttributeBenchmark benchmark;
+  benchmark.AddSectionName(GetMemPoolAttrSectionName(attribute));
+  benchmark.Run(attribute);
+}
+
+/**
+ * @warning **MemPool APIs are not fully implemented within current version
+ *          or HIP and therefore they cannot be appropriately executed on AMD and NVIDIA platforms.
+ *          Therefore, all tests related to MemPool APIs are implemented without formal
+ *          verification and will be verified once HIP fully supports MemPool APIs.**
+ * Test Description
+ * ------------------------
+ *  - Executes `hipMemPoolGetAttribute`:
+ *    -# Supported attributes:
+ *      - `hipMemPoolAttrReleaseThreshold`
+ *      - `hipMemPoolReuseFollowEventDependencies`
+ *      - `hipMemPoolReuseAllowOpportunistic`
+ *      - `hipMemPoolReuseAllowInternalDependencies`
+ * Test source
+ * ------------------------
+ *  - performance/stream/hipMemPoolGetAttribute.cc
+ * Test requirements
+ * ------------------------
+ *  - Device supports memory pools
+ *  - HIP_VERSION >= 5.2
+ */
+TEST_CASE("Performance_hipMemPoolGetAttribute") {
+  if (!AreMemPoolsSupported(0)) {
+    HipTest::HIP_SKIP_TEST("GPU 0 doesn't support hipDeviceAttributeMemoryPoolsSupported "
+                           "attribute. Hence skipping the testing with Pass result.\n");
+    return;
+  }
+  hipMemPoolAttr attribute = GENERATE(hipMemPoolAttrReleaseThreshold,
+                                      hipMemPoolReuseFollowEventDependencies,
+                                      hipMemPoolReuseAllowOpportunistic,
+                                      hipMemPoolReuseAllowInternalDependencies);
+  RunBenchmark(attribute);
+}

catch/performance/stream/hipMemPoolImportFromShareableHandle.cc

@@ -0,0 +1,75 @@
+/*
+Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved.
+Permission is hereby granted, free of charge, to any person obtaining a copy
+of this software and associated documentation files (the "Software"), to deal
+in the Software without restriction, including without limitation the rights
+to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+copies of the Software, and to permit persons to whom the Software is
+furnished to do so, subject to the following conditions:
+The above copyright notice and this permission notice shall be included in
+all copies or substantial portions of the Software.
+THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+THE SOFTWARE.
+*/
+
+#include "mem_pools_performance_common.hh"
+
+/**
+ * @addtogroup stream stream
+ * @{
+ * @ingroup PerformanceTest
+ */
+
+class MemPoolImportFromShareableHandleBenchmark : public Benchmark<MemPoolImportFromShareableHandleBenchmark> {
+ public:
+  void operator()() {
+    hipMemPool_t mem_pool{nullptr};
+    int share_handle;
+
+    hipMemPoolProps props = CreateMemPoolProps(0, kHandleType);
+    HIP_CHECK(hipMemPoolCreate(&mem_pool, &props));
+    HIP_CHECK(hipMemPoolExportToShareableHandle(&share_handle, mem_pool, kHandleType, 0));
+
+    TIMED_SECTION(kTimerTypeCpu) {
+      HIP_CHECK(hipMemPoolImportFromShareableHandle(&mem_pool, &share_handle, kHandleType, 0));
+    }
+
+    HIP_CHECK(hipMemPoolDestroy(mem_pool));
+  }
+};
+
+static void RunBenchmark() {
+  MemPoolImportFromShareableHandleBenchmark benchmark;
+  benchmark.Run();
+}
+
+/**
+ * @warning **MemPool APIs are not fully implemented within current version
+ *          or HIP and therefore they cannot be appropriately executed on AMD and NVIDIA platforms.
+ *          Therefore, all tests related to MemPool APIs are implemented without formal
+ *          verification and will be verified once HIP fully supports MemPool APIs.**
+ * Test Description
+ * ------------------------
+ *  - Executes `hipMemPoolImportFromShareableHandle`.
+ *  - Uses the same process for import and export operations.
+ * Test source
+ * ------------------------
+ *  - performance/stream/hipMemPoolImportFromShareableHandle.cc
+ * Test requirements
+ * ------------------------
+ *  - Device supports memory pools
+ *  - HIP_VERSION >= 5.2
+ */
+TEST_CASE("Performance_hipMemPoolImportFromShareableHandle") {
+  if (!AreMemPoolsSupported(0)) {
+    HipTest::HIP_SKIP_TEST("GPU 0 doesn't support hipDeviceAttributeMemoryPoolsSupported "
+                           "attribute. Hence skipping the testing with Pass result.\n");
+    return;
+  }
+  RunBenchmark();
+}

catch/performance/stream/hipMemPoolImportPointer.cc

@@ -0,0 +1,87 @@
+/*
+Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved.
+Permission is hereby granted, free of charge, to any person obtaining a copy
+of this software and associated documentation files (the "Software"), to deal
+in the Software without restriction, including without limitation the rights
+to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+copies of the Software, and to permit persons to whom the Software is
+furnished to do so, subject to the following conditions:
+The above copyright notice and this permission notice shall be included in
+all copies or substantial portions of the Software.
+THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+THE SOFTWARE.
+*/
+
+#include "mem_pools_performance_common.hh"
+
+/**
+ * @addtogroup stream stream
+ * @{
+ * @ingroup PerformanceTest
+ */
+
+class MemPoolImportPointerBenchmark : public Benchmark<MemPoolImportPointerBenchmark> {
+ public:
+  void operator()(const size_t array_size) {
+    float* device_ptr{nullptr};
+    float* device_ptr_import{nullptr};
+    hipMemPool_t mem_pool{nullptr};
+    hipMemPoolPtrExportData exp_data;
+
+    hipMemPoolProps props = CreateMemPoolProps(0, kHandleType);
+    HIP_CHECK(hipMemPoolCreate(&mem_pool, &props));
+    HIP_CHECK(hipMallocFromPoolAsync(&device_ptr, array_size * sizeof(float), mem_pool, nullptr));
+    HIP_CHECK(hipStreamSynchronize(nullptr));
+    HIP_CHECK(hipMemPoolExportPointer(&exp_data, device_ptr));
+
+    TIMED_SECTION(kTimerTypeCpu) {
+      HIP_CHECK(hipMemPoolImportPointer(reinterpret_cast<void**>(device_ptr_import), mem_pool, &exp_data));
+    }
+
+    HIP_CHECK(hipFree(device_ptr));
+    HIP_CHECK(hipFree(device_ptr_import));
+    HIP_CHECK(hipMemPoolDestroy(mem_pool));
+  }
+};
+
+static void RunBenchmark(const size_t array_size) {
+  MemPoolImportPointerBenchmark benchmark;
+  benchmark.AddSectionName(std::to_string(array_size));
+  benchmark.Run(array_size);
+}
+
+/**
+ * @warning **MemPool APIs are not fully implemented within current version
+ *          or HIP and therefore they cannot be appropriately executed on AMD and NVIDIA platforms.
+ *          Therefore, all tests related to MemPool APIs are implemented without formal
+ *          verification and will be verified once HIP fully supports MemPool APIs.**
+ * Test Description
+ * ------------------------
+ *  - Executes `hipMemPoolImportPointer`:
+ *    -# Allocation size:
+ *      - 4 KB
+ *      - 4 MB
+ *      - 16 MB
+ *  - Uses the same process for import and export operations.
+ * Test source
+ * ------------------------
+ *  - performance/stream/hipMemPoolImportPointer.cc
+ * Test requirements
+ * ------------------------
+ *  - Device supports memory pools
+ *  - HIP_VERSION >= 5.2
+ */
+TEST_CASE("Performance_hipMemPoolImportPointer") {
+  if (!AreMemPoolsSupported(0)) {
+    HipTest::HIP_SKIP_TEST("GPU 0 doesn't support hipDeviceAttributeMemoryPoolsSupported "
+                           "attribute. Hence skipping the testing with Pass result.\n");
+    return;
+  }
+  size_t array_size = GENERATE(4_KB, 4_MB, 16_MB);
+  RunBenchmark(array_size);
+}

catch/performance/stream/hipMemPoolSetAccess.cc

@@ -0,0 +1,79 @@
+/*
+Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved.
+Permission is hereby granted, free of charge, to any person obtaining a copy
+of this software and associated documentation files (the "Software"), to deal
+in the Software without restriction, including without limitation the rights
+to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+copies of the Software, and to permit persons to whom the Software is
+furnished to do so, subject to the following conditions:
+The above copyright notice and this permission notice shall be included in
+all copies or substantial portions of the Software.
+THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+THE SOFTWARE.
+*/
+
+#include "mem_pools_performance_common.hh"
+
+/**
+ * @addtogroup stream stream
+ * @{
+ * @ingroup PerformanceTest
+ */
+
+class MemPoolSetAccessBenchmark : public Benchmark<MemPoolSetAccessBenchmark> {
+ public:
+  void operator()() {
+    hipMemPool_t mem_pool{nullptr};
+    hipMemPoolProps pool_props = CreateMemPoolProps(0, hipMemHandleTypeNone);
+    HIP_CHECK(hipMemPoolCreate(&mem_pool, &pool_props));
+
+    hipMemAccessDesc desc_list = {
+      {
+        hipMemLocationTypeDevice,
+        0
+      },
+      hipMemAccessFlagsProtReadWrite
+    };
+
+    TIMED_SECTION(kTimerTypeCpu) {
+      HIP_CHECK(hipMemPoolSetAccess(mem_pool, &desc_list, 1));
+    }
+
+    HIP_CHECK(hipMemPoolDestroy(mem_pool));
+  }
+};
+
+static void RunBenchmark() {
+  MemPoolSetAccessBenchmark benchmark;
+  benchmark.Run();
+}
+
+/**
+ * @warning **MemPool APIs are not fully implemented within current version
+ *          or HIP and therefore they cannot be appropriately executed on AMD and NVIDIA platforms.
+ *          Therefore, all tests related to MemPool APIs are implemented without formal
+ *          verification and will be verified once HIP fully supports MemPool APIs.**
+ * Test Description
+ * ------------------------
+ *  - Executes `hipMemPoolSetAccess` with `hipMemAccessFlagsProtReadWrite`.
+ * Test source
+ * ------------------------
+ *  - performance/stream/hipMemPoolSetAccess.cc
+ * Test requirements
+ * ------------------------
+ *  - Device supports memory pools
+ *  - HIP_VERSION >= 5.2
+ */
+TEST_CASE("Performance_hipMemPoolSetAccess") {
+  if (!AreMemPoolsSupported(0)) {
+    HipTest::HIP_SKIP_TEST("GPU 0 doesn't support hipDeviceAttributeMemoryPoolsSupported "
+                           "attribute. Hence skipping the testing with Pass result.\n");
+    return;
+  }
+  RunBenchmark();
+}

catch/performance/stream/hipMemPoolSetAttribute.cc

@@ -0,0 +1,83 @@
+/*
+Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved.
+Permission is hereby granted, free of charge, to any person obtaining a copy
+of this software and associated documentation files (the "Software"), to deal
+in the Software without restriction, including without limitation the rights
+to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+copies of the Software, and to permit persons to whom the Software is
+furnished to do so, subject to the following conditions:
+The above copyright notice and this permission notice shall be included in
+all copies or substantial portions of the Software.
+THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+THE SOFTWARE.
+*/
+
+#include "mem_pools_performance_common.hh"
+
+/**
+ * @addtogroup stream stream
+ * @{
+ * @ingroup PerformanceTest
+ */
+
+class MemPoolSetAttributeBenchmark : public Benchmark<MemPoolSetAttributeBenchmark> {
+ public:
+  void operator()(const hipMemPoolAttr attribute) {
+    hipMemPool_t mem_pool{nullptr};
+    hipMemPoolProps pool_props = CreateMemPoolProps(0, hipMemHandleTypeNone);
+    HIP_CHECK(hipMemPoolCreate(&mem_pool, &pool_props));
+
+    int value{0};
+
+    TIMED_SECTION(kTimerTypeCpu) {
+      HIP_CHECK(hipMemPoolSetAttribute(mem_pool, attribute, &value));
+    }
+
+    HIP_CHECK(hipMemPoolDestroy(mem_pool));
+  }
+};
+
+static void RunBenchmark(const hipMemPoolAttr attribute) {
+  MemPoolSetAttributeBenchmark benchmark;
+  benchmark.AddSectionName(GetMemPoolAttrSectionName(attribute));
+  benchmark.Run(attribute);
+}
+
+/**
+ * @warning **MemPool APIs are not fully implemented within current version
+ *          or HIP and therefore they cannot be appropriately executed on AMD and NVIDIA platforms.
+ *          Therefore, all tests related to MemPool APIs are implemented without formal
+ *          verification and will be verified once HIP fully supports MemPool APIs.**
+ * Test Description
+ * ------------------------
+ *  - Executes `hipMemPoolSetAttribute`:
+ *    -# Supported attributes:
+ *      - `hipMemPoolAttrReleaseThreshold`
+ *      - `hipMemPoolReuseFollowEventDependencies`
+ *      - `hipMemPoolReuseAllowOpportunistic`
+ *      - `hipMemPoolReuseAllowInternalDependencies`
+ * Test source
+ * ------------------------
+ *  - performance/stream/hipMemPoolSetAttribute.cc
+ * Test requirements
+ * ------------------------
+ *  - Device supports memory pools
+ *  - HIP_VERSION >= 5.2
+ */
+TEST_CASE("Performance_hipMemPoolSetAttribute") {
+  if (!AreMemPoolsSupported(0)) {
+    HipTest::HIP_SKIP_TEST("GPU 0 doesn't support hipDeviceAttributeMemoryPoolsSupported "
+                           "attribute. Hence skipping the testing with Pass result.\n");
+    return;
+  }
+  hipMemPoolAttr attribute = GENERATE(hipMemPoolAttrReleaseThreshold,
+                                      hipMemPoolReuseFollowEventDependencies,
+                                      hipMemPoolReuseAllowOpportunistic,
+                                      hipMemPoolReuseAllowInternalDependencies);
+  RunBenchmark(attribute);
+}

catch/performance/stream/hipMemPoolTrimTo.cc

@@ -0,0 +1,77 @@
+/*
+Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved.
+Permission is hereby granted, free of charge, to any person obtaining a copy
+of this software and associated documentation files (the "Software"), to deal
+in the Software without restriction, including without limitation the rights
+to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+copies of the Software, and to permit persons to whom the Software is
+furnished to do so, subject to the following conditions:
+The above copyright notice and this permission notice shall be included in
+all copies or substantial portions of the Software.
+THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+THE SOFTWARE.
+*/
+
+#include "mem_pools_performance_common.hh"
+
+/**
+ * @addtogroup stream stream
+ * @{
+ * @ingroup PerformanceTest
+ */
+
+class MemPoolTrimToBenchmark : public Benchmark<MemPoolTrimToBenchmark> {
+ public:
+  void operator()(const size_t min_bytes_to_hold) {
+    hipMemPool_t mem_pool{nullptr};
+    hipMemPoolProps pool_props = CreateMemPoolProps(0, hipMemHandleTypeNone);
+    HIP_CHECK(hipMemPoolCreate(&mem_pool, &pool_props));
+
+    TIMED_SECTION(kTimerTypeCpu) {
+      HIP_CHECK(hipMemPoolTrimTo(mem_pool, min_bytes_to_hold));
+    }
+
+    HIP_CHECK(hipMemPoolDestroy(mem_pool));
+  }
+};
+
+static void RunBenchmark(const size_t min_bytes_to_hold) {
+  MemPoolTrimToBenchmark benchmark;
+  benchmark.AddSectionName(std::to_string(min_bytes_to_hold));
+  benchmark.Run(min_bytes_to_hold);
+}
+
+/**
+ * @warning **MemPool APIs are not fully implemented within current version
+ *          or HIP and therefore they cannot be appropriately executed on AMD and NVIDIA platforms.
+ *          Therefore, all tests related to MemPool APIs are implemented without formal
+ *          verification and will be verified once HIP fully supports MemPool APIs.**
+ * Test Description
+ * ------------------------
+ *  - Executes `hipMemPoolTrimTo`:
+ *    -# Minimum bytes to hold:
+ *      - 4 KB
+ *      - 4 MB
+ *      - 16 MB
+ * Test source
+ * ------------------------
+ *  - performance/stream/hipMemPoolTrimTo.cc
+ * Test requirements
+ * ------------------------
+ *  - Device supports memory pools
+ *  - HIP_VERSION >= 5.2
+ */
+TEST_CASE("Performance_hipMemPoolTrimTo") {
+  if (!AreMemPoolsSupported(0)) {
+    HipTest::HIP_SKIP_TEST("GPU 0 doesn't support hipDeviceAttributeMemoryPoolsSupported "
+                           "attribute. Hence skipping the testing with Pass result.\n");
+    return;
+  }
+  size_t min_bytes_to_hold = GENERATE(4_KB, 4_MB, 16_MB);
+  RunBenchmark(min_bytes_to_hold);
+}

catch/performance/stream/hipStreamAddCallback.cc

@@ -0,0 +1,61 @@
+/*
+Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved.
+Permission is hereby granted, free of charge, to any person obtaining a copy
+of this software and associated documentation files (the "Software"), to deal
+in the Software without restriction, including without limitation the rights
+to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+copies of the Software, and to permit persons to whom the Software is
+furnished to do so, subject to the following conditions:
+The above copyright notice and this permission notice shall be included in
+all copies or substantial portions of the Software.
+THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+THE SOFTWARE.
+*/
+
+#include <hip_test_common.hh>
+#include <performance_common.hh>
+
+/**
+ * @addtogroup stream stream
+ * @{
+ * @ingroup PerformanceTest
+ */
+
+void Callback(hipStream_t stream, hipError_t status, void* user_data) {}
+
+class StreamAddCallbackBenchmark : public Benchmark<StreamAddCallbackBenchmark> {
+ public:
+  void operator()() {
+    const StreamGuard stream_guard{Streams::created};
+    const hipStream_t stream = stream_guard.stream();
+
+    TIMED_SECTION(kTimerTypeCpu) {
+      HIP_CHECK(hipStreamAddCallback(stream, Callback, nullptr, 0));
+    }
+  }
+};
+
+static void RunBenchmark() {
+  StreamAddCallbackBenchmark benchmark;
+  benchmark.Run();
+}
+
+/**
+ * Test Description
+ * ------------------------
+ *  - Executes `hipStreamAddCallback` on the created stream.
+ * Test source
+ * ------------------------
+ *  - performance/stream/hipStreamAddCallback.cc
+ * Test requirements
+ * ------------------------
+ *  - HIP_VERSION >= 5.2
+ */
+TEST_CASE("Performance_hipStreamAddCallback") {
+  RunBenchmark();
+}

catch/performance/stream/hipStreamBasic.cc

@@ -0,0 +1,269 @@
+/*
+Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved.
+Permission is hereby granted, free of charge, to any person obtaining a copy
+of this software and associated documentation files (the "Software"), to deal
+in the Software without restriction, including without limitation the rights
+to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+copies of the Software, and to permit persons to whom the Software is
+furnished to do so, subject to the following conditions:
+The above copyright notice and this permission notice shall be included in
+all copies or substantial portions of the Software.
+THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+THE SOFTWARE.
+*/
+
+#include <hip_test_common.hh>
+#include <performance_common.hh>
+#include <resource_guards.hh>
+
+/**
+ * @addtogroup stream stream
+ * @{
+ * @ingroup PerformanceTest
+ * Contains performance tests for all hipStream related APIs
+ */
+ 
+ class HipDeviceGetStreamPriorityRangeBenchmark : public Benchmark<HipDeviceGetStreamPriorityRangeBenchmark> {
+ public:
+  void operator()() {
+    int priority_min, priority_max;
+    TIMED_SECTION(kTimerTypeCpu) { HIP_CHECK(hipDeviceGetStreamPriorityRange(&priority_min, &priority_max)); }
+  }
+};
+
+class HipStreamQueryBenchmark : public Benchmark<HipStreamQueryBenchmark> {
+ public:
+  void operator()(bool perform_work) {
+    hipError_t error;
+    hipStream_t stream;
+    HIP_CHECK(hipStreamCreate(&stream));
+    void *dptr;
+    
+    if(perform_work) {
+      HIP_CHECK(hipMallocAsync(&dptr, 2048 * 4, stream));
+    }
+
+    TIMED_SECTION(kTimerTypeCpu) { error = hipStreamQuery(stream); }
+    
+    if(perform_work) {
+      HIP_CHECK(hipFreeAsync(dptr, stream));
+      HIP_CHECK(hipStreamSynchronize(stream));
+    }
+    
+    HIP_CHECK(hipStreamDestroy(stream));
+  }
+};
+
+class HipStreamSynchronizeBenchmark : public Benchmark<HipStreamSynchronizeBenchmark> {
+ public:
+  void operator()() {
+    hipError_t error;
+    hipStream_t stream;
+    HIP_CHECK(hipStreamCreate(&stream));
+    
+    TIMED_SECTION(kTimerTypeCpu) { error = hipStreamSynchronize(stream); }
+      
+    HIP_CHECK(hipStreamDestroy(stream));
+  }
+};
+
+class HipStreamDestroyBenchmark : public Benchmark<HipStreamDestroyBenchmark> {
+ public:
+  void operator()() {
+    hipStream_t stream;
+    HIP_CHECK(hipStreamCreate(&stream));
+
+    TIMED_SECTION(kTimerTypeCpu) { HIP_CHECK(hipStreamDestroy(stream)); }
+  }
+};
+
+class HipStreamCreateBenchmark : public Benchmark<HipStreamCreateBenchmark> {
+ public:
+  void operator()() {
+    hipStream_t stream;
+
+    TIMED_SECTION(kTimerTypeCpu) { HIP_CHECK(hipStreamCreate(&stream)); }
+
+    HIP_CHECK(hipStreamDestroy(stream));
+  }
+};
+
+class HipStreamCreateWithPriorityBenchmark : public Benchmark<HipStreamCreateWithPriorityBenchmark> {
+ public:
+  void operator()(unsigned int flag) {
+    hipStream_t stream;
+    int priority_min, priority_max, priority_mid;
+    
+    HIP_CHECK(hipDeviceGetStreamPriorityRange(&priority_min, &priority_max));
+    priority_mid = (priority_max + priority_min) / 2;
+    
+    TIMED_SECTION(kTimerTypeCpu) { HIP_CHECK(hipStreamCreateWithPriority(&stream, flag, priority_mid)); }
+
+    HIP_CHECK(hipStreamDestroy(stream));
+  }
+};
+
+
+
+static std::string GetStreamCreateFlagName(unsigned flag) {
+  switch (flag) {
+    case hipStreamDefault:
+      return "hipStreamDefault";
+    case hipStreamNonBlocking:
+      return "hipStreamNonBlocking";
+    default:
+      return "flag combination";
+  }
+}
+
+class HipStreamCreateWithFlagsBenchmark : public Benchmark<HipStreamCreateWithFlagsBenchmark> {
+ public:
+  void operator()(unsigned int flag) {
+    hipStream_t stream;
+
+    TIMED_SECTION(kTimerTypeCpu) { HIP_CHECK(hipStreamCreateWithFlags(&stream, flag)); }
+
+    HIP_CHECK(hipStreamDestroy(stream));
+  }
+};
+
+/**
+ * Test Description
+ * ------------------------
+ *  - Executes `hipStreamCreate`:
+ * Test source
+ * ------------------------
+ *  - performance/stream/hipStreamBasic.cc
+ * Test requirements
+ * ------------------------
+ *  - HIP_VERSION >= 5.2
+ */
+TEST_CASE("Performance_hipStreamCreate") {
+  HipStreamCreateBenchmark benchmark;
+  benchmark.Run();
+}
+
+static void RunBenchmark(unsigned flag) {
+  HipStreamCreateWithFlagsBenchmark benchmark;
+  benchmark.AddSectionName(GetStreamCreateFlagName(flag));
+  benchmark.Run(flag);
+}
+
+static void RunBenchmarkWithPriority(unsigned flag) {
+  HipStreamCreateWithPriorityBenchmark benchmark;
+  benchmark.AddSectionName(GetStreamCreateFlagName(flag));
+  benchmark.Run(flag);
+}
+
+/**
+ * Test Description
+ * ------------------------
+ *  - Executes `hipStreamCreateWithFlags` with all flags:
+ *    -# Flags
+ *      - hipStreamDefault
+ *      - hipStreamNonBlocking
+ * Test source
+ * ------------------------
+ *  - performance/stream/hipStreamBasic.cc
+ * Test requirements
+ * ------------------------
+ *  - HIP_VERSION >= 5.2
+ */
+TEST_CASE("Performance_hipStreamCreateWithFlags") {
+  const auto flag = GENERATE(hipStreamDefault, hipStreamNonBlocking);
+  RunBenchmark(flag);
+}
+
+/**
+ * Test Description
+ * ------------------------
+ *  - Executes `hipStreamCreateWithPriority` with all flags:
+ *    -# Flags
+ *      - hipStreamDefault
+ *      - hipStreamNonBlocking
+ * Test source
+ * ------------------------
+ *  - performance/stream/hipStreamBasic.cc
+ * Test requirements
+ * ------------------------
+ *  - HIP_VERSION >= 5.2
+ */
+TEST_CASE("Performance_hipStreamCreateWithPriority") {
+  const auto flag = GENERATE(hipStreamDefault, hipStreamNonBlocking);
+  RunBenchmarkWithPriority(flag);
+}
+
+/**
+ * Test Description
+ * ------------------------
+ *  - Executes `hipStreamDestroy`:
+ * Test source
+ * ------------------------
+ *  - performance/stream/hipStreamBasic.cc
+ * Test requirements
+ * ------------------------
+ *  - HIP_VERSION >= 5.2
+ */
+TEST_CASE("Performance_hipStreamDestroy") {
+  HipStreamDestroyBenchmark benchmark;
+  benchmark.Run();
+}
+
+/**
+ * Test Description
+ * ------------------------
+ *  - Executes `hipDeviceGetStreamPriorityRange`:
+ * Test source
+ * ------------------------
+ *  - performance/stream/hipStreamBasic.cc
+ * Test requirements
+ * ------------------------
+ *  - HIP_VERSION >= 5.2
+ */
+TEST_CASE("Performance_hipDeviceGetStreamPriorityRange") {
+  HipDeviceGetStreamPriorityRangeBenchmark benchmark;
+  benchmark.Run();
+}
+
+/**
+ * Test Description
+ * ------------------------
+ *  - Executes `hipStreamQuery`:
+ * Test source
+ * ------------------------
+ *  - performance/stream/hipStreamBasic.cc
+ * Test requirements
+ * ------------------------
+ *  - HIP_VERSION >= 5.2
+ */
+TEST_CASE("Performance_hipStreamQuery") {
+  const auto perform_work = GENERATE(true, false);
+  HipStreamQueryBenchmark benchmark;
+  if(perform_work) {
+    benchmark.AddSectionName("stream with work");
+  } else {
+    benchmark.AddSectionName("stream without work");
+  }
+  benchmark.Run(perform_work);
+}
+
+/**
+ * Test Description
+ * ------------------------
+ *  - Executes `hipDeviceGetStreamPriorityRange`:
+ * Test source
+ * ------------------------
+ *  - performance/stream/hipStreamBasic.cc
+ * Test requirements
+ * ------------------------
+ *  - HIP_VERSION >= 5.2
+ */
+TEST_CASE("Performance_hipStreamSynchronize") {
+  HipStreamSynchronizeBenchmark benchmark;
+  benchmark.Run();
+}

catch/performance/stream/hipStreamGetFlags.cc

@@ -0,0 +1,75 @@
+/*
+Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved.
+Permission is hereby granted, free of charge, to any person obtaining a copy
+of this software and associated documentation files (the "Software"), to deal
+in the Software without restriction, including without limitation the rights
+to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+copies of the Software, and to permit persons to whom the Software is
+furnished to do so, subject to the following conditions:
+The above copyright notice and this permission notice shall be included in
+all copies or substantial portions of the Software.
+THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+THE SOFTWARE.
+*/
+
+#include <hip_test_common.hh>
+#include <performance_common.hh>
+
+/**
+ * @addtogroup stream stream
+ * @{
+ * @ingroup PerformanceTest
+ */
+
+class StreamGetFlagsBenchmark : public Benchmark<StreamGetFlagsBenchmark> {
+ public:
+  void operator()(unsigned int expected_flag) {
+    unsigned int returned_flags{};
+    hipStream_t stream;
+
+    HIP_CHECK(hipStreamCreateWithFlags(&stream, expected_flag));
+    TIMED_SECTION(kTimerTypeCpu) {
+      HIP_CHECK(hipStreamGetFlags(stream, &returned_flags))
+    }
+    HIP_CHECK(hipStreamDestroy(stream));
+  }
+};
+
+static void RunBenchmark(unsigned int expected_flag) {
+  StreamGetFlagsBenchmark benchmark;
+  switch (expected_flag) {
+    case hipStreamDefault:
+      benchmark.AddSectionName("hipStreamDefault");
+      break;
+    case hipStreamNonBlocking:
+      benchmark.AddSectionName("hipStreamNonBlocking");
+      break;
+    default:
+      benchmark.AddSectionName("unknown flag type");
+  }
+  benchmark.Run(expected_flag);
+}
+
+/**
+ * Test Description
+ * ------------------------
+ *  - Executes `hipStreamGetFlags`:
+ *    -# Flags:
+ *      - `hipStreamDefault`
+ *      - `hipStreamNonBlocking`
+ * Test source
+ * ------------------------
+ *  - performance/stream/hipStreamGetFlags.cc
+ * Test requirements
+ * ------------------------
+ *  - HIP_VERSION >= 5.2
+ */
+TEST_CASE("Performance_hipStreamGetFlags") {
+  unsigned int expected_flag = GENERATE(hipStreamDefault, hipStreamNonBlocking);
+  RunBenchmark(expected_flag);
+}

catch/performance/stream/hipStreamGetPriority.cc

@@ -0,0 +1,74 @@
+/*
+Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved.
+Permission is hereby granted, free of charge, to any person obtaining a copy
+of this software and associated documentation files (the "Software"), to deal
+in the Software without restriction, including without limitation the rights
+to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+copies of the Software, and to permit persons to whom the Software is
+furnished to do so, subject to the following conditions:
+The above copyright notice and this permission notice shall be included in
+all copies or substantial portions of the Software.
+THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+THE SOFTWARE.
+*/
+
+#include <hip_test_common.hh>
+#include <performance_common.hh>
+
+/**
+ * @addtogroup stream stream
+ * @{
+ * @ingroup PerformanceTest
+ */
+
+class StreamGetPriorityBenchmark : public Benchmark<StreamGetPriorityBenchmark> {
+ public:
+  void operator()(Streams stream_type) {
+    const StreamGuard stream_guard{stream_type};
+    const hipStream_t stream = stream_guard.stream();
+
+    int priority{};
+    TIMED_SECTION(kTimerTypeCpu) {
+      HIP_CHECK(hipStreamGetPriority(stream, &priority));
+    }
+  }
+};
+
+static void RunBenchmark(Streams stream_type) {
+  StreamGetPriorityBenchmark benchmark;
+  switch (stream_type) {
+    case Streams::nullstream:
+      benchmark.AddSectionName("null stream");
+      break;
+    case Streams::created:
+      benchmark.AddSectionName("created");
+      break;
+    default:
+      benchmark.AddSectionName("per thread stream");
+  }
+  benchmark.Run(stream_type);
+}
+
+/**
+ * Test Description
+ * ------------------------
+ *  - Executes `hipStreamGetPriority`:
+ *    -# Stream types:
+ *      - `null`
+ *      - created
+ * Test source
+ * ------------------------
+ *  - performance/stream/hipStreamGetPriority.cc
+ * Test requirements
+ * ------------------------
+ *  - HIP_VERSION >= 5.2
+ */
+TEST_CASE("Performance_hipStreamGetPriority") {
+  Streams stream_type = GENERATE(Streams::nullstream, Streams::created);
+  RunBenchmark(stream_type);
+}

catch/performance/stream/hipStreamWaitEvent.cc

@@ -0,0 +1,80 @@
+/*
+Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved.
+Permission is hereby granted, free of charge, to any person obtaining a copy
+of this software and associated documentation files (the "Software"), to deal
+in the Software without restriction, including without limitation the rights
+to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+copies of the Software, and to permit persons to whom the Software is
+furnished to do so, subject to the following conditions:
+The above copyright notice and this permission notice shall be included in
+all copies or substantial portions of the Software.
+THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+THE SOFTWARE.
+*/
+
+#include <hip_test_common.hh>
+#include <performance_common.hh>
+
+/**
+ * @addtogroup stream stream
+ * @{
+ * @ingroup PerformanceTest
+ */
+
+class StreamWaitEventBenchmark : public Benchmark<StreamWaitEventBenchmark> {
+ public:
+  void operator()(Streams stream_type) {
+    const StreamGuard stream_guard{stream_type};
+    const hipStream_t stream = stream_guard.stream();
+    hipEvent_t wait_event{nullptr};
+
+    HIP_CHECK(hipEventCreate(&wait_event));
+    REQUIRE(wait_event != nullptr);
+    HIP_CHECK(hipEventRecord(wait_event, stream));
+
+    TIMED_SECTION(kTimerTypeCpu) {
+      HIP_CHECK(hipStreamWaitEvent(stream, wait_event, 0));
+      HIP_CHECK(hipStreamSynchronize(stream));
+    }
+    HIP_CHECK(hipEventDestroy(wait_event));
+  }
+};
+
+static void RunBenchmark(Streams stream_type) {
+  StreamWaitEventBenchmark benchmark{};
+  switch (stream_type) {
+    case Streams::nullstream:
+      benchmark.AddSectionName("null stream");
+      break;
+    case Streams::created:
+      benchmark.AddSectionName("created");
+      break;
+    default:
+      benchmark.AddSectionName("per thread stream");
+  }
+  benchmark.Run(stream_type);
+}
+
+/**
+ * Test Description
+ * ------------------------
+ *  - Executes `hipStreamWaitEvent`:
+ *    -# Stream types:
+ *      - `null`
+ *      - created
+ * Test source
+ * ------------------------
+ *  - performance/stream/hipStreamWaitEvent.cc
+ * Test requirements
+ * ------------------------
+ *  - HIP_VERSION >= 5.2
+ */
+TEST_CASE("Performance_hipStreamWaitEvent") {
+  Streams stream_type = GENERATE(Streams::nullstream, Streams::created);
+  RunBenchmark(stream_type);
+}

catch/performance/stream/hipStreamWaitValue.cc

@@ -0,0 +1,171 @@
+/*
+Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved.
+Permission is hereby granted, free of charge, to any person obtaining a copy
+of this software and associated documentation files (the "Software"), to deal
+in the Software without restriction, including without limitation the rights
+to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+copies of the Software, and to permit persons to whom the Software is
+furnished to do so, subject to the following conditions:
+The above copyright notice and this permission notice shall be included in
+all copies or substantial portions of the Software.
+THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+THE SOFTWARE.
+*/
+
+#include <hip_test_common.hh>
+#include <performance_common.hh>
+
+/**
+ * @addtogroup stream stream
+ * @{
+ * @ingroup PerformanceTest
+ */
+
+static int IsStreamWaitValueSupported(int device_id) {
+  int wait_value_supported = 0;
+#if HT_AMD
+  HIP_CHECK(hipDeviceGetAttribute(&wait_value_supported, hipDeviceAttributeCanUseStreamWaitValue,
+                                  device_id));
+#else
+  cuDeviceGetAttribute(&wait_value_supported, CU_DEVICE_ATTRIBUTE_CAN_USE_64_BIT_STREAM_MEM_OPS,
+                       device_id);
+#endif
+  return wait_value_supported;
+}
+
+class StreamWaitValue32Benchmark : public Benchmark<StreamWaitValue32Benchmark> {
+ public:
+  void operator()(const size_t array_size, unsigned int flag) {
+    const StreamGuard stream_guard{Streams::created};
+    const hipStream_t stream = stream_guard.stream();
+    uint32_t* value_ptr;
+    uint32_t value{0};
+    if (flag == hipStreamWaitValueAnd) {
+      value = 1;
+    }
+    HIP_CHECK(hipMalloc(&value_ptr, sizeof(uint32_t) * array_size));
+    HIP_CHECK(hipMemset(value_ptr, value, sizeof(uint32_t) * array_size));
+
+    TIMED_SECTION(kTimerTypeCpu) {
+      HIP_CHECK(hipStreamWaitValue32(stream, value_ptr, value, flag));
+    }
+    HIP_CHECK(hipFree(value_ptr));
+  }
+};
+
+class StreamWaitValue64Benchmark : public Benchmark<StreamWaitValue64Benchmark> {
+ public:
+  void operator()(const size_t array_size, unsigned int flag) {
+    const StreamGuard stream_guard{Streams::created};
+    const hipStream_t stream = stream_guard.stream();
+    uint64_t* value_ptr;
+    uint64_t value{0};
+    if (flag == hipStreamWaitValueAnd) {
+      value = 1;
+    }
+    HIP_CHECK(hipMalloc(&value_ptr, sizeof(uint64_t) * array_size));
+    HIP_CHECK(hipMemset(value_ptr, value, sizeof(uint64_t) * array_size));
+
+    TIMED_SECTION(kTimerTypeCpu) {
+      HIP_CHECK(hipStreamWaitValue64(stream, value_ptr, value, flag));
+    }
+    HIP_CHECK(hipFree(value_ptr));
+  }
+};
+
+template <typename WaitValueBenchmark>
+static void RunBenchmark(const size_t array_size, unsigned int flag) {
+  WaitValueBenchmark benchmark;
+  benchmark.AddSectionName(std::to_string(array_size));
+  switch (flag) {
+    case hipStreamWaitValueGte:
+      benchmark.AddSectionName("greater than or equal");
+      break;
+    case hipStreamWaitValueEq:
+      benchmark.AddSectionName("equal");
+      break;
+    case hipStreamWaitValueAnd:
+      benchmark.AddSectionName("logical and");
+      break;
+    case hipStreamWaitValueNor:
+      benchmark.AddSectionName("logical nor");
+      break;
+    default:
+      benchmark.AddSectionName("unknown flag");
+  }
+  benchmark.Run(array_size, flag);
+}
+
+/**
+ * Test Description
+ * ------------------------
+ *  - Executes `hipStreamWaitValue32` for different array sizes:
+ *    -# 4 KB
+ *    -# 4 MB
+ *    -# 16 MB
+ *  - Uses different flag types for wait criteria:
+ *    -# Greater than or equal
+ *    -# Equal
+ *    -# Logical AND
+ *    -# Logical OR
+ * Test source
+ * ------------------------
+ *  - performance/stream/hipStreamWaitValue.cc
+ * Test requirements
+ * ------------------------
+ *  - Device supports Stream Wait Value operations
+ *  - HIP_VERSION >= 5.2
+ */
+TEST_CASE("Performance_hipStreamWaitValue32") {
+#if HT_AMD
+  if (!IsStreamWaitValueSupported(0)) {
+    HipTest::HIP_SKIP_TEST(
+        "GPU 0 doesn't support hipStreamWaitValue32() function. "
+        "Hence skipping the testing with Pass result.\n");
+    return;
+  }
+#endif
+  size_t array_size = GENERATE(4_KB, 4_MB, 16_MB);
+  unsigned int flag = GENERATE(hipStreamWaitValueGte, hipStreamWaitValueEq, hipStreamWaitValueAnd,
+                               hipStreamWaitValueNor);
+  RunBenchmark<StreamWaitValue32Benchmark>(array_size, flag);
+}
+
+/**
+ * Test Description
+ * ------------------------
+ *  - Executes `hipStreamWaitValue64`:
+ *    -# Allocation size:
+ *      - 4 KB
+ *      - 4 MB
+ *      - 16 MB
+ *    -# Wait type:
+ *      - Greater than or equal
+ *      - Equal
+ *      - Logical AND
+ *      - Logical OR
+ * Test source
+ * ------------------------
+ *  - performance/stream/hipStreamWaitValue.cc
+ * Test requirements
+ * ------------------------
+ *  - Device supports Stream Wait Value operations
+ *  - HIP_VERSION >= 5.2
+ */
+TEST_CASE("Performance_hipStreamWaitValue64") {
+  if (!IsStreamWaitValueSupported(0)) {
+    HipTest::HIP_SKIP_TEST(
+        "GPU 0 doesn't support hipStreamWaitValue64() function. "
+        "Hence skipping the testing with Pass result.\n");
+    return;
+  }
+  size_t array_size = GENERATE(4_KB, 4_MB, 16_MB);
+  unsigned int flag = GENERATE(hipStreamWaitValueGte, hipStreamWaitValueEq, hipStreamWaitValueAnd,
+                               hipStreamWaitValueNor);
+  RunBenchmark<StreamWaitValue64Benchmark>(array_size, flag);
+}

catch/performance/stream/hipStreamWriteValue.cc

@@ -0,0 +1,121 @@
+/*
+Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved.
+Permission is hereby granted, free of charge, to any person obtaining a copy
+of this software and associated documentation files (the "Software"), to deal
+in the Software without restriction, including without limitation the rights
+to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+copies of the Software, and to permit persons to whom the Software is
+furnished to do so, subject to the following conditions:
+The above copyright notice and this permission notice shall be included in
+all copies or substantial portions of the Software.
+THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+THE SOFTWARE.
+*/
+
+#include <hip_test_common.hh>
+#include <performance_common.hh>
+
+/**
+ * @addtogroup stream stream
+ * @{
+ * @ingroup PerformanceTest
+ */
+
+#if HT_NVIDIA
+static int IsStreamWriteValueSupported(int device_id) {
+  int write_value_supported = 0;
+
+  cuDeviceGetAttribute(&write_value_supported, CU_DEVICE_ATTRIBUTE_CAN_USE_64_BIT_STREAM_MEM_OPS,
+                       device_id);
+  return write_value_supported;
+}
+#endif
+
+class StreamWriteValue32Benchmark : public Benchmark<StreamWriteValue32Benchmark> {
+ public:
+  void operator()(const size_t array_size) {
+    const StreamGuard stream_guard{Streams::created};
+    const hipStream_t stream = stream_guard.stream();
+    uint32_t* value_ptr;
+    uint32_t value{0};
+    HIP_CHECK(hipMalloc(&value_ptr, sizeof(uint32_t) * array_size));
+    HIP_CHECK(hipMemset(value_ptr, value, sizeof(uint32_t) * array_size));
+
+    TIMED_SECTION(kTimerTypeCpu) { HIP_CHECK(hipStreamWriteValue32(stream, value_ptr, value, 0)); }
+    HIP_CHECK(hipFree(value_ptr));
+  }
+};
+
+class StreamWriteValue64Benchmark : public Benchmark<StreamWriteValue64Benchmark> {
+ public:
+  void operator()(const size_t array_size) {
+    const StreamGuard stream_guard{Streams::created};
+    const hipStream_t stream = stream_guard.stream();
+    uint64_t* value_ptr;
+    uint64_t value{0};
+    HIP_CHECK(hipMalloc(&value_ptr, sizeof(uint64_t) * array_size));
+    HIP_CHECK(hipMemset(value_ptr, value, sizeof(uint64_t) * array_size));
+
+    TIMED_SECTION(kTimerTypeCpu) { HIP_CHECK(hipStreamWriteValue64(stream, value_ptr, value, 0)); }
+    HIP_CHECK(hipFree(value_ptr));
+  }
+};
+
+template <typename WriteValueBenchmark> static void RunBenchmark(const size_t array_size) {
+  WriteValueBenchmark benchmark;
+  benchmark.AddSectionName(std::to_string(array_size));
+  benchmark.Run(array_size);
+}
+
+/**
+ * Test Description
+ * ------------------------
+ *  - Executes `hipStreamWriteValue32`:
+ *    -# Allocation size:
+ *      - 4 KB
+ *      - 4 MB
+ *      - 16 MB
+ * Test source
+ * ------------------------
+ *  - performance/stream/hipStreamWriteValue.cc
+ * Test requirements
+ * ------------------------
+ *  - HIP_VERSION >= 5.2
+ */
+TEST_CASE("Performance_hipStreamWriteValue32") {
+  size_t array_size = GENERATE(4_KB, 4_MB, 16_MB);
+  RunBenchmark<StreamWriteValue32Benchmark>(array_size);
+}
+
+/**
+ * Test Description
+ * ------------------------
+ *  - Executes `hipStreamWriteValue64`:
+ *    -# Allocation size:
+ *      - 4 KB
+ *      - 4 MB
+ *      - 16 MB
+ * Test source
+ * ------------------------
+ *  - performance/stream/hipStreamWriteValue.cc
+ * Test requirements
+ * ------------------------
+ *  - HIP_VERSION >= 5.2
+ */
+TEST_CASE("Performance_hipStreamWriteValue64") {
+#if HT_NVIDIA
+  if (!IsStreamWriteValueSupported(0)) {
+    HipTest::HIP_SKIP_TEST(
+        "GPU 0 doesn't support hipStreamWriteValue64() function. "
+        "Hence skipping the testing with Pass result.\n");
+    return;
+  }
+#endif
+  size_t array_size = GENERATE(4_KB, 4_MB, 16_MB);
+  RunBenchmark<StreamWriteValue64Benchmark>(array_size);
+}

catch/performance/stream/mem_pools_performance_common.hh

@@ -0,0 +1,74 @@
+/*
+Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved.
+Permission is hereby granted, free of charge, to any person obtaining a copy
+of this software and associated documentation files (the "Software"), to deal
+in the Software without restriction, including without limitation the rights
+to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+copies of the Software, and to permit persons to whom the Software is
+furnished to do so, subject to the following conditions:
+The above copyright notice and this permission notice shall be included in
+all copies or substantial portions of the Software.
+THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+THE SOFTWARE.
+*/
+
+#pragma once
+
+#include <hip_test_common.hh>
+#include <performance_common.hh>
+
+#if __linux__
+  static const hipMemAllocationHandleType kHandleType = hipMemHandleTypePosixFileDescriptor;
+#else
+  static const hipMemAllocationHandleType kHandleType = hipMemHandleTypeWin32;
+#endif
+
+static int AreMemPoolsSupported(int device_id) {
+  int mem_pools_supported = 0;
+  HIP_CHECK(hipDeviceGetAttribute(&mem_pools_supported,
+                                  hipDeviceAttributeMemoryPoolsSupported, 0));
+  return mem_pools_supported;
+}
+
+static hipMemPoolProps CreateMemPoolProps(const int device_id, const hipMemAllocationHandleType handle_type) {
+  hipMemPoolProps kPoolProps = {
+    hipMemAllocationTypePinned,
+    handle_type,
+    {
+      hipMemLocationTypeDevice,
+      device_id
+    },
+    nullptr,
+    {0}
+  };
+
+  return kPoolProps;
+}
+
+static std::string GetMemPoolAttrSectionName(const hipMemPoolAttr attribute) {
+  switch (attribute) {
+    case hipMemPoolReuseFollowEventDependencies:
+      return "ReuseFollowEventDependencies";
+    case hipMemPoolReuseAllowOpportunistic:
+      return "ReuseAllowOpportunistic";
+    case hipMemPoolReuseAllowInternalDependencies:
+      return "ReuseAllowInternalDependencies";
+    case hipMemPoolAttrReleaseThreshold:
+      return "AttrReleaseThreshold";
+    case hipMemPoolAttrReservedMemCurrent:
+      return "AttrReservedMemCurrent";
+    case hipMemPoolAttrReservedMemHigh:
+      return "AttrReservedMemHigh";
+    case hipMemPoolAttrUsedMemCurrent:
+      return "AttrUsedMemCurrent";
+    case hipMemPoolAttrUsedMemHigh:
+      return "AttrUsedMemHigh";
+    default:
+      return "unknown attribute";
+  }
+}