diff --git a/catch/hipTestMain/config/config_amd_linux_MI2xx.json b/catch/hipTestMain/config/config_amd_linux_MI2xx.json index bf00e2e8e1..5e27deaa80 100644 --- a/catch/hipTestMain/config/config_amd_linux_MI2xx.json +++ b/catch/hipTestMain/config/config_amd_linux_MI2xx.json @@ -7,7 +7,8 @@ "Unit_hipInit_Negative", "Unit_BuiltinAtomicsRTC_fmaxCoherentGlobalMem", "Unit_BuiltinAtomicsRTC__fminCoherentGlobalMem", - "Unit_BuiltInAtomicAdd_CoherentGlobalMemWithRtc" + "Unit_BuiltInAtomicAdd_CoherentGlobalMemWithRtc", + "Unit_hipMemGetAddressRange_Negative" ] } diff --git a/catch/hipTestMain/config/config_amd_linux_common.json b/catch/hipTestMain/config/config_amd_linux_common.json index 224e85d155..864c6aa5c4 100644 --- a/catch/hipTestMain/config/config_amd_linux_common.json +++ b/catch/hipTestMain/config/config_amd_linux_common.json @@ -17,6 +17,7 @@ "Unit_hipDeviceReset_Positive_Threaded", "Unit_hipGraphDestroyNode_Complx_ChkNumOfNodesNDep", "Unit_hipGraphDestroyNode_Complx_ChkNumOfNodesNDep_ClonedGrph", - "Unit_hipGraphDestroyNode_Complx_ChkNumOfNodesNDep_ChldNode" + "Unit_hipGraphDestroyNode_Complx_ChkNumOfNodesNDep_ChldNode", + "Unit_hipMemGetAddressRange_Negative" ] } diff --git a/catch/hipTestMain/config/config_amd_windows_MI2xx.json b/catch/hipTestMain/config/config_amd_windows_MI2xx.json index 89c484b79b..7898310ea5 100644 --- a/catch/hipTestMain/config/config_amd_windows_MI2xx.json +++ b/catch/hipTestMain/config/config_amd_windows_MI2xx.json @@ -92,6 +92,7 @@ "Unit_hipDeviceGetPCIBusId_Negative_PartialFill", "Unit_hipInit_Negative", "Unit_hipGraphAddEventRecordNode_Functional_ElapsedTime", - "Unit_hipStreamBeginCapture_captureComplexGraph" + "Unit_hipStreamBeginCapture_captureComplexGraph", + "Unit_hipMemGetAddressRange_Negative" ] } diff --git a/catch/hipTestMain/config/config_amd_windows_common.json b/catch/hipTestMain/config/config_amd_windows_common.json index 417985d83c..5586728b6f 100644 --- a/catch/hipTestMain/config/config_amd_windows_common.json +++ b/catch/hipTestMain/config/config_amd_windows_common.json @@ -96,6 +96,7 @@ "Unit_hipGraphAddEventRecordNode_Functional_WithoutFlags", "Unit_hipGraphDestroyNode_Complx_ChkNumOfNodesNDep", "Unit_hipGraphDestroyNode_Complx_ChkNumOfNodesNDep_ClonedGrph", - "Unit_hipGraphDestroyNode_Complx_ChkNumOfNodesNDep_ChldNode" + "Unit_hipGraphDestroyNode_Complx_ChkNumOfNodesNDep_ChldNode", + "Unit_hipMemGetAddressRange_Negative" ] } diff --git a/catch/unit/memory/CMakeLists.txt b/catch/unit/memory/CMakeLists.txt index c89fae9bfc..962f493a60 100644 --- a/catch/unit/memory/CMakeLists.txt +++ b/catch/unit/memory/CMakeLists.txt @@ -70,7 +70,9 @@ set(TEST_SRC hipPtrGetAttribute.cc hipMemPoolApi.cc hipMemcpyPeer.cc + hipMemcpyPeer_old.cc hipMemcpyPeerAsync.cc + hipMemcpyPeerAsync_old.cc hipMemcpyWithStream_old.cc hipMemcpyWithStream.cc hipMemcpyWithStreamMultiThread.cc @@ -109,6 +111,7 @@ set(TEST_SRC hipMemsetAsync.cc hipMemAdvise.cc hipMemRangeGetAttributes.cc + hipMemGetAddressRange.cc ) else() set(TEST_SRC @@ -158,7 +161,9 @@ set(TEST_SRC hipPtrGetAttribute.cc hipMemPoolApi.cc hipMemcpyPeer.cc + hipMemcpyPeer_old.cc hipMemcpyPeerAsync.cc + hipMemcpyPeerAsync_old.cc hipMemcpyWithStream_old.cc hipMemcpyWithStream.cc hipMemcpyWithStreamMultiThread.cc @@ -194,6 +199,7 @@ set(TEST_SRC hipMemAdvise.cc hipMemRangeGetAttributes.cc hipGetSymbolSizeAddress.cc + hipMemGetAddressRange.cc ) endif() diff --git a/catch/unit/memory/hipMemGetAddressRange.cc b/catch/unit/memory/hipMemGetAddressRange.cc new file mode 100644 index 0000000000..7491163bda --- /dev/null +++ b/catch/unit/memory/hipMemGetAddressRange.cc @@ -0,0 +1,88 @@ +/* +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. +*/ +/* +Testcase Scenarios : +Unit_hipMemGetAddressRange_Positive - Test hipMemGetAddressRange api for various memory allocation +types and offsets Unit_hipMemGetAddressRange_Negative - Test unsuccessful execution of +hipMemGetAddressRange api when parameters are invalid +*/ +#include +#include +#include +#include + +TEST_CASE("Unit_hipMemGetAddressRange_Positive") { + hipDeviceptr_t base_ptr; + size_t mem_size = 0; + const auto allocation_size = GENERATE(kPageSize / 2, kPageSize, kPageSize * 2); + const int offset = GENERATE(0, 20, 40, 60, 80); + + SECTION("Host address range") { + using LA = LinearAllocs; + LinearAllocGuard host_alloc(LA::hipHostMalloc, allocation_size); + + HIP_CHECK(hipMemGetAddressRange(&base_ptr, &mem_size, + reinterpret_cast(host_alloc.ptr() + offset))); + + REQUIRE(reinterpret_cast(host_alloc.ptr()) == base_ptr); + REQUIRE(mem_size == allocation_size); + } + SECTION("Device address range") { + using LA = LinearAllocs; + const auto device_allocation_type = GENERATE(LA::hipMalloc, LA::hipMallocManaged); + LinearAllocGuard device_alloc(device_allocation_type, allocation_size); + + HIP_CHECK(hipMemGetAddressRange(&base_ptr, &mem_size, + reinterpret_cast(device_alloc.ptr() + offset))); + + REQUIRE(reinterpret_cast(device_alloc.ptr()) == base_ptr); + REQUIRE(mem_size == allocation_size); + } + SECTION("Pitch address range") { + size_t width = 32; + size_t height = 32; + LinearAllocGuard2D device_alloc(width, height); + + HIP_CHECK(hipMemGetAddressRange(&base_ptr, &mem_size, + reinterpret_cast(device_alloc.ptr() + offset))); + + REQUIRE(reinterpret_cast(device_alloc.ptr()) == base_ptr); + REQUIRE(mem_size == (device_alloc.pitch() * height)); + } +} + +TEST_CASE("Unit_hipMemGetAddressRange_Negative") { + hipDeviceptr_t base_ptr; + size_t mem_size = 0; + const auto allocation_size = kPageSize / 2; + const int offset = kPageSize; + LinearAllocGuard host_alloc(LinearAllocs::hipHostMalloc, allocation_size); + + hipDeviceptr_t dummy_ptr; + + SECTION("Device pointer is invalid") { + HIP_CHECK_ERROR(hipMemGetAddressRange(&base_ptr, &mem_size, dummy_ptr), hipErrorNotFound); + } + SECTION("Offset is greater than allocated size") { + HIP_CHECK_ERROR( + hipMemGetAddressRange(&base_ptr, &mem_size, + reinterpret_cast(host_alloc.ptr() + offset)), + hipErrorNotFound); + } +} diff --git a/catch/unit/memory/hipMemcpyPeer.cc b/catch/unit/memory/hipMemcpyPeer.cc index 33ee2e9f70..245b3e0b2c 100644 --- a/catch/unit/memory/hipMemcpyPeer.cc +++ b/catch/unit/memory/hipMemcpyPeer.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,143 +16,201 @@ 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. */ - +/* +Testcase Scenarios : +Unit_hipMemcpyPeer_Positive_Default - Test basic P2P memcpy between two devices +with hipMemcpyPeer api Unit_hipMemcpyPeer_Positive_Synchronization_Behavior - +Test synchronization behavior for hipMemcpyPeer api +Unit_hipMemcpyPeer_Positive_ZeroSize - Test that no data is copied when +sizeBytes is set to 0 Unit_hipMemcpyPeer_Negative_Parameters - Test unsuccessful +execution of hipMemcpyPeer api when parameters are invalid +*/ +#include #include -#include -#include -/* -This testfile verifies the following scenarios of hipMemcpyPeer API -1. Negative Scenarios -2. Basic scenario of hipMemcpyPeer API -*/ +#include +#include -/*This testcase verifies the negative scenarios of hipmemcpypeer -*/ -TEST_CASE("Unit_hipMemcpyPeer_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}; - HIP_CHECK(hipSetDevice(0)); - HipTest::initArrays(&A_d, nullptr, nullptr, - &A_h, &B_h, nullptr, numElements*sizeof(int)); - HipTest::setDefaultData(numElements, A_h, B_h, nullptr); - HIP_CHECK(hipSetDevice(1)); - HipTest::initArrays(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(hipMemcpyPeer(nullptr, 1, A_d, 0, copy_bytes) != hipSuccess); - } +TEST_CASE("Unit_hipMemcpyPeer_Positive_Default") { + const auto device_count = HipTest::getDeviceCount(); + if (device_count < 2) { + HipTest::HIP_SKIP_TEST("Skipping because devices < 2"); + return; + } - SECTION("Nullptr to Source Pointer") { - REQUIRE(hipMemcpyPeer(B_d, 1, nullptr, 0, copy_bytes) != hipSuccess); - } + const auto allocation_size = GENERATE(kPageSize / 2, kPageSize, kPageSize * 2); - SECTION("Pass NumElements as 0") { - HIP_CHECK(hipMemcpy(A_d, A_h, numElements*sizeof(int), - hipMemcpyHostToDevice)); - HIP_CHECK(hipMemcpyPeer(B_d, 1, A_d, 0, 0)); - HIP_CHECK(hipMemcpy(A_h, B_d, numElements*sizeof(int), - hipMemcpyDeviceToHost)); - HipTest::checkTest(A_h, B_h, numElements); - } + int can_access_peer = 0; + 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); - SECTION("Passing more than allocated size") { - REQUIRE(hipMemcpyPeer(B_d, 1, A_d, 0, - ((numElements+40)*sizeof(int))) != hipSuccess); - } + 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)); - SECTION("Passing invalid Destination device ID") { - REQUIRE(hipMemcpyPeer(B_d, numDevices, A_d, 0, copy_bytes) != - hipSuccess); - } + LinearAllocGuard src_alloc(LinearAllocs::hipMalloc, allocation_size); + LinearAllocGuard result(LinearAllocs::hipHostMalloc, allocation_size); + HIP_CHECK(hipSetDevice(dst_device)); + LinearAllocGuard dst_alloc(LinearAllocs::hipMalloc, allocation_size); - SECTION("Passing invalid Source device ID") { - REQUIRE(hipMemcpyPeer(B_d, 1, A_d, numDevices, copy_bytes) != - hipSuccess); - } - HipTest::freeArrays(A_d, B_d, nullptr, A_h, B_h, nullptr, false); - } 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<<>>(src_alloc.ptr(), expected_value, element_count); + HIP_CHECK(hipGetLastError()); + + HIP_CHECK( + hipMemcpyPeer(dst_alloc.ptr(), dst_device, src_alloc.ptr(), src_device, allocation_size)); + + 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 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 hipMemcpyPeer API - * Then performs the addition and validates the sum - */ -TEST_CASE("Unit_hipMemcpyPeer_Basic") { - 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}; +TEST_CASE("Unit_hipMemcpyPeer_Positive_Synchronization_Behavior") { + HIP_CHECK(hipDeviceSynchronize()); - // Initialization of Variables on GPU-0 - HIP_CHECK(hipSetDevice(0)); - HipTest::initArrays(&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)); + const auto device_count = HipTest::getDeviceCount(); + if (device_count < 2) { + HipTest::HIP_SKIP_TEST("Skipping because devices < 2"); + return; + } - // Initialization of Variables on GPU-1 - HIP_CHECK(hipSetDevice(1)); - HipTest::initArrays(&X_d, &Y_d, &Z_d, nullptr, - nullptr, nullptr, numElements*sizeof(int)); + int can_access_peer = 0; + const auto src_device = 0; + const auto dst_device = 1; - // Launching kernel and performing vector addition on GPU-0 - HIP_CHECK(hipSetDevice(0)); - hipLaunchKernelGGL(HipTest::vectorADD, dim3(1), dim3(1), - 0, 0, static_cast(A_d), - static_cast(B_d), C_d, numElements*sizeof(int)); - HIP_CHECK(hipGetLastError()); - HIP_CHECK(hipMemcpy(C_h, C_d, numElements*sizeof(int), - hipMemcpyDeviceToHost)); - HipTest::checkVectorADD(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)); - HIP_CHECK(hipSetDevice(1)); - // Copying data from GPU-0 to GPU-1 and performing vector addition - HIP_CHECK(hipMemcpyPeer(X_d, 1, A_d, 0, copy_bytes)); - HIP_CHECK(hipMemcpyPeer(Y_d, 1, B_d, 0, copy_bytes)); - hipLaunchKernelGGL(HipTest::vectorADD, dim3(1), dim3(1), - 0, 0, static_cast(X_d), - static_cast(Y_d), Z_d, numElements*sizeof(int)); - HIP_CHECK(hipGetLastError()); - HIP_CHECK(hipMemcpy(C_h, Z_d, numElements*sizeof(int), - hipMemcpyDeviceToHost)); - HipTest::checkVectorADD(A_h, B_h, C_h, numElements); + LinearAllocGuard src_alloc(LinearAllocs::hipMalloc, kPageSize); + HIP_CHECK(hipSetDevice(dst_device)); + LinearAllocGuard dst_alloc(LinearAllocs::hipMalloc, kPageSize); - // Cleaning the memory - HipTest::freeArrays(A_d, B_d, C_d, A_h, B_h, C_h, false); - HipTest::freeArrays(X_d, Y_d, Z_d, nullptr, nullptr, nullptr, false); - } else { - SUCCEED("Machine Does not have P2P capability"); - } + HIP_CHECK(hipSetDevice(src_device)); + LaunchDelayKernel(std::chrono::milliseconds{100}, nullptr); + + HIP_CHECK(hipMemcpyPeer(dst_alloc.ptr(), dst_device, src_alloc.ptr(), src_device, kPageSize)); + 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_hipMemcpyPeer_Positive_ZeroSize") { + const auto device_count = HipTest::getDeviceCount(); + if (device_count < 2) { + HipTest::HIP_SKIP_TEST("Skipping because devices < 2"); + return; + } + + 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 src_alloc(LinearAllocs::hipMalloc, allocation_size); + LinearAllocGuard result(LinearAllocs::hipHostMalloc, allocation_size, + hipHostMallocPortable); + HIP_CHECK(hipSetDevice(dst_device)); + LinearAllocGuard 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<<>>(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(hipMemcpyPeer(dst_alloc.ptr(), dst_device, src_alloc.ptr(), src_device, 0)); + + 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_hipMemcpyPeer_Negative_Parameters") { + const auto device_count = HipTest::getDeviceCount(); + if (device_count < 2) { + HipTest::HIP_SKIP_TEST("Skipping because devices < 2"); + return; + } + + 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 src_alloc(LinearAllocs::hipMalloc, kPageSize); + HIP_CHECK(hipSetDevice(dst_device)); + LinearAllocGuard dst_alloc(LinearAllocs::hipMalloc, kPageSize); + + HIP_CHECK(hipSetDevice(src_device)); + + SECTION("Nullptr to Destination Pointer") { + HIP_CHECK_ERROR(hipMemcpyPeer(nullptr, dst_device, src_alloc.ptr(), src_device, kPageSize), + hipErrorInvalidValue); + } + + SECTION("Nullptr to Source Pointer") { + HIP_CHECK_ERROR(hipMemcpyPeer(dst_alloc.ptr(), dst_device, nullptr, src_device, kPageSize), + hipErrorInvalidValue); + } + + SECTION("Passing more than allocated size") { + HIP_CHECK_ERROR( + hipMemcpyPeer(dst_alloc.ptr(), dst_device, src_alloc.ptr(), src_device, kPageSize + 1), + hipErrorInvalidValue); + } + + SECTION("Passing invalid Destination device ID") { + HIP_CHECK_ERROR( + hipMemcpyPeer(dst_alloc.ptr(), device_count, src_alloc.ptr(), src_device, kPageSize), + hipErrorInvalidDevice); + } + + SECTION("Passing invalid Source device ID") { + HIP_CHECK_ERROR( + hipMemcpyPeer(dst_alloc.ptr(), dst_device, src_alloc.ptr(), device_count, kPageSize), + hipErrorInvalidDevice); + } + + HIP_CHECK(hipDeviceDisablePeerAccess(dst_device)); + } else { + INFO("Peer access cannot be enabled between devices " << src_device << " " << dst_device); } } diff --git a/catch/unit/memory/hipMemcpyPeerAsync.cc b/catch/unit/memory/hipMemcpyPeerAsync.cc index 159c8309e3..90f9fdb8f4 100644 --- a/catch/unit/memory/hipMemcpyPeerAsync.cc +++ b/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 #include -#include -#include -#include +#include +#include -/*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(&A_d, nullptr, nullptr, - &A_h, &B_h, nullptr, numElements*sizeof(int)); - HipTest::setDefaultData(numElements, A_h, B_h, nullptr); - HIP_CHECK(hipSetDevice(1)); - HipTest::initArrays(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(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(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(&A_d, &B_d, &C_d, - &A_h, &B_h, &C_h, numElements*sizeof(int)); - HipTest::setDefaultData(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(&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(A_d), - static_cast(B_d), C_d, numElements*sizeof(int)); - HIP_CHECK(hipGetLastError()); - HIP_CHECK(hipMemcpy(C_h, C_d, numElements*sizeof(int), - hipMemcpyDeviceToHost)); - HipTest::checkVectorADD(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(X_d), - static_cast(Y_d), Z_d, numElements*sizeof(int)); - HIP_CHECK(hipGetLastError()); - HIP_CHECK(hipMemcpy(C_h, Z_d, numElements*sizeof(int), - hipMemcpyDeviceToHost)); - HipTest::checkVectorADD(A_h, B_h, C_h, numElements); + LinearAllocGuard src_alloc(LinearAllocs::hipMalloc, allocation_size); + LinearAllocGuard result(LinearAllocs::hipHostMalloc, allocation_size, + hipHostMallocPortable); + HIP_CHECK(hipSetDevice(dst_device)); + LinearAllocGuard dst_alloc(LinearAllocs::hipMalloc, allocation_size); - // Cleaning the Memory - HipTest::freeArrays(A_d, B_d, C_d, A_h, B_h, C_h, false); - HipTest::freeArrays(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<<>>(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(&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(&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(A_d), - static_cast(B_d), C_d, numElements*sizeof(int)); - HIP_CHECK(hipGetLastError()); - HIP_CHECK(hipMemcpy(C_h, C_d, numElements*sizeof(int), - hipMemcpyDeviceToHost)); - HipTest::checkVectorADD(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(X_d), - static_cast(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(A_h, B_h, C_h, numElements); - HipTest::freeArrays(A_d, B_d, C_d, A_h, B_h, C_h, false); - HipTest::freeArrays(X_d, Y_d, Z_d, nullptr, nullptr, nullptr, false); - HIP_CHECK(hipStreamDestroy(stream)); - } else { - SUCCEED("Machine Does not have P2P capability"); - } + LinearAllocGuard src_alloc(LinearAllocs::hipMalloc, kPageSize); + HIP_CHECK(hipSetDevice(dst_device)); + LinearAllocGuard 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 src_alloc(LinearAllocs::hipMalloc, allocation_size); + LinearAllocGuard result(LinearAllocs::hipHostMalloc, allocation_size, + hipHostMallocPortable); + HIP_CHECK(hipSetDevice(dst_device)); + LinearAllocGuard 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<<>>(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 src_alloc(LinearAllocs::hipMalloc, kPageSize); + HIP_CHECK(hipSetDevice(dst_device)); + LinearAllocGuard 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); + } +} diff --git a/catch/unit/memory/hipMemcpyPeerAsync_old.cc b/catch/unit/memory/hipMemcpyPeerAsync_old.cc new file mode 100644 index 0000000000..159c8309e3 --- /dev/null +++ b/catch/unit/memory/hipMemcpyPeerAsync_old.cc @@ -0,0 +1,262 @@ +/* +Copyright (c) 2021 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. +*/ + +/* +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 +*/ + + +#include +#include +#include +#include + +/*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(&A_d, nullptr, nullptr, + &A_h, &B_h, nullptr, numElements*sizeof(int)); + HipTest::setDefaultData(numElements, A_h, B_h, nullptr); + HIP_CHECK(hipSetDevice(1)); + HipTest::initArrays(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(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(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"); + } +} +/* + * 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 + */ + +TEST_CASE("Unit_hipMemcpyPeerAsync_Basic") { + 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 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(&A_d, &B_d, &C_d, + &A_h, &B_h, &C_h, numElements*sizeof(int)); + HipTest::setDefaultData(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)); + + // Initialization of Variables in GPU-1 + HIP_CHECK(hipSetDevice(1)); + HipTest::initArrays(&X_d, &Y_d, &Z_d, nullptr, + nullptr, nullptr, numElements*sizeof(int)); + + // Launching kernel and performing vector addition in GPU-0 + HIP_CHECK(hipSetDevice(0)); + hipLaunchKernelGGL(HipTest::vectorADD, dim3(1), dim3(1), + 0, 0, static_cast(A_d), + static_cast(B_d), C_d, numElements*sizeof(int)); + HIP_CHECK(hipGetLastError()); + HIP_CHECK(hipMemcpy(C_h, C_d, numElements*sizeof(int), + hipMemcpyDeviceToHost)); + HipTest::checkVectorADD(A_h, B_h, C_h, numElements); + + // 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(X_d), + static_cast(Y_d), Z_d, numElements*sizeof(int)); + HIP_CHECK(hipGetLastError()); + HIP_CHECK(hipMemcpy(C_h, Z_d, numElements*sizeof(int), + hipMemcpyDeviceToHost)); + HipTest::checkVectorADD(A_h, B_h, C_h, numElements); + + // Cleaning the Memory + HipTest::freeArrays(A_d, B_d, C_d, A_h, B_h, C_h, false); + HipTest::freeArrays(X_d, Y_d, Z_d, nullptr, nullptr, nullptr, false); + HIP_CHECK(hipStreamDestroy(stream)); + } else { + SUCCEED("Machine Does not have P2P capability"); + } + } else { + SUCCEED("Number of devices are < 2"); + } +} + +/* + * 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)); + + // Initialization of all variables in GPU-0 + HipTest::initArrays(&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(&X_d, &Y_d, &Z_d, nullptr, + nullptr, nullptr, numElements*sizeof(int)); + + // Stream created in GPU-1 + HIP_CHECK(hipSetDevice(1)); + HIP_CHECK(hipStreamCreate(&stream)); + + // Performing vector addition and validate the data + HIP_CHECK(hipSetDevice(0)); + hipLaunchKernelGGL(HipTest::vectorADD, dim3(1), dim3(1), + 0, 0, static_cast(A_d), + static_cast(B_d), C_d, numElements*sizeof(int)); + HIP_CHECK(hipGetLastError()); + HIP_CHECK(hipMemcpy(C_h, C_d, numElements*sizeof(int), + hipMemcpyDeviceToHost)); + HipTest::checkVectorADD(A_h, B_h, C_h, numElements); + + // 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(X_d), + static_cast(Y_d), Z_d, numElements*sizeof(int)); + HIP_CHECK(hipGetLastError()); + HIP_CHECK(hipMemcpy(C_h, Z_d, numElements*sizeof(int), + hipMemcpyDeviceToHost)); + + // Cleaning the data + HipTest::checkVectorADD(A_h, B_h, C_h, numElements); + HipTest::freeArrays(A_d, B_d, C_d, A_h, B_h, C_h, false); + HipTest::freeArrays(X_d, Y_d, Z_d, nullptr, nullptr, nullptr, false); + HIP_CHECK(hipStreamDestroy(stream)); + } else { + SUCCEED("Machine Does not have P2P capability"); + } + } else { + SUCCEED("Number of devices are < 2"); + } +} + diff --git a/catch/unit/memory/hipMemcpyPeer_old.cc b/catch/unit/memory/hipMemcpyPeer_old.cc new file mode 100644 index 0000000000..33ee2e9f70 --- /dev/null +++ b/catch/unit/memory/hipMemcpyPeer_old.cc @@ -0,0 +1,158 @@ +/* +Copyright (c) 2021 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 +#include +#include +/* +This testfile verifies the following scenarios of hipMemcpyPeer API +1. Negative Scenarios +2. Basic scenario of hipMemcpyPeer API +*/ + +/*This testcase verifies the negative scenarios of hipmemcpypeer +*/ +TEST_CASE("Unit_hipMemcpyPeer_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}; + HIP_CHECK(hipSetDevice(0)); + HipTest::initArrays(&A_d, nullptr, nullptr, + &A_h, &B_h, nullptr, numElements*sizeof(int)); + HipTest::setDefaultData(numElements, A_h, B_h, nullptr); + HIP_CHECK(hipSetDevice(1)); + HipTest::initArrays(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(hipMemcpyPeer(nullptr, 1, A_d, 0, copy_bytes) != hipSuccess); + } + + SECTION("Nullptr to Source Pointer") { + REQUIRE(hipMemcpyPeer(B_d, 1, nullptr, 0, copy_bytes) != hipSuccess); + } + + SECTION("Pass NumElements as 0") { + HIP_CHECK(hipMemcpy(A_d, A_h, numElements*sizeof(int), + hipMemcpyHostToDevice)); + HIP_CHECK(hipMemcpyPeer(B_d, 1, A_d, 0, 0)); + HIP_CHECK(hipMemcpy(A_h, B_d, numElements*sizeof(int), + hipMemcpyDeviceToHost)); + HipTest::checkTest(A_h, B_h, numElements); + } + + SECTION("Passing more than allocated size") { + REQUIRE(hipMemcpyPeer(B_d, 1, A_d, 0, + ((numElements+40)*sizeof(int))) != hipSuccess); + } + + SECTION("Passing invalid Destination device ID") { + REQUIRE(hipMemcpyPeer(B_d, numDevices, A_d, 0, copy_bytes) != + hipSuccess); + } + + SECTION("Passing invalid Source device ID") { + REQUIRE(hipMemcpyPeer(B_d, 1, A_d, numDevices, copy_bytes) != + hipSuccess); + } + HipTest::freeArrays(A_d, B_d, nullptr, A_h, B_h, nullptr, false); + } else { + SUCCEED("Machine Does not have P2P capability"); + } + } else { + SUCCEED("Number of devices are < 2"); + } +} + +/* + * 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 hipMemcpyPeer API + * Then performs the addition and validates the sum + */ +TEST_CASE("Unit_hipMemcpyPeer_Basic") { + 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}; + + // Initialization of Variables on GPU-0 + HIP_CHECK(hipSetDevice(0)); + HipTest::initArrays(&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)); + + // Initialization of Variables on GPU-1 + HIP_CHECK(hipSetDevice(1)); + HipTest::initArrays(&X_d, &Y_d, &Z_d, nullptr, + nullptr, nullptr, numElements*sizeof(int)); + + // Launching kernel and performing vector addition on GPU-0 + HIP_CHECK(hipSetDevice(0)); + hipLaunchKernelGGL(HipTest::vectorADD, dim3(1), dim3(1), + 0, 0, static_cast(A_d), + static_cast(B_d), C_d, numElements*sizeof(int)); + HIP_CHECK(hipGetLastError()); + HIP_CHECK(hipMemcpy(C_h, C_d, numElements*sizeof(int), + hipMemcpyDeviceToHost)); + HipTest::checkVectorADD(A_h, B_h, C_h, numElements); + + HIP_CHECK(hipSetDevice(1)); + // Copying data from GPU-0 to GPU-1 and performing vector addition + HIP_CHECK(hipMemcpyPeer(X_d, 1, A_d, 0, copy_bytes)); + HIP_CHECK(hipMemcpyPeer(Y_d, 1, B_d, 0, copy_bytes)); + hipLaunchKernelGGL(HipTest::vectorADD, dim3(1), dim3(1), + 0, 0, static_cast(X_d), + static_cast(Y_d), Z_d, numElements*sizeof(int)); + HIP_CHECK(hipGetLastError()); + HIP_CHECK(hipMemcpy(C_h, Z_d, numElements*sizeof(int), + hipMemcpyDeviceToHost)); + HipTest::checkVectorADD(A_h, B_h, C_h, numElements); + + // Cleaning the memory + HipTest::freeArrays(A_d, B_d, C_d, A_h, B_h, C_h, false); + HipTest::freeArrays(X_d, Y_d, Z_d, nullptr, nullptr, nullptr, false); + } else { + SUCCEED("Machine Does not have P2P capability"); + } + } else { + SUCCEED("Number of devices are < 2"); + } +}