EXSWHTEC-224 - Test cases ID clean up and documentation for Stream Management

catch/include/hip_test_defgroups.hh

@@ -131,6 +131,26 @@ THE SOFTWARE.
  * @}
  */
 
+/**
+ * @defgroup StreamOTest Ordered Memory Allocator
+ * @{
+ * This section describes the tests for Stream Ordered Memory Allocator functions of HIP runtime
+ * API.
+ */
+
+/**
+ * @defgroup StreamTest Stream Management
+ * @{
+ * This section describes tests for the stream management functions of HIP runtime API.
+ * @}
+ */
+
+/**
+ * @defgroup StreamMTest Stream Memory Operations
+ * @{
+ * This section describes tests for the Stream Memory Wait and Write functions of HIP runtime API.
+ */
+
 /**
  * @defgroup ShflTest warp shuffle function Management
  * @{

catch/unit/memory/hipMemPoolApi.cc

@@ -17,11 +17,6 @@
    THE SOFTWARE.
  */
 
-/* Test Case Description:
-   1) This testcase verifies the  basic scenario - supported on
-     all devices
-*/
-
 #include <hip_test_common.hh>
 #include <hip_test_kernels.hh>
 #include <hip_test_checkers.hh>
@@ -31,6 +26,12 @@
 #include <thread>
 #include <chrono>
 
+/**
+ * @addtogroup hipMallocAsync hipMallocAsync
+ * @{
+ * @ingroup StreamOTest
+ */
+
 constexpr hipMemPoolProps kPoolProps = {
   hipMemAllocationTypePinned,
   hipMemHandleTypeNone,
@@ -42,10 +43,25 @@ constexpr hipMemPoolProps kPoolProps = {
   {0}
 };
 
-/*
-   This testcase verifies HIP Mem Pool API basic scenario - supported on all devices
+/**
+ * @addtogroup hipMallocAsync hipMallocAsync
+ * @{
+ * @ingroup StreamOTest
  */
 
+ /**
+ * Test Description
+ * ------------------------
+ *  - Allocates memory for the array.
+ *  - Checks basic functionalities.
+ * Test source
+ * ------------------------
+ *  - unit/memory/hipMemPoolApi.cc
+ * Test requirements
+ * ------------------------
+ *  - Runtime supports Memory Pools
+ *  - HIP_VERSION >= 5.2
+ */
 TEST_CASE("Unit_hipMemPoolApi_Basic") {
   int mem_pool_support = 0;
   HIP_CHECK(hipDeviceGetAttribute(&mem_pool_support, hipDeviceAttributeMemoryPoolsSupported, 0));
@@ -108,6 +124,100 @@ TEST_CASE("Unit_hipMemPoolApi_Basic") {
   HIP_CHECK(hipStreamDestroy(stream));
 }
 
+/**
+ * Test Description
+ * ------------------------
+ *  - Checks that the freed memory is used for allocation again.
+ *  - Launches kernel to create a realistic test case.
+ * Test source
+ * ------------------------
+ *  - unit/memory/hipMemPoolApi.cc
+ * Test requirements
+ * ------------------------
+ *  - Runtime supports Memory Pools
+ *  - HIP_VERSION >= 5.2
+ */
+TEST_CASE("Unit_hipMemPoolApi_Default") {
+  int mem_pool_support = 0;
+  HIP_CHECK(hipDeviceGetAttribute(&mem_pool_support, hipDeviceAttributeMemoryPoolsSupported, 0));
+  if (!mem_pool_support) {
+    SUCCEED("Runtime doesn't support Memory Pool. Skip the test case.");
+    return;
+  }
+
+  hipMemPool_t mem_pool;
+  HIP_CHECK(hipDeviceGetDefaultMemPool(&mem_pool, 0));
+
+  float *A, *B, *C;
+  hipStream_t stream;
+  HIP_CHECK(hipStreamCreate(&stream));
+
+  size_t numElements = 8 * 1024 * 1024;
+  HIP_CHECK(hipMallocAsync(reinterpret_cast<void**>(&A), numElements * sizeof(float), stream));
+
+  numElements = 1024;
+  HIP_CHECK(hipMallocAsync(reinterpret_cast<void**>(&C), numElements * sizeof(float), stream));
+
+  int blocks = 2;
+  int clkRate;
+
+  if (IsGfx11()) {
+    HIP_CHECK(hipDeviceGetAttribute(&clkRate, hipDeviceAttributeWallClockRate, 0));
+    kernel500ms_gfx11<<<32, blocks, 0, stream>>>(A, clkRate);
+  } else {
+    HIP_CHECK(hipDeviceGetAttribute(&clkRate, hipDeviceAttributeClockRate, 0));
+
+    kernel500ms<<<32, blocks, 0, stream>>>(A, clkRate);
+  }
+
+  hipMemPoolAttr attr;
+  // Not a real free, since kernel isn't done
+  HIP_CHECK(hipFreeAsync(reinterpret_cast<void*>(A), stream));
+
+  numElements = 8 * 1024 * 1024;
+  HIP_CHECK(hipMallocAsync(reinterpret_cast<void**>(&B), numElements * sizeof(float), stream));
+  // Runtime must reuse the pointer
+  REQUIRE(A == B);
+
+  // Make a sync before the second kernel launch to make sure memory B isn't gone
+  HIP_CHECK(hipStreamSynchronize(stream));
+
+  // Second kernel launch with new memory
+  if (IsGfx11()) {
+    kernel500ms_gfx11<<<32, blocks, 0, stream>>>(B, clkRate);
+  } else {
+    kernel500ms<<<32, blocks, 0, stream>>>(B, clkRate);
+  }
+
+  HIP_CHECK(hipFreeAsync(reinterpret_cast<void*>(B), stream));
+
+  HIP_CHECK(hipStreamSynchronize(stream));
+
+  std::uint64_t value64 = 0;
+  attr = hipMemPoolAttrReservedMemCurrent;
+  HIP_CHECK(hipMemPoolGetAttribute(mem_pool, attr, &value64));
+  // Make sure the current reserved is at least allocation size of buffer C (4KB)
+  REQUIRE(sizeof(float) * 1024 <= value64);
+
+  attr = hipMemPoolAttrUsedMemHigh;
+  HIP_CHECK(hipMemPoolGetAttribute(mem_pool, attr, &value64));
+  // Make sure the high watermark usage works - the both buffers must be reported
+  REQUIRE(sizeof(float) * (8 * 1024 * 1024 + 1024) == value64);
+
+  attr = hipMemPoolAttrUsedMemCurrent;
+  HIP_CHECK(hipMemPoolGetAttribute(mem_pool, attr, &value64));
+  // Make sure the current usage reports just one buffer, because the above free doesn't hold memory
+  REQUIRE(sizeof(float) * 1024 == value64);
+
+  HIP_CHECK(hipFreeAsync(reinterpret_cast<void*>(C), stream));
+  HIP_CHECK(hipStreamDestroy(stream));
+}
+
+/**
+ * End doxygen group hipMallocAsync.
+ * @}
+ */
+
 constexpr auto wait_ms = 500;
 
 __global__ void kernel500ms(float* hostRes, int clkRate) {
@@ -138,6 +248,25 @@ __global__ void kernel500ms_gfx11(float* hostRes, int clkRate) {
 #endif
 }
 
+/**
+ * @addtogroup hipFreeAsync hipFreeAsync
+ * @{
+ * @ingroup StreamOTest
+ */
+
+/**
+ * Test Description
+ * ------------------------
+ *  - Checks if memory usage is different before and after synchronization.
+ *  - Synchronization will force free to execute.
+ * Test source
+ * ------------------------
+ *  - unit/memory/hipMemPoolApi.cc
+ * Test requirements
+ * ------------------------
+ *  - Runtime supports Memory Pools
+ *  - HIP_VERSION >= 5.2
+ */
 TEST_CASE("Unit_hipMemPoolApi_BasicAlloc") {
   int mem_pool_support = 0;
   HIP_CHECK(hipDeviceGetAttribute(&mem_pool_support, hipDeviceAttributeMemoryPoolsSupported, 0));
@@ -226,6 +355,30 @@ TEST_CASE("Unit_hipMemPoolApi_BasicAlloc") {
   HIP_CHECK(hipStreamDestroy(stream));
 }
 
+/**
+ * End doxygen group hipFreeAsync.
+ * @}
+ */
+
+/**
+ * @addtogroup hipMemPoolTrimTo hipMemPoolTrimTo
+ * @{
+ * @ingroup StreamOTest
+ */
+
+/**
+ * Test Description
+ * ------------------------
+ *  - Check if a trim operation is no-op when memory is still in use.
+ *  - Check that trim works correctly once the memory is not in use.
+ * Test source
+ * ------------------------
+ *  - unit/memory/hipMemPoolApi.cc
+ * Test requirements
+ * ------------------------
+ *  - Runtime supports Memory Pools
+ *  - HIP_VERSION >= 5.2
+ */
 TEST_CASE("Unit_hipMemPoolApi_BasicTrim") {
   int mem_pool_support = 0;
   HIP_CHECK(hipDeviceGetAttribute(&mem_pool_support, hipDeviceAttributeMemoryPoolsSupported, 0));
@@ -314,6 +467,31 @@ TEST_CASE("Unit_hipMemPoolApi_BasicTrim") {
   HIP_CHECK(hipStreamDestroy(stream));
 }
 
+/**
+ * End doxygen group hipMemPoolTrimTo.
+ * @}
+ */
+
+/**
+ * @addtogroup hipMallocFromPoolAsync hipMallocFromPoolAsync
+ * @{
+ * @ingroup StreamOTest
+ */
+
+/**
+ * Test Description
+ * ------------------------
+ *  - Checks that memory from pool is reused when freed.
+ *  - Allocate the same array after the memory is freed.
+ *  - Verify that the old and new pointers are the same.
+ * Test source
+ * ------------------------
+ *  - unit/memory/hipMemPoolApi.cc
+ * Test requirements
+ * ------------------------
+ *  - Runtime supports Memory Pools
+ *  - HIP_VERSION >= 5.2
+ */
 TEST_CASE("Unit_hipMemPoolApi_BasicReuse") {
   int mem_pool_support = 0;
   HIP_CHECK(hipDeviceGetAttribute(&mem_pool_support, hipDeviceAttributeMemoryPoolsSupported, 0));
@@ -390,6 +568,21 @@ TEST_CASE("Unit_hipMemPoolApi_BasicReuse") {
   HIP_CHECK(hipStreamDestroy(stream));
 }
 
+/**
+ * Test Description
+ * ------------------------
+ *  - Verifies that an oportunistic flag behaves correctly with allocations.
+ *    -# When oportunistic is disallowed and no reuse
+ *    -# When oportunistic is allowed and reuse
+ *    -# When oportunistic is allowed and no reuse
+ * Test source
+ * ------------------------
+ *  - unit/memory/hipMemPoolApi.cc
+ * Test requirements
+ * ------------------------
+ *  - Runtime supports Memory Pools
+ *  - HIP_VERSION >= 5.2
+ */
 TEST_CASE("Unit_hipMemPoolApi_Opportunistic") {
   int mem_pool_support = 0;
   HIP_CHECK(hipDeviceGetAttribute(&mem_pool_support, hipDeviceAttributeMemoryPoolsSupported, 0));
@@ -553,78 +746,7 @@ TEST_CASE("Unit_hipMemPoolApi_Opportunistic") {
   HIP_CHECK(hipStreamDestroy(stream2));
 }
 
-TEST_CASE("Unit_hipMemPoolApi_Default") {
-  int mem_pool_support = 0;
-  HIP_CHECK(hipDeviceGetAttribute(&mem_pool_support, hipDeviceAttributeMemoryPoolsSupported, 0));
-  if (!mem_pool_support) {
-    SUCCEED("Runtime doesn't support Memory Pool. Skip the test case.");
-    return;
-  }
-
-  hipMemPool_t mem_pool;
-  HIP_CHECK(hipDeviceGetDefaultMemPool(&mem_pool, 0));
-
-  float *A, *B, *C;
-  hipStream_t stream;
-  HIP_CHECK(hipStreamCreate(&stream));
-
-  size_t numElements = 8 * 1024 * 1024;
-  HIP_CHECK(hipMallocAsync(reinterpret_cast<void**>(&A), numElements * sizeof(float), stream));
-
-  numElements = 1024;
-  HIP_CHECK(hipMallocAsync(reinterpret_cast<void**>(&C), numElements * sizeof(float), stream));
-
-  int blocks = 2;
-  int clkRate;
-
-  if (IsGfx11()) {
-    HIP_CHECK(hipDeviceGetAttribute(&clkRate, hipDeviceAttributeWallClockRate, 0));
-    kernel500ms_gfx11<<<32, blocks, 0, stream>>>(A, clkRate);
-  } else {
-    HIP_CHECK(hipDeviceGetAttribute(&clkRate, hipDeviceAttributeClockRate, 0));
-
-    kernel500ms<<<32, blocks, 0, stream>>>(A, clkRate);
-  }
-
-  hipMemPoolAttr attr;
-  // Not a real free, since kernel isn't done
-  HIP_CHECK(hipFreeAsync(reinterpret_cast<void*>(A), stream));
-
-  numElements = 8 * 1024 * 1024;
-  HIP_CHECK(hipMallocAsync(reinterpret_cast<void**>(&B), numElements * sizeof(float), stream));
-  // Runtime must reuse the pointer
-  REQUIRE(A == B);
-
-  // Make a sync before the second kernel launch to make sure memory B isn't gone
-  HIP_CHECK(hipStreamSynchronize(stream));
-
-  // Second kernel launch with new memory
-  if (IsGfx11()) {
-    kernel500ms_gfx11<<<32, blocks, 0, stream>>>(B, clkRate);
-  } else {
-    kernel500ms<<<32, blocks, 0, stream>>>(B, clkRate);
-  }
-
-  HIP_CHECK(hipFreeAsync(reinterpret_cast<void*>(B), stream));
-
-  HIP_CHECK(hipStreamSynchronize(stream));
-
-  std::uint64_t value64 = 0;
-  attr = hipMemPoolAttrReservedMemCurrent;
-  HIP_CHECK(hipMemPoolGetAttribute(mem_pool, attr, &value64));
-  // Make sure the current reserved is at least allocation size of buffer C (4KB)
-  REQUIRE(sizeof(float) * 1024 <= value64);
-
-  attr = hipMemPoolAttrUsedMemHigh;
-  HIP_CHECK(hipMemPoolGetAttribute(mem_pool, attr, &value64));
-  // Make sure the high watermark usage works - the both buffers must be reported
-  REQUIRE(sizeof(float) * (8 * 1024 * 1024 + 1024) == value64);
-
-  attr = hipMemPoolAttrUsedMemCurrent;
-  HIP_CHECK(hipMemPoolGetAttribute(mem_pool, attr, &value64));
-  // Make sure the current usage reports just one buffer, because the above free doesn't hold memory
-  REQUIRE(sizeof(float) * 1024 == value64);
-
-  HIP_CHECK(hipFreeAsync(reinterpret_cast<void*>(C), stream));
-  HIP_CHECK(hipStreamDestroy(stream));
-}
+/**
+ * End doxygen group hipMallocFromPoolAsync.
+ * @}
+ */