AI/rocm-systems
2
0 Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

SWDEV-422641 - [catch2][dtest] Added test for Annoying truncation of 'size_t' to 'unsigned int' inside hipMemcpy3D

Browse Source 
Change-Id: I0faef0278acbb82186839fcd556af508b6fc84ec
This commit is contained in:
mbhiutra authored and Rakesh Roy committed 2023-11-06 09:22:04 -05:00
1 parent 4d36246849
commit f729e892a2
6 files changed
+873 -206
Download Patch File Download Diff File Expand all files Collapse all files

No files matched your search

catch/unit/memory/hipMemcpy2DAsync.cc
+193 -54
View File
@@ -1,5 +1,5 @@
/*
Copyright (c) 2021 Advanced Micro Devices, Inc. All rights reserved.
Copyright (c) 2021-2023 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,6 +17,16 @@ OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE.
*/
/**
* @addtogroup hipMemcpy2DAsync hipMemcpy2DAsync
* @{
* @ingroup MemcpyTest
* `hipMemcpy2DAsync(void* dst, size_t dpitch, const void* src,
* size_t spitch, size_t width, size_t height,
* hipMemcpyKind kind, hipStream_t stream = 0 )` -
* Copies data between host and device.
*/
// Testcase Description:
// 1) Verifies the working of Memcpy2DAsync API negative scenarios by
// Pass NULL to destination pointer
@@ -44,21 +54,30 @@ static constexpr auto NUM_H{16};
static constexpr auto COLUMNS{6};
static constexpr auto ROWS{6};
/*
This performs the following scenarios of hipMemcpy2DAsync API on same GPU
1. H2D-D2D-D2H for Host Memory<-->Device Memory
2. H2D-D2D-D2H for Pinned Host Memory<-->Device Memory
/**
* Test Description
* ------------------------
* - This performs the following scenarios of hipMemcpy2DAsync API on same GPU
1. H2D-D2D-D2H for Host Memory<-->Device Memory
2. H2D-D2D-D2H for Pinned Host Memory<-->Device Memory
Input : "A_h" initialized based on data type
Input : "A_h" initialized based on data type
"A_h" --> "A_d" using H2D copy
"A_d" --> "B_d" using D2D copy
"B_d" --> "B_h" using D2H copy
Output: Validating A_h with B_h both should be equal for
Output: Validating A_h with B_h both should be equal for
the number of COLUMNS and ROWS copied
*/
TEMPLATE_TEST_CASE("Unit_hipMemcpy2DAsync_Host&PinnedMem", "", int, float, double) {
CHECK_IMAGE_SUPPORT
* Test source
* ------------------------
* - unit/memory/hipMemcpy2DAsync.cc
* Test requirements
* ------------------------
* - HIP_VERSION >= 5.2
*/
TEMPLATE_TEST_CASE("Unit_hipMemcpy2DAsync_Host&PinnedMem", ""
, int, float, double) {
CHECK_IMAGE_SUPPORT
// 1 refers to pinned host memory
auto mem_type = GENERATE(0, 1);
HIP_CHECK(hipSetDevice(0));
@@ -108,9 +127,9 @@ TEMPLATE_TEST_CASE("Unit_hipMemcpy2DAsync_Host&PinnedMem", "", int, float, doubl
hipMemcpyHostToDevice, hipStreamPerThread));
// Performs D2D on same GPU device
HIP_CHECK(hipMemcpy2DAsync(B_d, pitch_B, A_d,
pitch_A, COLUMNS*sizeof(TestType),
ROWS, hipMemcpyDeviceToDevice, hipStreamPerThread));
HIP_CHECK(hipMemcpy2DAsync(B_d, pitch_B, A_d, pitch_A,
COLUMNS*sizeof(TestType), ROWS,
hipMemcpyDeviceToDevice, hipStreamPerThread));
// hipMemcpy2DAsync Device to Host
HIP_CHECK(hipMemcpy2DAsync(B_h, COLUMNS*sizeof(TestType), B_d, pitch_B,
@@ -136,22 +155,30 @@ TEMPLATE_TEST_CASE("Unit_hipMemcpy2DAsync_Host&PinnedMem", "", int, float, doubl
HIP_CHECK(hipStreamDestroy(stream));
}
/*
This testcases performs the following scenarios of hipMemcpy2DAsync API
on Peer GPU
1. H2D-D2D-D2H for Host Memory<-->Device Memory
2. H2D-D2D-D2H for Pinned Host Memory<-->Device Memory
/**
* Test Description
* ------------------------
* - This testcases performs the following scenarios of hipMemcpy2DAsync API on Peer GPU
1. H2D-D2D-D2H for Host Memory<-->Device Memory
2. H2D-D2D-D2H for Pinned Host Memory<-->Device Memory
Input : "A_h" initialized based on data type
"A_h" --> "A_d" using H2D copy
"A_d" --> "X_d" using D2D copy
"X_d" --> "B_h" using D2H copy
Output: Validating A_h with B_h both should be equal for
the number of COLUMNS and ROWS copied
*/
TEMPLATE_TEST_CASE("Unit_hipMemcpy2DAsync_multiDevice-Host&PinnedMem", "", int, float, double) {
Input : "A_h" initialized based on data type
"A_h" --> "A_d" using H2D copy
"A_d" --> "X_d" using D2D copy
"X_d" --> "B_h" using D2H copy
Output: Validating A_h with B_h both should be equal for
the number of COLUMNS and ROWS copied
* Test source
* ------------------------
* - unit/memory/hipMemcpy2DAsync.cc
* Test requirements
* ------------------------
* - HIP_VERSION >= 5.2
*/
TEMPLATE_TEST_CASE("Unit_hipMemcpy2DAsync_multiDevice-Host&PinnedMem", ""
, int, float, double) {
CHECK_IMAGE_SUPPORT
auto mem_type = GENERATE(0, 1);
int numDevices = 0;
int canAccessPeer = 0;
@@ -226,23 +253,31 @@ TEMPLATE_TEST_CASE("Unit_hipMemcpy2DAsync_multiDevice-Host&PinnedMem", "", int,
}
}
/*
This testcases performs the following scenarios of hipMemcpy2DAsync API
on Peer GPU
1. H2D-D2D-D2H for Host Memory<-->Device Memory
2. H2D-D2D-D2H for Pinned Host Memory<-->Device Memory
Memory is allocated in GPU-0 and Stream is created in GPU-1
/**
* Test Description
* ------------------------
* - This testcases performs the following scenarios of hipMemcpy2DAsync API on Peer GPU
1. H2D-D2D-D2H for Host Memory<-->Device Memory
2. H2D-D2D-D2H for Pinned Host Memory<-->Device Memory
Memory is allocated in GPU-0 and Stream is created in GPU-1
Input : "A_h" initialized based on data type
"A_h" --> "A_d" using H2D copy
"A_d" --> "X_d" using D2D copy
"X_d" --> "B_h" using D2H copy
Output: Validating A_h with B_h both should be equal for
the number of COLUMNS and ROWS copied
*/
TEMPLATE_TEST_CASE("Unit_hipMemcpy2DAsync_multiDevice-StreamOnDiffDevice", "", int, float, double) {
Input : "A_h" initialized based on data type
"A_h" --> "A_d" using H2D copy
"A_d" --> "X_d" using D2D copy
"X_d" --> "B_h" using D2H copy
Output: Validating A_h with B_h both should be equal for
the number of COLUMNS and ROWS copied
* Test source
* ------------------------
* - unit/memory/hipMemcpy2DAsync.cc
* Test requirements
* ------------------------
* - HIP_VERSION >= 5.2
*/
TEMPLATE_TEST_CASE("Unit_hipMemcpy2DAsync_multiDevice-StreamOnDiffDevice", ""
, int, float, double) {
CHECK_IMAGE_SUPPORT
auto mem_type = GENERATE(0, 1);
int numDevices = 0;
int canAccessPeer = 0;
@@ -275,7 +310,6 @@ TEMPLATE_TEST_CASE("Unit_hipMemcpy2DAsync_multiDevice-StreamOnDiffDevice", "", i
// Initialize the data
HipTest::setDefaultData<TestType>(NUM_W*NUM_H, A_h, B_h, C_h);
// Change device
HIP_CHECK(hipSetDevice(1));
HIP_CHECK(hipStreamCreate(&stream));
@@ -284,7 +318,6 @@ TEMPLATE_TEST_CASE("Unit_hipMemcpy2DAsync_multiDevice-StreamOnDiffDevice", "", i
HIP_CHECK(hipMemcpy2DAsync(A_d, pitch_A, A_h, COLUMNS*sizeof(TestType),
COLUMNS*sizeof(TestType), ROWS, hipMemcpyHostToDevice, stream));
// Device to Device
HIP_CHECK(hipMemcpy2DAsync(X_d, pitch_X, A_d,
pitch_A, COLUMNS*sizeof(TestType),
@@ -318,12 +351,24 @@ TEMPLATE_TEST_CASE("Unit_hipMemcpy2DAsync_multiDevice-StreamOnDiffDevice", "", i
}
}
/*
This testcase verifies the null checks of hipMemcpy2DAsync API
*/
/**
* Test Description
* ------------------------
* - This testcase verifies the null checks of hipMemcpy2DAsync API
1. hipMemcpy2DAsync API where Source Pitch is zero
2. hipMemcpy2DAsync API where Destination Pitch is zero
3. hipMemcpy2DAsync API where height is zero
4. hipMemcpy2DAsync API where width is zero
* Test source
* ------------------------
* - unit/memory/hipMemcpy2DAsync.cc
* Test requirements
* ------------------------
* - HIP_VERSION >= 5.2
*/
TEST_CASE("Unit_hipMemcpy2DAsync_SizeCheck") {
CHECK_IMAGE_SUPPORT
HIP_CHECK(hipSetDevice(0));
int* A_h{nullptr}, *A_d{nullptr};
size_t pitch_A;
@@ -369,12 +414,23 @@ TEST_CASE("Unit_hipMemcpy2DAsync_SizeCheck") {
free(A_h);
}
/*
This testcase performs the negative scenarios of hipMemcpy2DAsync API
*/
/**
* Test Description
* ------------------------
* - This testcase performs the negative scenarios of hipMemcpy2DAsync API
1. hipMemcpy2DAsync API by Passing nullptr to destination
2. hipMemcpy2DAsync API by Passing nullptr to source
3. hipMemcpy2DAsync API where width is > destination pitch
* Test source
* ------------------------
* - unit/memory/hipMemcpy2DAsync.cc
* Test requirements
* ------------------------
* - HIP_VERSION >= 5.2
*/
TEST_CASE("Unit_hipMemcpy2DAsync_Negative") {
CHECK_IMAGE_SUPPORT
HIP_CHECK(hipSetDevice(0));
int* A_h{nullptr}, *A_d{nullptr};
size_t pitch_A;
@@ -397,8 +453,8 @@ TEST_CASE("Unit_hipMemcpy2DAsync_Negative") {
hipMemcpyDeviceToHost, stream) != hipSuccess);
}
SECTION("hipMemcpy2DAsync API by Passing nullptr to destination") {
REQUIRE(hipMemcpy2DAsync(nullptr, width, nullptr,
SECTION("hipMemcpy2DAsync API by Passing nullptr to source") {
REQUIRE(hipMemcpy2DAsync(A_h, width, nullptr,
pitch_A, COLUMNS*sizeof(int), ROWS,
hipMemcpyDeviceToHost, stream) != hipSuccess);
}
@@ -414,3 +470,86 @@ TEST_CASE("Unit_hipMemcpy2DAsync_Negative") {
HIP_CHECK(hipStreamDestroy(stream));
free(A_h);
}
static void hipMemcpy2DAsync_Basic_Size_Test(size_t inc) {
constexpr int defaultProgramSize = 256 * 1024 * 1024;
constexpr int N = 2;
constexpr int value = 42;
int *in, *out, *dev;
size_t newSize = 0, inp = 0;
size_t size = sizeof(int) * N * inc;
size_t free, total;
HIP_CHECK(hipMemGetInfo(&free, &total));
if ( free < 2 * size )
newSize = ( free - defaultProgramSize ) / 2;
else
newSize = size;
INFO("Array size: " << size/1024.0/1024.0 << " MB or " << size << " Bytes.");
INFO("Free memory: " << free/1024.0/1024.0 << " MB or " << free << " Bytes");
INFO("NewSize:" << newSize/1024.0/1024.0 << "MB or " << newSize << " Bytes");
HIP_CHECK(hipHostMalloc(&in, newSize));
HIP_CHECK(hipHostMalloc(&out, newSize));
HIP_CHECK(hipMalloc(&dev, newSize));
inp = newSize / (sizeof(int) * N);
for (size_t i=0; i < N; i++) {
in[i * inp] = value;
}
size_t pitch = sizeof(int) * inp;
hipStream_t stream;
HIP_CHECK(hipStreamCreate(&stream));
HIP_CHECK(hipMemcpy2DAsync(dev, pitch, in, pitch, sizeof(int),
N, hipMemcpyHostToDevice, stream));
HIP_CHECK(hipMemcpy2DAsync(out, pitch, dev, pitch, sizeof(int),
N, hipMemcpyDeviceToHost, stream));
HIP_CHECK(hipStreamSynchronize(stream));
for (size_t i=0; i < N; i++) {
REQUIRE(out[i * inp] == value);
}
HIP_CHECK(hipFree(dev));
HIP_CHECK(hipHostFree(in));
HIP_CHECK(hipHostFree(out));
HIP_CHECK(hipStreamDestroy(stream));
}
/**
* Test Description
* ------------------------
* - This testcase performs multidevice size check on hipMemcpy2DAsync API
1. Verify hipMemcpy2DAsync with 1 << 20 size
2. Verify hipMemcpy2DAsync with 1 << 21 size
* Test source
* ------------------------
* - unit/memory/hipMemcpy2DAsync.cc
* Test requirements
* ------------------------
* - HIP_VERSION >= 6.0
*/
TEST_CASE("Unit_hipMemcpy2DAsync_multiDevice_Basic_Size_Test") {
CHECK_IMAGE_SUPPORT
size_t input = 1 << 20;
int numDevices = 0;
HIP_CHECK(hipGetDeviceCount(&numDevices));
for (int i=0; i < numDevices; i++) {
HIP_CHECK(hipSetDevice(i));
SECTION("Verify hipMemcpy2DAsync with 1 << 20 size") {
hipMemcpy2DAsync_Basic_Size_Test(input);
}
SECTION("Verify hipMemcpy2DAsync with 1 << 21 size") {
input <<= 1;
hipMemcpy2DAsync_Basic_Size_Test(input);
}
}
}