EXSWHTEC-72 - Implement hipMemcpyPeer/hipMemcpyPeerAsync and hipMemGetAddressRange tests (#34)

- Reimplement and expand hipMemcpyPeer/hipMemcpyPeerAsync tests using resource guards
- Implement positive and negative test for hipMemGetAddressRange

catch/unit/memory/hipMemcpyPeerAsync.cc

@@ -1,5 +1,5 @@
 /*
-Copyright (c) 2021 Advanced Micro Devices, Inc. All rights reserved.
+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
@@ -16,247 +16,236 @@ 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.
 */
-
 /*
-This testfile verifies the following scenarios of hipMemcpyPeerAsync API
-1. Negative Scenarios
-2. Memory on one GPU and stream created on another GPU
-3. Basic scenario of hipMemcpyPeerAsync API
+Testcase Scenarios :
+Unit_hipMemcpyPeerAsync_Positive_Default - Test basic P2P async memcpy between
+two devices with hipMemcpyPeerAsync api
+Unit_hipMemcpyPeerAsync_Positive_Synchronization_Behavior - Test synchronization
+behavior for hipMemcpyPeerAsync api Unit_hipMemcpyPeerAsync_Positive_ZeroSize -
+Test that no data is copied when sizeBytes is set to 0
+Unit_hipMemcpyPeerAsync_Negative_Parameters - Test unsuccessful execution of
+hipMemcpyPeerAsync api when parameters are invalid
 */
-
-
+#include <hip/hip_runtime_api.h>
 #include <hip_test_common.hh>
-#include <hip_test_kernels.hh>
-#include <hip_test_checkers.hh>
-#include <iostream>
+#include <resource_guards.hh>
+#include <utils.hh>
 
-/*This testcase verifies the negative scenarios of hipmemcpypeerAsync
-*/
-TEST_CASE("Unit_hipMemcpyPeerAsync_Negative") {
-  constexpr auto numElements{10};
-  constexpr auto copy_bytes{numElements*sizeof(int)};
-  int numDevices = 0;
-  HIP_CHECK(hipGetDeviceCount(&numDevices));
-  if (numDevices > 1) {
-    int canAccessPeer = 0;
-    HIP_CHECK(hipDeviceCanAccessPeer(&canAccessPeer, 0, 1));
-    if (canAccessPeer) {
-      // Initialization of variables
-      int *A_d{nullptr}, *B_d{nullptr};
-      int *A_h{nullptr}, *B_h{nullptr};
-      hipStream_t stream;
-      HIP_CHECK(hipSetDevice(0));
-      HIP_CHECK(hipStreamCreate(&stream));
-      HipTest::initArrays<int>(&A_d, nullptr, nullptr,
-          &A_h, &B_h, nullptr, numElements*sizeof(int));
-      HipTest::setDefaultData<int>(numElements, A_h, B_h, nullptr);
-      HIP_CHECK(hipSetDevice(1));
-      HipTest::initArrays<int>(nullptr, &B_d, nullptr,
-          nullptr, nullptr, nullptr, numElements*sizeof(int));
-      HIP_CHECK(hipMemcpy(B_d, B_h, numElements*sizeof(int),
-                          hipMemcpyHostToDevice));
 
-      SECTION("Nullptr to Destination Pointer") {
-        REQUIRE(hipMemcpyPeerAsync(nullptr, 1, A_d, 0, copy_bytes,
-              stream) != hipSuccess);
-      }
-
-      SECTION("Nullptr to Source Pointer") {
-        REQUIRE(hipMemcpyPeerAsync(B_d, 1, nullptr, 0, copy_bytes,
-              stream) != hipSuccess);
-      }
-
-      SECTION("Pass NumElements as 0") {
-        HIP_CHECK(hipMemcpy(A_d, A_h, numElements*sizeof(int),
-                            hipMemcpyHostToDevice));
-        HIP_CHECK(hipMemcpyPeerAsync(B_d, 1, A_d, 0, 0,
-                                     stream));
-        HIP_CHECK(hipMemcpy(A_h, B_d, numElements*sizeof(int),
-                            hipMemcpyDeviceToHost));
-        HipTest::checkTest<int>(A_h, B_h, numElements);
-      }
-
-      SECTION("Passing more than allocated size") {
-        REQUIRE(hipMemcpyPeerAsync(B_d, 1, A_d, 0,
-              100*sizeof(int), stream) != hipSuccess);
-      }
-
-      SECTION("Passing invalid Destination device ID") {
-        REQUIRE(hipMemcpyPeerAsync(B_d, numDevices, A_d, 0, copy_bytes,
-              stream) != hipSuccess);
-      }
-
-      SECTION("Passing invalid Source device ID") {
-        REQUIRE(hipMemcpyPeerAsync(B_d, 0, A_d, numDevices, copy_bytes,
-              stream) != hipSuccess);
-      }
-
-      HipTest::freeArrays<int>(A_d, B_d, nullptr, A_h, B_h, nullptr, false);
-      HIP_CHECK(hipStreamDestroy(stream));
-    } else {
-      SUCCEED("Machine Does not have P2P capability");
-    }
-  } else {
-    SUCCEED("Number of devices are < 2");
+TEST_CASE("Unit_hipMemcpyPeerAsync_Positive_Default") {
+  const auto device_count = HipTest::getDeviceCount();
+  if (device_count < 2) {
+    HipTest::HIP_SKIP_TEST("Skipping because devices < 2");
+    return;
   }
-}
-/*
- * This test case verifies the basic scenario of hipMemcpyPeer API
- * Initializes data in GPU-0
- * Launches the kernel and performs addition in GPU-0
- * Copies the data from GPU-0 to GPU-1 using hipMemcpyPeerAsync API
- * Then performs the addition and validates the sum
- */
+  const auto stream_type = GENERATE(Streams::nullstream, Streams::perThread, Streams::created);
+  const StreamGuard stream_guard(stream_type);
+  const hipStream_t stream = stream_guard.stream();
 
-TEST_CASE("Unit_hipMemcpyPeerAsync_Basic") {
-  constexpr auto numElements{10};
-  constexpr auto copy_bytes{numElements*sizeof(int)};
+  const auto allocation_size = GENERATE(kPageSize / 2, kPageSize, kPageSize * 2);
 
-  int numDevices = 0;
-  HIP_CHECK(hipGetDeviceCount(&numDevices));
-  if (numDevices > 1) {
-    int canAccessPeer = 0;
-    HIP_CHECK(hipDeviceCanAccessPeer(&canAccessPeer, 0, 1));
-    if (canAccessPeer) {
-      // Initialization of Variables on GPU-0
-      int *A_d{nullptr}, *B_d{nullptr}, *C_d{nullptr};
-      int *X_d{nullptr}, *Y_d{nullptr}, *Z_d{nullptr};
-      int *A_h{nullptr}, *B_h{nullptr}, *C_h{nullptr};
-      hipStream_t stream;
-      HIP_CHECK(hipSetDevice(0));
-      HipTest::initArrays<int>(&A_d, &B_d, &C_d,
-          &A_h, &B_h, &C_h, numElements*sizeof(int));
-      HipTest::setDefaultData<int>(numElements, A_h, B_h, nullptr);
-      HIP_CHECK(hipMemcpy(A_d, A_h, numElements*sizeof(int),
-                          hipMemcpyHostToDevice));
-      HIP_CHECK(hipMemcpy(B_d, B_h, numElements*sizeof(int),
-                          hipMemcpyHostToDevice));
-      HIP_CHECK(hipStreamCreate(&stream));
+  int can_access_peer = 0;
 
-      // Initialization of Variables in GPU-1
-      HIP_CHECK(hipSetDevice(1));
-      HipTest::initArrays<int>(&X_d, &Y_d, &Z_d, nullptr,
-          nullptr, nullptr, numElements*sizeof(int));
+  const auto src_device = GENERATE(range(0, HipTest::getDeviceCount()));
+  const auto dst_device = GENERATE(range(0, HipTest::getDeviceCount()));
+  INFO("Src device: " << src_device << ", Dst device: " << dst_device);
 
-      // Launching kernel and performing vector addition in GPU-0
-      HIP_CHECK(hipSetDevice(0));
-      hipLaunchKernelGGL(HipTest::vectorADD, dim3(1), dim3(1),
-          0, 0, static_cast<const int*>(A_d),
-          static_cast<const int*>(B_d), C_d, numElements*sizeof(int));
-      HIP_CHECK(hipGetLastError());
-      HIP_CHECK(hipMemcpy(C_h, C_d, numElements*sizeof(int),
-               hipMemcpyDeviceToHost));
-      HipTest::checkVectorADD<int>(A_h, B_h, C_h, numElements);
+  HIP_CHECK(hipSetDevice(src_device));
+  HIP_CHECK(hipDeviceCanAccessPeer(&can_access_peer, src_device, dst_device));
+  if (can_access_peer) {
+    HIP_CHECK(hipDeviceEnablePeerAccess(dst_device, 0));
 
-      // Copying data from GPU-0 to GPU-1 and performing vector addition
-      HIP_CHECK(hipSetDevice(1));
-      SECTION("Calling hipMemcpyPerAsync() using user defined stream obj") {
-        HIP_CHECK(hipMemcpyPeerAsync(X_d, 1, A_d, 0, copy_bytes,
-                                     stream));
-        HIP_CHECK(hipMemcpyPeerAsync(Y_d, 1, B_d, 0, copy_bytes,
-                                     stream));
-        HIP_CHECK(hipStreamSynchronize(stream));
-      }
-      SECTION("Calling hipMemcpyPerAsync() using hipStreamPerThread") {
-        HIP_CHECK(hipMemcpyPeerAsync(X_d, 1, A_d, 0, copy_bytes,
-                                     hipStreamPerThread));
-        HIP_CHECK(hipMemcpyPeerAsync(Y_d, 1, B_d, 0, copy_bytes,
-                                     hipStreamPerThread));
-        HIP_CHECK(hipStreamSynchronize(hipStreamPerThread));
-      }
-      hipLaunchKernelGGL(HipTest::vectorADD, dim3(1), dim3(1),
-          0, 0, static_cast<const int*>(X_d),
-          static_cast<const int*>(Y_d), Z_d, numElements*sizeof(int));
-      HIP_CHECK(hipGetLastError());
-      HIP_CHECK(hipMemcpy(C_h, Z_d, numElements*sizeof(int),
-               hipMemcpyDeviceToHost));
-      HipTest::checkVectorADD<int>(A_h, B_h, C_h, numElements);
+    LinearAllocGuard<int> src_alloc(LinearAllocs::hipMalloc, allocation_size);
+    LinearAllocGuard<int> result(LinearAllocs::hipHostMalloc, allocation_size,
+                                 hipHostMallocPortable);
+    HIP_CHECK(hipSetDevice(dst_device));
+    LinearAllocGuard<int> dst_alloc(LinearAllocs::hipMalloc, allocation_size);
 
-      // Cleaning the Memory
-      HipTest::freeArrays<int>(A_d, B_d, C_d, A_h, B_h, C_h, false);
-      HipTest::freeArrays<int>(X_d, Y_d, Z_d, nullptr, nullptr, nullptr, false);
-      HIP_CHECK(hipStreamDestroy(stream));
-    } else {
-      SUCCEED("Machine Does not have P2P capability");
-    }
+    const auto element_count = allocation_size / sizeof(*src_alloc.ptr());
+    constexpr auto thread_count = 1024;
+    const auto block_count = element_count / thread_count + 1;
+    constexpr int expected_value = 22;
+    HIP_CHECK(hipSetDevice(src_device));
+    VectorSet<<<block_count, thread_count, 0, stream>>>(src_alloc.ptr(), expected_value,
+                                                        element_count);
+    HIP_CHECK(hipGetLastError());
+
+    HIP_CHECK(hipMemcpyPeerAsync(dst_alloc.ptr(), dst_device, src_alloc.ptr(), src_device,
+                                 allocation_size, stream));
+
+    HIP_CHECK(hipStreamSynchronize(stream));
+
+    HIP_CHECK(
+        hipMemcpy(result.host_ptr(), dst_alloc.ptr(), allocation_size, hipMemcpyDeviceToHost));
+
+    HIP_CHECK(hipDeviceDisablePeerAccess(dst_device));
+
+    ArrayFindIfNot(result.host_ptr(), expected_value, element_count);
   } else {
-    SUCCEED("Number of devices are < 2");
+    INFO("Peer access cannot be enabled between devices " << src_device << " " << dst_device);
   }
 }
 
-/*
- * This test case verifies the following functionality where
-   Memory is allocated in One GPU and
-   stream created on another GPU
- * Initializes all the data in GPU-0
- * Creating stream in GPU-1
- * Launches the kernel and performs addition in GPU-0
- * Copies the data from GPU-0 to GPU-1 using hipMemcpyPeerAsync API
- * where stream is created in GPU-1
- * Then performs the addition and validates the sum
- */
-TEST_CASE("Unit_hipMemcpyPeerAsync_StreamOnDiffDevice") {
-  constexpr auto numElements{10};
-  constexpr auto copy_bytes{numElements*sizeof(int)};
-  int numDevices = 0;
-  HIP_CHECK(hipGetDeviceCount(&numDevices));
-  if (numDevices > 1) {
-    int canAccessPeer = 0;
-    HIP_CHECK(hipDeviceCanAccessPeer(&canAccessPeer, 0, 1));
-    if (canAccessPeer) {
-      int *A_d{nullptr}, *B_d{nullptr}, *C_d{nullptr};
-      int *X_d{nullptr}, *Y_d{nullptr}, *Z_d{nullptr};
-      int *A_h{nullptr}, *B_h{nullptr}, *C_h{nullptr};
-      hipStream_t stream;
-      HIP_CHECK(hipSetDevice(0));
+TEST_CASE("Unit_hipMemcpyPeerAsync_Positive_Synchronization_Behavior") {
+  HIP_CHECK(hipDeviceSynchronize());
 
-      // Initialization of all variables in GPU-0
-      HipTest::initArrays<int>(&A_d, &B_d, &C_d,
-          &A_h, &B_h, &C_h, numElements*sizeof(int));
-      HIP_CHECK(hipMemcpy(A_d, A_h, numElements*sizeof(int),
-               hipMemcpyHostToDevice));
-      HIP_CHECK(hipMemcpy(B_d, B_h, numElements*sizeof(int),
-            hipMemcpyHostToDevice));
-      HipTest::initArrays<int>(&X_d, &Y_d, &Z_d, nullptr,
-          nullptr, nullptr, numElements*sizeof(int));
+  const auto device_count = HipTest::getDeviceCount();
+  if (device_count < 2) {
+    HipTest::HIP_SKIP_TEST("Skipping because devices < 2");
+    return;
+  }
 
-      // Stream created in GPU-1
-      HIP_CHECK(hipSetDevice(1));
-      HIP_CHECK(hipStreamCreate(&stream));
+  const StreamGuard stream_guard(Streams::created);
+  const hipStream_t stream = stream_guard.stream();
 
-      // Performing vector addition and validate the data
-      HIP_CHECK(hipSetDevice(0));
-      hipLaunchKernelGGL(HipTest::vectorADD, dim3(1), dim3(1),
-          0, 0, static_cast<const int*>(A_d),
-          static_cast<const int*>(B_d), C_d, numElements*sizeof(int));
-      HIP_CHECK(hipGetLastError());
-      HIP_CHECK(hipMemcpy(C_h, C_d, numElements*sizeof(int),
-               hipMemcpyDeviceToHost));
-      HipTest::checkVectorADD<int>(A_h, B_h, C_h, numElements);
+  int can_access_peer = 0;
+  const auto src_device = 0;
+  const auto dst_device = 1;
 
-      // Copying the data from GPU-0 to GPU-1 where stream is from diff device
-      HIP_CHECK(hipMemcpyPeerAsync(X_d, 1, A_d, 0, copy_bytes,
-               stream));
-      HIP_CHECK(hipMemcpyPeerAsync(Y_d, 1, B_d, 0, copy_bytes,
-               stream));
-      HIP_CHECK(hipStreamSynchronize(stream));
-      hipLaunchKernelGGL(HipTest::vectorADD, dim3(1), dim3(1),
-          0, 0, static_cast<const int*>(X_d),
-          static_cast<const int*>(Y_d), Z_d, numElements*sizeof(int));
-      HIP_CHECK(hipGetLastError());
-      HIP_CHECK(hipMemcpy(C_h, Z_d, numElements*sizeof(int),
-               hipMemcpyDeviceToHost));
+  HIP_CHECK(hipSetDevice(src_device));
+  HIP_CHECK(hipDeviceCanAccessPeer(&can_access_peer, src_device, dst_device));
+  if (can_access_peer) {
+    HIP_CHECK(hipDeviceEnablePeerAccess(dst_device, 0));
 
-      // Cleaning the data
-      HipTest::checkVectorADD<int>(A_h, B_h, C_h, numElements);
-      HipTest::freeArrays<int>(A_d, B_d, C_d, A_h, B_h, C_h, false);
-      HipTest::freeArrays<int>(X_d, Y_d, Z_d, nullptr, nullptr, nullptr, false);
-      HIP_CHECK(hipStreamDestroy(stream));
-    } else {
-      SUCCEED("Machine Does not have P2P capability");
-    }
+    LinearAllocGuard<int> src_alloc(LinearAllocs::hipMalloc, kPageSize);
+    HIP_CHECK(hipSetDevice(dst_device));
+    LinearAllocGuard<int> dst_alloc(LinearAllocs::hipMalloc, kPageSize);
+
+    HIP_CHECK(hipSetDevice(src_device));
+    LaunchDelayKernel(std::chrono::milliseconds{100}, nullptr);
+
+    HIP_CHECK(hipMemcpyPeerAsync(dst_alloc.ptr(), dst_device, src_alloc.ptr(), src_device,
+                                 kPageSize, stream));
+    HIP_CHECK_ERROR(hipStreamQuery(nullptr), hipErrorNotReady);
+
+    HIP_CHECK(hipDeviceDisablePeerAccess(dst_device));
   } else {
-    SUCCEED("Number of devices are < 2");
+    INFO("Peer access cannot be enabled between devices " << src_device << " " << dst_device);
   }
 }
 
+TEST_CASE("Unit_hipMemcpyPeerAsync_Positive_ZeroSize") {
+  const auto device_count = HipTest::getDeviceCount();
+  if (device_count < 2) {
+    HipTest::HIP_SKIP_TEST("Skipping because devices < 2");
+    return;
+  }
+
+  const StreamGuard stream_guard(Streams::created);
+  const hipStream_t stream = stream_guard.stream();
+
+  const auto allocation_size = kPageSize;
+
+  int can_access_peer = 0;
+  const auto src_device = 0;
+  const auto dst_device = 1;
+
+  HIP_CHECK(hipSetDevice(src_device));
+  HIP_CHECK(hipDeviceCanAccessPeer(&can_access_peer, src_device, dst_device));
+  if (can_access_peer) {
+    HIP_CHECK(hipDeviceEnablePeerAccess(dst_device, 0));
+
+    LinearAllocGuard<int> src_alloc(LinearAllocs::hipMalloc, allocation_size);
+    LinearAllocGuard<int> result(LinearAllocs::hipHostMalloc, allocation_size,
+                                 hipHostMallocPortable);
+    HIP_CHECK(hipSetDevice(dst_device));
+    LinearAllocGuard<int> dst_alloc(LinearAllocs::hipMalloc, allocation_size);
+
+    const auto element_count = allocation_size / sizeof(*src_alloc.ptr());
+    constexpr auto thread_count = 1024;
+    const auto block_count = element_count / thread_count + 1;
+    constexpr int set_value = 22;
+    HIP_CHECK(hipSetDevice(src_device));
+    VectorSet<<<block_count, thread_count, 0, stream>>>(src_alloc.ptr(), set_value, element_count);
+    HIP_CHECK(hipGetLastError());
+
+    constexpr int expected_value = 21;
+    std::fill_n(src_alloc.host_ptr(), element_count, expected_value);
+
+    HIP_CHECK(
+        hipMemcpyPeerAsync(dst_alloc.ptr(), dst_device, src_alloc.ptr(), src_device, 0, stream));
+
+    HIP_CHECK(hipStreamSynchronize(stream));
+
+    HIP_CHECK(
+        hipMemcpy(result.host_ptr(), dst_alloc.ptr(), allocation_size, hipMemcpyDeviceToHost));
+
+    HIP_CHECK(hipDeviceDisablePeerAccess(dst_device));
+
+    ArrayFindIfNot(result.host_ptr(), expected_value, element_count);
+  } else {
+    INFO("Peer access cannot be enabled between devices " << src_device << " " << dst_device);
+  }
+}
+
+TEST_CASE("Unit_hipMemcpyPeerAsync_Negative_Parameters") {
+  const auto device_count = HipTest::getDeviceCount();
+  if (device_count < 2) {
+    HipTest::HIP_SKIP_TEST("Skipping because devices < 2");
+    return;
+  }
+
+  const StreamGuard stream_guard(Streams::created);
+  const hipStream_t stream = stream_guard.stream();
+
+  constexpr auto InvalidStream = [] {
+    StreamGuard sg(Streams::created);
+    return sg.stream();
+  };
+
+  int can_access_peer = 0;
+  const auto src_device = 0;
+  const auto dst_device = 1;
+
+  HIP_CHECK(hipSetDevice(src_device));
+  HIP_CHECK(hipDeviceCanAccessPeer(&can_access_peer, src_device, dst_device));
+  if (can_access_peer) {
+    HIP_CHECK(hipDeviceEnablePeerAccess(dst_device, 0));
+
+    LinearAllocGuard<int> src_alloc(LinearAllocs::hipMalloc, kPageSize);
+    HIP_CHECK(hipSetDevice(dst_device));
+    LinearAllocGuard<int> dst_alloc(LinearAllocs::hipMalloc, kPageSize);
+
+    HIP_CHECK(hipSetDevice(src_device));
+
+    SECTION("Nullptr to Destination Pointer") {
+      HIP_CHECK_ERROR(
+          hipMemcpyPeerAsync(nullptr, dst_device, src_alloc.ptr(), src_device, kPageSize, stream),
+          hipErrorInvalidValue);
+    }
+
+    SECTION("Nullptr to Source Pointer") {
+      HIP_CHECK_ERROR(
+          hipMemcpyPeerAsync(dst_alloc.ptr(), dst_device, nullptr, src_device, kPageSize, stream),
+          hipErrorInvalidValue);
+    }
+
+    SECTION("Passing more than allocated size") {
+      HIP_CHECK_ERROR(hipMemcpyPeerAsync(dst_alloc.ptr(), dst_device, src_alloc.ptr(), src_device,
+                                         kPageSize + 1, stream),
+                      hipErrorInvalidValue);
+    }
+
+    SECTION("Passing invalid Destination device ID") {
+      HIP_CHECK_ERROR(hipMemcpyPeerAsync(dst_alloc.ptr(), device_count, src_alloc.ptr(), src_device,
+                                         kPageSize, stream),
+                      hipErrorInvalidDevice);
+    }
+
+    SECTION("Passing invalid Source device ID") {
+      HIP_CHECK_ERROR(hipMemcpyPeerAsync(dst_alloc.ptr(), dst_device, src_alloc.ptr(), device_count,
+                                         kPageSize, stream),
+                      hipErrorInvalidDevice);
+    }
+
+    SECTION("Passing invalid Stream") {
+      HIP_CHECK_ERROR(hipMemcpyPeerAsync(dst_alloc.ptr(), dst_device, src_alloc.ptr(), src_device,
+                                         kPageSize, InvalidStream()),
+                      hipErrorContextIsDestroyed);
+    }
+
+    HIP_CHECK(hipDeviceDisablePeerAccess(dst_device));
+  } else {
+    INFO("Peer access cannot be enabled between devices " << src_device << " " << dst_device);
+  }
+}