diff --git a/projects/hip-tests/catch/unit/memory/CMakeLists.txt b/projects/hip-tests/catch/unit/memory/CMakeLists.txt index 6554caec00..d0dba71e6c 100644 --- a/projects/hip-tests/catch/unit/memory/CMakeLists.txt +++ b/projects/hip-tests/catch/unit/memory/CMakeLists.txt @@ -26,6 +26,7 @@ set(TEST_SRC malloc.cc hipMemcpy2DToArray.cc hipMemcpy2DToArrayAsync.cc + hipMemcpy2DToArrayAsync_old.cc hipMemcpy3D.cc hipMemcpy3DAsync.cc hipMemcpyParam2D.cc @@ -34,6 +35,7 @@ set(TEST_SRC hipMemcpy2DAsync.cc hipMemcpy2DFromArray.cc hipMemcpy2DFromArrayAsync.cc + hipMemcpy2DFromArrayAsync_old.cc hipMemcpyAtoH.cc hipMemcpyHtoA.cc hipMemcpyAllApiNegative.cc @@ -111,6 +113,7 @@ set(TEST_SRC malloc.cc hipMemcpy2DToArray.cc hipMemcpy2DToArrayAsync.cc + hipMemcpy2DToArrayAsync_old.cc hipMemcpy3D.cc hipMemcpy3DAsync.cc hipMemcpyParam2D.cc @@ -119,6 +122,7 @@ set(TEST_SRC hipMemcpy2DAsync.cc hipMemcpy2DFromArray.cc hipMemcpy2DFromArrayAsync.cc + hipMemcpy2DFromArrayAsync_old.cc hipMemcpyAtoH.cc hipMemcpyHtoA.cc hipMemcpyAllApiNegative.cc diff --git a/projects/hip-tests/catch/unit/memory/array_memcpy_tests_common.hh b/projects/hip-tests/catch/unit/memory/array_memcpy_tests_common.hh new file mode 100644 index 0000000000..71a02a8aca --- /dev/null +++ b/projects/hip-tests/catch/unit/memory/array_memcpy_tests_common.hh @@ -0,0 +1,442 @@ +/* +Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved. +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +#pragma once + +#include + +#include +#include +#include +#include +#include "hipArrayCommon.hh" + +/* Array -> Host */ +template +void MemcpyAtoHShell(F memcpy_func, size_t width, const hipStream_t kernel_stream = nullptr) { + const unsigned int flag = hipArrayDefault; + + size_t allocation_size = width * sizeof(T); + + LinearAllocGuard host_allocation(LinearAllocs::hipHostMalloc, allocation_size); + ArrayAllocGuard array_allocation(make_hipExtent(width, 0, 0), flag); + + const auto element_count = allocation_size / sizeof(T); + constexpr int fill_value = 42; + std::fill_n(host_allocation.host_ptr(), element_count, fill_value); + + HIP_CHECK(hipMemcpy2DToArray(array_allocation.ptr(), 0, 0, host_allocation.host_ptr(), + sizeof(T) * width, sizeof(T) * width, 1, hipMemcpyHostToDevice)); + std::fill_n(host_allocation.host_ptr(), element_count, 0); + + HIP_CHECK(memcpy_func(host_allocation.host_ptr(), array_allocation.ptr())); + if (should_synchronize) { + HIP_CHECK(hipStreamSynchronize(kernel_stream)); + } + + ArrayFindIfNot(host_allocation.host_ptr(), fill_value, element_count); +} + +template +void Memcpy2DHostFromAShell(F memcpy_func, size_t width, size_t height, + const hipStream_t kernel_stream = nullptr) { + const unsigned int flag = hipArrayDefault; + + size_t allocation_size = width * height * sizeof(T); + + LinearAllocGuard host_allocation(LinearAllocs::hipHostMalloc, allocation_size); + ArrayAllocGuard array_allocation(make_hipExtent(width, height, 0), flag); + + const auto element_count = allocation_size / sizeof(T); + constexpr int fill_value = 42; + std::fill_n(host_allocation.host_ptr(), element_count, fill_value); + + HIP_CHECK(hipMemcpy2DToArray(array_allocation.ptr(), 0, 0, host_allocation.host_ptr(), + sizeof(T) * width, sizeof(T) * width, height, + hipMemcpyHostToDevice)); + std::fill_n(host_allocation.host_ptr(), element_count, 0); + + HIP_CHECK(memcpy_func(host_allocation.host_ptr(), sizeof(T) * width, array_allocation.ptr())); + if (should_synchronize) { + HIP_CHECK(hipStreamSynchronize(kernel_stream)); + } + + ArrayFindIfNot(host_allocation.host_ptr(), fill_value, element_count); +} + +/* Array -> Device */ +template +void Memcpy2DDeviceFromAShell(F memcpy_func, size_t width, size_t height, + const hipStream_t kernel_stream = nullptr) { + const unsigned int flag = hipArrayDefault; + size_t allocation_size = width * height * sizeof(T); + + const auto device_count = HipTest::getDeviceCount(); + const auto src_device = GENERATE_COPY(range(0, device_count)); + const auto dst_device = GENERATE_COPY(range(0, device_count)); + INFO("Src device: " << src_device << ", Dst device: " << dst_device); + + HIP_CHECK(hipSetDevice(src_device)); + if constexpr (enable_peer_access) { + if (src_device == dst_device) { + return; + } + int can_access_peer = 0; + HIP_CHECK(hipDeviceCanAccessPeer(&can_access_peer, src_device, dst_device)); + if (!can_access_peer) { + INFO("Peer access cannot be enabled between devices " << src_device << " " << dst_device); + REQUIRE(can_access_peer); + } + HIP_CHECK(hipDeviceEnablePeerAccess(dst_device, 0)); + } + + LinearAllocGuard host_allocation(LinearAllocs::hipHostMalloc, allocation_size); + ArrayAllocGuard array_allocation(make_hipExtent(width, height, 0), flag); + HIP_CHECK(hipSetDevice(dst_device)); + LinearAllocGuard2D device_allocation(width, height); + + HIP_CHECK(hipSetDevice(src_device)); + const auto element_count = allocation_size / sizeof(T); + constexpr int fill_value = 42; + std::fill_n(host_allocation.host_ptr(), element_count, fill_value); + + HIP_CHECK(hipMemcpy2DToArray(array_allocation.ptr(), 0, 0, host_allocation.host_ptr(), + sizeof(T) * width, sizeof(T) * width, height, + hipMemcpyHostToDevice)); + std::fill_n(host_allocation.host_ptr(), element_count, 0); + + HIP_CHECK( + memcpy_func(device_allocation.ptr(), device_allocation.pitch(), array_allocation.ptr())); + if (should_synchronize) { + HIP_CHECK(hipStreamSynchronize(kernel_stream)); + } + + HIP_CHECK(hipMemcpy2D(host_allocation.host_ptr(), sizeof(T) * width, device_allocation.ptr(), + device_allocation.pitch(), sizeof(T) * width, height, + hipMemcpyDeviceToHost)); + + if constexpr (enable_peer_access) { + // If we've gotten this far, EnablePeerAccess must have succeeded, so we only need to check this + // condition + HIP_CHECK(hipDeviceDisablePeerAccess(dst_device)); + } + + ArrayFindIfNot(host_allocation.host_ptr(), fill_value, element_count); +} + +/* Host -> Array */ +template +void MemcpyHtoAShell(F memcpy_func, size_t width, const hipStream_t kernel_stream = nullptr) { + const unsigned int flag = hipArrayDefault; + + size_t allocation_size = width * sizeof(T); + + LinearAllocGuard host_allocation(LinearAllocs::hipHostMalloc, allocation_size); + ArrayAllocGuard array_allocation(make_hipExtent(width, 0, 0), flag); + + const auto element_count = allocation_size / sizeof(T); + constexpr int fill_value = 41; + std::fill_n(host_allocation.host_ptr(), element_count, fill_value); + + HIP_CHECK(memcpy_func(array_allocation.ptr(), host_allocation.host_ptr())); + if (should_synchronize) { + HIP_CHECK(hipStreamSynchronize(kernel_stream)); + } + + std::fill_n(host_allocation.host_ptr(), element_count, 0); + + HIP_CHECK(hipMemcpy2DFromArray(host_allocation.host_ptr(), sizeof(T) * width, + array_allocation.ptr(), 0, 0, sizeof(T) * width, 1, + hipMemcpyDeviceToHost)); + + ArrayFindIfNot(host_allocation.host_ptr(), fill_value, element_count); +} + +template +void Memcpy2DHosttoAShell(F memcpy_func, size_t width, size_t height, + const hipStream_t kernel_stream = nullptr) { + const unsigned int flag = hipArrayDefault; + ; + + size_t allocation_size = width * height * sizeof(T); + + LinearAllocGuard host_allocation(LinearAllocs::hipHostMalloc, allocation_size); + ArrayAllocGuard array_allocation(make_hipExtent(width, height, 0), flag); + + const auto element_count = allocation_size / sizeof(T); + constexpr int fill_value = 41; + std::fill_n(host_allocation.host_ptr(), element_count, fill_value); + + HIP_CHECK(memcpy_func(array_allocation.ptr(), host_allocation.host_ptr(), sizeof(T) * width)); + if (should_synchronize) { + HIP_CHECK(hipStreamSynchronize(kernel_stream)); + } + + std::fill_n(host_allocation.host_ptr(), element_count, 0); + + HIP_CHECK(hipMemcpy2DFromArray(host_allocation.host_ptr(), sizeof(T) * width, + array_allocation.ptr(), 0, 0, sizeof(T) * width, height, + hipMemcpyDeviceToHost)); + + ArrayFindIfNot(host_allocation.host_ptr(), fill_value, element_count); +} + +/* Device -> Array */ +template +void Memcpy2DDevicetoAShell(F memcpy_func, size_t width, size_t height, + const hipStream_t kernel_stream = nullptr) { + const unsigned int flag = hipArrayDefault; + size_t allocation_size = width * height * sizeof(T); + + const auto device_count = HipTest::getDeviceCount(); + const auto src_device = GENERATE_COPY(range(0, device_count)); + const auto dst_device = GENERATE_COPY(range(0, device_count)); + INFO("Src device: " << src_device << ", Dst device: " << dst_device); + + HIP_CHECK(hipSetDevice(src_device)); + if constexpr (enable_peer_access) { + if (src_device == dst_device) { + return; + } + int can_access_peer = 0; + HIP_CHECK(hipDeviceCanAccessPeer(&can_access_peer, src_device, dst_device)); + if (!can_access_peer) { + INFO("Peer access cannot be enabled between devices " << src_device << " " << dst_device); + REQUIRE(can_access_peer); + } + HIP_CHECK(hipDeviceEnablePeerAccess(dst_device, 0)); + } + + LinearAllocGuard host_allocation(LinearAllocs::hipHostMalloc, allocation_size); + LinearAllocGuard2D device_allocation(width, height); + HIP_CHECK(hipSetDevice(dst_device)); + ArrayAllocGuard array_allocation(make_hipExtent(width, height, 0), flag); + + HIP_CHECK(hipSetDevice(src_device)); + const auto element_count = allocation_size / sizeof(T); + constexpr int fill_value = 41; + std::fill_n(host_allocation.host_ptr(), element_count, fill_value); + + HIP_CHECK(hipMemcpy2D(device_allocation.ptr(), device_allocation.pitch(), + host_allocation.host_ptr(), sizeof(T) * width, sizeof(T) * width, height, + hipMemcpyHostToDevice)); + + HIP_CHECK( + memcpy_func(array_allocation.ptr(), device_allocation.ptr(), device_allocation.pitch())); + if (should_synchronize) { + HIP_CHECK(hipStreamSynchronize(kernel_stream)); + } + + std::fill_n(host_allocation.host_ptr(), element_count, 0); + + HIP_CHECK(hipMemcpy2DFromArray(host_allocation.host_ptr(), sizeof(T) * width, + array_allocation.ptr(), 0, 0, sizeof(T) * width, height, + hipMemcpyDeviceToHost)); + + if constexpr (enable_peer_access) { + // If we've gotten this far, EnablePeerAccess must have succeeded, so we only need to check this + // condition + HIP_CHECK(hipDeviceDisablePeerAccess(dst_device)); + } + + ArrayFindIfNot(host_allocation.host_ptr(), fill_value, element_count); +} + +// Synchronization behavior checks +template +void MemcpyArraySyncBehaviorCheck(F memcpy_func, const bool should_sync, + const hipStream_t kernel_stream) { + LaunchDelayKernel(std::chrono::milliseconds{100}, kernel_stream); + HIP_CHECK(memcpy_func()); + if (should_sync) { + HIP_CHECK(hipStreamQuery(kernel_stream)); + } else { + HIP_CHECK_ERROR(hipStreamQuery(kernel_stream), hipErrorNotReady); + } +} + +/* Host -> Array Sync check */ +template +void MemcpyHtoASyncBehavior(F memcpy_func, size_t width, size_t height, const bool should_sync, + const hipStream_t kernel_stream = nullptr) { + const unsigned int flag = hipArrayDefault; + size_t num_h = (height == 0) ? 1 : height; + size_t allocation_size = width * num_h * sizeof(int); + + LinearAllocGuard host_alloc(LinearAllocs::hipHostMalloc, allocation_size); + ArrayAllocGuard array_allocation(make_hipExtent(width, height, 0), flag); + MemcpyArraySyncBehaviorCheck(std::bind(memcpy_func, array_allocation.ptr(), host_alloc.ptr()), + should_sync, kernel_stream); +} + +/* Array -> Host sync check */ +template +void MemcpyAtoHPageableSyncBehavior(F memcpy_func, size_t width, size_t height, + const bool should_sync, + const hipStream_t kernel_stream = nullptr) { + const unsigned int flag = hipArrayDefault; + size_t num_h = (height == 0) ? 1 : height; + size_t allocation_size = width * num_h * sizeof(int); + + LinearAllocGuard host_alloc(LinearAllocs::hipHostMalloc, allocation_size); + ArrayAllocGuard array_allocation(make_hipExtent(width, height, 0), flag); + MemcpyArraySyncBehaviorCheck(std::bind(memcpy_func, host_alloc.ptr(), array_allocation.ptr()), + should_sync, kernel_stream); +} + +template +void MemcpyAtoHPinnedSyncBehavior(F memcpy_func, size_t width, size_t height, + const bool should_sync, + const hipStream_t kernel_stream = nullptr) { + const unsigned int flag = hipArrayDefault; + size_t num_h = (height == 0) ? 1 : height; + size_t allocation_size = width * num_h * sizeof(int); + + LinearAllocGuard host_alloc(LinearAllocs::hipHostMalloc, allocation_size); + ArrayAllocGuard array_allocation(make_hipExtent(width, height, 0), flag); + MemcpyArraySyncBehaviorCheck(std::bind(memcpy_func, host_alloc.ptr(), array_allocation.ptr()), + should_sync, kernel_stream); +} + +/* Device -> Array sync check */ +template +void MemcpyDtoASyncBehavior(F memcpy_func, size_t width, size_t height, const bool should_sync, + const hipStream_t kernel_stream = nullptr) { + const unsigned int flag = hipArrayDefault; + + ArrayAllocGuard array_allocation(make_hipExtent(width, height, 0), flag); + LinearAllocGuard2D device_allocation(width, height); + + MemcpyArraySyncBehaviorCheck(std::bind(memcpy_func, array_allocation.ptr(), + device_allocation.ptr(), device_allocation.pitch()), + should_sync, kernel_stream); +} + +/* Array -> Device sync check */ +template +void MemcpyAtoDSyncBehavior(F memcpy_func, size_t width, size_t height, const bool should_sync, + const hipStream_t kernel_stream = nullptr) { + const unsigned int flag = hipArrayDefault; + + LinearAllocGuard2D device_allocation(width, height); + ArrayAllocGuard array_allocation(make_hipExtent(width, height, 0), flag); + + MemcpyArraySyncBehaviorCheck(std::bind(memcpy_func, device_allocation.ptr(), + device_allocation.pitch(), array_allocation.ptr()), + should_sync, kernel_stream); +} + +/* Array -> Host/Device zero copy */ +template +void Memcpy2DFromArrayZeroWidthHeight(F memcpy_func, size_t width, size_t height, + const hipStream_t stream = nullptr) { + const unsigned int flag = hipArrayDefault; + const auto element_count = width * height; + + SECTION("Device to Host") { + ArrayAllocGuard array_alloc(make_hipExtent(width, height, 0), flag); + LinearAllocGuard host_alloc(LinearAllocs::hipHostMalloc, width * height * sizeof(int)); + int fill_value = 42; + std::fill_n(host_alloc.host_ptr(), width * height, fill_value); + HIP_CHECK(hipMemcpy2DToArray(array_alloc.ptr(), 0, 0, host_alloc.host_ptr(), + sizeof(int) * width, sizeof(int) * width, height, + hipMemcpyHostToDevice)); + fill_value = 41; + std::fill_n(host_alloc.host_ptr(), width * height, fill_value); + + HIP_CHECK(memcpy_func(host_alloc.host_ptr(), sizeof(int) * width, array_alloc.ptr())); + if (should_synchronize) { + HIP_CHECK(hipStreamSynchronize(stream)); + } + ArrayFindIfNot(host_alloc.host_ptr(), fill_value, element_count); + } + SECTION("Device to Device") { + ArrayAllocGuard array_alloc(make_hipExtent(width, height, 0), flag); + LinearAllocGuard2D device_alloc(width, height); + LinearAllocGuard host_alloc(LinearAllocs::hipHostMalloc, width * height * sizeof(int)); + int fill_value = 42; + std::fill_n(host_alloc.host_ptr(), width * height, fill_value); + HIP_CHECK(hipMemcpy2DToArray(array_alloc.ptr(), 0, 0, host_alloc.host_ptr(), + sizeof(int) * width, sizeof(int) * width, height, + hipMemcpyHostToDevice)); + fill_value = 41; + std::fill_n(host_alloc.host_ptr(), width * height, fill_value); + HIP_CHECK(hipMemcpy2D(device_alloc.ptr(), device_alloc.pitch(), host_alloc.host_ptr(), + sizeof(int) * width, sizeof(int) * width, height, hipMemcpyHostToDevice)); + + HIP_CHECK(memcpy_func(device_alloc.ptr(), device_alloc.pitch(), array_alloc.ptr())); + if constexpr (should_synchronize) { + HIP_CHECK(hipStreamSynchronize(stream)); + } + HIP_CHECK(hipMemcpy2D(host_alloc.host_ptr(), sizeof(int) * width, device_alloc.ptr(), + device_alloc.pitch(), sizeof(int) * width, height, + hipMemcpyDeviceToHost)); + ArrayFindIfNot(host_alloc.host_ptr(), fill_value, element_count); + } +} + +/* Host/Device -> Array zero copy */ +template +void Memcpy2DToArrayZeroWidthHeight(F memcpy_func, size_t width, size_t height, + const hipStream_t stream = nullptr) { + const unsigned int flag = hipArrayDefault; + const auto element_count = width * height; + + SECTION("Host to Device") { + ArrayAllocGuard array_alloc(make_hipExtent(width, height, 0), flag); + LinearAllocGuard host_alloc(LinearAllocs::hipHostMalloc, width * height * sizeof(int)); + int fill_value = 42; + std::fill_n(host_alloc.host_ptr(), width * height, fill_value); + HIP_CHECK(hipMemcpy2DToArray(array_alloc.ptr(), 0, 0, host_alloc.host_ptr(), + sizeof(int) * width, sizeof(int) * width, height, + hipMemcpyHostToDevice)); + fill_value = 41; + std::fill_n(host_alloc.host_ptr(), width * height, fill_value); + + HIP_CHECK(memcpy_func(array_alloc.ptr(), host_alloc.host_ptr(), sizeof(int) * width)); + if (should_synchronize) { + HIP_CHECK(hipStreamSynchronize(stream)); + } + HIP_CHECK(hipMemcpy2DFromArray(host_alloc.host_ptr(), sizeof(int) * width, array_alloc.ptr(), 0, + 0, sizeof(int) * width, height, hipMemcpyDeviceToHost)); + ArrayFindIfNot(host_alloc.host_ptr(), 42, element_count); + } + SECTION("Device to Device") { + ArrayAllocGuard array_alloc(make_hipExtent(width, height, 0), flag); + LinearAllocGuard2D device_alloc(width, height); + LinearAllocGuard host_alloc(LinearAllocs::hipHostMalloc, width * height * sizeof(int)); + int fill_value = 42; + std::fill_n(host_alloc.host_ptr(), width * height, fill_value); + HIP_CHECK(hipMemcpy2DToArray(array_alloc.ptr(), 0, 0, host_alloc.host_ptr(), + sizeof(int) * width, sizeof(int) * width, height, + hipMemcpyHostToDevice)); + fill_value = 41; + std::fill_n(host_alloc.host_ptr(), width * height, fill_value); + HIP_CHECK(hipMemcpy2D(device_alloc.ptr(), device_alloc.pitch(), host_alloc.host_ptr(), + sizeof(int) * width, sizeof(int) * width, height, hipMemcpyHostToDevice)); + + HIP_CHECK(memcpy_func(array_alloc.ptr(), device_alloc.ptr(), device_alloc.pitch())); + if constexpr (should_synchronize) { + HIP_CHECK(hipStreamSynchronize(stream)); + } + HIP_CHECK(hipMemcpy2DFromArray(host_alloc.host_ptr(), sizeof(int) * width, array_alloc.ptr(), 0, + 0, sizeof(int) * width, height, hipMemcpyDeviceToHost)); + ArrayFindIfNot(host_alloc.host_ptr(), 42, element_count); + } +} diff --git a/projects/hip-tests/catch/unit/memory/hipMemcpy2DFromArrayAsync.cc b/projects/hip-tests/catch/unit/memory/hipMemcpy2DFromArrayAsync.cc index f7c8e75b71..5dcd4085c6 100644 --- a/projects/hip-tests/catch/unit/memory/hipMemcpy2DFromArrayAsync.cc +++ b/projects/hip-tests/catch/unit/memory/hipMemcpy2DFromArrayAsync.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,357 +16,270 @@ 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 file verifies the following scenarios of hipMemcpy2DFromArrayAsync API -1. Negative Scenarios -2. Extent Validation Scenarios -3. hipMemcpy2DFromArrayAsync Basic Scenario -4. Pinned Memory scenarios on same and peer GPU -5. Device Context change scenario where memory is allocated in - one GPU and stream is created in peer GPU. +Testcase Scenarios : +Unit_hipMemcpy2DFromArrayAsync_Positive_Default - Test basic async memcpy +between 2D array and host/device with hipMemcpy2DFromArrayAsync api +Unit_hipMemcpy2DFromArrayAsync_Positive_Synchronization_Behavior - Test +synchronization behavior for hipMemcpy2DFromArrayAsync api +Unit_hipMemcpy2DFromArrayAsync_Positive_ZeroWidthHeight - Test that no data is +copied when width/height is set to 0 +Unit_hipMemcpy2DFromArrayAsync_Negative_Parameters - Test unsuccessful execution +of hipMemcpy2DFromArrayAsync api when parameters are invalid */ +#include "array_memcpy_tests_common.hh" +#include #include -#include +#include +#include +TEST_CASE("Unit_hipMemcpy2DFromArrayAsync_Positive_Default") { + using namespace std::placeholders; -static constexpr auto NUM_W{10}; -static constexpr auto NUM_H{10}; + const auto stream_type = GENERATE(Streams::nullstream, Streams::perThread, Streams::created); + const StreamGuard stream_guard(stream_type); + const hipStream_t stream = stream_guard.stream(); -/* - * This testcase copies the data from host to device of - hipMemcpy2DFromArrayAsync API - * INPUT: Copying Host variable hData(Initialized with value Phi(1.618)) - * --> A_d device variable - * OUTPUT: For validating the result,Copying A_d device variable - * --> A_h host variable - * and verifying A_h with Phi - */ -TEST_CASE("Unit_hipMemcpy2DFromArrayAsync_Basic") { - HIP_CHECK(hipSetDevice(0)); - hipArray *A_d{nullptr}; - size_t width{sizeof(float)*NUM_W}; - float *A_h{nullptr}, *hData{nullptr}; - hipStream_t stream; + const auto width = GENERATE(16, 32, 48); + const auto height = GENERATE(1, 16, 32, 48); - // Initialization of variables - HipTest::initArrays(nullptr, nullptr, nullptr, - &A_h, &hData, nullptr, - width*NUM_H, false); - hipChannelFormatDesc desc = hipCreateChannelDesc(); - HIP_CHECK(hipMallocArray(&A_d, &desc, NUM_W, NUM_H, hipArrayDefault)); - HipTest::setDefaultData(width*NUM_H, A_h, hData, nullptr); - HIP_CHECK(hipStreamCreate(&stream)); - - HIP_CHECK(hipMemcpy2DToArray(A_d, 0, 0, hData, width, - width, NUM_H, - hipMemcpyHostToDevice)); - SECTION("Calling hipMemcpy2DFromArrayAsync() with user declared stream obj") { - HIP_CHECK(hipMemcpy2DFromArrayAsync(A_h, width, A_d, - 0, 0, width, NUM_H, - hipMemcpyDeviceToHost, stream)); - HIP_CHECK(hipStreamSynchronize(stream)); - } - SECTION("Calling hipMemcpy2DFromArrayAsync() with hipStreamPerThread") { - HIP_CHECK(hipMemcpy2DFromArrayAsync(A_h, width, A_d, - 0, 0, width, NUM_H, - hipMemcpyDeviceToHost, hipStreamPerThread)); - HIP_CHECK(hipStreamSynchronize(hipStreamPerThread)); - } - REQUIRE(HipTest::checkArray(A_h, hData, NUM_W, NUM_H) == true); - - // Cleaning the memory - HIP_CHECK(hipFreeArray(A_d)); - HIP_CHECK(hipStreamDestroy(stream)); - HipTest::freeArrays(nullptr, nullptr, nullptr, - A_h, hData, nullptr, false); -} - -/* - * This testcase verifies the extent validation scenarios - * of hipMemcpy2DFromArrayAsync API - */ -TEST_CASE("Unit_hipMemcpy2DFromArrayAsync_ExtentValidation") { - HIP_CHECK(hipSetDevice(0)); - hipArray *A_d{nullptr}; - size_t width{sizeof(float)*NUM_W}; - float *A_h{nullptr}, *hData{nullptr}, *valData{nullptr}; - hipStream_t stream; - - // Initialization of variables - HipTest::initArrays(nullptr, nullptr, nullptr, - &A_h, &hData, nullptr, - width*NUM_H, false); - HipTest::initArrays(nullptr, nullptr, nullptr, - nullptr, &valData, nullptr, - width*NUM_H, false); - hipChannelFormatDesc desc = hipCreateChannelDesc(); - HIP_CHECK(hipMallocArray(&A_d, &desc, NUM_W, NUM_H, hipArrayDefault)); - HIP_CHECK(hipStreamCreate(&stream)); - - SECTION("Destination width is 0") { - REQUIRE(hipMemcpy2DFromArrayAsync(A_h, 0, A_d, - 0, 0, NUM_W*sizeof(float), - NUM_H, hipMemcpyDeviceToHost, stream) - != hipSuccess); - } - // hipMemcpy2DFromArrayAsync API would return success for - // width and height as 0 - // and does not perform any copy - // Validating the result with the initialized value - // 1.Initializing A_d with Pi value - // 2.copying A_d-->hData variable - // with height 0(copy will not be performed) - // 3 validating hData<-->A_h which will not be equal as copy is not done. - SECTION("Height is 0") { - HIP_CHECK(hipMemcpy2DToArray(A_d, 0, 0, - A_h, width, width, - NUM_H, hipMemcpyHostToDevice)); - HIP_CHECK(hipMemcpy2DFromArrayAsync(hData, width, A_d, - 0, 0, NUM_W*sizeof(float), - 0, hipMemcpyDeviceToHost, stream)); - HIP_CHECK(hipStreamSynchronize(stream)); - REQUIRE(HipTest::checkArray(hData, valData, NUM_W, NUM_H) == true); - } - // hipMemcpy2DFromArrayAsync API would return success for - // width and height as 0 - // and does not perform any copy - // Validating the result with the initialized value - // 1.Initializing A_d with Pi value - // 2.copying A_d-->hData variable - // with width 0(copy will not be performed) - // 3 validating hData<-->A_h which will not be equal as copy is not done. - - SECTION("Width is 0") { - HIP_CHECK(hipMemcpy2DToArray(A_d, 0, 0, - A_h, width, width, - NUM_H, hipMemcpyHostToDevice)); - HIP_CHECK(hipMemcpy2DFromArrayAsync(hData, width, A_d, - 0, 0, 0, - NUM_H, hipMemcpyDeviceToHost, stream)); - HIP_CHECK(hipStreamSynchronize(stream)); - REQUIRE(HipTest::checkArray(hData, valData, NUM_W, NUM_H) == true); + SECTION("Array to host") { + Memcpy2DHostFromAShell( + std::bind(hipMemcpy2DFromArrayAsync, _1, _2, _3, 0, 0, width * sizeof(int), height, + hipMemcpyDeviceToHost, stream), + width, height, stream); } - // Cleaning the memory - HIP_CHECK(hipFreeArray(A_d)); - HIP_CHECK(hipStreamDestroy(stream)); - HipTest::freeArrays(nullptr, nullptr, nullptr, - A_h, hData, nullptr, false); - HipTest::freeArrays(nullptr, nullptr, nullptr, - nullptr, valData, nullptr, false); -} -/* - * This Scenario Verifies hipMemcpy2DFromArrayAsync API by copying the - * data from pinned host memory to device on same GPU - * INPUT: Copying Host variable PinnMem(Initialized with value "10" ) - * --> A_d device variable - * OUTPUT: For validating the result,Copying A_d device variable - * --> A_h host variable - * and verifying A_h with PinnedMem[0](i.e., 10) - */ -TEST_CASE("Unit_hipMemcpy2DFromArrayAsync_PinnedHostMemSameGpu") { - HIP_CHECK(hipSetDevice(0)); - hipArray *A_d{nullptr}; - constexpr auto def_val{10}; - size_t width{sizeof(float)*NUM_W}; - float *A_h{nullptr}, *PinnMem{nullptr}; - hipStream_t stream; - - // Initialization of variables - HipTest::initArrays(nullptr, nullptr, nullptr, - &A_h, nullptr, nullptr, - width*NUM_H, false); - HIP_CHECK(hipHostMalloc(reinterpret_cast(&PinnMem), width * NUM_H)); - hipChannelFormatDesc desc = hipCreateChannelDesc(); - HIP_CHECK(hipMallocArray(&A_d, &desc, NUM_W, NUM_H, hipArrayDefault)); - HipTest::setDefaultData(width*NUM_H, A_h, nullptr, nullptr); - for (int i = 0; i < NUM_W*NUM_H; i++) { - PinnMem[i] = def_val + i; + SECTION("Array to host with default kind") { + Memcpy2DHostFromAShell( + std::bind(hipMemcpy2DFromArrayAsync, _1, _2, _3, 0, 0, width * sizeof(int), height, + hipMemcpyDefault, stream), + width, height, stream); } - HIP_CHECK(hipStreamCreate(&stream)); - - HIP_CHECK(hipMemcpy2DToArray(A_d, 0, 0, PinnMem, - width, width, NUM_H, hipMemcpyHostToDevice)); - HIP_CHECK(hipMemcpy2DFromArrayAsync(A_h, width, A_d, - 0, 0, width, NUM_H, - hipMemcpyDeviceToHost, stream)); - HIP_CHECK(hipStreamSynchronize(stream)); - REQUIRE(HipTest::checkArray(A_h, PinnMem, NUM_W, NUM_H) == true); - - // Cleaning the memory - HIP_CHECK(hipFreeArray(A_d)); - HIP_CHECK(hipHostFree(PinnMem)); - HIP_CHECK(hipStreamDestroy(stream)); - HipTest::freeArrays(nullptr, nullptr, nullptr, - A_h, nullptr, nullptr, false); -} -/* - * This Scenario Verifies hipMemcpy2DFromArrayAsync API by copying the - * data from pinned host memory to device from Peer GPU. - * Device Memory is allocated in GPU 0 and the API is trigerred from GPU1 - * INPUT: Initialize data, A_h --> A_d device variable - * whose memory is allocated in GPU 0 - then A_d-->E_h in GPU1 - * OUTPUT: validating the result by comparing A_h and E_h - */ -TEST_CASE("Unit_hipMemcpy2DFromArrayAsync_multiDevicePinnedHostMem") { - int numDevices = 0; - constexpr auto def_val{10}; - HIP_CHECK(hipGetDeviceCount(&numDevices)); - if (numDevices > 1) { - int canAccessPeer = 0; - HIP_CHECK(hipDeviceCanAccessPeer(&canAccessPeer, 0, 1)); - if (canAccessPeer) { - HIP_CHECK(hipSetDevice(0)); - hipArray *A_d{nullptr}; - size_t width{sizeof(float)*NUM_W}; - float *A_h{nullptr}, *E_h{nullptr}; - hipStream_t stream; - HIP_CHECK(hipStreamCreate(&stream)); - - // Initialization of variables - HipTest::initArrays(nullptr, nullptr, nullptr, - &A_h, nullptr, nullptr, - width*NUM_H, false); - hipChannelFormatDesc desc = hipCreateChannelDesc(); - HIP_CHECK(hipMallocArray(&A_d, &desc, NUM_W, NUM_H, hipArrayDefault)); - HipTest::setDefaultData(width*NUM_H, A_h, nullptr, nullptr); - HIP_CHECK(hipHostMalloc(reinterpret_cast(&E_h), width * NUM_H)); - for (int i = 0; i < NUM_W*NUM_H; i++) { - E_h[i] = def_val + i; - } - - HIP_CHECK(hipMemcpy2DToArray(A_d, 0, 0, A_h, width, - width, NUM_H, hipMemcpyHostToDevice)); - HIP_CHECK(hipSetDevice(1)); - HIP_CHECK(hipMemcpy2DFromArrayAsync(E_h, width, A_d, - 0, 0, width, NUM_H, - hipMemcpyDeviceToHost, stream)); - HIP_CHECK(hipStreamSynchronize(stream)); - REQUIRE(HipTest::checkArray(A_h, E_h, NUM_W, NUM_H) == true); - - // Cleaning the memory - HIP_CHECK(hipFreeArray(A_d)); - HIP_CHECK(hipHostFree(E_h)); - HIP_CHECK(hipStreamDestroy(stream)); - HipTest::freeArrays(nullptr, nullptr, nullptr, - A_h, nullptr, nullptr, false); - } else { - SUCCEED("Device Does not have P2P capability"); +#if HT_NVIDIA // EXSWHTEC-213 + SECTION("Array to device") { + SECTION("Peer access disabled") { + Memcpy2DDeviceFromAShell( + std::bind(hipMemcpy2DFromArrayAsync, _1, _2, _3, 0, 0, width * sizeof(int), height, + hipMemcpyDeviceToDevice, stream), + width, height, stream); } - } else { - SUCCEED("Number of devices are < 2"); - } -} - -/* - * This scenario verifies the hipMemcpy2DFromArrayAsync API in case of device - * context change. - * Memory is allocated in GPU-0 and the API is triggered from GPU-1 - * INPUT: Copying Host variable hData(Initial value Phi) - * --> A_d device variable - * whose memory is allocated in GPU 0 - * OUTPUT: For validating the result,Copying A_d device variable - * --> A_h host variable - * and verifying A_h with Phi - * */ -TEST_CASE("Unit_hipMemcpy2DFromArrayAsync_multiDeviceContextChange") { - int numDevices = 0; - HIP_CHECK(hipGetDeviceCount(&numDevices)); - if (numDevices > 1) { - int canAccessPeer = 0; - HIP_CHECK(hipDeviceCanAccessPeer(&canAccessPeer, 0, 1)); - if (canAccessPeer) { - HIP_CHECK(hipSetDevice(0)); - hipArray *A_d{nullptr}; - size_t width{sizeof(float)*NUM_W}; - float *A_h{nullptr}, *hData{nullptr}; - hipStream_t stream; - - // Initialization of variables - HipTest::initArrays(nullptr, nullptr, nullptr, - &A_h, &hData, nullptr, - width*NUM_H, false); - hipChannelFormatDesc desc = hipCreateChannelDesc(); - HIP_CHECK(hipMallocArray(&A_d, &desc, NUM_W, NUM_H, hipArrayDefault)); - HipTest::setDefaultData(width*NUM_H, A_h, hData, nullptr); - - HIP_CHECK(hipSetDevice(1)); - HIP_CHECK(hipStreamCreate(&stream)); - HIP_CHECK(hipMemcpy2DToArray(A_d, 0, 0, hData, width, width, - NUM_H, hipMemcpyHostToDevice)); - - HIP_CHECK(hipMemcpy2DFromArrayAsync(A_h, width, A_d, - 0, 0, width, NUM_H, - hipMemcpyDeviceToHost, stream)); - HIP_CHECK(hipStreamSynchronize(stream)); - REQUIRE(HipTest::checkArray(A_h, hData, NUM_W, NUM_H) == true); - - // Cleaning the memory - HIP_CHECK(hipFreeArray(A_d)); - HIP_CHECK(hipStreamDestroy(stream)); - HipTest::freeArrays(nullptr, nullptr, nullptr, - A_h, hData, nullptr, false); - } else { - SUCCEED("Device Does not have P2P capability"); + SECTION("Peer access enabled") { + Memcpy2DDeviceFromAShell( + std::bind(hipMemcpy2DFromArrayAsync, _1, _2, _3, 0, 0, width * sizeof(int), height, + hipMemcpyDeviceToDevice, stream), + width, height, stream); } - } else { - SUCCEED("Number of devices are < 2"); - } -} -/* This testcase verifies the negative scenarios - * of hipMemcpy2DFromArrayAsync API - */ -TEST_CASE("Unit_hipMemcpy2DFromArrayAsync_Negative") { - HIP_CHECK(hipSetDevice(0)); - hipArray *A_d{nullptr}; - size_t width{sizeof(float)*NUM_W}; - float *A_h{nullptr}, *hData{nullptr}; - hipStream_t stream; - HIP_CHECK(hipStreamCreate(&stream)); - - // Initialization of variables - HipTest::initArrays(nullptr, nullptr, nullptr, - &A_h, &hData, nullptr, - width*NUM_H, false); - HipTest::setDefaultData(width*NUM_H, A_h, hData, nullptr); - hipChannelFormatDesc desc = hipCreateChannelDesc(); - HIP_CHECK(hipMallocArray(&A_d, &desc, NUM_W, NUM_H, hipArrayDefault)); - - SECTION("Nullptr to destination") { - REQUIRE(hipMemcpy2DFromArrayAsync(nullptr, width, A_d, - 0, 0, width, NUM_H, - hipMemcpyDeviceToHost, - stream) != hipSuccess); } - SECTION("Nullptr to source") { - REQUIRE(hipMemcpy2DFromArrayAsync(A_h, width, nullptr, - 0, 0, width, NUM_H, - hipMemcpyDeviceToHost, - stream) != hipSuccess); + SECTION("Array to device with default kind") { + SECTION("Peer access disabled") { + Memcpy2DDeviceFromAShell( + std::bind(hipMemcpy2DFromArrayAsync, _1, _2, _3, 0, 0, width * sizeof(int), height, + hipMemcpyDefault, stream), + width, height, stream); + } + SECTION("Peer access enabled") { + Memcpy2DDeviceFromAShell( + std::bind(hipMemcpy2DFromArrayAsync, _1, _2, _3, 0, 0, width * sizeof(int), height, + hipMemcpyDefault, stream), + width, height, stream); + } } - - SECTION("Passing offset more than 0") { - REQUIRE(hipMemcpy2DFromArrayAsync(A_h, width, A_d, 1, - 1, width, NUM_H, - hipMemcpyDeviceToHost, - stream) != hipSuccess); - } - - SECTION("Passing array more than allocated") { - REQUIRE(hipMemcpy2DFromArrayAsync(A_h, width, A_d, 0, - 0, width+2, NUM_H+2, - hipMemcpyDeviceToHost, - stream) != hipSuccess); - } - - // Cleaning of Memory - HIP_CHECK(hipFreeArray(A_d)); - HIP_CHECK(hipStreamDestroy(stream)); - HipTest::freeArrays(nullptr, nullptr, nullptr, - A_h, hData, nullptr, false); +#endif } +TEST_CASE("Unit_hipMemcpy2DFromArrayAsync_Positive_Synchronization_Behavior") { + using namespace std::placeholders; + HIP_CHECK(hipDeviceSynchronize()); + + SECTION("Array to host") { + const auto width = GENERATE(16, 32, 48); + const auto height = GENERATE(16, 32, 48); + + MemcpyAtoHPageableSyncBehavior( + std::bind(hipMemcpy2DFromArrayAsync, _1, width * sizeof(int), _2, 0, 0, width * sizeof(int), + height, hipMemcpyDeviceToHost, nullptr), + width, height, false); + MemcpyAtoHPinnedSyncBehavior( + std::bind(hipMemcpy2DFromArrayAsync, _1, width * sizeof(int), _2, 0, 0, width * sizeof(int), + height, hipMemcpyDeviceToHost, nullptr), + width, height, false); + } + + SECTION("Array to device") { + const auto width = GENERATE(16, 32, 48); + const auto height = GENERATE(16, 32, 48); + + MemcpyAtoDSyncBehavior(std::bind(hipMemcpy2DFromArrayAsync, _1, _2, _3, 0, 0, + width * sizeof(int), height, hipMemcpyDeviceToDevice, nullptr), + width, height, false); + } +} + +TEST_CASE("Unit_hipMemcpy2DFromArrayAsync_Positive_ZeroWidthHeight") { + using namespace std::placeholders; + + const auto stream_type = GENERATE(Streams::nullstream, Streams::perThread, Streams::created); + const StreamGuard stream_guard(stream_type); + const hipStream_t stream = stream_guard.stream(); + + const auto width = 16; + const auto height = 16; + + SECTION("Array to host") { + SECTION("Height is 0") { + Memcpy2DFromArrayZeroWidthHeight( + std::bind(hipMemcpy2DFromArrayAsync, _1, _2, _3, 0, 0, width * sizeof(int), 0, + hipMemcpyDeviceToHost, stream), + width, height, stream); + } + SECTION("Width is 0") { + Memcpy2DFromArrayZeroWidthHeight(std::bind(hipMemcpy2DFromArrayAsync, _1, _2, _3, 0, 0, + 0, height, hipMemcpyDeviceToHost, stream), + width, height, stream); + } + } + SECTION("Array to device") { + SECTION("Height is 0") { + Memcpy2DFromArrayZeroWidthHeight( + std::bind(hipMemcpy2DFromArrayAsync, _1, _2, _3, 0, 0, width * sizeof(int), 0, + hipMemcpyDeviceToDevice, stream), + width, height, stream); + } + SECTION("Width is 0") { + Memcpy2DFromArrayZeroWidthHeight(std::bind(hipMemcpy2DFromArrayAsync, _1, _2, _3, 0, 0, + 0, height, hipMemcpyDeviceToDevice, stream), + width, height, stream); + } + } +} + +TEST_CASE("Unit_hipMemcpy2DFromArrayAsync_Negative_Parameters") { + using namespace std::placeholders; + + const auto width = 32; + const auto height = 32; + const auto allocation_size = 2 * width * height * sizeof(int); + + const unsigned int flag = hipArrayDefault; + + constexpr auto InvalidStream = [] { + StreamGuard sg(Streams::created); + return sg.stream(); + }; + + ArrayAllocGuard array_alloc(make_hipExtent(width, height, 0), flag); + LinearAllocGuard2D device_alloc(width, height); + LinearAllocGuard host_alloc(LinearAllocs::hipHostMalloc, allocation_size); + + SECTION("Array to host") { + SECTION("dst == nullptr") { + HIP_CHECK_ERROR( + hipMemcpy2DFromArrayAsync(nullptr, 2 * width * sizeof(int), array_alloc.ptr(), 0, 0, + width * sizeof(int), height, hipMemcpyDeviceToHost, nullptr), + hipErrorInvalidValue); + } + SECTION("src == nullptr") { + HIP_CHECK_ERROR( + hipMemcpy2DFromArrayAsync(host_alloc.ptr(), 2 * width * sizeof(int), nullptr, 0, 0, + width * sizeof(int), height, hipMemcpyDeviceToHost, nullptr), + hipErrorInvalidHandle); + } +#if HT_NVIDIA // EXSWHTEC-212 + SECTION("dpitch < width") { + HIP_CHECK_ERROR(hipMemcpy2DFromArrayAsync(host_alloc.ptr(), width * sizeof(int) - 10, + array_alloc.ptr(), 0, 0, width * sizeof(int), + height, hipMemcpyDeviceToHost, nullptr), + hipErrorInvalidPitchValue); + } + SECTION("Offset + width/height overflows") { + HIP_CHECK_ERROR( + hipMemcpy2DFromArrayAsync(host_alloc.ptr(), 2 * width * sizeof(int), array_alloc.ptr(), 1, + 0, width * sizeof(int), height, hipMemcpyDeviceToHost, nullptr), + hipErrorInvalidValue); + HIP_CHECK_ERROR( + hipMemcpy2DFromArrayAsync(host_alloc.ptr(), 2 * width * sizeof(int), array_alloc.ptr(), 0, + 1, width * sizeof(int), height, hipMemcpyDeviceToHost, nullptr), + hipErrorInvalidValue); + } + SECTION("Width/height overflows") { + HIP_CHECK_ERROR(hipMemcpy2DFromArrayAsync(host_alloc.ptr(), 2 * width * sizeof(int), + array_alloc.ptr(), 0, 0, width * sizeof(int) + 1, + height, hipMemcpyDeviceToHost, nullptr), + hipErrorInvalidValue); + HIP_CHECK_ERROR(hipMemcpy2DFromArrayAsync(host_alloc.ptr(), 2 * width * sizeof(int), + array_alloc.ptr(), 0, 0, width * sizeof(int), + height + 1, hipMemcpyDeviceToHost, nullptr), + hipErrorInvalidValue); + } + SECTION("Memcpy kind is invalid") { + HIP_CHECK_ERROR(hipMemcpy2DFromArrayAsync(host_alloc.ptr(), 2 * width * sizeof(int), + array_alloc.ptr(), 0, 0, width * sizeof(int), + height, static_cast(-1), nullptr), + hipErrorInvalidMemcpyDirection); + } + SECTION("Invalid stream") { + HIP_CHECK_ERROR(hipMemcpy2DFromArrayAsync(host_alloc.ptr(), 2 * width * sizeof(int), + array_alloc.ptr(), 0, 0, width * sizeof(int), + height, hipMemcpyDeviceToHost, InvalidStream()), + hipErrorContextIsDestroyed); + } +#endif + } + SECTION("Array to device") { + SECTION("dst == nullptr") { + HIP_CHECK_ERROR( + hipMemcpy2DFromArrayAsync(nullptr, device_alloc.pitch(), array_alloc.ptr(), 0, 0, + width * sizeof(int), height, hipMemcpyDeviceToDevice, nullptr), + hipErrorInvalidValue); + } + SECTION("src == nullptr") { + HIP_CHECK_ERROR( + hipMemcpy2DFromArrayAsync(device_alloc.ptr(), device_alloc.pitch(), nullptr, 0, 0, + width * sizeof(int), height, hipMemcpyDeviceToDevice, nullptr), + hipErrorInvalidHandle); + } +#if HT_NVIDIA // EXSWHTEC-212 + SECTION("dpitch < width") { + HIP_CHECK_ERROR(hipMemcpy2DFromArrayAsync(device_alloc.ptr(), width * sizeof(int) - 10, + array_alloc.ptr(), 0, 0, width * sizeof(int), + height, hipMemcpyDeviceToDevice, nullptr), + hipErrorInvalidPitchValue); + } + SECTION("Offset + width/height overflows") { + HIP_CHECK_ERROR(hipMemcpy2DFromArrayAsync(device_alloc.ptr(), device_alloc.pitch(), + array_alloc.ptr(), 1, 0, width * sizeof(int), + height, hipMemcpyDeviceToDevice, nullptr), + hipErrorInvalidValue); + HIP_CHECK_ERROR(hipMemcpy2DFromArrayAsync(device_alloc.ptr(), device_alloc.pitch(), + array_alloc.ptr(), 0, 1, width * sizeof(int), + height, hipMemcpyDeviceToDevice, nullptr), + hipErrorInvalidValue); + } + SECTION("Width/height overflows") { + HIP_CHECK_ERROR(hipMemcpy2DFromArrayAsync(device_alloc.ptr(), device_alloc.pitch(), + array_alloc.ptr(), 0, 0, width * sizeof(int) + 1, + height, hipMemcpyDeviceToDevice, nullptr), + hipErrorInvalidValue); + HIP_CHECK_ERROR(hipMemcpy2DFromArrayAsync(device_alloc.ptr(), device_alloc.pitch(), + array_alloc.ptr(), 0, 0, width * sizeof(int), + height + 1, hipMemcpyDeviceToDevice, nullptr), + hipErrorInvalidValue); + } + SECTION("Memcpy kind is invalid") { + HIP_CHECK_ERROR(hipMemcpy2DFromArrayAsync(device_alloc.ptr(), device_alloc.pitch(), + array_alloc.ptr(), 0, 0, width * sizeof(int), + height, static_cast(-1), nullptr), + hipErrorInvalidMemcpyDirection); + } + SECTION("Invalid stream") { + HIP_CHECK_ERROR(hipMemcpy2DFromArrayAsync(device_alloc.ptr(), device_alloc.pitch(), + array_alloc.ptr(), 0, 0, width * sizeof(int), + height, hipMemcpyDeviceToDevice, InvalidStream()), + hipErrorContextIsDestroyed); + } +#endif + } +} diff --git a/projects/hip-tests/catch/unit/memory/hipMemcpy2DFromArrayAsync_old.cc b/projects/hip-tests/catch/unit/memory/hipMemcpy2DFromArrayAsync_old.cc new file mode 100644 index 0000000000..f7c8e75b71 --- /dev/null +++ b/projects/hip-tests/catch/unit/memory/hipMemcpy2DFromArrayAsync_old.cc @@ -0,0 +1,372 @@ +/* +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 file verifies the following scenarios of hipMemcpy2DFromArrayAsync API +1. Negative Scenarios +2. Extent Validation Scenarios +3. hipMemcpy2DFromArrayAsync Basic Scenario +4. Pinned Memory scenarios on same and peer GPU +5. Device Context change scenario where memory is allocated in + one GPU and stream is created in peer GPU. +*/ + +#include +#include + + +static constexpr auto NUM_W{10}; +static constexpr auto NUM_H{10}; + +/* + * This testcase copies the data from host to device of + hipMemcpy2DFromArrayAsync API + * INPUT: Copying Host variable hData(Initialized with value Phi(1.618)) + * --> A_d device variable + * OUTPUT: For validating the result,Copying A_d device variable + * --> A_h host variable + * and verifying A_h with Phi + */ +TEST_CASE("Unit_hipMemcpy2DFromArrayAsync_Basic") { + HIP_CHECK(hipSetDevice(0)); + hipArray *A_d{nullptr}; + size_t width{sizeof(float)*NUM_W}; + float *A_h{nullptr}, *hData{nullptr}; + hipStream_t stream; + + // Initialization of variables + HipTest::initArrays(nullptr, nullptr, nullptr, + &A_h, &hData, nullptr, + width*NUM_H, false); + hipChannelFormatDesc desc = hipCreateChannelDesc(); + HIP_CHECK(hipMallocArray(&A_d, &desc, NUM_W, NUM_H, hipArrayDefault)); + HipTest::setDefaultData(width*NUM_H, A_h, hData, nullptr); + HIP_CHECK(hipStreamCreate(&stream)); + + HIP_CHECK(hipMemcpy2DToArray(A_d, 0, 0, hData, width, + width, NUM_H, + hipMemcpyHostToDevice)); + SECTION("Calling hipMemcpy2DFromArrayAsync() with user declared stream obj") { + HIP_CHECK(hipMemcpy2DFromArrayAsync(A_h, width, A_d, + 0, 0, width, NUM_H, + hipMemcpyDeviceToHost, stream)); + HIP_CHECK(hipStreamSynchronize(stream)); + } + SECTION("Calling hipMemcpy2DFromArrayAsync() with hipStreamPerThread") { + HIP_CHECK(hipMemcpy2DFromArrayAsync(A_h, width, A_d, + 0, 0, width, NUM_H, + hipMemcpyDeviceToHost, hipStreamPerThread)); + HIP_CHECK(hipStreamSynchronize(hipStreamPerThread)); + } + REQUIRE(HipTest::checkArray(A_h, hData, NUM_W, NUM_H) == true); + + // Cleaning the memory + HIP_CHECK(hipFreeArray(A_d)); + HIP_CHECK(hipStreamDestroy(stream)); + HipTest::freeArrays(nullptr, nullptr, nullptr, + A_h, hData, nullptr, false); +} + +/* + * This testcase verifies the extent validation scenarios + * of hipMemcpy2DFromArrayAsync API + */ +TEST_CASE("Unit_hipMemcpy2DFromArrayAsync_ExtentValidation") { + HIP_CHECK(hipSetDevice(0)); + hipArray *A_d{nullptr}; + size_t width{sizeof(float)*NUM_W}; + float *A_h{nullptr}, *hData{nullptr}, *valData{nullptr}; + hipStream_t stream; + + // Initialization of variables + HipTest::initArrays(nullptr, nullptr, nullptr, + &A_h, &hData, nullptr, + width*NUM_H, false); + HipTest::initArrays(nullptr, nullptr, nullptr, + nullptr, &valData, nullptr, + width*NUM_H, false); + hipChannelFormatDesc desc = hipCreateChannelDesc(); + HIP_CHECK(hipMallocArray(&A_d, &desc, NUM_W, NUM_H, hipArrayDefault)); + HIP_CHECK(hipStreamCreate(&stream)); + + SECTION("Destination width is 0") { + REQUIRE(hipMemcpy2DFromArrayAsync(A_h, 0, A_d, + 0, 0, NUM_W*sizeof(float), + NUM_H, hipMemcpyDeviceToHost, stream) + != hipSuccess); + } + // hipMemcpy2DFromArrayAsync API would return success for + // width and height as 0 + // and does not perform any copy + // Validating the result with the initialized value + // 1.Initializing A_d with Pi value + // 2.copying A_d-->hData variable + // with height 0(copy will not be performed) + // 3 validating hData<-->A_h which will not be equal as copy is not done. + SECTION("Height is 0") { + HIP_CHECK(hipMemcpy2DToArray(A_d, 0, 0, + A_h, width, width, + NUM_H, hipMemcpyHostToDevice)); + HIP_CHECK(hipMemcpy2DFromArrayAsync(hData, width, A_d, + 0, 0, NUM_W*sizeof(float), + 0, hipMemcpyDeviceToHost, stream)); + HIP_CHECK(hipStreamSynchronize(stream)); + REQUIRE(HipTest::checkArray(hData, valData, NUM_W, NUM_H) == true); + } + // hipMemcpy2DFromArrayAsync API would return success for + // width and height as 0 + // and does not perform any copy + // Validating the result with the initialized value + // 1.Initializing A_d with Pi value + // 2.copying A_d-->hData variable + // with width 0(copy will not be performed) + // 3 validating hData<-->A_h which will not be equal as copy is not done. + + SECTION("Width is 0") { + HIP_CHECK(hipMemcpy2DToArray(A_d, 0, 0, + A_h, width, width, + NUM_H, hipMemcpyHostToDevice)); + HIP_CHECK(hipMemcpy2DFromArrayAsync(hData, width, A_d, + 0, 0, 0, + NUM_H, hipMemcpyDeviceToHost, stream)); + HIP_CHECK(hipStreamSynchronize(stream)); + REQUIRE(HipTest::checkArray(hData, valData, NUM_W, NUM_H) == true); + } + + // Cleaning the memory + HIP_CHECK(hipFreeArray(A_d)); + HIP_CHECK(hipStreamDestroy(stream)); + HipTest::freeArrays(nullptr, nullptr, nullptr, + A_h, hData, nullptr, false); + HipTest::freeArrays(nullptr, nullptr, nullptr, + nullptr, valData, nullptr, false); +} +/* + * This Scenario Verifies hipMemcpy2DFromArrayAsync API by copying the + * data from pinned host memory to device on same GPU + * INPUT: Copying Host variable PinnMem(Initialized with value "10" ) + * --> A_d device variable + * OUTPUT: For validating the result,Copying A_d device variable + * --> A_h host variable + * and verifying A_h with PinnedMem[0](i.e., 10) + */ +TEST_CASE("Unit_hipMemcpy2DFromArrayAsync_PinnedHostMemSameGpu") { + HIP_CHECK(hipSetDevice(0)); + hipArray *A_d{nullptr}; + constexpr auto def_val{10}; + size_t width{sizeof(float)*NUM_W}; + float *A_h{nullptr}, *PinnMem{nullptr}; + hipStream_t stream; + + // Initialization of variables + HipTest::initArrays(nullptr, nullptr, nullptr, + &A_h, nullptr, nullptr, + width*NUM_H, false); + HIP_CHECK(hipHostMalloc(reinterpret_cast(&PinnMem), width * NUM_H)); + hipChannelFormatDesc desc = hipCreateChannelDesc(); + HIP_CHECK(hipMallocArray(&A_d, &desc, NUM_W, NUM_H, hipArrayDefault)); + HipTest::setDefaultData(width*NUM_H, A_h, nullptr, nullptr); + for (int i = 0; i < NUM_W*NUM_H; i++) { + PinnMem[i] = def_val + i; + } + HIP_CHECK(hipStreamCreate(&stream)); + + HIP_CHECK(hipMemcpy2DToArray(A_d, 0, 0, PinnMem, + width, width, NUM_H, hipMemcpyHostToDevice)); + HIP_CHECK(hipMemcpy2DFromArrayAsync(A_h, width, A_d, + 0, 0, width, NUM_H, + hipMemcpyDeviceToHost, stream)); + HIP_CHECK(hipStreamSynchronize(stream)); + REQUIRE(HipTest::checkArray(A_h, PinnMem, NUM_W, NUM_H) == true); + + // Cleaning the memory + HIP_CHECK(hipFreeArray(A_d)); + HIP_CHECK(hipHostFree(PinnMem)); + HIP_CHECK(hipStreamDestroy(stream)); + HipTest::freeArrays(nullptr, nullptr, nullptr, + A_h, nullptr, nullptr, false); +} +/* + * This Scenario Verifies hipMemcpy2DFromArrayAsync API by copying the + * data from pinned host memory to device from Peer GPU. + * Device Memory is allocated in GPU 0 and the API is trigerred from GPU1 + * INPUT: Initialize data, A_h --> A_d device variable + * whose memory is allocated in GPU 0 + then A_d-->E_h in GPU1 + * OUTPUT: validating the result by comparing A_h and E_h + */ +TEST_CASE("Unit_hipMemcpy2DFromArrayAsync_multiDevicePinnedHostMem") { + int numDevices = 0; + constexpr auto def_val{10}; + HIP_CHECK(hipGetDeviceCount(&numDevices)); + if (numDevices > 1) { + int canAccessPeer = 0; + HIP_CHECK(hipDeviceCanAccessPeer(&canAccessPeer, 0, 1)); + if (canAccessPeer) { + HIP_CHECK(hipSetDevice(0)); + hipArray *A_d{nullptr}; + size_t width{sizeof(float)*NUM_W}; + float *A_h{nullptr}, *E_h{nullptr}; + hipStream_t stream; + HIP_CHECK(hipStreamCreate(&stream)); + + // Initialization of variables + HipTest::initArrays(nullptr, nullptr, nullptr, + &A_h, nullptr, nullptr, + width*NUM_H, false); + hipChannelFormatDesc desc = hipCreateChannelDesc(); + HIP_CHECK(hipMallocArray(&A_d, &desc, NUM_W, NUM_H, hipArrayDefault)); + HipTest::setDefaultData(width*NUM_H, A_h, nullptr, nullptr); + HIP_CHECK(hipHostMalloc(reinterpret_cast(&E_h), width * NUM_H)); + for (int i = 0; i < NUM_W*NUM_H; i++) { + E_h[i] = def_val + i; + } + + HIP_CHECK(hipMemcpy2DToArray(A_d, 0, 0, A_h, width, + width, NUM_H, hipMemcpyHostToDevice)); + HIP_CHECK(hipSetDevice(1)); + HIP_CHECK(hipMemcpy2DFromArrayAsync(E_h, width, A_d, + 0, 0, width, NUM_H, + hipMemcpyDeviceToHost, stream)); + HIP_CHECK(hipStreamSynchronize(stream)); + REQUIRE(HipTest::checkArray(A_h, E_h, NUM_W, NUM_H) == true); + + // Cleaning the memory + HIP_CHECK(hipFreeArray(A_d)); + HIP_CHECK(hipHostFree(E_h)); + HIP_CHECK(hipStreamDestroy(stream)); + HipTest::freeArrays(nullptr, nullptr, nullptr, + A_h, nullptr, nullptr, false); + } else { + SUCCEED("Device Does not have P2P capability"); + } + } else { + SUCCEED("Number of devices are < 2"); + } +} + +/* + * This scenario verifies the hipMemcpy2DFromArrayAsync API in case of device + * context change. + * Memory is allocated in GPU-0 and the API is triggered from GPU-1 + * INPUT: Copying Host variable hData(Initial value Phi) + * --> A_d device variable + * whose memory is allocated in GPU 0 + * OUTPUT: For validating the result,Copying A_d device variable + * --> A_h host variable + * and verifying A_h with Phi + * */ +TEST_CASE("Unit_hipMemcpy2DFromArrayAsync_multiDeviceContextChange") { + int numDevices = 0; + HIP_CHECK(hipGetDeviceCount(&numDevices)); + if (numDevices > 1) { + int canAccessPeer = 0; + HIP_CHECK(hipDeviceCanAccessPeer(&canAccessPeer, 0, 1)); + if (canAccessPeer) { + HIP_CHECK(hipSetDevice(0)); + hipArray *A_d{nullptr}; + size_t width{sizeof(float)*NUM_W}; + float *A_h{nullptr}, *hData{nullptr}; + hipStream_t stream; + + // Initialization of variables + HipTest::initArrays(nullptr, nullptr, nullptr, + &A_h, &hData, nullptr, + width*NUM_H, false); + hipChannelFormatDesc desc = hipCreateChannelDesc(); + HIP_CHECK(hipMallocArray(&A_d, &desc, NUM_W, NUM_H, hipArrayDefault)); + HipTest::setDefaultData(width*NUM_H, A_h, hData, nullptr); + + HIP_CHECK(hipSetDevice(1)); + HIP_CHECK(hipStreamCreate(&stream)); + HIP_CHECK(hipMemcpy2DToArray(A_d, 0, 0, hData, width, width, + NUM_H, hipMemcpyHostToDevice)); + + HIP_CHECK(hipMemcpy2DFromArrayAsync(A_h, width, A_d, + 0, 0, width, NUM_H, + hipMemcpyDeviceToHost, stream)); + HIP_CHECK(hipStreamSynchronize(stream)); + REQUIRE(HipTest::checkArray(A_h, hData, NUM_W, NUM_H) == true); + + // Cleaning the memory + HIP_CHECK(hipFreeArray(A_d)); + HIP_CHECK(hipStreamDestroy(stream)); + HipTest::freeArrays(nullptr, nullptr, nullptr, + A_h, hData, nullptr, false); + } else { + SUCCEED("Device Does not have P2P capability"); + } + } else { + SUCCEED("Number of devices are < 2"); + } +} +/* This testcase verifies the negative scenarios + * of hipMemcpy2DFromArrayAsync API + */ +TEST_CASE("Unit_hipMemcpy2DFromArrayAsync_Negative") { + HIP_CHECK(hipSetDevice(0)); + hipArray *A_d{nullptr}; + size_t width{sizeof(float)*NUM_W}; + float *A_h{nullptr}, *hData{nullptr}; + hipStream_t stream; + HIP_CHECK(hipStreamCreate(&stream)); + + // Initialization of variables + HipTest::initArrays(nullptr, nullptr, nullptr, + &A_h, &hData, nullptr, + width*NUM_H, false); + HipTest::setDefaultData(width*NUM_H, A_h, hData, nullptr); + hipChannelFormatDesc desc = hipCreateChannelDesc(); + HIP_CHECK(hipMallocArray(&A_d, &desc, NUM_W, NUM_H, hipArrayDefault)); + + SECTION("Nullptr to destination") { + REQUIRE(hipMemcpy2DFromArrayAsync(nullptr, width, A_d, + 0, 0, width, NUM_H, + hipMemcpyDeviceToHost, + stream) != hipSuccess); + } + + SECTION("Nullptr to source") { + REQUIRE(hipMemcpy2DFromArrayAsync(A_h, width, nullptr, + 0, 0, width, NUM_H, + hipMemcpyDeviceToHost, + stream) != hipSuccess); + } + + SECTION("Passing offset more than 0") { + REQUIRE(hipMemcpy2DFromArrayAsync(A_h, width, A_d, 1, + 1, width, NUM_H, + hipMemcpyDeviceToHost, + stream) != hipSuccess); + } + + SECTION("Passing array more than allocated") { + REQUIRE(hipMemcpy2DFromArrayAsync(A_h, width, A_d, 0, + 0, width+2, NUM_H+2, + hipMemcpyDeviceToHost, + stream) != hipSuccess); + } + + // Cleaning of Memory + HIP_CHECK(hipFreeArray(A_d)); + HIP_CHECK(hipStreamDestroy(stream)); + HipTest::freeArrays(nullptr, nullptr, nullptr, + A_h, hData, nullptr, false); +} + diff --git a/projects/hip-tests/catch/unit/memory/hipMemcpy2DToArrayAsync.cc b/projects/hip-tests/catch/unit/memory/hipMemcpy2DToArrayAsync.cc index 990b9c651e..2cf02a3e01 100644 --- a/projects/hip-tests/catch/unit/memory/hipMemcpy2DToArrayAsync.cc +++ b/projects/hip-tests/catch/unit/memory/hipMemcpy2DToArrayAsync.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,354 +16,265 @@ 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 file verifies the following scenarios of hipMemcpy2DToArrayAsync API -1. Negative Scenarios -2. Extent Validation Scenarios -3. hipMemcpy2DToArrayAsync Basic Scenario -4. Pinned Memory scenarios on same and peer GPU -5. Device Context change scenario where memory is allocated in - one GPU and stream is created in peer GPU. +Testcase Scenarios : +Unit_hipMemcpy2DToArrayAsync_Positive_Default - Test basic async memcpy between +host/device and 2D array with hipMemcpy2DToArrayAsync api +Unit_hipMemcpy2DToArrayAsync_Positive_Synchronization_Behavior - Test +synchronization behavior for hipMemcpy2DToArrayAsync api +Unit_hipMemcpy2DToArrayAsync_Positive_ZeroWidthHeight - Test that no data is +copied when width/height is set to 0 +Unit_hipMemcpy2DToArrayAsync_Negative_Parameters - Test unsuccessful execution +of hipMemcpy2DToArrayAsync api when parameters are invalid */ +#include "array_memcpy_tests_common.hh" +#include #include -#include -#include +#include +#include -static constexpr auto NUM_W{10}; -static constexpr auto NUM_H{10}; +TEST_CASE("Unit_hipMemcpy2DToArrayAsync_Positive_Default") { + using namespace std::placeholders; -/* - * This Scenario copies the data from host to device - * INPUT: Copying Host variable hData(Initialized with value Phi(1.618)) - * --> A_d device variable - * OUTPUT: For validating the result,Copying A_d device variable - * --> A_h host variable - * and verifying A_h with Phi - */ -TEST_CASE("Unit_hipMemcpy2DToArrayAsync_Basic") { - HIP_CHECK(hipSetDevice(0)); - hipArray *A_d{nullptr}; - size_t width{sizeof(float)*NUM_W}; - float *A_h{nullptr}, *hData{nullptr}; - hipStream_t stream; + const auto stream_type = GENERATE(Streams::nullstream, Streams::perThread, Streams::created); + const StreamGuard stream_guard(stream_type); + const hipStream_t stream = stream_guard.stream(); - // Initialization of variables - HipTest::initArrays(nullptr, nullptr, nullptr, - &A_h, &hData, nullptr, - width*NUM_H, false); - hipChannelFormatDesc desc = hipCreateChannelDesc(); - HIP_CHECK(hipMallocArray(&A_d, &desc, NUM_W, NUM_H, hipArrayDefault)); - HipTest::setDefaultData(width*NUM_H, A_h, hData, nullptr); - HIP_CHECK(hipStreamCreate(&stream)); - SECTION("Calling hipMemcpy2DToArrayAsync() with user declared stream obj") { - HIP_CHECK(hipMemcpy2DToArrayAsync(A_d, 0, 0, hData, width, - width, NUM_H, - hipMemcpyHostToDevice, stream)); - HIP_CHECK(hipStreamSynchronize(stream)); - } - SECTION("Calling hipMemcpy2DToArrayAsync() with hipStreamPerThread") { - HIP_CHECK(hipMemcpy2DToArrayAsync(A_d, 0, 0, hData, width, - width, NUM_H, - hipMemcpyHostToDevice, hipStreamPerThread)); - HIP_CHECK(hipStreamSynchronize(hipStreamPerThread)); - } - HIP_CHECK(hipMemcpy2DFromArray(A_h, width, A_d, - 0, 0, width, NUM_H, - hipMemcpyDeviceToHost)); - REQUIRE(HipTest::checkArray(A_h, hData, NUM_W, NUM_H) == true); + const auto width = GENERATE(16, 32, 48); + const auto height = GENERATE(1, 16, 32, 48); - // Cleaning the memory - HIP_CHECK(hipFreeArray(A_d)); - HIP_CHECK(hipStreamDestroy(stream)); - HipTest::freeArrays(nullptr, nullptr, nullptr, - A_h, hData, nullptr, false); -} - -/* - * This testcase verifies the extent validation scenarios - */ -TEST_CASE("Unit_hipMemcpy2DToArrayAsync_ExtentValidation") { - HIP_CHECK(hipSetDevice(0)); - hipArray *A_d{nullptr}; - size_t width{sizeof(float)*NUM_W}; - float *A_h{nullptr}, *hData{nullptr}; - hipStream_t stream; - - // Initialization of variables - HipTest::initArrays(nullptr, nullptr, nullptr, - &A_h, &hData, nullptr, - width*NUM_H, false); - hipChannelFormatDesc desc = hipCreateChannelDesc(); - HIP_CHECK(hipMallocArray(&A_d, &desc, NUM_W, NUM_H, hipArrayDefault)); - HIP_CHECK(hipStreamCreate(&stream)); - - SECTION("Source width is 0") { - REQUIRE(hipMemcpy2DToArrayAsync(A_d, 0, 0, hData, 0, - width, NUM_H, hipMemcpyHostToDevice, - stream) != hipSuccess); - HIP_CHECK(hipStreamSynchronize(stream)); - } - // hipMemcpy2DToArray API would return success for width and height as 0 - // and does not perform any copy - // Validating the result with the initialized value - // 1.Initializing A_d with Pi value - // 2.copying hData(Phi)-->A_d device variable - // with height 0(copy will not be performed) - // 3.copying A_d-->hData and validating it with A_h data - SECTION("Height is 0") { - HIP_CHECK(hipMemcpy2DToArrayAsync(A_d, 0, 0, A_h, width, - width, NUM_H, hipMemcpyHostToDevice, - stream)); - HIP_CHECK(hipMemcpy2DToArrayAsync(A_d, 0, 0, hData, width, - width, 0, hipMemcpyHostToDevice, - stream)); - HIP_CHECK(hipStreamSynchronize(stream)); - HIP_CHECK(hipMemcpy2DFromArray(hData, width, A_d, - 0, 0, width, NUM_H, - hipMemcpyDeviceToHost)); - REQUIRE(HipTest::checkArray(hData, A_h, NUM_W, NUM_H) == true); - } - // hipMemcpy2DToArray API would return success for width and height as 0 - // and does not perform any copy - // Validating the result with the initialized value - // 1.Initializing A_d with Pi value - // 2.copying hData(Phi)-->A_d device variable - // with width 0(copy will not be performed) - // 3.copying A_d-->hData and validating it with A_h data - SECTION("Width is 0") { - HIP_CHECK(hipMemcpy2DToArrayAsync(A_d, 0, 0, A_h, width, - width, NUM_H, hipMemcpyHostToDevice, - stream)); - HIP_CHECK(hipMemcpy2DToArrayAsync(A_d, 0, 0, hData, width, - 0, NUM_H, hipMemcpyHostToDevice, - stream)); - HIP_CHECK(hipStreamSynchronize(stream)); - HIP_CHECK(hipMemcpy2DFromArray(hData, width, A_d, - 0, 0, width, NUM_H, - hipMemcpyDeviceToHost)); - REQUIRE(HipTest::checkArray(hData, A_h, NUM_W, NUM_H) == true); + SECTION("Host to Array") { + Memcpy2DHosttoAShell( + std::bind(hipMemcpy2DToArrayAsync, _1, 0, 0, _2, _3, width * sizeof(int), height, + hipMemcpyHostToDevice, stream), + width, height, stream); } - // Cleaning the memory - HIP_CHECK(hipFreeArray(A_d)); - HIP_CHECK(hipStreamDestroy(stream)); - HipTest::freeArrays(nullptr, nullptr, nullptr, - A_h, hData, nullptr, false); -} -/* - * This Scenario Verifies hipMemcpy2DToArray API by copying the - * data from pinned host memory to device on same GPU - * INPUT: Copying Host variable PinnMem(Initialized with value "10" ) - * --> A_d device variable - * OUTPUT: For validating the result,Copying A_d device variable - * --> A_h host variable - * and verifying A_h with PinnedMem[0](i.e., 10) - */ -TEST_CASE("Unit_hipMemcpy2DToArrayAsync_PinnedHostMemSameGpu") { - HIP_CHECK(hipSetDevice(0)); - hipArray *A_d{nullptr}; - constexpr auto def_val{10}; - size_t width{sizeof(float)*NUM_W}; - float *A_h{nullptr}, *PinnMem{nullptr}; - hipStream_t stream; - - // Initialization of variables - HipTest::initArrays(nullptr, nullptr, nullptr, - &A_h, nullptr, nullptr, - width*NUM_H, false); - HIP_CHECK(hipHostMalloc(reinterpret_cast(&PinnMem), width * NUM_H)); - hipChannelFormatDesc desc = hipCreateChannelDesc(); - HIP_CHECK(hipMallocArray(&A_d, &desc, NUM_W, NUM_H, hipArrayDefault)); - HipTest::setDefaultData(width*NUM_H, A_h, nullptr, nullptr); - for (int i = 0; i < NUM_W*NUM_H; i++) { - PinnMem[i] = def_val + i; + SECTION("Host to Array with default kind") { + Memcpy2DHosttoAShell( + std::bind(hipMemcpy2DToArrayAsync, _1, 0, 0, _2, _3, width * sizeof(int), height, + hipMemcpyDefault, stream), + width, height, stream); } - HIP_CHECK(hipStreamCreate(&stream)); - - HIP_CHECK(hipMemcpy2DToArrayAsync(A_d, 0, 0, PinnMem, - width, width, NUM_H, hipMemcpyHostToDevice, - stream)); - HIP_CHECK(hipStreamSynchronize(stream)); - HIP_CHECK(hipMemcpy2DFromArray(A_h, width, A_d, - 0, 0, width, NUM_H, - hipMemcpyDeviceToHost)); - REQUIRE(HipTest::checkArray(A_h, PinnMem, NUM_W, NUM_H) == true); - - // Cleaning the memory - HIP_CHECK(hipFreeArray(A_d)); - HIP_CHECK(hipHostFree(PinnMem)); - HIP_CHECK(hipStreamDestroy(stream)); - HipTest::freeArrays(nullptr, nullptr, nullptr, - A_h, nullptr, nullptr, false); -} -/* - * This Scenario Verifies hipMemcpy2DToArray API by copying the - * data from pinned host memory to device from Peer GPU. - * Device Memory is allocated in GPU 0 and the API is trigerred from GPU1 - * INPUT: Copying Host variable E_h(Initialized with value 10+i(numelements)) - * --> A_d device variable - * whose memory is allocated in GPU 0 - * OUTPUT: For validating the result,Copying A_d device variable - * --> A_h host variable - * and verifying A_h with E_h[0]+i(i.e., 10+i) - */ -TEST_CASE("Unit_hipMemcpy2DToArrayAsync_multiDevicePinnedHostMem") { - int numDevices = 0; - constexpr auto def_val{10}; - HIP_CHECK(hipGetDeviceCount(&numDevices)); - if (numDevices > 1) { - int canAccessPeer = 0; - HIP_CHECK(hipDeviceCanAccessPeer(&canAccessPeer, 0, 1)); - if (canAccessPeer) { - HIP_CHECK(hipSetDevice(0)); - hipArray *A_d{nullptr}; - size_t width{sizeof(float)*NUM_W}; - float *A_h{nullptr}, *E_h{nullptr}; - hipStream_t stream; - HIP_CHECK(hipStreamCreate(&stream)); - - // Initialization of variables - HipTest::initArrays(nullptr, nullptr, nullptr, - &A_h, nullptr, nullptr, - width*NUM_H, false); - hipChannelFormatDesc desc = hipCreateChannelDesc(); - HIP_CHECK(hipMallocArray(&A_d, &desc, NUM_W, NUM_H, hipArrayDefault)); - HipTest::setDefaultData(width*NUM_H, A_h, nullptr, nullptr); - HIP_CHECK(hipSetDevice(1)); - HIP_CHECK(hipHostMalloc(reinterpret_cast(&E_h), width * NUM_H)); - for (int i = 0; i < NUM_W*NUM_H; i++) { - E_h[i] = def_val + i; - } - - HIP_CHECK(hipMemcpy2DToArrayAsync(A_d, 0, 0, E_h, width, - width, NUM_H, hipMemcpyHostToDevice, - stream)); - HIP_CHECK(hipStreamSynchronize(stream)); - HIP_CHECK(hipSetDevice(0)); - HIP_CHECK(hipMemcpy2DFromArray(A_h, width, A_d, - 0, 0, width, NUM_H, - hipMemcpyDeviceToHost)); - REQUIRE(HipTest::checkArray(A_h, E_h, NUM_W, NUM_H) == true); - - // Cleaning the memory - HIP_CHECK(hipFreeArray(A_d)); - HIP_CHECK(hipHostFree(E_h)); - HIP_CHECK(hipStreamDestroy(stream)); - HipTest::freeArrays(nullptr, nullptr, nullptr, - A_h, nullptr, nullptr, false); - } else { - SUCCEED("Machine Does not have P2P capability"); +#if HT_NVIDIA // EXSWHTEC-213 + SECTION("Device to Array") { + SECTION("Peer access disabled") { + Memcpy2DDevicetoAShell( + std::bind(hipMemcpy2DToArrayAsync, _1, 0, 0, _2, _3, width * sizeof(int), height, + hipMemcpyDeviceToDevice, stream), + width, height, stream); } - } else { - SUCCEED("Number of devices are < 2"); - } -} - -/* - * This scenario verifies the hipMemcpy2DToArray API in case of device - * context change. - * Memory is allocated in GPU-0 and the API is triggered from GPU-1 - * INPUT: Copying Host variable hData(Initial value Phi) - * --> A_d device variable - * whose memory is allocated in GPU 0 - * OUTPUT: For validating the result,Copying A_d device variable - * --> A_h host variable - * and verifying A_h with Phi - * */ -TEST_CASE("Unit_hipMemcpy2DToArrayAsync_multiDeviceDeviceContextChange") { - int numDevices = 0; - HIP_CHECK(hipGetDeviceCount(&numDevices)); - if (numDevices > 1) { - int canAccessPeer = 0; - HIP_CHECK(hipDeviceCanAccessPeer(&canAccessPeer, 0, 1)); - if (canAccessPeer) { - HIP_CHECK(hipSetDevice(0)); - hipArray *A_d{nullptr}; - size_t width{sizeof(float)*NUM_W}; - float *A_h{nullptr}, *hData{nullptr}; - hipStream_t stream; - - // Initialization of variables - HipTest::initArrays(nullptr, nullptr, nullptr, - &A_h, &hData, nullptr, - width*NUM_H, false); - hipChannelFormatDesc desc = hipCreateChannelDesc(); - HIP_CHECK(hipMallocArray(&A_d, &desc, NUM_W, NUM_H, hipArrayDefault)); - HipTest::setDefaultData(width*NUM_H, A_h, hData, nullptr); - - HIP_CHECK(hipSetDevice(1)); - HIP_CHECK(hipStreamCreate(&stream)); - HIP_CHECK(hipMemcpy2DToArrayAsync(A_d, 0, 0, hData, width, width, - NUM_H, hipMemcpyHostToDevice, - stream)); - HIP_CHECK(hipStreamSynchronize(stream)); - - HIP_CHECK(hipMemcpy2DFromArray(A_h, width, A_d, - 0, 0, width, NUM_H, - hipMemcpyDeviceToHost)); - REQUIRE(HipTest::checkArray(A_h, hData, NUM_W, NUM_H) == true); - - // Cleaning the memory - HIP_CHECK(hipFreeArray(A_d)); - HIP_CHECK(hipStreamDestroy(stream)); - HipTest::freeArrays(nullptr, nullptr, nullptr, - A_h, hData, nullptr, false); - } else { - SUCCEED("Machine Does not have P2P capability"); + SECTION("Peer access enabled") { + Memcpy2DDevicetoAShell( + std::bind(hipMemcpy2DToArrayAsync, _1, 0, 0, _2, _3, width * sizeof(int), height, + hipMemcpyDeviceToDevice, stream), + width, height, stream); } - } else { - SUCCEED("Number of devices are < 2"); - } -} -/* This testcase verifies the negative scenarios - */ -TEST_CASE("Unit_hipMemcpy2DToArrayAsync_Negative") { - HIP_CHECK(hipSetDevice(0)); - hipArray *A_d{nullptr}; - size_t width{sizeof(float)*NUM_W}; - float *A_h{nullptr}, *hData{nullptr}; - hipStream_t stream; - HIP_CHECK(hipStreamCreate(&stream)); - - // Initialization of variables - HipTest::initArrays(nullptr, nullptr, nullptr, - &A_h, &hData, nullptr, - width*NUM_H, false); - HipTest::setDefaultData(width*NUM_H, A_h, hData, nullptr); - hipChannelFormatDesc desc = hipCreateChannelDesc(); - HIP_CHECK(hipMallocArray(&A_d, &desc, NUM_W, NUM_H, hipArrayDefault)); - - SECTION("Nullptr to destination") { - REQUIRE(hipMemcpy2DToArrayAsync(nullptr, 0, 0, hData, width, - width, NUM_H, hipMemcpyHostToDevice, - stream) != hipSuccess); } - SECTION("Nullptr to source") { - REQUIRE(hipMemcpy2DToArrayAsync(A_d, 0, 0, nullptr, - width, width, NUM_H, hipMemcpyHostToDevice, - stream) != hipSuccess); + SECTION("Device to Array with default kind") { + SECTION("Peer access disabled") { + Memcpy2DDevicetoAShell( + std::bind(hipMemcpy2DToArrayAsync, _1, 0, 0, _2, _3, width * sizeof(int), height, + hipMemcpyDefault, stream), + width, height, stream); + } + SECTION("Peer access enabled") { + Memcpy2DDevicetoAShell( + std::bind(hipMemcpy2DToArrayAsync, _1, 0, 0, _2, _3, width * sizeof(int), height, + hipMemcpyDefault, stream), + width, height, stream); + } } - - SECTION("Passing offset more than 0") { - REQUIRE(hipMemcpy2DToArrayAsync(A_d, 1, 1, hData, width, - width, NUM_H, hipMemcpyHostToDevice, - stream) != hipSuccess); - } - - SECTION("Passing array more than allocated") { - REQUIRE(hipMemcpy2DToArrayAsync(A_d, 0, 0, hData, width, - width+2, NUM_H+2, hipMemcpyHostToDevice, - stream) != hipSuccess); - } - - // Cleaning of Memory - HIP_CHECK(hipFreeArray(A_d)); - HIP_CHECK(hipStreamDestroy(stream)); - HipTest::freeArrays(nullptr, nullptr, nullptr, - A_h, hData, nullptr, false); +#endif } +TEST_CASE("Unit_hipMemcpy2DToArrayAsync_Positive_Synchronization_Behavior") { + using namespace std::placeholders; + HIP_CHECK(hipDeviceSynchronize()); + + SECTION("Host to Array") { + const auto width = GENERATE(16, 32, 48); + const auto height = GENERATE(16, 32, 48); + + MemcpyHtoASyncBehavior(std::bind(hipMemcpy2DToArrayAsync, _1, 0, 0, _2, width * sizeof(int), + width * sizeof(int), height, hipMemcpyHostToDevice, nullptr), + width, height, false); + } + + SECTION("Device to Array") { + const auto width = GENERATE(16, 32, 48); + const auto height = GENERATE(16, 32, 48); + + MemcpyDtoASyncBehavior(std::bind(hipMemcpy2DToArrayAsync, _1, 0, 0, _2, _3, width * sizeof(int), + height, hipMemcpyDeviceToDevice, nullptr), + width, height, false); + } +} + +TEST_CASE("Unit_hipMemcpy2DToArrayAsync_Positive_ZeroWidthHeight") { + using namespace std::placeholders; + const auto width = 16; + const auto height = 16; + + const auto stream_type = GENERATE(Streams::nullstream, Streams::perThread, Streams::created); + const StreamGuard stream_guard(stream_type); + const hipStream_t stream = stream_guard.stream(); + + SECTION("Array to host") { + SECTION("Height is 0") { + Memcpy2DToArrayZeroWidthHeight( + std::bind(hipMemcpy2DToArrayAsync, _1, 0, 0, _2, _3, width * sizeof(int), 0, + hipMemcpyHostToDevice, stream), + width, height, stream); + } + SECTION("Width is 0") { + Memcpy2DToArrayZeroWidthHeight(std::bind(hipMemcpy2DToArrayAsync, _1, 0, 0, _2, _3, 0, + height, hipMemcpyHostToDevice, stream), + width, height, stream); + } + } + SECTION("Array to device") { + SECTION("Height is 0") { + Memcpy2DToArrayZeroWidthHeight( + std::bind(hipMemcpy2DToArrayAsync, _1, 0, 0, _2, _3, width * sizeof(int), 0, + hipMemcpyDeviceToDevice, stream), + width, height, stream); + } + SECTION("Width is 0") { + Memcpy2DToArrayZeroWidthHeight(std::bind(hipMemcpy2DToArrayAsync, _1, 0, 0, _2, _3, 0, + height, hipMemcpyDeviceToDevice, stream), + width, height, stream); + } + } +} + +TEST_CASE("Unit_hipMemcpy2DToArrayAsync_Negative_Parameters") { + using namespace std::placeholders; + + const auto width = 32; + const auto height = 32; + const auto allocation_size = 2 * width * height * sizeof(int); + + const unsigned int flag = hipArrayDefault; + + constexpr auto InvalidStream = [] { + StreamGuard sg(Streams::created); + return sg.stream(); + }; + + ArrayAllocGuard array_alloc(make_hipExtent(width, height, 0), flag); + LinearAllocGuard2D device_alloc(width, height); + LinearAllocGuard host_alloc(LinearAllocs::hipHostMalloc, allocation_size); + + SECTION("Host to Array") { + SECTION("dst == nullptr") { + HIP_CHECK_ERROR( + hipMemcpy2DToArrayAsync(nullptr, 0, 0, host_alloc.ptr(), 2 * width * sizeof(int), + width * sizeof(int), height, hipMemcpyHostToDevice, nullptr), + hipErrorInvalidHandle); + } + SECTION("src == nullptr") { + HIP_CHECK_ERROR( + hipMemcpy2DToArrayAsync(array_alloc.ptr(), 0, 0, nullptr, 2 * width * sizeof(int), + width * sizeof(int), height, hipMemcpyHostToDevice, nullptr), + hipErrorInvalidValue); + } +#if HT_NVIDIA // EXSWHTEC-212 + SECTION("spitch < width") { + HIP_CHECK_ERROR(hipMemcpy2DToArrayAsync(array_alloc.ptr(), 0, 0, host_alloc.ptr(), + width * sizeof(int) - 10, width * sizeof(int), height, + hipMemcpyHostToDevice, nullptr), + hipErrorInvalidPitchValue); + } + SECTION("Offset + width/height overflows") { + HIP_CHECK_ERROR(hipMemcpy2DToArrayAsync(array_alloc.ptr(), 1, 0, host_alloc.ptr(), + 2 * width * sizeof(int), width * sizeof(int), height, + hipMemcpyHostToDevice, nullptr), + hipErrorInvalidValue); + HIP_CHECK_ERROR(hipMemcpy2DToArrayAsync(array_alloc.ptr(), 0, 1, host_alloc.ptr(), + 2 * width * sizeof(int), width * sizeof(int), height, + hipMemcpyHostToDevice, nullptr), + hipErrorInvalidValue); + } + SECTION("Width/height overflows") { + HIP_CHECK_ERROR(hipMemcpy2DToArrayAsync(array_alloc.ptr(), 0, 0, host_alloc.ptr(), + 2 * width * sizeof(int), width * sizeof(int) + 1, + height, hipMemcpyHostToDevice, nullptr), + hipErrorInvalidValue); + HIP_CHECK_ERROR(hipMemcpy2DToArrayAsync(array_alloc.ptr(), 0, 0, host_alloc.ptr(), + 2 * width * sizeof(int), width * sizeof(int), + height + 1, hipMemcpyHostToDevice, nullptr), + hipErrorInvalidValue); + } + SECTION("Memcpy kind is invalid") { + HIP_CHECK_ERROR(hipMemcpy2DToArrayAsync(array_alloc.ptr(), 0, 0, host_alloc.ptr(), + 2 * width * sizeof(int), width * sizeof(int), height, + static_cast(-1), nullptr), + hipErrorInvalidMemcpyDirection); + } + SECTION("Invalid stream") { + HIP_CHECK_ERROR(hipMemcpy2DToArrayAsync(array_alloc.ptr(), 0, 0, host_alloc.ptr(), + 2 * width * sizeof(int), width * sizeof(int), height, + hipMemcpyHostToDevice, InvalidStream()), + hipErrorContextIsDestroyed); + } +#endif + } + SECTION("Device to Array") { + SECTION("dst == nullptr") { + HIP_CHECK_ERROR( + hipMemcpy2DToArrayAsync(nullptr, 0, 0, device_alloc.ptr(), device_alloc.pitch(), + width * sizeof(int), height, hipMemcpyDeviceToDevice, nullptr), + hipErrorInvalidHandle); + } + SECTION("src == nullptr") { + HIP_CHECK_ERROR( + hipMemcpy2DToArrayAsync(array_alloc.ptr(), 0, 0, nullptr, device_alloc.pitch(), + width * sizeof(int), height, hipMemcpyDeviceToDevice, nullptr), + hipErrorInvalidValue); + } +#if HT_NVIDIA // EXSWHTEC-212 + SECTION("spitch < width") { + HIP_CHECK_ERROR(hipMemcpy2DToArrayAsync(array_alloc.ptr(), 0, 0, device_alloc.ptr(), + width * sizeof(int) - 10, width * sizeof(int), height, + hipMemcpyDeviceToDevice, nullptr), + hipErrorInvalidPitchValue); + } + SECTION("Offset + width/height overflows") { + HIP_CHECK_ERROR( + hipMemcpy2DToArrayAsync(array_alloc.ptr(), 1, 0, device_alloc.ptr(), device_alloc.pitch(), + width * sizeof(int), height, hipMemcpyDeviceToDevice, nullptr), + hipErrorInvalidValue); + HIP_CHECK_ERROR( + hipMemcpy2DToArrayAsync(array_alloc.ptr(), 0, 1, device_alloc.ptr(), device_alloc.pitch(), + width * sizeof(int), height, hipMemcpyDeviceToDevice, nullptr), + hipErrorInvalidValue); + } + SECTION("Width/height overflows") { + HIP_CHECK_ERROR(hipMemcpy2DToArrayAsync(array_alloc.ptr(), 0, 0, device_alloc.ptr(), + device_alloc.pitch(), width * sizeof(int) + 1, height, + hipMemcpyDeviceToDevice, nullptr), + hipErrorInvalidValue); + HIP_CHECK_ERROR(hipMemcpy2DToArrayAsync(array_alloc.ptr(), 0, 0, device_alloc.ptr(), + device_alloc.pitch(), width * sizeof(int), height + 1, + hipMemcpyDeviceToDevice, nullptr), + hipErrorInvalidValue); + } + SECTION("Memcpy kind is invalid") { + HIP_CHECK_ERROR(hipMemcpy2DToArrayAsync(array_alloc.ptr(), 0, 0, device_alloc.ptr(), + device_alloc.pitch(), width * sizeof(int), height, + static_cast(-1), nullptr), + hipErrorInvalidMemcpyDirection); + } + SECTION("Invalid stream") { + HIP_CHECK_ERROR(hipMemcpy2DToArrayAsync(array_alloc.ptr(), 0, 0, device_alloc.ptr(), + device_alloc.pitch(), width * sizeof(int), height, + hipMemcpyDeviceToDevice, InvalidStream()), + hipErrorContextIsDestroyed); + } +#endif + } +} diff --git a/projects/hip-tests/catch/unit/memory/hipMemcpy2DToArrayAsync_old.cc b/projects/hip-tests/catch/unit/memory/hipMemcpy2DToArrayAsync_old.cc new file mode 100644 index 0000000000..990b9c651e --- /dev/null +++ b/projects/hip-tests/catch/unit/memory/hipMemcpy2DToArrayAsync_old.cc @@ -0,0 +1,369 @@ +/* +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 file verifies the following scenarios of hipMemcpy2DToArrayAsync API +1. Negative Scenarios +2. Extent Validation Scenarios +3. hipMemcpy2DToArrayAsync Basic Scenario +4. Pinned Memory scenarios on same and peer GPU +5. Device Context change scenario where memory is allocated in + one GPU and stream is created in peer GPU. +*/ + +#include +#include +#include + + +static constexpr auto NUM_W{10}; +static constexpr auto NUM_H{10}; + +/* + * This Scenario copies the data from host to device + * INPUT: Copying Host variable hData(Initialized with value Phi(1.618)) + * --> A_d device variable + * OUTPUT: For validating the result,Copying A_d device variable + * --> A_h host variable + * and verifying A_h with Phi + */ +TEST_CASE("Unit_hipMemcpy2DToArrayAsync_Basic") { + HIP_CHECK(hipSetDevice(0)); + hipArray *A_d{nullptr}; + size_t width{sizeof(float)*NUM_W}; + float *A_h{nullptr}, *hData{nullptr}; + hipStream_t stream; + + // Initialization of variables + HipTest::initArrays(nullptr, nullptr, nullptr, + &A_h, &hData, nullptr, + width*NUM_H, false); + hipChannelFormatDesc desc = hipCreateChannelDesc(); + HIP_CHECK(hipMallocArray(&A_d, &desc, NUM_W, NUM_H, hipArrayDefault)); + HipTest::setDefaultData(width*NUM_H, A_h, hData, nullptr); + HIP_CHECK(hipStreamCreate(&stream)); + SECTION("Calling hipMemcpy2DToArrayAsync() with user declared stream obj") { + HIP_CHECK(hipMemcpy2DToArrayAsync(A_d, 0, 0, hData, width, + width, NUM_H, + hipMemcpyHostToDevice, stream)); + HIP_CHECK(hipStreamSynchronize(stream)); + } + SECTION("Calling hipMemcpy2DToArrayAsync() with hipStreamPerThread") { + HIP_CHECK(hipMemcpy2DToArrayAsync(A_d, 0, 0, hData, width, + width, NUM_H, + hipMemcpyHostToDevice, hipStreamPerThread)); + HIP_CHECK(hipStreamSynchronize(hipStreamPerThread)); + } + HIP_CHECK(hipMemcpy2DFromArray(A_h, width, A_d, + 0, 0, width, NUM_H, + hipMemcpyDeviceToHost)); + REQUIRE(HipTest::checkArray(A_h, hData, NUM_W, NUM_H) == true); + + // Cleaning the memory + HIP_CHECK(hipFreeArray(A_d)); + HIP_CHECK(hipStreamDestroy(stream)); + HipTest::freeArrays(nullptr, nullptr, nullptr, + A_h, hData, nullptr, false); +} + +/* + * This testcase verifies the extent validation scenarios + */ +TEST_CASE("Unit_hipMemcpy2DToArrayAsync_ExtentValidation") { + HIP_CHECK(hipSetDevice(0)); + hipArray *A_d{nullptr}; + size_t width{sizeof(float)*NUM_W}; + float *A_h{nullptr}, *hData{nullptr}; + hipStream_t stream; + + // Initialization of variables + HipTest::initArrays(nullptr, nullptr, nullptr, + &A_h, &hData, nullptr, + width*NUM_H, false); + hipChannelFormatDesc desc = hipCreateChannelDesc(); + HIP_CHECK(hipMallocArray(&A_d, &desc, NUM_W, NUM_H, hipArrayDefault)); + HIP_CHECK(hipStreamCreate(&stream)); + + SECTION("Source width is 0") { + REQUIRE(hipMemcpy2DToArrayAsync(A_d, 0, 0, hData, 0, + width, NUM_H, hipMemcpyHostToDevice, + stream) != hipSuccess); + HIP_CHECK(hipStreamSynchronize(stream)); + } + // hipMemcpy2DToArray API would return success for width and height as 0 + // and does not perform any copy + // Validating the result with the initialized value + // 1.Initializing A_d with Pi value + // 2.copying hData(Phi)-->A_d device variable + // with height 0(copy will not be performed) + // 3.copying A_d-->hData and validating it with A_h data + SECTION("Height is 0") { + HIP_CHECK(hipMemcpy2DToArrayAsync(A_d, 0, 0, A_h, width, + width, NUM_H, hipMemcpyHostToDevice, + stream)); + HIP_CHECK(hipMemcpy2DToArrayAsync(A_d, 0, 0, hData, width, + width, 0, hipMemcpyHostToDevice, + stream)); + HIP_CHECK(hipStreamSynchronize(stream)); + HIP_CHECK(hipMemcpy2DFromArray(hData, width, A_d, + 0, 0, width, NUM_H, + hipMemcpyDeviceToHost)); + REQUIRE(HipTest::checkArray(hData, A_h, NUM_W, NUM_H) == true); + } + // hipMemcpy2DToArray API would return success for width and height as 0 + // and does not perform any copy + // Validating the result with the initialized value + // 1.Initializing A_d with Pi value + // 2.copying hData(Phi)-->A_d device variable + // with width 0(copy will not be performed) + // 3.copying A_d-->hData and validating it with A_h data + SECTION("Width is 0") { + HIP_CHECK(hipMemcpy2DToArrayAsync(A_d, 0, 0, A_h, width, + width, NUM_H, hipMemcpyHostToDevice, + stream)); + HIP_CHECK(hipMemcpy2DToArrayAsync(A_d, 0, 0, hData, width, + 0, NUM_H, hipMemcpyHostToDevice, + stream)); + HIP_CHECK(hipStreamSynchronize(stream)); + HIP_CHECK(hipMemcpy2DFromArray(hData, width, A_d, + 0, 0, width, NUM_H, + hipMemcpyDeviceToHost)); + REQUIRE(HipTest::checkArray(hData, A_h, NUM_W, NUM_H) == true); + } + + // Cleaning the memory + HIP_CHECK(hipFreeArray(A_d)); + HIP_CHECK(hipStreamDestroy(stream)); + HipTest::freeArrays(nullptr, nullptr, nullptr, + A_h, hData, nullptr, false); +} +/* + * This Scenario Verifies hipMemcpy2DToArray API by copying the + * data from pinned host memory to device on same GPU + * INPUT: Copying Host variable PinnMem(Initialized with value "10" ) + * --> A_d device variable + * OUTPUT: For validating the result,Copying A_d device variable + * --> A_h host variable + * and verifying A_h with PinnedMem[0](i.e., 10) + */ +TEST_CASE("Unit_hipMemcpy2DToArrayAsync_PinnedHostMemSameGpu") { + HIP_CHECK(hipSetDevice(0)); + hipArray *A_d{nullptr}; + constexpr auto def_val{10}; + size_t width{sizeof(float)*NUM_W}; + float *A_h{nullptr}, *PinnMem{nullptr}; + hipStream_t stream; + + // Initialization of variables + HipTest::initArrays(nullptr, nullptr, nullptr, + &A_h, nullptr, nullptr, + width*NUM_H, false); + HIP_CHECK(hipHostMalloc(reinterpret_cast(&PinnMem), width * NUM_H)); + hipChannelFormatDesc desc = hipCreateChannelDesc(); + HIP_CHECK(hipMallocArray(&A_d, &desc, NUM_W, NUM_H, hipArrayDefault)); + HipTest::setDefaultData(width*NUM_H, A_h, nullptr, nullptr); + for (int i = 0; i < NUM_W*NUM_H; i++) { + PinnMem[i] = def_val + i; + } + HIP_CHECK(hipStreamCreate(&stream)); + + HIP_CHECK(hipMemcpy2DToArrayAsync(A_d, 0, 0, PinnMem, + width, width, NUM_H, hipMemcpyHostToDevice, + stream)); + HIP_CHECK(hipStreamSynchronize(stream)); + HIP_CHECK(hipMemcpy2DFromArray(A_h, width, A_d, + 0, 0, width, NUM_H, + hipMemcpyDeviceToHost)); + REQUIRE(HipTest::checkArray(A_h, PinnMem, NUM_W, NUM_H) == true); + + // Cleaning the memory + HIP_CHECK(hipFreeArray(A_d)); + HIP_CHECK(hipHostFree(PinnMem)); + HIP_CHECK(hipStreamDestroy(stream)); + HipTest::freeArrays(nullptr, nullptr, nullptr, + A_h, nullptr, nullptr, false); +} +/* + * This Scenario Verifies hipMemcpy2DToArray API by copying the + * data from pinned host memory to device from Peer GPU. + * Device Memory is allocated in GPU 0 and the API is trigerred from GPU1 + * INPUT: Copying Host variable E_h(Initialized with value 10+i(numelements)) + * --> A_d device variable + * whose memory is allocated in GPU 0 + * OUTPUT: For validating the result,Copying A_d device variable + * --> A_h host variable + * and verifying A_h with E_h[0]+i(i.e., 10+i) + */ +TEST_CASE("Unit_hipMemcpy2DToArrayAsync_multiDevicePinnedHostMem") { + int numDevices = 0; + constexpr auto def_val{10}; + HIP_CHECK(hipGetDeviceCount(&numDevices)); + if (numDevices > 1) { + int canAccessPeer = 0; + HIP_CHECK(hipDeviceCanAccessPeer(&canAccessPeer, 0, 1)); + if (canAccessPeer) { + HIP_CHECK(hipSetDevice(0)); + hipArray *A_d{nullptr}; + size_t width{sizeof(float)*NUM_W}; + float *A_h{nullptr}, *E_h{nullptr}; + hipStream_t stream; + HIP_CHECK(hipStreamCreate(&stream)); + + // Initialization of variables + HipTest::initArrays(nullptr, nullptr, nullptr, + &A_h, nullptr, nullptr, + width*NUM_H, false); + hipChannelFormatDesc desc = hipCreateChannelDesc(); + HIP_CHECK(hipMallocArray(&A_d, &desc, NUM_W, NUM_H, hipArrayDefault)); + HipTest::setDefaultData(width*NUM_H, A_h, nullptr, nullptr); + HIP_CHECK(hipSetDevice(1)); + HIP_CHECK(hipHostMalloc(reinterpret_cast(&E_h), width * NUM_H)); + for (int i = 0; i < NUM_W*NUM_H; i++) { + E_h[i] = def_val + i; + } + + HIP_CHECK(hipMemcpy2DToArrayAsync(A_d, 0, 0, E_h, width, + width, NUM_H, hipMemcpyHostToDevice, + stream)); + HIP_CHECK(hipStreamSynchronize(stream)); + HIP_CHECK(hipSetDevice(0)); + HIP_CHECK(hipMemcpy2DFromArray(A_h, width, A_d, + 0, 0, width, NUM_H, + hipMemcpyDeviceToHost)); + REQUIRE(HipTest::checkArray(A_h, E_h, NUM_W, NUM_H) == true); + + // Cleaning the memory + HIP_CHECK(hipFreeArray(A_d)); + HIP_CHECK(hipHostFree(E_h)); + HIP_CHECK(hipStreamDestroy(stream)); + HipTest::freeArrays(nullptr, nullptr, nullptr, + A_h, nullptr, nullptr, false); + } else { + SUCCEED("Machine Does not have P2P capability"); + } + } else { + SUCCEED("Number of devices are < 2"); + } +} + +/* + * This scenario verifies the hipMemcpy2DToArray API in case of device + * context change. + * Memory is allocated in GPU-0 and the API is triggered from GPU-1 + * INPUT: Copying Host variable hData(Initial value Phi) + * --> A_d device variable + * whose memory is allocated in GPU 0 + * OUTPUT: For validating the result,Copying A_d device variable + * --> A_h host variable + * and verifying A_h with Phi + * */ +TEST_CASE("Unit_hipMemcpy2DToArrayAsync_multiDeviceDeviceContextChange") { + int numDevices = 0; + HIP_CHECK(hipGetDeviceCount(&numDevices)); + if (numDevices > 1) { + int canAccessPeer = 0; + HIP_CHECK(hipDeviceCanAccessPeer(&canAccessPeer, 0, 1)); + if (canAccessPeer) { + HIP_CHECK(hipSetDevice(0)); + hipArray *A_d{nullptr}; + size_t width{sizeof(float)*NUM_W}; + float *A_h{nullptr}, *hData{nullptr}; + hipStream_t stream; + + // Initialization of variables + HipTest::initArrays(nullptr, nullptr, nullptr, + &A_h, &hData, nullptr, + width*NUM_H, false); + hipChannelFormatDesc desc = hipCreateChannelDesc(); + HIP_CHECK(hipMallocArray(&A_d, &desc, NUM_W, NUM_H, hipArrayDefault)); + HipTest::setDefaultData(width*NUM_H, A_h, hData, nullptr); + + HIP_CHECK(hipSetDevice(1)); + HIP_CHECK(hipStreamCreate(&stream)); + HIP_CHECK(hipMemcpy2DToArrayAsync(A_d, 0, 0, hData, width, width, + NUM_H, hipMemcpyHostToDevice, + stream)); + HIP_CHECK(hipStreamSynchronize(stream)); + + HIP_CHECK(hipMemcpy2DFromArray(A_h, width, A_d, + 0, 0, width, NUM_H, + hipMemcpyDeviceToHost)); + REQUIRE(HipTest::checkArray(A_h, hData, NUM_W, NUM_H) == true); + + // Cleaning the memory + HIP_CHECK(hipFreeArray(A_d)); + HIP_CHECK(hipStreamDestroy(stream)); + HipTest::freeArrays(nullptr, nullptr, nullptr, + A_h, hData, nullptr, false); + } else { + SUCCEED("Machine Does not have P2P capability"); + } + } else { + SUCCEED("Number of devices are < 2"); + } +} +/* This testcase verifies the negative scenarios + */ +TEST_CASE("Unit_hipMemcpy2DToArrayAsync_Negative") { + HIP_CHECK(hipSetDevice(0)); + hipArray *A_d{nullptr}; + size_t width{sizeof(float)*NUM_W}; + float *A_h{nullptr}, *hData{nullptr}; + hipStream_t stream; + HIP_CHECK(hipStreamCreate(&stream)); + + // Initialization of variables + HipTest::initArrays(nullptr, nullptr, nullptr, + &A_h, &hData, nullptr, + width*NUM_H, false); + HipTest::setDefaultData(width*NUM_H, A_h, hData, nullptr); + hipChannelFormatDesc desc = hipCreateChannelDesc(); + HIP_CHECK(hipMallocArray(&A_d, &desc, NUM_W, NUM_H, hipArrayDefault)); + + SECTION("Nullptr to destination") { + REQUIRE(hipMemcpy2DToArrayAsync(nullptr, 0, 0, hData, width, + width, NUM_H, hipMemcpyHostToDevice, + stream) != hipSuccess); + } + + SECTION("Nullptr to source") { + REQUIRE(hipMemcpy2DToArrayAsync(A_d, 0, 0, nullptr, + width, width, NUM_H, hipMemcpyHostToDevice, + stream) != hipSuccess); + } + + SECTION("Passing offset more than 0") { + REQUIRE(hipMemcpy2DToArrayAsync(A_d, 1, 1, hData, width, + width, NUM_H, hipMemcpyHostToDevice, + stream) != hipSuccess); + } + + SECTION("Passing array more than allocated") { + REQUIRE(hipMemcpy2DToArrayAsync(A_d, 0, 0, hData, width, + width+2, NUM_H+2, hipMemcpyHostToDevice, + stream) != hipSuccess); + } + + // Cleaning of Memory + HIP_CHECK(hipFreeArray(A_d)); + HIP_CHECK(hipStreamDestroy(stream)); + HipTest::freeArrays(nullptr, nullptr, nullptr, + A_h, hData, nullptr, false); +} +