EXSWHTEC-92 - Implement tests for async memcpy of 2D hipArray (#31)
- Implement tests for hipMemcpy2DFromArrayAsync using resource guards and templates
- Implement tests for hipMemcpy2DToArrayAsync using resource guards and templates
[ROCm/hip-tests commit: 818923bfbc]
This commit is contained in:
nives-vukovic
GitHub
@@ -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
|
||||
|
||||
@@ -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 <functional>
|
||||
|
||||
#include <hip_test_common.hh>
|
||||
#include <hip/hip_runtime_api.h>
|
||||
#include <utils.hh>
|
||||
#include <resource_guards.hh>
|
||||
#include "hipArrayCommon.hh"
|
||||
|
||||
/* Array -> Host */
|
||||
template <bool should_synchronize, typename T, typename F>
|
||||
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<T> host_allocation(LinearAllocs::hipHostMalloc, allocation_size);
|
||||
ArrayAllocGuard<T> 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 <bool should_synchronize, typename T, typename F>
|
||||
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<T> host_allocation(LinearAllocs::hipHostMalloc, allocation_size);
|
||||
ArrayAllocGuard<T> 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 <bool should_synchronize, bool enable_peer_access, typename T, typename F>
|
||||
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<T> host_allocation(LinearAllocs::hipHostMalloc, allocation_size);
|
||||
ArrayAllocGuard<T> array_allocation(make_hipExtent(width, height, 0), flag);
|
||||
HIP_CHECK(hipSetDevice(dst_device));
|
||||
LinearAllocGuard2D<T> 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 <bool should_synchronize, typename T, typename F>
|
||||
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<T> host_allocation(LinearAllocs::hipHostMalloc, allocation_size);
|
||||
ArrayAllocGuard<T> 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 <bool should_synchronize, typename T, typename F>
|
||||
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<T> host_allocation(LinearAllocs::hipHostMalloc, allocation_size);
|
||||
ArrayAllocGuard<T> 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 <bool should_synchronize, bool enable_peer_access, typename T, typename F>
|
||||
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<T> host_allocation(LinearAllocs::hipHostMalloc, allocation_size);
|
||||
LinearAllocGuard2D<T> device_allocation(width, height);
|
||||
HIP_CHECK(hipSetDevice(dst_device));
|
||||
ArrayAllocGuard<T> 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 <typename F>
|
||||
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 <typename F>
|
||||
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<int> host_alloc(LinearAllocs::hipHostMalloc, allocation_size);
|
||||
ArrayAllocGuard<int> 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 <typename F>
|
||||
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<int> host_alloc(LinearAllocs::hipHostMalloc, allocation_size);
|
||||
ArrayAllocGuard<int> array_allocation(make_hipExtent(width, height, 0), flag);
|
||||
MemcpyArraySyncBehaviorCheck(std::bind(memcpy_func, host_alloc.ptr(), array_allocation.ptr()),
|
||||
should_sync, kernel_stream);
|
||||
}
|
||||
|
||||
template <typename F>
|
||||
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<int> host_alloc(LinearAllocs::hipHostMalloc, allocation_size);
|
||||
ArrayAllocGuard<int> 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 <typename F>
|
||||
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<int> array_allocation(make_hipExtent(width, height, 0), flag);
|
||||
LinearAllocGuard2D<int> 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 <typename F>
|
||||
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<int> device_allocation(width, height);
|
||||
ArrayAllocGuard<int> 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 <bool should_synchronize, typename F>
|
||||
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<int> array_alloc(make_hipExtent(width, height, 0), flag);
|
||||
LinearAllocGuard<int> 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<int> array_alloc(make_hipExtent(width, height, 0), flag);
|
||||
LinearAllocGuard2D<int> device_alloc(width, height);
|
||||
LinearAllocGuard<int> 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 <bool should_synchronize, typename F>
|
||||
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<int> array_alloc(make_hipExtent(width, height, 0), flag);
|
||||
LinearAllocGuard<int> 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<int> array_alloc(make_hipExtent(width, height, 0), flag);
|
||||
LinearAllocGuard2D<int> device_alloc(width, height);
|
||||
LinearAllocGuard<int> 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);
|
||||
}
|
||||
}
|
||||
@@ -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 <hip/hip_runtime_api.h>
|
||||
#include <hip_test_common.hh>
|
||||
#include <hip_test_checkers.hh>
|
||||
#include <resource_guards.hh>
|
||||
#include <utils.hh>
|
||||
|
||||
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<float>(nullptr, nullptr, nullptr,
|
||||
&A_h, &hData, nullptr,
|
||||
width*NUM_H, false);
|
||||
hipChannelFormatDesc desc = hipCreateChannelDesc<float>();
|
||||
HIP_CHECK(hipMallocArray(&A_d, &desc, NUM_W, NUM_H, hipArrayDefault));
|
||||
HipTest::setDefaultData<float>(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<float>(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<float>(nullptr, nullptr, nullptr,
|
||||
&A_h, &hData, nullptr,
|
||||
width*NUM_H, false);
|
||||
HipTest::initArrays<float>(nullptr, nullptr, nullptr,
|
||||
nullptr, &valData, nullptr,
|
||||
width*NUM_H, false);
|
||||
hipChannelFormatDesc desc = hipCreateChannelDesc<float>();
|
||||
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<true, int>(
|
||||
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<float>(nullptr, nullptr, nullptr,
|
||||
A_h, hData, nullptr, false);
|
||||
HipTest::freeArrays<float>(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<float>(nullptr, nullptr, nullptr,
|
||||
&A_h, nullptr, nullptr,
|
||||
width*NUM_H, false);
|
||||
HIP_CHECK(hipHostMalloc(reinterpret_cast<void**>(&PinnMem), width * NUM_H));
|
||||
hipChannelFormatDesc desc = hipCreateChannelDesc<float>();
|
||||
HIP_CHECK(hipMallocArray(&A_d, &desc, NUM_W, NUM_H, hipArrayDefault));
|
||||
HipTest::setDefaultData<float>(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<true, int>(
|
||||
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<float>(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<float>(nullptr, nullptr, nullptr,
|
||||
&A_h, nullptr, nullptr,
|
||||
width*NUM_H, false);
|
||||
hipChannelFormatDesc desc = hipCreateChannelDesc<float>();
|
||||
HIP_CHECK(hipMallocArray(&A_d, &desc, NUM_W, NUM_H, hipArrayDefault));
|
||||
HipTest::setDefaultData<float>(width*NUM_H, A_h, nullptr, nullptr);
|
||||
HIP_CHECK(hipHostMalloc(reinterpret_cast<void**>(&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<float>(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<true, false, int>(
|
||||
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<float>(nullptr, nullptr, nullptr,
|
||||
&A_h, &hData, nullptr,
|
||||
width*NUM_H, false);
|
||||
hipChannelFormatDesc desc = hipCreateChannelDesc<float>();
|
||||
HIP_CHECK(hipMallocArray(&A_d, &desc, NUM_W, NUM_H, hipArrayDefault));
|
||||
HipTest::setDefaultData<float>(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<float>(nullptr, nullptr, nullptr,
|
||||
A_h, hData, nullptr, false);
|
||||
} else {
|
||||
SUCCEED("Device Does not have P2P capability");
|
||||
SECTION("Peer access enabled") {
|
||||
Memcpy2DDeviceFromAShell<true, true, int>(
|
||||
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<float>(nullptr, nullptr, nullptr,
|
||||
&A_h, &hData, nullptr,
|
||||
width*NUM_H, false);
|
||||
HipTest::setDefaultData<float>(width*NUM_H, A_h, hData, nullptr);
|
||||
hipChannelFormatDesc desc = hipCreateChannelDesc<float>();
|
||||
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<true, false, int>(
|
||||
std::bind(hipMemcpy2DFromArrayAsync, _1, _2, _3, 0, 0, width * sizeof(int), height,
|
||||
hipMemcpyDefault, stream),
|
||||
width, height, stream);
|
||||
}
|
||||
SECTION("Peer access enabled") {
|
||||
Memcpy2DDeviceFromAShell<true, true, int>(
|
||||
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<float>(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<true>(
|
||||
std::bind(hipMemcpy2DFromArrayAsync, _1, _2, _3, 0, 0, width * sizeof(int), 0,
|
||||
hipMemcpyDeviceToHost, stream),
|
||||
width, height, stream);
|
||||
}
|
||||
SECTION("Width is 0") {
|
||||
Memcpy2DFromArrayZeroWidthHeight<true>(std::bind(hipMemcpy2DFromArrayAsync, _1, _2, _3, 0, 0,
|
||||
0, height, hipMemcpyDeviceToHost, stream),
|
||||
width, height, stream);
|
||||
}
|
||||
}
|
||||
SECTION("Array to device") {
|
||||
SECTION("Height is 0") {
|
||||
Memcpy2DFromArrayZeroWidthHeight<true>(
|
||||
std::bind(hipMemcpy2DFromArrayAsync, _1, _2, _3, 0, 0, width * sizeof(int), 0,
|
||||
hipMemcpyDeviceToDevice, stream),
|
||||
width, height, stream);
|
||||
}
|
||||
SECTION("Width is 0") {
|
||||
Memcpy2DFromArrayZeroWidthHeight<true>(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<int> array_alloc(make_hipExtent(width, height, 0), flag);
|
||||
LinearAllocGuard2D<int> device_alloc(width, height);
|
||||
LinearAllocGuard<int> 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<hipMemcpyKind>(-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<hipMemcpyKind>(-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
|
||||
}
|
||||
}
|
||||
|
||||
@@ -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 <hip_test_common.hh>
|
||||
#include <hip_test_checkers.hh>
|
||||
|
||||
|
||||
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<float>(nullptr, nullptr, nullptr,
|
||||
&A_h, &hData, nullptr,
|
||||
width*NUM_H, false);
|
||||
hipChannelFormatDesc desc = hipCreateChannelDesc<float>();
|
||||
HIP_CHECK(hipMallocArray(&A_d, &desc, NUM_W, NUM_H, hipArrayDefault));
|
||||
HipTest::setDefaultData<float>(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<float>(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<float>(nullptr, nullptr, nullptr,
|
||||
&A_h, &hData, nullptr,
|
||||
width*NUM_H, false);
|
||||
HipTest::initArrays<float>(nullptr, nullptr, nullptr,
|
||||
nullptr, &valData, nullptr,
|
||||
width*NUM_H, false);
|
||||
hipChannelFormatDesc desc = hipCreateChannelDesc<float>();
|
||||
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<float>(nullptr, nullptr, nullptr,
|
||||
A_h, hData, nullptr, false);
|
||||
HipTest::freeArrays<float>(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<float>(nullptr, nullptr, nullptr,
|
||||
&A_h, nullptr, nullptr,
|
||||
width*NUM_H, false);
|
||||
HIP_CHECK(hipHostMalloc(reinterpret_cast<void**>(&PinnMem), width * NUM_H));
|
||||
hipChannelFormatDesc desc = hipCreateChannelDesc<float>();
|
||||
HIP_CHECK(hipMallocArray(&A_d, &desc, NUM_W, NUM_H, hipArrayDefault));
|
||||
HipTest::setDefaultData<float>(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<float>(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<float>(nullptr, nullptr, nullptr,
|
||||
&A_h, nullptr, nullptr,
|
||||
width*NUM_H, false);
|
||||
hipChannelFormatDesc desc = hipCreateChannelDesc<float>();
|
||||
HIP_CHECK(hipMallocArray(&A_d, &desc, NUM_W, NUM_H, hipArrayDefault));
|
||||
HipTest::setDefaultData<float>(width*NUM_H, A_h, nullptr, nullptr);
|
||||
HIP_CHECK(hipHostMalloc(reinterpret_cast<void**>(&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<float>(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<float>(nullptr, nullptr, nullptr,
|
||||
&A_h, &hData, nullptr,
|
||||
width*NUM_H, false);
|
||||
hipChannelFormatDesc desc = hipCreateChannelDesc<float>();
|
||||
HIP_CHECK(hipMallocArray(&A_d, &desc, NUM_W, NUM_H, hipArrayDefault));
|
||||
HipTest::setDefaultData<float>(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<float>(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<float>(nullptr, nullptr, nullptr,
|
||||
&A_h, &hData, nullptr,
|
||||
width*NUM_H, false);
|
||||
HipTest::setDefaultData<float>(width*NUM_H, A_h, hData, nullptr);
|
||||
hipChannelFormatDesc desc = hipCreateChannelDesc<float>();
|
||||
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<float>(nullptr, nullptr, nullptr,
|
||||
A_h, hData, nullptr, false);
|
||||
}
|
||||
|
||||
@@ -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 <hip/hip_runtime_api.h>
|
||||
#include <hip_test_common.hh>
|
||||
#include <hip_test_checkers.hh>
|
||||
#include <iostream>
|
||||
#include <resource_guards.hh>
|
||||
#include <utils.hh>
|
||||
|
||||
|
||||
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<float>(nullptr, nullptr, nullptr,
|
||||
&A_h, &hData, nullptr,
|
||||
width*NUM_H, false);
|
||||
hipChannelFormatDesc desc = hipCreateChannelDesc<float>();
|
||||
HIP_CHECK(hipMallocArray(&A_d, &desc, NUM_W, NUM_H, hipArrayDefault));
|
||||
HipTest::setDefaultData<float>(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<float>(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<float>(nullptr, nullptr, nullptr,
|
||||
&A_h, &hData, nullptr,
|
||||
width*NUM_H, false);
|
||||
hipChannelFormatDesc desc = hipCreateChannelDesc<float>();
|
||||
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<true, int>(
|
||||
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<float>(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<float>(nullptr, nullptr, nullptr,
|
||||
&A_h, nullptr, nullptr,
|
||||
width*NUM_H, false);
|
||||
HIP_CHECK(hipHostMalloc(reinterpret_cast<void**>(&PinnMem), width * NUM_H));
|
||||
hipChannelFormatDesc desc = hipCreateChannelDesc<float>();
|
||||
HIP_CHECK(hipMallocArray(&A_d, &desc, NUM_W, NUM_H, hipArrayDefault));
|
||||
HipTest::setDefaultData<float>(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<true, int>(
|
||||
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<float>(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<float>(nullptr, nullptr, nullptr,
|
||||
&A_h, nullptr, nullptr,
|
||||
width*NUM_H, false);
|
||||
hipChannelFormatDesc desc = hipCreateChannelDesc<float>();
|
||||
HIP_CHECK(hipMallocArray(&A_d, &desc, NUM_W, NUM_H, hipArrayDefault));
|
||||
HipTest::setDefaultData<float>(width*NUM_H, A_h, nullptr, nullptr);
|
||||
HIP_CHECK(hipSetDevice(1));
|
||||
HIP_CHECK(hipHostMalloc(reinterpret_cast<void**>(&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<float>(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<true, false, int>(
|
||||
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<float>(nullptr, nullptr, nullptr,
|
||||
&A_h, &hData, nullptr,
|
||||
width*NUM_H, false);
|
||||
hipChannelFormatDesc desc = hipCreateChannelDesc<float>();
|
||||
HIP_CHECK(hipMallocArray(&A_d, &desc, NUM_W, NUM_H, hipArrayDefault));
|
||||
HipTest::setDefaultData<float>(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<float>(nullptr, nullptr, nullptr,
|
||||
A_h, hData, nullptr, false);
|
||||
} else {
|
||||
SUCCEED("Machine Does not have P2P capability");
|
||||
SECTION("Peer access enabled") {
|
||||
Memcpy2DDevicetoAShell<true, true, int>(
|
||||
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<float>(nullptr, nullptr, nullptr,
|
||||
&A_h, &hData, nullptr,
|
||||
width*NUM_H, false);
|
||||
HipTest::setDefaultData<float>(width*NUM_H, A_h, hData, nullptr);
|
||||
hipChannelFormatDesc desc = hipCreateChannelDesc<float>();
|
||||
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<true, false, int>(
|
||||
std::bind(hipMemcpy2DToArrayAsync, _1, 0, 0, _2, _3, width * sizeof(int), height,
|
||||
hipMemcpyDefault, stream),
|
||||
width, height, stream);
|
||||
}
|
||||
SECTION("Peer access enabled") {
|
||||
Memcpy2DDevicetoAShell<true, true, int>(
|
||||
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<float>(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<false>(
|
||||
std::bind(hipMemcpy2DToArrayAsync, _1, 0, 0, _2, _3, width * sizeof(int), 0,
|
||||
hipMemcpyHostToDevice, stream),
|
||||
width, height, stream);
|
||||
}
|
||||
SECTION("Width is 0") {
|
||||
Memcpy2DToArrayZeroWidthHeight<false>(std::bind(hipMemcpy2DToArrayAsync, _1, 0, 0, _2, _3, 0,
|
||||
height, hipMemcpyHostToDevice, stream),
|
||||
width, height, stream);
|
||||
}
|
||||
}
|
||||
SECTION("Array to device") {
|
||||
SECTION("Height is 0") {
|
||||
Memcpy2DToArrayZeroWidthHeight<false>(
|
||||
std::bind(hipMemcpy2DToArrayAsync, _1, 0, 0, _2, _3, width * sizeof(int), 0,
|
||||
hipMemcpyDeviceToDevice, stream),
|
||||
width, height, stream);
|
||||
}
|
||||
SECTION("Width is 0") {
|
||||
Memcpy2DToArrayZeroWidthHeight<false>(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<int> array_alloc(make_hipExtent(width, height, 0), flag);
|
||||
LinearAllocGuard2D<int> device_alloc(width, height);
|
||||
LinearAllocGuard<int> 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<hipMemcpyKind>(-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<hipMemcpyKind>(-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
|
||||
}
|
||||
}
|
||||
|
||||
@@ -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 <hip_test_common.hh>
|
||||
#include <hip_test_checkers.hh>
|
||||
#include <iostream>
|
||||
|
||||
|
||||
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<float>(nullptr, nullptr, nullptr,
|
||||
&A_h, &hData, nullptr,
|
||||
width*NUM_H, false);
|
||||
hipChannelFormatDesc desc = hipCreateChannelDesc<float>();
|
||||
HIP_CHECK(hipMallocArray(&A_d, &desc, NUM_W, NUM_H, hipArrayDefault));
|
||||
HipTest::setDefaultData<float>(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<float>(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<float>(nullptr, nullptr, nullptr,
|
||||
&A_h, &hData, nullptr,
|
||||
width*NUM_H, false);
|
||||
hipChannelFormatDesc desc = hipCreateChannelDesc<float>();
|
||||
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<float>(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<float>(nullptr, nullptr, nullptr,
|
||||
&A_h, nullptr, nullptr,
|
||||
width*NUM_H, false);
|
||||
HIP_CHECK(hipHostMalloc(reinterpret_cast<void**>(&PinnMem), width * NUM_H));
|
||||
hipChannelFormatDesc desc = hipCreateChannelDesc<float>();
|
||||
HIP_CHECK(hipMallocArray(&A_d, &desc, NUM_W, NUM_H, hipArrayDefault));
|
||||
HipTest::setDefaultData<float>(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<float>(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<float>(nullptr, nullptr, nullptr,
|
||||
&A_h, nullptr, nullptr,
|
||||
width*NUM_H, false);
|
||||
hipChannelFormatDesc desc = hipCreateChannelDesc<float>();
|
||||
HIP_CHECK(hipMallocArray(&A_d, &desc, NUM_W, NUM_H, hipArrayDefault));
|
||||
HipTest::setDefaultData<float>(width*NUM_H, A_h, nullptr, nullptr);
|
||||
HIP_CHECK(hipSetDevice(1));
|
||||
HIP_CHECK(hipHostMalloc(reinterpret_cast<void**>(&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<float>(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<float>(nullptr, nullptr, nullptr,
|
||||
&A_h, &hData, nullptr,
|
||||
width*NUM_H, false);
|
||||
hipChannelFormatDesc desc = hipCreateChannelDesc<float>();
|
||||
HIP_CHECK(hipMallocArray(&A_d, &desc, NUM_W, NUM_H, hipArrayDefault));
|
||||
HipTest::setDefaultData<float>(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<float>(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<float>(nullptr, nullptr, nullptr,
|
||||
&A_h, &hData, nullptr,
|
||||
width*NUM_H, false);
|
||||
HipTest::setDefaultData<float>(width*NUM_H, A_h, hData, nullptr);
|
||||
hipChannelFormatDesc desc = hipCreateChannelDesc<float>();
|
||||
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<float>(nullptr, nullptr, nullptr,
|
||||
A_h, hData, nullptr, false);
|
||||
}
|
||||
|
||||
Reference in New Issue
Block a user