EXSWHTEC-99 - Reimplement tests for hipStreamAttachMemAsync (#52)

- Negative parameter tests - Validate basic behavior - Validate the behavior when pageable memory access is supported - Validate the behavior for hipMemAttachGlobal - Validate the behavior for hipMemAttachHost - Validate the behavior for hipMemAttachSingle
2023-01-17 12:58:57 +01:00
parent 203b994230
commit 8ae538a341
4 changed files with 372 additions and 3 deletions
@@ -112,6 +112,7 @@ set(TEST_SRC
    hipMemsetAsync.cc
    hipMemAdvise.cc
    hipMemRangeGetAttributes.cc
+    hipStreamAttachMemAsync.cc
    hipMemRangeGetAttributes_old.cc
    hipMemGetAddressRange.cc
 )
@@ -203,6 +204,7 @@ set(TEST_SRC
    hipMemRangeGetAttributes.cc
    hipMemRangeGetAttributes_old.cc
    hipGetSymbolSizeAddress.cc
+    hipStreamAttachMemAsync.cc
    hipMemGetAddressRange.cc
 )
 endif()
@@ -0,0 +1,224 @@
+/*
+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/hip_runtime_api.h>
+#include <hip_test_common.hh>
+#include <kernels.hh>
+#include <resource_guards.hh>
+#include <utils.hh>
+
+TEST_CASE("Unit_hipStreamAttachMemAsync_Positive_Basic") {
+  if (!DeviceAttributesSupport(0, hipDeviceAttributeManagedMemory)) {
+    HipTest::HIP_SKIP_TEST("Managed memory is not supported");
+    return;
+  }
+
+  StreamGuard stream(Streams::created);
+  LinearAllocGuard<hipDeviceptr_t> managed(LinearAllocs::hipMallocManaged, kPageSize,
+                                           hipMemAttachHost);
+
+  HIP_CHECK(hipStreamAttachMemAsync(stream.stream(), managed.ptr(), 0));
+  HIP_CHECK(hipStreamSynchronize(stream.stream()));
+}
+
+TEST_CASE("Unit_hipStreamAttachMemAsync_Positive_Pageable") {
+  if (!DeviceAttributesSupport(0, hipDeviceAttributeManagedMemory)) {
+    HipTest::HIP_SKIP_TEST("Managed memory is not supported");
+    return;
+  }
+
+  if (!DeviceAttributesSupport(0, hipDeviceAttributePageableMemoryAccess)) {
+    HipTest::HIP_SKIP_TEST("Pageable memory access is not supported");
+    return;
+  }
+
+  StreamGuard stream(Streams::created);
+  LinearAllocGuard<hipDeviceptr_t> pageable(LinearAllocs::malloc, kPageSize);
+
+  HIP_CHECK(hipStreamAttachMemAsync(stream.stream(), pageable.ptr(), kPageSize));
+  HIP_CHECK(hipStreamSynchronize(stream.stream()));
+}
+
+// CUDA docs:
+// If the cudaMemAttachGlobal flag is specified, the memory can be accessed by any stream on any
+// device.
+TEST_CASE("Unit_hipStreamAttachMemAsync_Positive_AttachGlobal") {
+  if (!DeviceAttributesSupport(0, hipDeviceAttributeManagedMemory)) {
+    HipTest::HIP_SKIP_TEST("Managed memory is not supported");
+    return;
+  }
+
+  const auto device_count = HipTest::getDeviceCount();
+  const auto stream_count = device_count < 2 ? 8 : device_count;
+
+  std::vector<std::unique_ptr<StreamGuard>> streams;
+  streams.reserve(stream_count);
+  for (int i = 0; i < stream_count; ++i) {
+    if (device_count > 1) {
+      HIP_CHECK(hipSetDevice(i));
+    }
+    streams.push_back(std::make_unique<StreamGuard>(Streams::created));
+  }
+
+  LinearAllocGuard<int> managed_global(LinearAllocs::hipMallocManaged, sizeof(int) * stream_count,
+                                       hipMemAttachHost);
+
+  HIP_CHECK(hipStreamAttachMemAsync(
+      nullptr, reinterpret_cast<hipDeviceptr_t*>(managed_global.ptr()), 0, hipMemAttachGlobal));
+  HIP_CHECK(hipStreamSynchronize(nullptr));
+
+  for (int i = 0; i < stream_count; ++i) {
+    HipTest::launchKernel(Set, 1, 1, 0, streams.at(i)->stream(), managed_global.ptr() + i, i);
+  }
+
+  for (auto&& stream : streams) {
+    HIP_CHECK(hipStreamSynchronize(stream->stream()));
+  }
+
+  for (int i = 0; i < stream_count; ++i) {
+    REQUIRE(managed_global.ptr()[i] == i);
+  }
+}
+
+// CUDA docs:
+// If the cudaMemAttachHost flag is specified, the program makes a guarantee that it won't access
+// the memory on the device from any stream on a device that has a zero value for the device
+// attribute cudaDevAttrConcurrentManagedAccess.
+TEST_CASE("Unit_hipStreamAttachMemAsync_Positive_AttachHost") {
+  if (!DeviceAttributesSupport(0, hipDeviceAttributeManagedMemory)) {
+    HipTest::HIP_SKIP_TEST("Managed memory is not supported");
+    return;
+  }
+
+  if (DeviceAttributesSupport(0, hipDeviceAttributeConcurrentManagedAccess)) {
+    HipTest::HIP_SKIP_TEST("Device supports concurrent managed access");
+    return;
+  }
+
+  StreamGuard stream(Streams::created);
+  LinearAllocGuard<int> managed_global(LinearAllocs::hipMallocManaged, sizeof(int));
+  LinearAllocGuard<int> managed_host(LinearAllocs::hipMallocManaged, sizeof(int));
+
+  HIP_CHECK(hipStreamAttachMemAsync(
+      stream.stream(), reinterpret_cast<hipDeviceptr_t*>(managed_host.ptr()), 0, hipMemAttachHost));
+  HIP_CHECK(hipStreamSynchronize(stream.stream()));
+
+  HipTest::launchKernel(Set, 1, 1, 0, stream.stream(), managed_global.ptr(), 32);
+  *managed_host.ptr() = 64;
+  HIP_CHECK(hipStreamSynchronize(stream.stream()));
+
+  REQUIRE(*managed_global.ptr() == 32);
+  REQUIRE(*managed_host.ptr() == 64);
+}
+
+// CUDA docs:
+// If the cudaMemAttachSingle flag is specified and stream is associated with a device that has a
+// zero value for the device attribute cudaDevAttrConcurrentManagedAccess, the program makes a
+// guarantee that it will only access the memory on the device from stream.
+TEST_CASE("Unit_hipStreamAttachMemAsync_Positive_AttachSingle") {
+  if (!DeviceAttributesSupport(0, hipDeviceAttributeManagedMemory)) {
+    HipTest::HIP_SKIP_TEST("Managed memory is not supported");
+    return;
+  }
+
+  if (DeviceAttributesSupport(0, hipDeviceAttributeConcurrentManagedAccess)) {
+    HipTest::HIP_SKIP_TEST("Device supports concurrent managed access");
+    return;
+  }
+
+  StreamGuard stream1(Streams::created);
+  StreamGuard stream2(Streams::created);
+
+  LinearAllocGuard<int> managed_global(LinearAllocs::hipMallocManaged, sizeof(int));
+  LinearAllocGuard<int> managed_single(LinearAllocs::hipMallocManaged, sizeof(int),
+                                       hipMemAttachHost);
+
+  HIP_CHECK(hipStreamAttachMemAsync(stream1.stream(),
+                                    reinterpret_cast<hipDeviceptr_t*>(managed_single.ptr()), 0,
+                                    hipMemAttachSingle));
+  HIP_CHECK(hipStreamSynchronize(stream1.stream()));
+
+  HipTest::launchKernel(Set, 1, 1, 0, stream1.stream(), managed_single.ptr(), 64);
+  HIP_CHECK(hipStreamSynchronize(stream1.stream()));
+
+  HipTest::launchKernel(Set, 1, 1, 0, stream2.stream(), managed_global.ptr(), 32);
+
+  REQUIRE(*managed_single.ptr() == 64);
+  *managed_single.ptr() = 128;
+
+  HIP_CHECK(hipStreamSynchronize(stream2.stream()));
+
+  REQUIRE(*managed_global.ptr() == 32);
+  REQUIRE(*managed_single.ptr() == 128);
+}
+
+TEST_CASE("Unit_hipStreamAttachMemAsync_Negative_Parameters") {
+  if (!DeviceAttributesSupport(0, hipDeviceAttributeManagedMemory)) {
+    HipTest::HIP_SKIP_TEST("Managed memory is not supported");
+    return;
+  }
+
+  StreamGuard stream(Streams::created);
+  LinearAllocGuard<hipDeviceptr_t> managed(LinearAllocs::hipMallocManaged, kPageSize,
+                                           hipMemAttachHost);
+
+  SECTION("invalid stream") {
+    HIP_CHECK(hipStreamDestroy(stream.stream()));
+    HIP_CHECK_ERROR(hipStreamAttachMemAsync(stream.stream(), managed.ptr()),
+                    hipErrorContextIsDestroyed);
+  }
+
+  SECTION("dev_ptr == nullptr") {
+    HIP_CHECK_ERROR(hipStreamAttachMemAsync(stream.stream(), nullptr), hipErrorInvalidValue);
+  }
+
+  SECTION("length is not zero nor entire allocation size") {
+    HIP_CHECK_ERROR(hipStreamAttachMemAsync(stream.stream(), managed.ptr(), kPageSize / 2),
+                    hipErrorInvalidValue);
+  }
+
+  SECTION("invalid flags") {
+    HIP_CHECK_ERROR(
+        hipStreamAttachMemAsync(stream.stream(), managed.ptr(), 0,
+                                hipMemAttachGlobal | hipMemAttachHost | hipMemAttachSingle),
+        hipErrorInvalidValue);
+  }
+
+  SECTION("attach single to nullstream") {
+    HIP_CHECK_ERROR(hipStreamAttachMemAsync(nullptr, managed.ptr(), 0, hipMemAttachSingle),
+                    hipErrorInvalidValue);
+  }
+
+  LinearAllocGuard<hipDeviceptr_t> pageable(LinearAllocs::malloc, kPageSize);
+
+  if (!DeviceAttributesSupport(0, hipDeviceAttributePageableMemoryAccess)) {
+    SECTION("dev_ptr is pageable memory") {
+      HIP_CHECK_ERROR(hipStreamAttachMemAsync(stream.stream(), pageable.ptr(), kPageSize),
+                      hipErrorInvalidValue);
+    }
+  } else {
+    SECTION("length is zero for pageable memory") {
+      HIP_CHECK_ERROR(hipStreamAttachMemAsync(stream.stream(), pageable.ptr(), 0),
+                      hipErrorInvalidValue);
+    }
+  }
+}
@@ -32,8 +32,8 @@ set(TEST_SRC
    hipStreamCreateWithPriority.cc
    hipStreamDestroy.cc
    hipAPIStreamDisable.cc
-    # hipStreamAttachMemAsync.cc   # Disabling it on nvidia due to issue in function definition of hipStreamAttachMemAsync
-                                   # Fixing would break ABI, to be re-enabled when the fix is made.
+    # hipStreamAttachMemAsync_old.cc   # Disabling it on nvidia due to issue in function definition of hipStreamAttachMemAsync
+                                       # Fixing would break ABI, to be re-enabled when the fix is made.
    streamCommon.cc
    hipStreamValue.cc
    hipStreamSynchronize.cc
@@ -42,7 +42,7 @@ set(TEST_SRC
    hipStreamACb_StrmSyncTiming.cc
    )

-    # set_source_files_properties(hipStreamAttachMemAsync.cc PROPERTIES COMPILE_FLAGS -std=c++17)
+    # set_source_files_properties(hipStreamAttachMemAsync_old.cc PROPERTIES COMPILE_FLAGS -std=c++17)
 endif()

 hip_add_exe_to_target(NAME StreamTest
@@ -0,0 +1,143 @@
+/*
+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.
+*/
+
+// TODO Enable it after hipStreamAttachMemAsync is feature complete on HIP
+
+#include <hip_test_common.hh>
+#include <memory>
+
+__device__ __managed__ int var = 0;
+
+enum class StreamAttachTestType { NullStream = 0, StreamPerThread, CreatedStream };
+
+TEST_CASE("Unit_hipStreamAttachMemAsync_Negative") {
+  hipStream_t stream{nullptr};
+
+  auto streamType =
+      GENERATE(StreamAttachTestType::NullStream, StreamAttachTestType::StreamPerThread,
+               StreamAttachTestType::CreatedStream);
+
+  if (streamType == StreamAttachTestType::StreamPerThread) {
+    stream = hipStreamPerThread;
+  } else if (streamType == StreamAttachTestType::CreatedStream) {
+    HIP_CHECK(hipStreamCreate(&stream));
+    REQUIRE(stream != nullptr);
+  }
+
+  SECTION("Invalid Resource Handle") {
+    int definitelyNotAManagedVariable = 0;
+    HIP_CHECK_ERROR(
+        hipStreamAttachMemAsync(stream, reinterpret_cast<void*>(&definitelyNotAManagedVariable),
+                                sizeof(int), hipMemAttachSingle),
+        hipErrorInvalidValue);
+  }
+
+  SECTION("Invalid devptr") {
+    HIP_CHECK_ERROR(hipStreamAttachMemAsync(stream, nullptr, sizeof(int), hipMemAttachSingle),
+                    hipErrorInvalidValue);
+  }
+
+  SECTION("Invalid Resource Size") {
+    HIP_CHECK_ERROR(hipStreamAttachMemAsync(stream, reinterpret_cast<void*>(&var), sizeof(int) - 1,
+                                            hipMemAttachSingle),
+                    hipErrorInvalidValue);
+  }
+
+  SECTION("Invalid Flags") {
+    HIP_CHECK_ERROR(
+        hipStreamAttachMemAsync(stream, reinterpret_cast<void*>(&var), sizeof(int) - 1,
+                                hipMemAttachSingle | hipMemAttachHost | hipMemAttachGlobal),
+        hipErrorInvalidValue);
+  }
+
+  if (streamType == StreamAttachTestType::CreatedStream) {
+    HIP_CHECK(hipStreamDestroy(stream));
+  }
+}
+
+__global__ void kernel(int* ptr, size_t size) {
+  auto i = threadIdx.x;
+  if (i < size) {
+    ptr[i] = 1024;
+  }
+}
+
+constexpr size_t size = 1024;
+__device__ __managed__ int m_memory[size];
+
+TEST_CASE("Unit_hipStreamAttachMemAsync_UseCase") {
+  hipStream_t stream{nullptr};
+
+  auto streamType =
+      GENERATE(StreamAttachTestType::NullStream, StreamAttachTestType::StreamPerThread,
+               StreamAttachTestType::CreatedStream);
+
+  if (streamType == StreamAttachTestType::CreatedStream) {
+    HIP_CHECK(hipStreamCreate(&stream));
+    REQUIRE(stream != nullptr);
+  }
+
+  SECTION("Size zero is valid") {
+    int* d_memory{nullptr};
+    HIP_CHECK(hipMallocManaged(&d_memory, sizeof(int) * size, hipMemAttachHost));
+    HIP_CHECK(
+        hipStreamAttachMemAsync(stream, reinterpret_cast<void*>(d_memory), 0, hipMemAttachHost));
+    HIP_CHECK(hipStreamSynchronize(stream));  // Wait for command to complete
+    HIP_CHECK(hipFree(d_memory));
+  }
+
+  SECTION("Access from device and host") {
+    int* d_memory{nullptr};
+
+    HIP_CHECK(hipMallocManaged(&d_memory, sizeof(int) * size, hipMemAttachHost));
+    HIP_CHECK(hipMemset(d_memory, 0, sizeof(int) * size));
+    HIP_CHECK(
+        hipStreamAttachMemAsync(stream, reinterpret_cast<void*>(d_memory), 0, hipMemAttachHost));
+    HIP_CHECK(hipStreamSynchronize(stream));  // Wait for the command to complete
+
+    kernel<<<1, size, 0, stream>>>(d_memory, size);
+    HIP_CHECK(hipStreamSynchronize(stream));  // Wait for the kernel to complete
+
+    auto ptr = std::make_unique<int[]>(size);
+    std::copy(d_memory, d_memory + size, ptr.get());
+
+    HIP_CHECK(hipFree(d_memory));
+
+    REQUIRE(std::all_of(ptr.get(), ptr.get() + size, [](int n) { return n == size; }));
+  }
+
+  SECTION("Access ManagedMemory") {
+    HIP_CHECK(hipMemset(m_memory, 0, sizeof(int) * size));
+    HIP_CHECK(
+        hipStreamAttachMemAsync(stream, reinterpret_cast<void*>(m_memory), 0, hipMemAttachHost));
+    HIP_CHECK(hipStreamSynchronize(stream));  // Wait for the command to complete
+
+    kernel<<<1, size, 0, stream>>>(m_memory, size);
+    HIP_CHECK(hipStreamSynchronize(stream));  // Wait for the kernel to complete
+
+    auto ptr = std::make_unique<int[]>(size);
+    std::copy(m_memory, m_memory + size, ptr.get());
+
+    REQUIRE(std::all_of(ptr.get(), ptr.get() + size, [](int n) { return n == size; }));
+  }
+
+  if (streamType == StreamAttachTestType::CreatedStream) {
+    HIP_CHECK(hipStreamDestroy(stream));
+  }
+}