AI/rocm-systems
2
0
Forc 0
Cód Saincheisteanna Iarratais Tarraingthe Gníomhartha Pacáistí Tionscadail Scaoileann Vicí Gníomhaíocht

EXSWHTEC-91 - Implement tests for memcpy of 1D/2D hipArray (#32)

Brabhsáil Foinse 
- Implement tests for hipMemcpyAtoH using resource guards and templates
- Implement tests for hipMemcpyHtoA using resource guards and templates
- Implement tests for hipMemcpy2DFromArray using resource guards and templates
- Implement tests for hipMemcpy2DToArray using resource guards and templates
Tá an tiomantas seo le fáil i:
nives-vukovic
2023-01-17 12:56:17 +01:00
tiomanta ag GitHub
tuismitheoir 818923bfbc
tiomantas 9d62df85a1
D'athraigh 9 comhad le 1720 breiseanna agus 956 scriosta
Íoslódáil an comhad paiste Íoslódáil Comhad Diff Leathnaigh gach comhaid Laghdaigh gach comhaid
catch/unit/memory/CMakeLists.txt
+8
Féach ar an gComhad
@@ -25,6 +25,7 @@ set(TEST_SRC
memset.cc
malloc.cc
hipMemcpy2DToArray.cc
hipMemcpy2DToArray_old.cc
hipMemcpy2DToArrayAsync.cc
hipMemcpy2DToArrayAsync_old.cc
hipMemcpy3D.cc
@@ -34,10 +35,13 @@ set(TEST_SRC
hipMemcpy2D.cc
hipMemcpy2DAsync.cc
hipMemcpy2DFromArray.cc
hipMemcpy2DFromArray_old.cc
hipMemcpy2DFromArrayAsync.cc
hipMemcpy2DFromArrayAsync_old.cc
hipMemcpyAtoH.cc
hipMemcpyAtoH_old.cc
hipMemcpyHtoA.cc
hipMemcpyHtoA_old.cc
hipMemcpyAllApiNegative.cc
hipMemcpy_MultiThread.cc
hipHostRegister.cc
@@ -112,6 +116,7 @@ set(TEST_SRC
memset.cc
malloc.cc
hipMemcpy2DToArray.cc
hipMemcpy2DToArray_old.cc
hipMemcpy2DToArrayAsync.cc
hipMemcpy2DToArrayAsync_old.cc
hipMemcpy3D.cc
@@ -121,10 +126,13 @@ set(TEST_SRC
hipMemcpy2D.cc
hipMemcpy2DAsync.cc
hipMemcpy2DFromArray.cc
hipMemcpy2DFromArray_old.cc
hipMemcpy2DFromArrayAsync.cc
hipMemcpy2DFromArrayAsync_old.cc
hipMemcpyAtoH.cc
hipMemcpyAtoH_old.cc
hipMemcpyHtoA.cc
hipMemcpyHtoA_old.cc
hipMemcpyAllApiNegative.cc
hipMemcpy_MultiThread.cc
hipHostRegister.cc
catch/unit/memory/hipMemcpy2DFromArray.cc
+219 -296
Féach ar an gComhad
@@ -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,316 +16,239 @@ 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 hipMemcpy2DFromArray API
1. Negative Scenarios
2. Extent Validation Scenarios
3. hipMemcpy2DFromArray Basic Scenario
4. Pinned Memory scenarios on same and peer GPU
5. Device Context change scenario where memory is allocated in
one GPU and API is triggered from peer GPU.
Testcase Scenarios :
Unit_hipMemcpy2DFromArray_Positive_Default - Test basic memcpy between 2D array
and host/device with hipMemcpy2DFromArray api
Unit_hipMemcpy2DFromArray_Positive_Synchronization_Behavior - Test
synchronization behavior for hipMemcpy2DFromArray api
Unit_hipMemcpy2DFromArray_Positive_ZeroWidthHeight - Test that no data is copied
when width/height is set to 0 Unit_hipMemcpy2DFromArray_Negative_Parameters -
Test unsuccessful execution of hipMemcpy2DFromArray 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>
static constexpr auto NUM_W{10};
static constexpr auto NUM_H{10};
/*
* This testcase verifies device to host copy for hipMemcpy2DFromArray 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_hipMemcpy2DFromArray_Basic") {
HIP_CHECK(hipSetDevice(0));
hipArray *A_d{nullptr};
size_t width{sizeof(float)*NUM_W};
float *A_h{nullptr}, *hData{nullptr};
// 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(hipMemcpy2DToArray(A_d, 0, 0, hData, width,
width, NUM_H,
hipMemcpyHostToDevice));
TEST_CASE("Unit_hipMemcpy2DFromArray_Positive_Default") {
using namespace std::placeholders;
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));
HipTest::freeArrays<float>(nullptr, nullptr, nullptr,
A_h, hData, nullptr, false);
}
/*
* This testcase verifies the extent validation scenarios
* of hipMemcpy2DFromArray API
*/
TEST_CASE("Unit_hipMemcpy2DFromArray_ExtentValidation") {
HIP_CHECK(hipSetDevice(0));
hipArray *A_d{nullptr};
size_t width{sizeof(float)*NUM_W};
float *A_h{nullptr}, *hData{nullptr}, *valData{nullptr};
// 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));
SECTION("Destination width is 0") {
REQUIRE(hipMemcpy2DFromArray(A_h, 0, A_d,
0, 0, NUM_W*sizeof(float),
NUM_H, hipMemcpyDeviceToHost) != hipSuccess);
}
// hipMemcpy2DFromArray 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(hipMemcpy2DFromArray(hData, width, A_d,
0, 0, width,
0, hipMemcpyDeviceToHost));
REQUIRE(HipTest::checkArray(hData, valData, NUM_W, NUM_H) == true);
}
// hipMemcpy2DFromArray 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(hipMemcpy2DFromArray(hData, width, A_d,
0, 0, 0,
NUM_H, hipMemcpyDeviceToHost));
REQUIRE(HipTest::checkArray(hData, valData, NUM_W, NUM_H) == true);
SECTION("Array to host") {
Memcpy2DHostFromAShell<false, int>(
std::bind(hipMemcpy2DFromArray, _1, _2, _3, 0, 0, width * sizeof(int), height,
hipMemcpyDeviceToHost),
width, height);
}
// Cleaning the memory
HIP_CHECK(hipFreeArray(A_d));
HipTest::freeArrays<float>(nullptr, nullptr, nullptr,
A_h, hData, nullptr, false);
HipTest::freeArrays<float>(nullptr, nullptr, nullptr,
nullptr, valData, nullptr, false);
}
/*
* This Scenario Verifies hipMemcpy2DFromArray 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_hipMemcpy2DFromArray_PinnedMemSameGPU") {
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};
// 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<false, int>(std::bind(hipMemcpy2DFromArray, _1, _2, _3, 0, 0,
width * sizeof(int), height, hipMemcpyDefault),
width, height);
}
HIP_CHECK(hipMemcpy2DToArray(A_d, 0, 0, PinnMem,
width, width, NUM_H,
hipMemcpyHostToDevice));
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));
HipTest::freeArrays<float>(nullptr, nullptr, nullptr,
A_h, nullptr, nullptr, false);
}
/*
* This Scenario Verifies hipMemcpy2DFromArray 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: Intializa A_d with A_h
* Copy A_d->E_h which is a pinned host memory
* OUTPUT: For validating the result,Copying A_d device variable
* --> E_h host variable
* and verifying A_h with E_h
*/
TEST_CASE("Unit_hipMemcpy2DFromArray_multiDevicePinnedMemPeerGpu") {
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};
// 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(hipMemcpy2DToArray(A_d, 0, 0, A_h,
width, width, NUM_H,
hipMemcpyHostToDevice));
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(hipSetDevice(1));
HIP_CHECK(hipMemcpy2DFromArray(E_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));
HipTest::freeArrays<float>(nullptr, nullptr, nullptr,
A_h, nullptr, nullptr, false);
} else {
SUCCEED("Device Does not have P2P capability");
#if HT_NVIDIA // EXSWHTEC-120
SECTION("Array to device") {
SECTION("Peer access disabled") {
Memcpy2DDeviceFromAShell<false, false, int>(
std::bind(hipMemcpy2DFromArray, _1, _2, _3, 0, 0, width * sizeof(int), height,
hipMemcpyDeviceToDevice),
width, height);
}
} else {
SUCCEED("Number of devices are < 2");
}
}
/*
* This scenario verifies the hipMemcpy2DFromArray 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_hipMemcpy2DFromArray_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};
// 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(hipMemcpy2DToArray(A_d, 0, 0, hData, width,
width, NUM_H,
hipMemcpyHostToDevice));
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));
HipTest::freeArrays<float>(nullptr, nullptr, nullptr,
A_h, hData, nullptr, false);
} else {
SUCCEED("Device Does not have P2P capability");
SECTION("Peer access enabled") {
Memcpy2DDeviceFromAShell<false, true, int>(
std::bind(hipMemcpy2DFromArray, _1, _2, _3, 0, 0, width * sizeof(int), height,
hipMemcpyDeviceToDevice),
width, height);
}
} else {
SUCCEED("Number of devices are < 2");
}
}
/* This testcase verifies the negative scenarios of
* hipMemcpy2DFromArray API
*/
TEST_CASE("Unit_hipMemcpy2DFromArray_Negative") {
HIP_CHECK(hipSetDevice(0));
hipArray *A_d{nullptr};
size_t width{sizeof(float)*NUM_W};
float *A_h{nullptr}, *hData{nullptr};
// 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(hipMemcpy2DFromArray(nullptr, width, A_d,
0, 0, width, NUM_H,
hipMemcpyDeviceToHost) != hipSuccess);
}
SECTION("Nullptr to source") {
REQUIRE(hipMemcpy2DFromArray(A_h, width, nullptr,
0, 0, width, NUM_H,
hipMemcpyDeviceToHost) != hipSuccess);
SECTION("Array to device with default kind") {
SECTION("Peer access disabled") {
Memcpy2DDeviceFromAShell<false, false, int>(
std::bind(hipMemcpy2DFromArray, _1, _2, _3, 0, 0, width * sizeof(int), height,
hipMemcpyDefault),
width, height);
}
SECTION("Peer access enabled") {
Memcpy2DDeviceFromAShell<false, true, int>(
std::bind(hipMemcpy2DFromArray, _1, _2, _3, 0, 0, width * sizeof(int), height,
hipMemcpyDefault),
width, height);
}
}
SECTION("Passing offset more than 0") {
REQUIRE(hipMemcpy2DFromArray(A_h, width, A_d, 1,
1, width, NUM_H,
hipMemcpyDeviceToHost) != hipSuccess);
}
SECTION("Passing array more than allocated") {
REQUIRE(hipMemcpy2DFromArray(A_h, width, A_d, 0,
0, width+2, NUM_H+2,
hipMemcpyDeviceToHost) != hipSuccess);
}
// Cleaning of memory
HIP_CHECK(hipFreeArray(A_d));
HipTest::freeArrays<float>(nullptr, nullptr, nullptr,
A_h, hData, nullptr, false);
#endif
}
TEST_CASE("Unit_hipMemcpy2DFromArray_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(hipMemcpy2DFromArray, _1, width * sizeof(int), _2, 0,
0, width * sizeof(int), height, hipMemcpyDeviceToHost),
width, height, true);
MemcpyAtoHPinnedSyncBehavior(std::bind(hipMemcpy2DFromArray, _1, width * sizeof(int), _2, 0, 0,
width * sizeof(int), height, hipMemcpyDeviceToHost),
width, height, true);
}
#if HT_NVIDIA // EXSWHTEC-214
SECTION("Array to device") {
const auto width = GENERATE(16, 32, 48);
const auto height = GENERATE(16, 32, 48);
MemcpyAtoDSyncBehavior(std::bind(hipMemcpy2DFromArray, _1, _2, _3, 0, 0, width * sizeof(int),
height, hipMemcpyDeviceToDevice),
width, height, false);
}
#endif
}
TEST_CASE("Unit_hipMemcpy2DFromArray_Positive_ZeroWidthHeight") {
using namespace std::placeholders;
const auto width = 16;
const auto height = 16;
SECTION("Array to host") {
SECTION("Height is 0") {
Memcpy2DFromArrayZeroWidthHeight<false>(
std::bind(hipMemcpy2DFromArray, _1, _2, _3, 0, 0, width * sizeof(int), 0,
hipMemcpyDeviceToHost),
width, height);
}
SECTION("Width is 0") {
Memcpy2DFromArrayZeroWidthHeight<false>(
std::bind(hipMemcpy2DFromArray, _1, _2, _3, 0, 0, 0, height, hipMemcpyDeviceToHost),
width, height);
}
}
SECTION("Array to device") {
SECTION("Height is 0") {
Memcpy2DFromArrayZeroWidthHeight<false>(
std::bind(hipMemcpy2DFromArray, _1, _2, _3, 0, 0, width * sizeof(int), 0,
hipMemcpyDeviceToDevice),
width, height);
}
SECTION("Width is 0") {
Memcpy2DFromArrayZeroWidthHeight<false>(
std::bind(hipMemcpy2DFromArray, _1, _2, _3, 0, 0, 0, height, hipMemcpyDeviceToDevice),
width, height);
}
}
}
TEST_CASE("Unit_hipMemcpy2DFromArray_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;
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(hipMemcpy2DFromArray(nullptr, 2 * width * sizeof(int), array_alloc.ptr(), 0,
0, width * sizeof(int), height, hipMemcpyDeviceToHost),
hipErrorInvalidValue);
}
SECTION("src == nullptr") {
HIP_CHECK_ERROR(hipMemcpy2DFromArray(host_alloc.ptr(), 2 * width * sizeof(int), nullptr, 0, 0,
width * sizeof(int), height, hipMemcpyDeviceToHost),
hipErrorInvalidHandle);
}
#if HT_NVIDIA // EXSWHTEC-119
SECTION("dpitch < width") {
HIP_CHECK_ERROR(
hipMemcpy2DFromArray(host_alloc.ptr(), width * sizeof(int) - 10, array_alloc.ptr(), 0, 0,
width * sizeof(int), height, hipMemcpyDeviceToHost),
hipErrorInvalidPitchValue);
}
SECTION("Offset + width/height overflows") {
HIP_CHECK_ERROR(
hipMemcpy2DFromArray(host_alloc.ptr(), 2 * width * sizeof(int), array_alloc.ptr(), 1, 0,
width * sizeof(int), height, hipMemcpyDeviceToHost),
hipErrorInvalidValue);
HIP_CHECK_ERROR(
hipMemcpy2DFromArray(host_alloc.ptr(), 2 * width * sizeof(int), array_alloc.ptr(), 0, 1,
width * sizeof(int), height, hipMemcpyDeviceToHost),
hipErrorInvalidValue);
}
SECTION("Width/height overflows") {
HIP_CHECK_ERROR(
hipMemcpy2DFromArray(host_alloc.ptr(), 2 * width * sizeof(int), array_alloc.ptr(), 0, 0,
width * sizeof(int) + 1, height, hipMemcpyDeviceToHost),
hipErrorInvalidValue);
HIP_CHECK_ERROR(
hipMemcpy2DFromArray(host_alloc.ptr(), 2 * width * sizeof(int), array_alloc.ptr(), 0, 0,
width * sizeof(int), height + 1, hipMemcpyDeviceToHost),
hipErrorInvalidValue);
}
SECTION("Memcpy kind is invalid") {
HIP_CHECK_ERROR(
hipMemcpy2DFromArray(host_alloc.ptr(), 2 * width * sizeof(int), array_alloc.ptr(), 0, 0,
width * sizeof(int), height, static_cast<hipMemcpyKind>(-1)),
hipErrorInvalidMemcpyDirection);
}
#endif
}
SECTION("Array to device") {
SECTION("dst == nullptr") {
HIP_CHECK_ERROR(hipMemcpy2DFromArray(nullptr, device_alloc.pitch(), array_alloc.ptr(), 0, 0,
width * sizeof(int), height, hipMemcpyDeviceToDevice),
hipErrorInvalidValue);
}
SECTION("src == nullptr") {
HIP_CHECK_ERROR(hipMemcpy2DFromArray(device_alloc.ptr(), device_alloc.pitch(), nullptr, 0, 0,
width * sizeof(int), height, hipMemcpyDeviceToDevice),
hipErrorInvalidHandle);
}
#if HT_NVIDIA // EXSWHTEC-119
SECTION("dpitch < width") {
HIP_CHECK_ERROR(
hipMemcpy2DFromArray(device_alloc.ptr(), width * sizeof(int) - 10, array_alloc.ptr(), 0,
0, width * sizeof(int), height, hipMemcpyDeviceToDevice),
hipErrorInvalidPitchValue);
}
SECTION("Offset + width/height overflows") {
HIP_CHECK_ERROR(
hipMemcpy2DFromArray(device_alloc.ptr(), device_alloc.pitch(), array_alloc.ptr(), 1, 0,
width * sizeof(int), height, hipMemcpyDeviceToDevice),
hipErrorInvalidValue);
HIP_CHECK_ERROR(
hipMemcpy2DFromArray(device_alloc.ptr(), device_alloc.pitch(), array_alloc.ptr(), 0, 1,
width * sizeof(int), height, hipMemcpyDeviceToDevice),
hipErrorInvalidValue);
}
SECTION("Width/height overflows") {
HIP_CHECK_ERROR(
hipMemcpy2DFromArray(device_alloc.ptr(), device_alloc.pitch(), array_alloc.ptr(), 0, 0,
width * sizeof(int) + 1, height, hipMemcpyDeviceToDevice),
hipErrorInvalidValue);
HIP_CHECK_ERROR(
hipMemcpy2DFromArray(device_alloc.ptr(), device_alloc.pitch(), array_alloc.ptr(), 0, 0,
width * sizeof(int), height + 1, hipMemcpyDeviceToDevice),
hipErrorInvalidValue);
}
SECTION("Memcpy kind is invalid") {
HIP_CHECK_ERROR(
hipMemcpy2DFromArray(device_alloc.ptr(), device_alloc.pitch(), array_alloc.ptr(), 0, 0,
width * sizeof(int), height, static_cast<hipMemcpyKind>(-1)),
hipErrorInvalidMemcpyDirection);
}
#endif
}
}
catch/unit/memory/hipMemcpy2DFromArray_old.cc
+331
Féach ar an gComhad
@@ -0,0 +1,331 @@
/*
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 hipMemcpy2DFromArray API
1. Negative Scenarios
2. Extent Validation Scenarios
3. hipMemcpy2DFromArray Basic Scenario
4. Pinned Memory scenarios on same and peer GPU
5. Device Context change scenario where memory is allocated in
one GPU and API is triggered from 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 verifies device to host copy for hipMemcpy2DFromArray 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_hipMemcpy2DFromArray_Basic") {
HIP_CHECK(hipSetDevice(0));
hipArray *A_d{nullptr};
size_t width{sizeof(float)*NUM_W};
float *A_h{nullptr}, *hData{nullptr};
// 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(hipMemcpy2DToArray(A_d, 0, 0, hData, width,
width, NUM_H,
hipMemcpyHostToDevice));
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));
HipTest::freeArrays<float>(nullptr, nullptr, nullptr,
A_h, hData, nullptr, false);
}
/*
* This testcase verifies the extent validation scenarios
* of hipMemcpy2DFromArray API
*/
TEST_CASE("Unit_hipMemcpy2DFromArray_ExtentValidation") {
HIP_CHECK(hipSetDevice(0));
hipArray *A_d{nullptr};
size_t width{sizeof(float)*NUM_W};
float *A_h{nullptr}, *hData{nullptr}, *valData{nullptr};
// 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));
SECTION("Destination width is 0") {
REQUIRE(hipMemcpy2DFromArray(A_h, 0, A_d,
0, 0, NUM_W*sizeof(float),
NUM_H, hipMemcpyDeviceToHost) != hipSuccess);
}
// hipMemcpy2DFromArray 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(hipMemcpy2DFromArray(hData, width, A_d,
0, 0, width,
0, hipMemcpyDeviceToHost));
REQUIRE(HipTest::checkArray(hData, valData, NUM_W, NUM_H) == true);
}
// hipMemcpy2DFromArray 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(hipMemcpy2DFromArray(hData, width, A_d,
0, 0, 0,
NUM_H, hipMemcpyDeviceToHost));
REQUIRE(HipTest::checkArray(hData, valData, NUM_W, NUM_H) == true);
}
// Cleaning the memory
HIP_CHECK(hipFreeArray(A_d));
HipTest::freeArrays<float>(nullptr, nullptr, nullptr,
A_h, hData, nullptr, false);
HipTest::freeArrays<float>(nullptr, nullptr, nullptr,
nullptr, valData, nullptr, false);
}
/*
* This Scenario Verifies hipMemcpy2DFromArray 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_hipMemcpy2DFromArray_PinnedMemSameGPU") {
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};
// 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(hipMemcpy2DToArray(A_d, 0, 0, PinnMem,
width, width, NUM_H,
hipMemcpyHostToDevice));
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));
HipTest::freeArrays<float>(nullptr, nullptr, nullptr,
A_h, nullptr, nullptr, false);
}
/*
* This Scenario Verifies hipMemcpy2DFromArray 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: Intializa A_d with A_h
* Copy A_d->E_h which is a pinned host memory
* OUTPUT: For validating the result,Copying A_d device variable
* --> E_h host variable
* and verifying A_h with E_h
*/
TEST_CASE("Unit_hipMemcpy2DFromArray_multiDevicePinnedMemPeerGpu") {
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};
// 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(hipMemcpy2DToArray(A_d, 0, 0, A_h,
width, width, NUM_H,
hipMemcpyHostToDevice));
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(hipSetDevice(1));
HIP_CHECK(hipMemcpy2DFromArray(E_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));
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 hipMemcpy2DFromArray 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_hipMemcpy2DFromArray_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};
// 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(hipMemcpy2DToArray(A_d, 0, 0, hData, width,
width, NUM_H,
hipMemcpyHostToDevice));
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));
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
* hipMemcpy2DFromArray API
*/
TEST_CASE("Unit_hipMemcpy2DFromArray_Negative") {
HIP_CHECK(hipSetDevice(0));
hipArray *A_d{nullptr};
size_t width{sizeof(float)*NUM_W};
float *A_h{nullptr}, *hData{nullptr};
// 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(hipMemcpy2DFromArray(nullptr, width, A_d,
0, 0, width, NUM_H,
hipMemcpyDeviceToHost) != hipSuccess);
}
SECTION("Nullptr to source") {
REQUIRE(hipMemcpy2DFromArray(A_h, width, nullptr,
0, 0, width, NUM_H,
hipMemcpyDeviceToHost) != hipSuccess);
}
SECTION("Passing offset more than 0") {
REQUIRE(hipMemcpy2DFromArray(A_h, width, A_d, 1,
1, width, NUM_H,
hipMemcpyDeviceToHost) != hipSuccess);
}
SECTION("Passing array more than allocated") {
REQUIRE(hipMemcpy2DFromArray(A_h, width, A_d, 0,
0, width+2, NUM_H+2,
hipMemcpyDeviceToHost) != hipSuccess);
}
// Cleaning of memory
HIP_CHECK(hipFreeArray(A_d));
HipTest::freeArrays<float>(nullptr, nullptr, nullptr,
A_h, hData, nullptr, false);
}
catch/unit/memory/hipMemcpy2DToArray.cc
+214 -297
Féach ar an gComhad
@@ -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,317 +16,234 @@ 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 hipMemcpy2DToArray API
1. Negative Scenarios
2. Extent Validation Scenarios
3. hipMemcpy2DToArray Basic Scenario
4. Pinned Memory scenarios on same and peer GPU
5. Device Context change scenario where memory is allocated in
one GPU and API is triggered from peer GPU.
Testcase Scenarios :
Unit_hipMemcpy2DToArray_Positive_Default - Test basic memcpy between host/device
and 2D array with hipMemcpy2DToArray api
Unit_hipMemcpy2DToArray_Positive_Synchronization_Behavior - Test synchronization
behavior for hipMemcpy2DToArray api
Unit_hipMemcpy2DToArray_Positive_ZeroWidthHeight - Test that no data is copied
when width/height is set to 0 Unit_hipMemcpy2DToArray_Negative_Parameters - Test
unsuccessful execution of hipMemcpy2DToArray 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};
/*
* 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_hipMemcpy2DToArray_Basic") {
HIP_CHECK(hipSetDevice(0));
hipArray *A_d{nullptr};
size_t width{sizeof(float)*NUM_W};
float *A_h{nullptr}, *hData{nullptr};
// 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(hipMemcpy2DToArray(A_d, 0, 0, hData, width,
width, NUM_H,
hipMemcpyHostToDevice));
TEST_CASE("Unit_hipMemcpy2DToArray_Positive_Default") {
using namespace std::placeholders;
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));
HipTest::freeArrays<float>(nullptr, nullptr, nullptr,
A_h, hData, nullptr, false);
}
/*
* This testcase verifies the extent validation scenarios
*/
TEST_CASE("Unit_hipMemcpy2DToArray_ExtentValidation") {
HIP_CHECK(hipSetDevice(0));
hipArray *A_d{nullptr};
size_t width{sizeof(float)*NUM_W};
float *A_h{nullptr}, *hData{nullptr};
// 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));
SECTION("Source width is 0") {
REQUIRE(hipMemcpy2DToArray(A_d, 0, 0, hData, 0,
width, NUM_H,
hipMemcpyHostToDevice) != hipSuccess);
}
// 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(hipMemcpy2DToArray(A_d, 0, 0,
A_h, width, width,
NUM_H, hipMemcpyHostToDevice));
HIP_CHECK(hipMemcpy2DToArray(A_d, 0, 0,
hData, width,
width, 0, hipMemcpyHostToDevice));
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(hipMemcpy2DToArray(A_d, 0, 0,
A_h, width, width,
NUM_H, hipMemcpyHostToDevice));
HIP_CHECK(hipMemcpy2DToArray(A_d, 0, 0,
hData, width,
0, NUM_H, hipMemcpyHostToDevice));
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<false, int>(std::bind(hipMemcpy2DToArray, _1, 0, 0, _2, _3,
width * sizeof(int), height, hipMemcpyHostToDevice),
width, height);
}
// Cleaning the memory
HIP_CHECK(hipFreeArray(A_d));
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_hipMemcpy2DToArray_PinnedMemSameGPU") {
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};
// 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<false, int>(std::bind(hipMemcpy2DToArray, _1, 0, 0, _2, _3,
width * sizeof(int), height, hipMemcpyDefault),
width, height);
}
HIP_CHECK(hipMemcpy2DToArray(A_d, 0, 0, PinnMem,
width, width, NUM_H,
hipMemcpyHostToDevice));
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));
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_hipMemcpy2DToArray_multiDevicePinnedMemPeerGpu") {
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};
// 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(hipMemcpy2DToArray(A_d, 0, 0, E_h,
width, width, NUM_H,
hipMemcpyHostToDevice));
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));
HipTest::freeArrays<float>(nullptr, nullptr, nullptr,
A_h, nullptr, nullptr, false);
} else {
SUCCEED("Machine Does not have P2P capability");
#if HT_NVIDIA // EXSWHTEC-120
SECTION("Device to Array") {
SECTION("Peer access disabled") {
Memcpy2DDevicetoAShell<false, false, int>(
std::bind(hipMemcpy2DToArray, _1, 0, 0, _2, _3, width * sizeof(int), height,
hipMemcpyDeviceToDevice),
width, height);
}
} 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_hipMemcpy2DToArray_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};
// 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(hipMemcpy2DToArray(A_d, 0, 0, hData, width,
width, NUM_H,
hipMemcpyHostToDevice));
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));
HipTest::freeArrays<float>(nullptr, nullptr, nullptr,
A_h, hData, nullptr, false);
} else {
SUCCEED("Machine Does not have P2P capability");
SECTION("Peer access enabled") {
Memcpy2DDevicetoAShell<false, true, int>(
std::bind(hipMemcpy2DToArray, _1, 0, 0, _2, _3, width * sizeof(int), height,
hipMemcpyDeviceToDevice),
width, height);
}
} else {
SUCCEED("Number of devices are < 2");
}
}
/* This testcase verifies the negative scenarios
*/
TEST_CASE("Unit_hipMemcpy2DToArray_Negative") {
HIP_CHECK(hipSetDevice(0));
hipArray *A_d{nullptr};
size_t width{sizeof(float)*NUM_W};
float *A_h{nullptr}, *hData{nullptr};
// 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(hipMemcpy2DToArray(nullptr, 0, 0, hData, width,
width, NUM_H,
hipMemcpyHostToDevice) != hipSuccess);
}
SECTION("Nullptr to source") {
REQUIRE(hipMemcpy2DToArray(A_d, 0, 0,
nullptr, width, width,
NUM_H, hipMemcpyHostToDevice) != hipSuccess);
SECTION("Device to Array with default kind") {
SECTION("Peer access disabled") {
Memcpy2DDevicetoAShell<false, false, int>(
std::bind(hipMemcpy2DToArray, _1, 0, 0, _2, _3, width * sizeof(int), height,
hipMemcpyDefault),
width, height);
}
SECTION("Peer access enabled") {
Memcpy2DDevicetoAShell<false, true, int>(
std::bind(hipMemcpy2DToArray, _1, 0, 0, _2, _3, width * sizeof(int), height,
hipMemcpyDefault),
width, height);
}
}
SECTION("Passing offset more than 0") {
REQUIRE(hipMemcpy2DToArray(A_d, 1, 1,
hData, width, width,
NUM_H, hipMemcpyHostToDevice) != hipSuccess);
}
SECTION("Passing array more than allocated") {
REQUIRE(hipMemcpy2DToArray(A_d, 0, 0,
hData, width, width+2,
NUM_H+2, hipMemcpyHostToDevice) != hipSuccess);
}
// Cleaning of memory
HIP_CHECK(hipFreeArray(A_d));
HipTest::freeArrays<float>(nullptr, nullptr, nullptr,
A_h, hData, nullptr, false);
#endif
}
TEST_CASE("Unit_hipMemcpy2DToArray_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(hipMemcpy2DToArray, _1, 0, 0, _2, width * sizeof(int),
width * sizeof(int), height, hipMemcpyHostToDevice),
width, height, true);
}
#if HT_NVIDIA // EXSWHTEC-214
SECTION("Device to Array") {
const auto width = GENERATE(16, 32, 48);
const auto height = GENERATE(16, 32, 48);
MemcpyDtoASyncBehavior(std::bind(hipMemcpy2DToArray, _1, 0, 0, _2, _3, width * sizeof(int),
height, hipMemcpyDeviceToDevice),
width, height, false);
}
#endif
}
TEST_CASE("Unit_hipMemcpy2DToArray_Positive_ZeroWidthHeight") {
using namespace std::placeholders;
const auto width = 16;
const auto height = 16;
SECTION("Array to host") {
SECTION("Height is 0") {
Memcpy2DToArrayZeroWidthHeight<false>(
std::bind(hipMemcpy2DToArray, _1, 0, 0, _2, _3, width * sizeof(int), 0,
hipMemcpyHostToDevice),
width, height);
}
SECTION("Width is 0") {
Memcpy2DToArrayZeroWidthHeight<false>(
std::bind(hipMemcpy2DToArray, _1, 0, 0, _2, _3, 0, height, hipMemcpyHostToDevice), width,
height);
}
}
SECTION("Array to device") {
SECTION("Height is 0") {
Memcpy2DToArrayZeroWidthHeight<false>(
std::bind(hipMemcpy2DToArray, _1, 0, 0, _2, _3, width * sizeof(int), 0,
hipMemcpyDeviceToDevice),
width, height);
}
SECTION("Width is 0") {
Memcpy2DToArrayZeroWidthHeight<false>(
std::bind(hipMemcpy2DToArray, _1, 0, 0, _2, _3, 0, height, hipMemcpyDeviceToDevice),
width, height);
}
}
}
TEST_CASE("Unit_hipMemcpy2DToArray_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;
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(hipMemcpy2DToArray(nullptr, 0, 0, host_alloc.ptr(), 2 * width * sizeof(int),
width * sizeof(int), height, hipMemcpyHostToDevice),
hipErrorInvalidHandle);
}
SECTION("src == nullptr") {
HIP_CHECK_ERROR(hipMemcpy2DToArray(array_alloc.ptr(), 0, 0, nullptr, 2 * width * sizeof(int),
width * sizeof(int), height, hipMemcpyHostToDevice),
hipErrorInvalidValue);
}
#if HT_NVIDIA // EXSWHTEC-119
SECTION("spitch < width") {
HIP_CHECK_ERROR(
hipMemcpy2DToArray(array_alloc.ptr(), 0, 0, host_alloc.ptr(), width * sizeof(int) - 10,
width * sizeof(int), height, hipMemcpyHostToDevice),
hipErrorInvalidPitchValue);
}
SECTION("Offset + width/height overflows") {
HIP_CHECK_ERROR(
hipMemcpy2DToArray(array_alloc.ptr(), 1, 0, host_alloc.ptr(), 2 * width * sizeof(int),
width * sizeof(int), height, hipMemcpyHostToDevice),
hipErrorInvalidValue);
HIP_CHECK_ERROR(
hipMemcpy2DToArray(array_alloc.ptr(), 0, 1, host_alloc.ptr(), 2 * width * sizeof(int),
width * sizeof(int), height, hipMemcpyHostToDevice),
hipErrorInvalidValue);
}
SECTION("Width/height overflows") {
HIP_CHECK_ERROR(
hipMemcpy2DToArray(array_alloc.ptr(), 0, 0, host_alloc.ptr(), 2 * width * sizeof(int),
width * sizeof(int) + 1, height, hipMemcpyHostToDevice),
hipErrorInvalidValue);
HIP_CHECK_ERROR(
hipMemcpy2DToArray(array_alloc.ptr(), 0, 0, host_alloc.ptr(), 2 * width * sizeof(int),
width * sizeof(int), height + 1, hipMemcpyHostToDevice),
hipErrorInvalidValue);
}
SECTION("Memcpy kind is invalid") {
HIP_CHECK_ERROR(
hipMemcpy2DToArray(array_alloc.ptr(), 0, 0, host_alloc.ptr(), 2 * width * sizeof(int),
width * sizeof(int), height, static_cast<hipMemcpyKind>(-1)),
hipErrorInvalidMemcpyDirection);
}
#endif
}
SECTION("Device to Array") {
SECTION("dst == nullptr") {
HIP_CHECK_ERROR(hipMemcpy2DToArray(nullptr, 0, 0, device_alloc.ptr(), device_alloc.pitch(),
width * sizeof(int), height, hipMemcpyDeviceToDevice),
hipErrorInvalidHandle);
}
SECTION("src == nullptr") {
HIP_CHECK_ERROR(hipMemcpy2DToArray(array_alloc.ptr(), 0, 0, nullptr, device_alloc.pitch(),
width * sizeof(int), height, hipMemcpyDeviceToDevice),
hipErrorInvalidValue);
}
#if HT_NVIDIA // EXSWHTEC-119
SECTION("spitch < width") {
HIP_CHECK_ERROR(
hipMemcpy2DToArray(array_alloc.ptr(), 0, 0, device_alloc.ptr(), width * sizeof(int) - 10,
width * sizeof(int), height, hipMemcpyDeviceToDevice),
hipErrorInvalidPitchValue);
}
SECTION("Offset + width/height overflows") {
HIP_CHECK_ERROR(
hipMemcpy2DToArray(array_alloc.ptr(), 1, 0, device_alloc.ptr(), device_alloc.pitch(),
width * sizeof(int), height, hipMemcpyDeviceToDevice),
hipErrorInvalidValue);
HIP_CHECK_ERROR(
hipMemcpy2DToArray(array_alloc.ptr(), 0, 1, device_alloc.ptr(), device_alloc.pitch(),
width * sizeof(int), height, hipMemcpyDeviceToDevice),
hipErrorInvalidValue);
}
SECTION("Width/height overflows") {
HIP_CHECK_ERROR(
hipMemcpy2DToArray(array_alloc.ptr(), 0, 0, device_alloc.ptr(), device_alloc.pitch(),
width * sizeof(int) + 1, height, hipMemcpyDeviceToDevice),
hipErrorInvalidValue);
HIP_CHECK_ERROR(
hipMemcpy2DToArray(array_alloc.ptr(), 0, 0, device_alloc.ptr(), device_alloc.pitch(),
width * sizeof(int), height + 1, hipMemcpyDeviceToDevice),
hipErrorInvalidValue);
}
SECTION("Memcpy kind is invalid") {
HIP_CHECK_ERROR(
hipMemcpy2DToArray(array_alloc.ptr(), 0, 0, device_alloc.ptr(), device_alloc.pitch(),
width * sizeof(int), height, static_cast<hipMemcpyKind>(-1)),
hipErrorInvalidMemcpyDirection);
}
#endif
}
}
catch/unit/memory/hipMemcpy2DToArray_old.cc
+332
Féach ar an gComhad
@@ -0,0 +1,332 @@
/*
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 hipMemcpy2DToArray API
1. Negative Scenarios
2. Extent Validation Scenarios
3. hipMemcpy2DToArray Basic Scenario
4. Pinned Memory scenarios on same and peer GPU
5. Device Context change scenario where memory is allocated in
one GPU and API is triggered from 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_hipMemcpy2DToArray_Basic") {
HIP_CHECK(hipSetDevice(0));
hipArray *A_d{nullptr};
size_t width{sizeof(float)*NUM_W};
float *A_h{nullptr}, *hData{nullptr};
// 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(hipMemcpy2DToArray(A_d, 0, 0, hData, width,
width, NUM_H,
hipMemcpyHostToDevice));
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));
HipTest::freeArrays<float>(nullptr, nullptr, nullptr,
A_h, hData, nullptr, false);
}
/*
* This testcase verifies the extent validation scenarios
*/
TEST_CASE("Unit_hipMemcpy2DToArray_ExtentValidation") {
HIP_CHECK(hipSetDevice(0));
hipArray *A_d{nullptr};
size_t width{sizeof(float)*NUM_W};
float *A_h{nullptr}, *hData{nullptr};
// 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));
SECTION("Source width is 0") {
REQUIRE(hipMemcpy2DToArray(A_d, 0, 0, hData, 0,
width, NUM_H,
hipMemcpyHostToDevice) != hipSuccess);
}
// 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(hipMemcpy2DToArray(A_d, 0, 0,
A_h, width, width,
NUM_H, hipMemcpyHostToDevice));
HIP_CHECK(hipMemcpy2DToArray(A_d, 0, 0,
hData, width,
width, 0, hipMemcpyHostToDevice));
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(hipMemcpy2DToArray(A_d, 0, 0,
A_h, width, width,
NUM_H, hipMemcpyHostToDevice));
HIP_CHECK(hipMemcpy2DToArray(A_d, 0, 0,
hData, width,
0, NUM_H, hipMemcpyHostToDevice));
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));
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_hipMemcpy2DToArray_PinnedMemSameGPU") {
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};
// 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(hipMemcpy2DToArray(A_d, 0, 0, PinnMem,
width, width, NUM_H,
hipMemcpyHostToDevice));
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));
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_hipMemcpy2DToArray_multiDevicePinnedMemPeerGpu") {
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};
// 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(hipMemcpy2DToArray(A_d, 0, 0, E_h,
width, width, NUM_H,
hipMemcpyHostToDevice));
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));
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_hipMemcpy2DToArray_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};
// 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(hipMemcpy2DToArray(A_d, 0, 0, hData, width,
width, NUM_H,
hipMemcpyHostToDevice));
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));
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_hipMemcpy2DToArray_Negative") {
HIP_CHECK(hipSetDevice(0));
hipArray *A_d{nullptr};
size_t width{sizeof(float)*NUM_W};
float *A_h{nullptr}, *hData{nullptr};
// 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(hipMemcpy2DToArray(nullptr, 0, 0, hData, width,
width, NUM_H,
hipMemcpyHostToDevice) != hipSuccess);
}
SECTION("Nullptr to source") {
REQUIRE(hipMemcpy2DToArray(A_d, 0, 0,
nullptr, width, width,
NUM_H, hipMemcpyHostToDevice) != hipSuccess);
}
SECTION("Passing offset more than 0") {
REQUIRE(hipMemcpy2DToArray(A_d, 1, 1,
hData, width, width,
NUM_H, hipMemcpyHostToDevice) != hipSuccess);
}
SECTION("Passing array more than allocated") {
REQUIRE(hipMemcpy2DToArray(A_d, 0, 0,
hData, width, width+2,
NUM_H+2, hipMemcpyHostToDevice) != hipSuccess);
}
// Cleaning of memory
HIP_CHECK(hipFreeArray(A_d));
HipTest::freeArrays<float>(nullptr, nullptr, nullptr,
A_h, hData, nullptr, false);
}
catch/unit/memory/hipMemcpyAtoH.cc
+85 -178
Féach ar an gComhad
@@ -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
@@ -17,164 +17,43 @@ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE.
*/
/*
* Test Scenarios:
* 1. Perform host and Pinned host Memory
* 2. Perform bytecount 0 validation for hipMemcpyAtoH API
* 3. Allocate Memory from one GPU device and call hipMemcpyAtoH from Peer
* GPU device
* 4. Perform hipMemcpyAtoH Negative Scenarios
*/
Testcase Scenarios :
Unit_hipMemcpy2DFromArray_Positive_Default - Test basic memcpy between 2D array
and host/device with hipMemcpy2DFromArray api
Unit_hipMemcpy2DFromArray_Positive_Synchronization_Behavior - Test
synchronization behavior for hipMemcpy2DFromArray api
Unit_hipMemcpy2DFromArray_Positive_ZeroWidthHeight - Test that no data is copied
when width/height is set to 0 Unit_hipMemcpy2DFromArray_Negative_Parameters -
Test unsuccessful execution of hipMemcpy2DFromArray 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_hipMemcpyAtoH_Positive_Default") {
using namespace std::placeholders;
static constexpr auto NUM_W{10};
static constexpr auto NUM_H{1};
static constexpr auto copy_bytes{2};
const auto width = GENERATE(512, 1024, 2048);
const auto allocation_size = width * sizeof(int);
/*
This testcase performs the basic and pinned host memory scenarios
of hipMemcpyAtoH API
Input: "A_d" initialized with "hData" Pi value
Output:"B_h" host variable output of hipMemcpyAtoH API
is then validated with "hData"
The same scenario is then verified with pinned host memory
*/
TEMPLATE_TEST_CASE("Unit_hipMemcpyAtoH_Basic", "[hipMemcpyAtoH]",
char, int, float) {
HIP_CHECK(hipSetDevice(0));
// 1 refers to pinned host memory scenario
auto memtype_check = GENERATE(0, 1);
hipArray *A_d;
TestType *hData{nullptr}, *B_h{nullptr};
size_t width{NUM_W * sizeof(TestType)};
// Initialization of data
if (memtype_check) {
HipTest::initArrays<TestType>(nullptr, nullptr, nullptr,
&hData, &B_h, nullptr, NUM_W, true);
} else {
HipTest::initArrays<TestType>(nullptr, nullptr, nullptr,
&hData, &B_h, nullptr, NUM_W);
}
HipTest::setDefaultData<TestType>(NUM_W, hData, B_h, nullptr);
hipChannelFormatDesc desc = hipCreateChannelDesc<TestType>();
HIP_CHECK(hipMallocArray(&A_d, &desc, NUM_W, NUM_H, hipArrayDefault));
HIP_CHECK(hipMemcpy2DToArray(A_d, 0, 0, hData, width,
width, NUM_H, hipMemcpyHostToDevice));
// Performing API call
REQUIRE(hipMemcpyAtoH(B_h, A_d, 0, copy_bytes*sizeof(TestType))
== hipSuccess);
// Validating the result
REQUIRE(HipTest::checkArray(B_h, hData, copy_bytes, NUM_H) == true);
// DeAllocating the memory
HIP_CHECK(hipFreeArray(A_d));
if (memtype_check) {
REQUIRE(HipTest::freeArrays<TestType>(nullptr, nullptr, nullptr, hData, B_h,
nullptr, true) == true);
} else {
REQUIRE(HipTest::freeArrays<TestType>(nullptr, nullptr, nullptr, hData, B_h,
nullptr, false) == true);
}
MemcpyAtoHShell<false, int>(std::bind(hipMemcpyAtoH, _1, _2, 0, allocation_size), width);
}
/*
This testcase performs the basic and pinned host memory scenarios
of hipMemcpyAtoH API
Memory is allocated in GPU-0 and the API is triggered from GPU-1
Input: "A_d" initialized with "hData" Pi value
Output:"B_h" host variable output of hipMemcpyAtoH API
is then validated with "hData"
*/
#if HT_AMD
TEMPLATE_TEST_CASE("Unit_hipMemcpyAtoH_multiDevice-PeerDeviceContext",
"[hipMemcpyAtoH]",
char, int, float) {
int numDevices = 0;
HIP_CHECK(hipGetDeviceCount(&numDevices));
if (numDevices > 1) {
int peerAccess = 0;
HIP_CHECK(hipDeviceCanAccessPeer(&peerAccess, 1, 0));
if (!peerAccess) {
SUCCEED("Skipped the test as there is no peer access");
} else {
HIP_CHECK(hipSetDevice(0));
hipArray *A_d;
TestType *hData{nullptr}, *B_h{nullptr};
size_t width{NUM_W * sizeof(TestType)};
TEST_CASE("Unit_hipMemcpyAtoH_Positive_Synchronization_Behavior") {
using namespace std::placeholders;
// Initialization of data
HipTest::initArrays<TestType>(nullptr, nullptr, nullptr,
&hData, &B_h, nullptr, NUM_W);
HipTest::setDefaultData<TestType>(NUM_W, hData, B_h, nullptr);
hipChannelFormatDesc desc = hipCreateChannelDesc<TestType>();
HIP_CHECK(hipMallocArray(&A_d, &desc, NUM_W, NUM_H, hipArrayDefault));
HIP_CHECK(hipMemcpy2DToArray(A_d, 0, 0, hData, width,
width, NUM_H, hipMemcpyHostToDevice));
const auto width = GENERATE(512, 1024, 2048);
const auto height = 0;
const auto allocation_size = width * sizeof(int);
HIP_CHECK(hipDeviceSynchronize());
// Changing the device context
HIP_CHECK(hipSetDevice(1));
// Performing API call
REQUIRE(hipMemcpyAtoH(B_h, A_d, 0, copy_bytes*sizeof(TestType))
== hipSuccess);
// Validating the result
REQUIRE(HipTest::checkArray(B_h, hData, copy_bytes, NUM_H) == true);
// DeAllocating the memory
HIP_CHECK(hipFreeArray(A_d));
REQUIRE(HipTest::freeArrays<TestType>(nullptr, nullptr, nullptr,
hData, B_h,
nullptr, false) == true);
}
} else {
SUCCEED("skipping the testcases as numDevices < 2");
}
}
#endif
/*
This testcase verifies the negative scenarios of hipMemcpyAtoH API
*/
TEST_CASE("Unit_hipMemcpyAtoH_Negative") {
HIP_CHECK(hipSetDevice(0));
hipArray *A_d;
float *hData{nullptr}, *B_h{nullptr};
size_t width{NUM_W * sizeof(float)};
// Initialization of data
HipTest::initArrays<float>(nullptr, nullptr, nullptr,
&hData, &B_h, nullptr, NUM_W);
HipTest::setDefaultData<float>(NUM_W, hData, B_h, nullptr);
hipChannelFormatDesc desc = hipCreateChannelDesc<float>();
HIP_CHECK(hipMallocArray(&A_d, &desc, NUM_W, NUM_H, hipArrayDefault));
HIP_CHECK(hipMemcpy2DToArray(A_d, 0, 0, hData, width,
width, NUM_H, hipMemcpyHostToDevice));
SECTION("Destination pointer is nullptr") {
REQUIRE(hipMemcpyAtoH(nullptr, A_d, 0, copy_bytes*sizeof(float))
!= hipSuccess);
}
SECTION("Source offset is more than allocated size") {
REQUIRE(hipMemcpyAtoH(B_h, A_d, 100, copy_bytes*sizeof(float))
!= hipSuccess);
}
SECTION("ByteCount is greater than allocated size") {
REQUIRE(hipMemcpyAtoH(B_h, A_d, 0, 12*sizeof(float)) != hipSuccess);
}
// DeAllocating the memory
HIP_CHECK(hipFreeArray(A_d));
REQUIRE(HipTest::freeArrays<float>(nullptr, nullptr, nullptr, hData, B_h,
nullptr, false) == true);
MemcpyAtoHPageableSyncBehavior(std::bind(hipMemcpyAtoH, _1, _2, 0, allocation_size), width,
height, true);
MemcpyAtoHPinnedSyncBehavior(std::bind(hipMemcpyAtoH, _1, _2, 0, allocation_size), width, height,
true);
}
/*
@@ -183,37 +62,65 @@ Excluded the testcase for amd,as there is already a bug raised
SWDEV-274683
*/
#if HT_NVIDIA
TEST_CASE("Unit_hipMemcpyAtoH_SizeCheck") {
HIP_CHECK(hipSetDevice(0));
hipArray *A_d;
float *hData{nullptr}, *B_h{nullptr}, *def_data{nullptr};
size_t width{NUM_W * sizeof(float)};
TEST_CASE("Unit_hipMemcpyAtoH_Positive_ZeroCount") {
const auto width = 1024;
const auto height = 0;
const auto allocation_size = width * sizeof(int);
// Initialization of data
HipTest::initArrays<float>(nullptr, nullptr, nullptr,
nullptr, &def_data, nullptr, NUM_W);
HipTest::initArrays<float>(nullptr, nullptr, nullptr,
&hData, &B_h, nullptr, NUM_W);
HipTest::setDefaultData<float>(NUM_W, hData, B_h, nullptr);
HipTest::setDefaultData<float>(NUM_W, nullptr, def_data, nullptr);
hipChannelFormatDesc desc = hipCreateChannelDesc<float>();
HIP_CHECK(hipMallocArray(&A_d, &desc, NUM_W, NUM_H, hipArrayDefault));
HIP_CHECK(hipMemcpy2DToArray(A_d, 0, 0, hData, width,
width, NUM_H, hipMemcpyHostToDevice));
const unsigned int flag = hipArrayDefault;
SECTION("Passing 0 to copy bytes") {
REQUIRE(hipMemcpyAtoH(B_h, A_d, 0, 0) == hipSuccess);
REQUIRE(HipTest::checkArray(B_h, def_data, NUM_W, NUM_H) == true);
}
ArrayAllocGuard<int> array_alloc(make_hipExtent(width, height, 0), flag);
LinearAllocGuard<uint8_t> host_alloc(LinearAllocs::hipHostMalloc, allocation_size);
SECTION(" Source Array is nullptr") {
REQUIRE(hipMemcpyAtoH(B_h, nullptr, 0, copy_bytes*sizeof(float))
!= hipSuccess);
}
int fill_value = 42;
std::fill_n(host_alloc.host_ptr(), width, fill_value);
HIP_CHECK(hipMemcpy2DToArray(array_alloc.ptr(), 0, 0, host_alloc.host_ptr(), sizeof(int) * width,
sizeof(int) * width, 1, hipMemcpyHostToDevice));
fill_value = 41;
std::fill_n(host_alloc.host_ptr(), width, fill_value);
HIP_CHECK(hipMemcpyAtoH(host_alloc.ptr(), array_alloc.ptr(), 0, 0));
// DeAllocating the memory
HIP_CHECK(hipFreeArray(A_d));
REQUIRE(HipTest::freeArrays<float>(nullptr, nullptr, nullptr, hData, B_h,
def_data, false) == true);
ArrayFindIfNot(host_alloc.host_ptr(), static_cast<uint8_t>(fill_value), width);
}
#endif
TEST_CASE("Unit_hipMemcpyAtoH_Negative_Parameters") {
using namespace std::placeholders;
const auto width = 1024;
const auto height = 0;
const auto allocation_size = width * sizeof(int);
const unsigned int flag = hipArrayDefault;
ArrayAllocGuard<int> array_alloc(make_hipExtent(width, height, 0), flag);
LinearAllocGuard<int> host_alloc(LinearAllocs::hipHostMalloc, allocation_size);
SECTION("dst == nullptr") {
HIP_CHECK_ERROR(hipMemcpyAtoH(nullptr, array_alloc.ptr(), 0, allocation_size),
hipErrorInvalidValue);
}
SECTION("src == nullptr") {
HIP_CHECK_ERROR(hipMemcpyAtoH(host_alloc.ptr(), nullptr, 0, allocation_size),
hipErrorInvalidValue);
}
SECTION("Offset is greater than allocated size") {
HIP_CHECK_ERROR(
hipMemcpyAtoH(host_alloc.ptr(), array_alloc.ptr(), allocation_size + 10, allocation_size),
hipErrorInvalidValue);
}
SECTION("Count is greater than allocated size") {
HIP_CHECK_ERROR(hipMemcpyAtoH(host_alloc.ptr(), array_alloc.ptr(), 0, allocation_size + 10),
hipErrorInvalidValue);
}
SECTION("2D array is allocated") {
const auto width_2d = 32;
const auto height_2d = width_2d;
const auto allocation_size_2d = width_2d * height_2d * sizeof(int);
ArrayAllocGuard<int> array_alloc_2d(make_hipExtent(width_2d, height_2d, 0), flag);
LinearAllocGuard<int> host_alloc_2d(LinearAllocs::hipHostMalloc, allocation_size_2d);
HIP_CHECK_ERROR(hipMemcpyAtoH(host_alloc_2d.ptr(), array_alloc_2d.ptr(), 0, allocation_size_2d),
hipErrorInvalidValue);
}
}
catch/unit/memory/hipMemcpyAtoH_old.cc
+219
Féach ar an gComhad
@@ -0,0 +1,219 @@
/*
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.
*/
/*
* Test Scenarios:
* 1. Perform host and Pinned host Memory
* 2. Perform bytecount 0 validation for hipMemcpyAtoH API
* 3. Allocate Memory from one GPU device and call hipMemcpyAtoH from Peer
* GPU device
* 4. Perform hipMemcpyAtoH Negative Scenarios
*/
#include <hip_test_common.hh>
#include <hip_test_checkers.hh>
static constexpr auto NUM_W{10};
static constexpr auto NUM_H{1};
static constexpr auto copy_bytes{2};
/*
This testcase performs the basic and pinned host memory scenarios
of hipMemcpyAtoH API
Input: "A_d" initialized with "hData" Pi value
Output:"B_h" host variable output of hipMemcpyAtoH API
is then validated with "hData"
The same scenario is then verified with pinned host memory
*/
TEMPLATE_TEST_CASE("Unit_hipMemcpyAtoH_Basic", "[hipMemcpyAtoH]",
char, int, float) {
HIP_CHECK(hipSetDevice(0));
// 1 refers to pinned host memory scenario
auto memtype_check = GENERATE(0, 1);
hipArray *A_d;
TestType *hData{nullptr}, *B_h{nullptr};
size_t width{NUM_W * sizeof(TestType)};
// Initialization of data
if (memtype_check) {
HipTest::initArrays<TestType>(nullptr, nullptr, nullptr,
&hData, &B_h, nullptr, NUM_W, true);
} else {
HipTest::initArrays<TestType>(nullptr, nullptr, nullptr,
&hData, &B_h, nullptr, NUM_W);
}
HipTest::setDefaultData<TestType>(NUM_W, hData, B_h, nullptr);
hipChannelFormatDesc desc = hipCreateChannelDesc<TestType>();
HIP_CHECK(hipMallocArray(&A_d, &desc, NUM_W, NUM_H, hipArrayDefault));
HIP_CHECK(hipMemcpy2DToArray(A_d, 0, 0, hData, width,
width, NUM_H, hipMemcpyHostToDevice));
// Performing API call
REQUIRE(hipMemcpyAtoH(B_h, A_d, 0, copy_bytes*sizeof(TestType))
== hipSuccess);
// Validating the result
REQUIRE(HipTest::checkArray(B_h, hData, copy_bytes, NUM_H) == true);
// DeAllocating the memory
HIP_CHECK(hipFreeArray(A_d));
if (memtype_check) {
REQUIRE(HipTest::freeArrays<TestType>(nullptr, nullptr, nullptr, hData, B_h,
nullptr, true) == true);
} else {
REQUIRE(HipTest::freeArrays<TestType>(nullptr, nullptr, nullptr, hData, B_h,
nullptr, false) == true);
}
}
/*
This testcase performs the basic and pinned host memory scenarios
of hipMemcpyAtoH API
Memory is allocated in GPU-0 and the API is triggered from GPU-1
Input: "A_d" initialized with "hData" Pi value
Output:"B_h" host variable output of hipMemcpyAtoH API
is then validated with "hData"
*/
#if HT_AMD
TEMPLATE_TEST_CASE("Unit_hipMemcpyAtoH_multiDevice-PeerDeviceContext",
"[hipMemcpyAtoH]",
char, int, float) {
int numDevices = 0;
HIP_CHECK(hipGetDeviceCount(&numDevices));
if (numDevices > 1) {
int peerAccess = 0;
HIP_CHECK(hipDeviceCanAccessPeer(&peerAccess, 1, 0));
if (!peerAccess) {
SUCCEED("Skipped the test as there is no peer access");
} else {
HIP_CHECK(hipSetDevice(0));
hipArray *A_d;
TestType *hData{nullptr}, *B_h{nullptr};
size_t width{NUM_W * sizeof(TestType)};
// Initialization of data
HipTest::initArrays<TestType>(nullptr, nullptr, nullptr,
&hData, &B_h, nullptr, NUM_W);
HipTest::setDefaultData<TestType>(NUM_W, hData, B_h, nullptr);
hipChannelFormatDesc desc = hipCreateChannelDesc<TestType>();
HIP_CHECK(hipMallocArray(&A_d, &desc, NUM_W, NUM_H, hipArrayDefault));
HIP_CHECK(hipMemcpy2DToArray(A_d, 0, 0, hData, width,
width, NUM_H, hipMemcpyHostToDevice));
HIP_CHECK(hipDeviceSynchronize());
// Changing the device context
HIP_CHECK(hipSetDevice(1));
// Performing API call
REQUIRE(hipMemcpyAtoH(B_h, A_d, 0, copy_bytes*sizeof(TestType))
== hipSuccess);
// Validating the result
REQUIRE(HipTest::checkArray(B_h, hData, copy_bytes, NUM_H) == true);
// DeAllocating the memory
HIP_CHECK(hipFreeArray(A_d));
REQUIRE(HipTest::freeArrays<TestType>(nullptr, nullptr, nullptr,
hData, B_h,
nullptr, false) == true);
}
} else {
SUCCEED("skipping the testcases as numDevices < 2");
}
}
#endif
/*
This testcase verifies the negative scenarios of hipMemcpyAtoH API
*/
TEST_CASE("Unit_hipMemcpyAtoH_Negative") {
HIP_CHECK(hipSetDevice(0));
hipArray *A_d;
float *hData{nullptr}, *B_h{nullptr};
size_t width{NUM_W * sizeof(float)};
// Initialization of data
HipTest::initArrays<float>(nullptr, nullptr, nullptr,
&hData, &B_h, nullptr, NUM_W);
HipTest::setDefaultData<float>(NUM_W, hData, B_h, nullptr);
hipChannelFormatDesc desc = hipCreateChannelDesc<float>();
HIP_CHECK(hipMallocArray(&A_d, &desc, NUM_W, NUM_H, hipArrayDefault));
HIP_CHECK(hipMemcpy2DToArray(A_d, 0, 0, hData, width,
width, NUM_H, hipMemcpyHostToDevice));
SECTION("Destination pointer is nullptr") {
REQUIRE(hipMemcpyAtoH(nullptr, A_d, 0, copy_bytes*sizeof(float))
!= hipSuccess);
}
SECTION("Source offset is more than allocated size") {
REQUIRE(hipMemcpyAtoH(B_h, A_d, 100, copy_bytes*sizeof(float))
!= hipSuccess);
}
SECTION("ByteCount is greater than allocated size") {
REQUIRE(hipMemcpyAtoH(B_h, A_d, 0, 12*sizeof(float)) != hipSuccess);
}
// DeAllocating the memory
HIP_CHECK(hipFreeArray(A_d));
REQUIRE(HipTest::freeArrays<float>(nullptr, nullptr, nullptr, hData, B_h,
nullptr, false) == true);
}
/*
This testcase verifies size 0 check of hipMemcpyAtoH API
Excluded the testcase for amd,as there is already a bug raised
SWDEV-274683
*/
#if HT_NVIDIA
TEST_CASE("Unit_hipMemcpyAtoH_SizeCheck") {
HIP_CHECK(hipSetDevice(0));
hipArray *A_d;
float *hData{nullptr}, *B_h{nullptr}, *def_data{nullptr};
size_t width{NUM_W * sizeof(float)};
// Initialization of data
HipTest::initArrays<float>(nullptr, nullptr, nullptr,
nullptr, &def_data, nullptr, NUM_W);
HipTest::initArrays<float>(nullptr, nullptr, nullptr,
&hData, &B_h, nullptr, NUM_W);
HipTest::setDefaultData<float>(NUM_W, hData, B_h, nullptr);
HipTest::setDefaultData<float>(NUM_W, nullptr, def_data, nullptr);
hipChannelFormatDesc desc = hipCreateChannelDesc<float>();
HIP_CHECK(hipMallocArray(&A_d, &desc, NUM_W, NUM_H, hipArrayDefault));
HIP_CHECK(hipMemcpy2DToArray(A_d, 0, 0, hData, width,
width, NUM_H, hipMemcpyHostToDevice));
SECTION("Passing 0 to copy bytes") {
REQUIRE(hipMemcpyAtoH(B_h, A_d, 0, 0) == hipSuccess);
REQUIRE(HipTest::checkArray(B_h, def_data, NUM_W, NUM_H) == true);
}
SECTION(" Source Array is nullptr") {
REQUIRE(hipMemcpyAtoH(B_h, nullptr, 0, copy_bytes*sizeof(float))
!= hipSuccess);
}
// DeAllocating the memory
HIP_CHECK(hipFreeArray(A_d));
REQUIRE(HipTest::freeArrays<float>(nullptr, nullptr, nullptr, hData, B_h,
def_data, false) == true);
}
#endif
catch/unit/memory/hipMemcpyHtoA.cc
+83 -185
Féach ar an gComhad
@@ -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,171 +16,41 @@ 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.
*/
/*
* Test Scenarios:
* 1. Perform simple and pinned host memory of hipMemcpyHtoA API
* 2. Allocate Memory from one GPU device and call hipMemcpyHtoA from Peer
* GPU device
* 3. Perform hipMemcpyHtoA Negative Scenarios
* 4. Perform bytecount 0 validation for hipMemcpyHtoA API
Testcase Scenarios :
Unit_hipMemcpyHtoA_Positive_Default - Test basic memcpy between host and 1D
array with hipMemcpyHtoA api
Unit_hipMemcpyHtoA_Positive_Synchronization_Behavior - Test synchronization
behavior for hipMemcpyHtoA api Unit_hipMemcpyHtoA_Positive_ZeroCount - Test that
no data is copied when allocation_size is set to 0
Unit_hipMemcpyHtoA_Negative_Parameters - Test unsuccessful execution of
hipMemcpyHtoA 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>
static constexpr auto NUM_W{10};
static constexpr auto NUM_H{1};
static constexpr auto copy_bytes{2};
TEST_CASE("Unit_hipMemcpyHtoA_Positive_Default") {
using namespace std::placeholders;
/*
This testcase performs the basic and pinned host memory scenarios
of hipMemcpyHtoA API
Input: "B_h" which is initialized with 1.6
Output: "A_d" output of hipMemcpyHtoA is copied to "hData" host variable
validated the result with "B_h"
const auto width = GENERATE(512, 1024, 2048);
const auto allocation_size = width * sizeof(int);
The same scenario is then verified with pinned host memory
*/
TEMPLATE_TEST_CASE("Unit_hipMemcpyHtoA_Basic", "[hipMemcpyHtoA]",
char, int, float) {
HIP_CHECK(hipSetDevice(0));
auto memtype_check = GENERATE(0, 1);
hipArray *A_d;
TestType *hData{nullptr}, *B_h{nullptr};
size_t width{NUM_W * sizeof(TestType)};
// Initialization of data
if (memtype_check) {
HipTest::initArrays<TestType>(nullptr, nullptr, nullptr,
&hData, &B_h, nullptr, NUM_W, true);
} else {
HipTest::initArrays<TestType>(nullptr, nullptr, nullptr,
&hData, &B_h, nullptr, NUM_W);
}
HipTest::setDefaultData<TestType>(NUM_W, hData, B_h, nullptr);
hipChannelFormatDesc desc = hipCreateChannelDesc<TestType>();
HIP_CHECK(hipMallocArray(&A_d, &desc, NUM_W, NUM_H, hipArrayDefault));
HIP_CHECK(hipMemcpy2DToArray(A_d, 0, 0, hData, width,
width, NUM_H, hipMemcpyHostToDevice));
// Performing API call
HIP_CHECK(hipMemcpyHtoA(A_d, 0, B_h, copy_bytes*sizeof(TestType)));
HIP_CHECK(hipMemcpy2DFromArray(hData, sizeof(TestType)*NUM_W, A_d,
0, 0, sizeof(TestType)*NUM_W, 1, hipMemcpyDeviceToHost));
// Validating the result
REQUIRE(HipTest::checkArray(B_h, hData, copy_bytes, NUM_H) == true);
// DeAllocating the memory
HIP_CHECK(hipFreeArray(A_d));
if (memtype_check) {
REQUIRE(HipTest::freeArrays<TestType>(nullptr, nullptr, nullptr, hData, B_h,
nullptr, true) == true);
} else {
REQUIRE(HipTest::freeArrays<TestType>(nullptr, nullptr, nullptr, hData, B_h,
nullptr, false) == true);
}
MemcpyHtoAShell<false, int>(std::bind(hipMemcpyHtoA, _1, 0, _2, allocation_size), width);
}
TEST_CASE("Unit_hipMemcpyHtoA_Positive_Synchronization_Behavior") {
using namespace std::placeholders;
/*
This testcase performs the peer device context scenario
of hipMemcpyHtoA API
Memory is allocated in GPU-0 and the API is triggered from GPU-1
Input: "B_h" which is initialized with 1.6
Output: "A_d" output of hipMemcpyHtoA is copied to "hData" host variable
validated the result with "B_h"
*/
#if HT_AMD
TEMPLATE_TEST_CASE("Unit_hipMemcpyHtoA_multiDevice-PeerDeviceContext",
"[hipMemcpyHtoA]",
char, int, float) {
int numDevices = 0;
HIP_CHECK(hipGetDeviceCount(&numDevices));
if (numDevices > 1) {
int peerAccess = 0;
HIP_CHECK(hipDeviceCanAccessPeer(&peerAccess, 1, 0));
if (!peerAccess) {
SUCCEED("Skipped the test as there is no peer access");
} else {
HIP_CHECK(hipSetDevice(0));
hipArray *A_d;
TestType *hData{nullptr}, *B_h{nullptr};
size_t width{NUM_W * sizeof(TestType)};
const auto width = GENERATE(512, 1024, 2048);
const auto height = 0;
const auto allocation_size = width * sizeof(int);
// Initialization of data
HipTest::initArrays<TestType>(nullptr, nullptr, nullptr,
&hData, &B_h, nullptr, NUM_W);
HipTest::setDefaultData<TestType>(NUM_W, hData, B_h, nullptr);
hipChannelFormatDesc desc = hipCreateChannelDesc<TestType>();
HIP_CHECK(hipMallocArray(&A_d, &desc, NUM_W, NUM_H, hipArrayDefault));
HIP_CHECK(hipMemcpy2DToArray(A_d, 0, 0, hData, width,
width, NUM_H, hipMemcpyHostToDevice));
// Changing the device context
HIP_CHECK(hipSetDevice(1));
// Performing API call
HIP_CHECK(hipMemcpyHtoA(A_d, 0, B_h, copy_bytes*sizeof(TestType)));
HIP_CHECK(hipMemcpy2DFromArray(hData, sizeof(TestType)*NUM_W, A_d,
0, 0, sizeof(TestType)*NUM_W, 1,
hipMemcpyDeviceToHost));
// Validating the result
REQUIRE(HipTest::checkArray(B_h, hData, copy_bytes, NUM_H) == true);
// DeAllocating the memory
HIP_CHECK(hipFreeArray(A_d));
REQUIRE(HipTest::freeArrays<TestType>(nullptr, nullptr, nullptr,
hData, B_h,
nullptr, false) == true);
}
} else {
SUCCEED("skipping the testcases as numDevices < 2");
}
}
#endif
/*
This testcase verifies the negative scenarios of hipMemcpyHtoA API
*/
TEST_CASE("Unit_hipMemcpyHtoA_Negative") {
HIP_CHECK(hipSetDevice(0));
hipArray *A_d;
float *hData{nullptr}, *B_h{nullptr};
size_t width{NUM_W * sizeof(float)};
// Initialization of data
HipTest::initArrays<float>(nullptr, nullptr, nullptr,
&hData, &B_h, nullptr, NUM_W);
HipTest::setDefaultData<float>(NUM_W, hData, B_h, nullptr);
hipChannelFormatDesc desc = hipCreateChannelDesc<float>();
HIP_CHECK(hipMallocArray(&A_d, &desc, NUM_W, NUM_H, hipArrayDefault));
HIP_CHECK(hipMemcpy2DToArray(A_d, 0, 0, hData, width,
width, NUM_H, hipMemcpyHostToDevice));
SECTION("Source pointer is nullptr") {
REQUIRE(hipMemcpyHtoA(A_d, 0, nullptr, copy_bytes*sizeof(float))
!= hipSuccess);
}
SECTION("Source offset is more than allocated size") {
REQUIRE(hipMemcpyHtoA(A_d, 100, B_h, copy_bytes*sizeof(float))
!= hipSuccess);
}
SECTION("ByteCount is greater than allocated size") {
REQUIRE(hipMemcpyHtoA(A_d, 0, B_h, 12*sizeof(float)) != hipSuccess);
}
// DeAllocating the memory
HIP_CHECK(hipFreeArray(A_d));
REQUIRE(HipTest::freeArrays<float>(nullptr, nullptr, nullptr, hData, B_h,
nullptr, false) == true);
MemcpyHtoASyncBehavior(std::bind(hipMemcpyHtoA, _1, 0, _2, allocation_size), width, height, true);
}
/*
@@ -189,41 +59,69 @@ This is excluded for AMD as we have a bug already raised
SWDEV-274683
*/
#if HT_NVIDIA
TEST_CASE("Unit_hipMemcpyHtoA_SizeCheck") {
HIP_CHECK(hipSetDevice(0));
hipArray *A_d;
float *hData{nullptr}, *B_h{nullptr}, *def_data{nullptr};
size_t width{NUM_W * sizeof(float)};
TEST_CASE("Unit_hipMemcpyHtoA_Positive_ZeroCount") {
const auto width = 1024;
const auto height = 0;
const auto allocation_size = width * sizeof(int);
// Initialization of data
HipTest::initArrays<float>(nullptr, nullptr, nullptr,
nullptr, &def_data, nullptr, NUM_W);
HipTest::initArrays<float>(nullptr, nullptr, nullptr,
&hData, &B_h, nullptr, NUM_W);
HipTest::setDefaultData<float>(NUM_W, hData, B_h, nullptr);
HipTest::setDefaultData<float>(NUM_W, nullptr, def_data, nullptr);
hipChannelFormatDesc desc = hipCreateChannelDesc<float>();
HIP_CHECK(hipMallocArray(&A_d, &desc, NUM_W, NUM_H, hipArrayDefault));
HIP_CHECK(hipMemcpy2DToArray(A_d, 0, 0, hData, width,
width, NUM_H, hipMemcpyHostToDevice));
const unsigned int flag = hipArrayDefault;
SECTION("Passing 0 to copy bytes") {
REQUIRE(hipMemcpyHtoA(A_d, 0, B_h, 0) == hipSuccess);
HIP_CHECK(hipMemcpy2DFromArray(def_data, sizeof(float)*NUM_W, A_d,
0, 0, sizeof(float)*NUM_W, 1,
hipMemcpyDeviceToHost));
ArrayAllocGuard<int> array_alloc(make_hipExtent(width, height, 0), flag);
LinearAllocGuard<uint8_t> host_alloc(LinearAllocs::hipHostMalloc, allocation_size);
REQUIRE(HipTest::checkArray(hData, def_data, NUM_W, NUM_H) == true);
}
int fill_value = 42;
std::fill_n(host_alloc.host_ptr(), width, fill_value);
HIP_CHECK(hipMemcpy2DToArray(array_alloc.ptr(), 0, 0, host_alloc.host_ptr(), sizeof(int) * width,
sizeof(int) * width, 1, hipMemcpyHostToDevice));
fill_value = 41;
std::fill_n(host_alloc.host_ptr(), width, fill_value);
HIP_CHECK(hipMemcpyHtoA(array_alloc.ptr(), 0, host_alloc.ptr(), 0));
SECTION(" Source Array is nullptr") {
REQUIRE(hipMemcpyHtoA(nullptr, 0, B_h, copy_bytes*sizeof(float))
!= hipSuccess);
}
HIP_CHECK(hipMemcpy2DFromArray(host_alloc.host_ptr(), sizeof(int) * width, array_alloc.ptr(), 0,
0, sizeof(int) * width, 1, hipMemcpyDeviceToHost));
// DeAllocating the memory
HIP_CHECK(hipFreeArray(A_d));
REQUIRE(HipTest::freeArrays<float>(nullptr, nullptr, nullptr, hData, B_h,
def_data, false) == true);
ArrayFindIfNot(host_alloc.host_ptr(), static_cast<uint8_t>(42), width);
}
#endif
TEST_CASE("Unit_hipMemcpyHtoA_Negative_Parameters") {
using namespace std::placeholders;
const auto width = 1024;
const auto height = 0;
const auto allocation_size = width * sizeof(int);
const unsigned int flag = hipArrayDefault;
ArrayAllocGuard<int> array_alloc(make_hipExtent(width, height, 0), flag);
LinearAllocGuard<int> host_alloc(LinearAllocs::hipHostMalloc, allocation_size);
SECTION("dst == nullptr") {
HIP_CHECK_ERROR(hipMemcpyHtoA(nullptr, 0, host_alloc.ptr(), allocation_size),
hipErrorInvalidValue);
}
SECTION("src == nullptr") {
HIP_CHECK_ERROR(hipMemcpyHtoA(array_alloc.ptr(), 0, nullptr, allocation_size),
hipErrorInvalidValue);
}
SECTION("Offset is greater than allocated size") {
HIP_CHECK_ERROR(
hipMemcpyHtoA(array_alloc.ptr(), allocation_size + 10, host_alloc.ptr(), allocation_size),
hipErrorInvalidValue);
}
SECTION("Count is greater than allocated size") {
HIP_CHECK_ERROR(hipMemcpyHtoA(array_alloc.ptr(), 0, host_alloc.ptr(), allocation_size + 10),
hipErrorInvalidValue);
}
SECTION("2D array is allocated") {
const auto width_2d = 32;
const auto height_2d = width_2d;
const auto allocation_size_2d = width_2d * height_2d * sizeof(int);
ArrayAllocGuard<int> array_alloc_2d(make_hipExtent(width_2d, height_2d, 0), flag);
LinearAllocGuard<int> host_alloc_2d(LinearAllocs::hipHostMalloc, allocation_size_2d);
HIP_CHECK_ERROR(hipMemcpyHtoA(array_alloc_2d.ptr(), 0, host_alloc_2d.ptr(), allocation_size_2d),
hipErrorInvalidValue);
}
}
catch/unit/memory/hipMemcpyHtoA_old.cc
+229
Féach ar an gComhad
@@ -0,0 +1,229 @@
/*
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.
*/
/*
* Test Scenarios:
* 1. Perform simple and pinned host memory of hipMemcpyHtoA API
* 2. Allocate Memory from one GPU device and call hipMemcpyHtoA from Peer
* GPU device
* 3. Perform hipMemcpyHtoA Negative Scenarios
* 4. Perform bytecount 0 validation for hipMemcpyHtoA API
*/
#include <hip_test_common.hh>
#include <hip_test_checkers.hh>
static constexpr auto NUM_W{10};
static constexpr auto NUM_H{1};
static constexpr auto copy_bytes{2};
/*
This testcase performs the basic and pinned host memory scenarios
of hipMemcpyHtoA API
Input: "B_h" which is initialized with 1.6
Output: "A_d" output of hipMemcpyHtoA is copied to "hData" host variable
validated the result with "B_h"
The same scenario is then verified with pinned host memory
*/
TEMPLATE_TEST_CASE("Unit_hipMemcpyHtoA_Basic", "[hipMemcpyHtoA]",
char, int, float) {
HIP_CHECK(hipSetDevice(0));
auto memtype_check = GENERATE(0, 1);
hipArray *A_d;
TestType *hData{nullptr}, *B_h{nullptr};
size_t width{NUM_W * sizeof(TestType)};
// Initialization of data
if (memtype_check) {
HipTest::initArrays<TestType>(nullptr, nullptr, nullptr,
&hData, &B_h, nullptr, NUM_W, true);
} else {
HipTest::initArrays<TestType>(nullptr, nullptr, nullptr,
&hData, &B_h, nullptr, NUM_W);
}
HipTest::setDefaultData<TestType>(NUM_W, hData, B_h, nullptr);
hipChannelFormatDesc desc = hipCreateChannelDesc<TestType>();
HIP_CHECK(hipMallocArray(&A_d, &desc, NUM_W, NUM_H, hipArrayDefault));
HIP_CHECK(hipMemcpy2DToArray(A_d, 0, 0, hData, width,
width, NUM_H, hipMemcpyHostToDevice));
// Performing API call
HIP_CHECK(hipMemcpyHtoA(A_d, 0, B_h, copy_bytes*sizeof(TestType)));
HIP_CHECK(hipMemcpy2DFromArray(hData, sizeof(TestType)*NUM_W, A_d,
0, 0, sizeof(TestType)*NUM_W, 1, hipMemcpyDeviceToHost));
// Validating the result
REQUIRE(HipTest::checkArray(B_h, hData, copy_bytes, NUM_H) == true);
// DeAllocating the memory
HIP_CHECK(hipFreeArray(A_d));
if (memtype_check) {
REQUIRE(HipTest::freeArrays<TestType>(nullptr, nullptr, nullptr, hData, B_h,
nullptr, true) == true);
} else {
REQUIRE(HipTest::freeArrays<TestType>(nullptr, nullptr, nullptr, hData, B_h,
nullptr, false) == true);
}
}
/*
This testcase performs the peer device context scenario
of hipMemcpyHtoA API
Memory is allocated in GPU-0 and the API is triggered from GPU-1
Input: "B_h" which is initialized with 1.6
Output: "A_d" output of hipMemcpyHtoA is copied to "hData" host variable
validated the result with "B_h"
*/
#if HT_AMD
TEMPLATE_TEST_CASE("Unit_hipMemcpyHtoA_multiDevice-PeerDeviceContext",
"[hipMemcpyHtoA]",
char, int, float) {
int numDevices = 0;
HIP_CHECK(hipGetDeviceCount(&numDevices));
if (numDevices > 1) {
int peerAccess = 0;
HIP_CHECK(hipDeviceCanAccessPeer(&peerAccess, 1, 0));
if (!peerAccess) {
SUCCEED("Skipped the test as there is no peer access");
} else {
HIP_CHECK(hipSetDevice(0));
hipArray *A_d;
TestType *hData{nullptr}, *B_h{nullptr};
size_t width{NUM_W * sizeof(TestType)};
// Initialization of data
HipTest::initArrays<TestType>(nullptr, nullptr, nullptr,
&hData, &B_h, nullptr, NUM_W);
HipTest::setDefaultData<TestType>(NUM_W, hData, B_h, nullptr);
hipChannelFormatDesc desc = hipCreateChannelDesc<TestType>();
HIP_CHECK(hipMallocArray(&A_d, &desc, NUM_W, NUM_H, hipArrayDefault));
HIP_CHECK(hipMemcpy2DToArray(A_d, 0, 0, hData, width,
width, NUM_H, hipMemcpyHostToDevice));
// Changing the device context
HIP_CHECK(hipSetDevice(1));
// Performing API call
HIP_CHECK(hipMemcpyHtoA(A_d, 0, B_h, copy_bytes*sizeof(TestType)));
HIP_CHECK(hipMemcpy2DFromArray(hData, sizeof(TestType)*NUM_W, A_d,
0, 0, sizeof(TestType)*NUM_W, 1,
hipMemcpyDeviceToHost));
// Validating the result
REQUIRE(HipTest::checkArray(B_h, hData, copy_bytes, NUM_H) == true);
// DeAllocating the memory
HIP_CHECK(hipFreeArray(A_d));
REQUIRE(HipTest::freeArrays<TestType>(nullptr, nullptr, nullptr,
hData, B_h,
nullptr, false) == true);
}
} else {
SUCCEED("skipping the testcases as numDevices < 2");
}
}
#endif
/*
This testcase verifies the negative scenarios of hipMemcpyHtoA API
*/
TEST_CASE("Unit_hipMemcpyHtoA_Negative") {
HIP_CHECK(hipSetDevice(0));
hipArray *A_d;
float *hData{nullptr}, *B_h{nullptr};
size_t width{NUM_W * sizeof(float)};
// Initialization of data
HipTest::initArrays<float>(nullptr, nullptr, nullptr,
&hData, &B_h, nullptr, NUM_W);
HipTest::setDefaultData<float>(NUM_W, hData, B_h, nullptr);
hipChannelFormatDesc desc = hipCreateChannelDesc<float>();
HIP_CHECK(hipMallocArray(&A_d, &desc, NUM_W, NUM_H, hipArrayDefault));
HIP_CHECK(hipMemcpy2DToArray(A_d, 0, 0, hData, width,
width, NUM_H, hipMemcpyHostToDevice));
SECTION("Source pointer is nullptr") {
REQUIRE(hipMemcpyHtoA(A_d, 0, nullptr, copy_bytes*sizeof(float))
!= hipSuccess);
}
SECTION("Source offset is more than allocated size") {
REQUIRE(hipMemcpyHtoA(A_d, 100, B_h, copy_bytes*sizeof(float))
!= hipSuccess);
}
SECTION("ByteCount is greater than allocated size") {
REQUIRE(hipMemcpyHtoA(A_d, 0, B_h, 12*sizeof(float)) != hipSuccess);
}
// DeAllocating the memory
HIP_CHECK(hipFreeArray(A_d));
REQUIRE(HipTest::freeArrays<float>(nullptr, nullptr, nullptr, hData, B_h,
nullptr, false) == true);
}
/*
This testcase verifies the size 0 check of hipMemcpyHtoA API
This is excluded for AMD as we have a bug already raised
SWDEV-274683
*/
#if HT_NVIDIA
TEST_CASE("Unit_hipMemcpyHtoA_SizeCheck") {
HIP_CHECK(hipSetDevice(0));
hipArray *A_d;
float *hData{nullptr}, *B_h{nullptr}, *def_data{nullptr};
size_t width{NUM_W * sizeof(float)};
// Initialization of data
HipTest::initArrays<float>(nullptr, nullptr, nullptr,
nullptr, &def_data, nullptr, NUM_W);
HipTest::initArrays<float>(nullptr, nullptr, nullptr,
&hData, &B_h, nullptr, NUM_W);
HipTest::setDefaultData<float>(NUM_W, hData, B_h, nullptr);
HipTest::setDefaultData<float>(NUM_W, nullptr, def_data, nullptr);
hipChannelFormatDesc desc = hipCreateChannelDesc<float>();
HIP_CHECK(hipMallocArray(&A_d, &desc, NUM_W, NUM_H, hipArrayDefault));
HIP_CHECK(hipMemcpy2DToArray(A_d, 0, 0, hData, width,
width, NUM_H, hipMemcpyHostToDevice));
SECTION("Passing 0 to copy bytes") {
REQUIRE(hipMemcpyHtoA(A_d, 0, B_h, 0) == hipSuccess);
HIP_CHECK(hipMemcpy2DFromArray(def_data, sizeof(float)*NUM_W, A_d,
0, 0, sizeof(float)*NUM_W, 1,
hipMemcpyDeviceToHost));
REQUIRE(HipTest::checkArray(hData, def_data, NUM_W, NUM_H) == true);
}
SECTION(" Source Array is nullptr") {
REQUIRE(hipMemcpyHtoA(nullptr, 0, B_h, copy_bytes*sizeof(float))
!= hipSuccess);
}
// DeAllocating the memory
HIP_CHECK(hipFreeArray(A_d));
REQUIRE(HipTest::freeArrays<float>(nullptr, nullptr, nullptr, hData, B_h,
def_data, false) == true);
}
#endif