AI/rocm-systems
2
0
Vork 0
Code Kwesties Pull-aanvragen Actions Paketten Projecten Publicaties Wiki Activiteit

SWDEV-436376 - [catch2][dtest] Adding new tests for hipMemcpyDeviceToDeviceNoCU

Bladeren bron 
Change-Id: Ibe338af731a4f0e21d5c8def8ac125532b75969e


[ROCm/hip-tests commit: 0c11f5920f]
This commit is contained in:
Rupam Chetia
2024-01-03 14:21:28 +05:30
gecommit door Rakesh Roy
bovenliggende e937d07a0f
commit ff6a16be9e
5 gewijzigde bestanden met toevoegingen van 333 en 14 verwijderingen
Download Patch-bestand Download Diff-bestand Expand all files Collapse all files
projects/hip-tests/catch/unit/graph/hipGraphAddMemcpyNode1D_old.cc
+7 -4
Bestand weergeven
@@ -1,5 +1,5 @@
/*
Copyright (c) 2023 Advanced Micro Devices, Inc. All rights reserved.
Copyright (c) 2023 - 2024 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
@@ -44,7 +44,8 @@ Negative -
#include <vector>
#include <numeric>
static void validateMemcpyNode1DArray(bool peerAccess) {
static void validateMemcpyNode1DArray(bool peerAccess,
hipMemcpyKind d2d_type = hipMemcpyDeviceToDevice) {
constexpr int SIZE{32};
int harray1D[SIZE]{};
int harray1Dres[SIZE]{};
@@ -79,7 +80,7 @@ static void validateMemcpyNode1DArray(bool peerAccess) {
// Device to Device (devArray1 -> devArray2)
HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyD2D, graph, &memcpyH2D, 1,
devArray2, devArray1, numBytes, hipMemcpyDeviceToDevice));
devArray2, devArray1, numBytes, d2d_type));
// Device to host (devArray2 -> harray1Dres)
HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyD2H, graph, &memcpyD2D, 1,
@@ -119,7 +120,9 @@ TEST_CASE("Unit_hipGraphAddMemcpyNode1D_Functional") {
SECTION("Memcpy with 1D array on default device") {
validateMemcpyNode1DArray(false);
}
SECTION("Memcpy with 1D array using DeviceToDeviceNoCU") {
validateMemcpyNode1DArray(false, hipMemcpyDeviceToDeviceNoCU);
}
SECTION("Memcpy with 1D array on peer device") {
int numDevices{}, peerAccess{};
HIP_CHECK(hipGetDeviceCount(&numDevices));
projects/hip-tests/catch/unit/memory/CMakeLists.txt
+2 -1
Bestand weergeven
@@ -94,7 +94,8 @@ if(HIP_PLATFORM MATCHES "amd")
hipMemAddressReserve.cc
hipMemRelease.cc
hipMemGetAllocationPropertiesFromHandle.cc
hipArray.cc)
hipArray.cc
hipMemcpyDeviceToDeviceNoCU.cc)
if(UNIX)
# Should be compiled for NVIDIA as well after EXSWHTEC-346 is addressed
# For windows build error occurs undefined symbol: hipPointerSetAttribute
projects/hip-tests/catch/unit/memory/hipMemcpy2DAsync_old.cc
+11 -4
Bestand weergeven
@@ -1,5 +1,5 @@
/*
Copyright (c) 2021-2023 Advanced Micro Devices, Inc. All rights reserved.
Copyright (c) 2021-2024 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
@@ -72,7 +72,7 @@ static constexpr auto ROWS{6};
* - unit/memory/hipMemcpy2DAsync.cc
* Test requirements
* ------------------------
* - HIP_VERSION >= 5.2
* - HIP_VERSION >= 6.1
*/
TEMPLATE_TEST_CASE("Unit_hipMemcpy2DAsync_Host&PinnedMem", ""
@@ -80,6 +80,7 @@ TEMPLATE_TEST_CASE("Unit_hipMemcpy2DAsync_Host&PinnedMem", ""
CHECK_IMAGE_SUPPORT
// 1 refers to pinned host memory
auto mem_type = GENERATE(0, 1);
auto memcpy_d2d_type = GENERATE(0, 1);
HIP_CHECK(hipSetDevice(0));
TestType *A_h{nullptr}, *B_h{nullptr}, *C_h{nullptr}, *A_d{nullptr},
*B_d{nullptr};
@@ -96,6 +97,12 @@ TEMPLATE_TEST_CASE("Unit_hipMemcpy2DAsync_Host&PinnedMem", ""
HipTest::initArrays<TestType>(nullptr, nullptr, nullptr,
&A_h, &B_h, &C_h, NUM_W*NUM_H, false);
}
hipMemcpyKind d2d_type;
if (memcpy_d2d_type) {
d2d_type = hipMemcpyDeviceToDevice;
} else {
d2d_type = hipMemcpyDeviceToDeviceNoCU;
}
HIP_CHECK(hipMallocPitch(reinterpret_cast<void**>(&A_d),
&pitch_A, width, NUM_H));
HIP_CHECK(hipMallocPitch(reinterpret_cast<void**>(&B_d),
@@ -112,7 +119,7 @@ TEMPLATE_TEST_CASE("Unit_hipMemcpy2DAsync_Host&PinnedMem", ""
// Performs D2D on same GPU device
HIP_CHECK(hipMemcpy2DAsync(B_d, pitch_B, A_d,
pitch_A, COLUMNS*sizeof(TestType),
ROWS, hipMemcpyDeviceToDevice, stream));
ROWS, d2d_type, stream));
// hipMemcpy2DAsync Device to Host
HIP_CHECK(hipMemcpy2DAsync(B_h, COLUMNS*sizeof(TestType), B_d, pitch_B,
@@ -129,7 +136,7 @@ TEMPLATE_TEST_CASE("Unit_hipMemcpy2DAsync_Host&PinnedMem", ""
// Performs D2D on same GPU device
HIP_CHECK(hipMemcpy2DAsync(B_d, pitch_B, A_d, pitch_A,
COLUMNS*sizeof(TestType), ROWS,
hipMemcpyDeviceToDevice, hipStreamPerThread));
d2d_type, hipStreamPerThread));
// hipMemcpy2DAsync Device to Host
HIP_CHECK(hipMemcpy2DAsync(B_h, COLUMNS*sizeof(TestType), B_d, pitch_B,
projects/hip-tests/catch/unit/memory/hipMemcpy2D_old.cc
+19 -5
Bestand weergeven
@@ -1,5 +1,5 @@
/*
Copyright (c) 2021-2023 Advanced Micro Devices, Inc. All rights reserved.
Copyright (c) 2021-2024 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
@@ -71,7 +71,7 @@ static constexpr auto ROWS{8};
* - unit/memory/hipMemcpy2D.cc
* Test requirements
* ------------------------
* - HIP_VERSION >= 6.0
* - HIP_VERSION >= 6.1
*/
TEMPLATE_TEST_CASE("Unit_hipMemcpy2D_H2D-D2D-D2H", ""
@@ -79,6 +79,7 @@ TEMPLATE_TEST_CASE("Unit_hipMemcpy2D_H2D-D2D-D2H", ""
CHECK_IMAGE_SUPPORT
// 1 refers to pinned host memory
auto mem_type = GENERATE(0, 1);
auto memcpy_d2d_type = GENERATE(0, 1);
HIP_CHECK(hipSetDevice(0));
TestType *A_h{nullptr}, *B_h{nullptr}, *C_h{nullptr}, *A_d{nullptr},
*B_d{nullptr};
@@ -93,6 +94,12 @@ TEMPLATE_TEST_CASE("Unit_hipMemcpy2D_H2D-D2D-D2H", ""
HipTest::initArrays<TestType>(nullptr, nullptr, nullptr,
&A_h, &B_h, &C_h, NUM_W*NUM_H, false);
}
hipMemcpyKind d2d_type;
if (memcpy_d2d_type) {
d2d_type = hipMemcpyDeviceToDevice;
} else {
d2d_type = hipMemcpyDeviceToDeviceNoCU;
}
HIP_CHECK(hipMallocPitch(reinterpret_cast<void**>(&A_d),
&pitch_A, width, NUM_H));
HIP_CHECK(hipMallocPitch(reinterpret_cast<void**>(&B_d),
@@ -109,7 +116,7 @@ TEMPLATE_TEST_CASE("Unit_hipMemcpy2D_H2D-D2D-D2H", ""
// Performs D2D on same GPU device
HIP_CHECK(hipMemcpy2D(B_d, pitch_B, A_d,
pitch_A, COLUMNS*sizeof(TestType),
ROWS, hipMemcpyDeviceToDevice));
ROWS, d2d_type));
// hipMemcpy2D Device to Host
HIP_CHECK(hipMemcpy2D(B_h, COLUMNS*sizeof(TestType), B_d, pitch_B,
@@ -151,7 +158,7 @@ TEMPLATE_TEST_CASE("Unit_hipMemcpy2D_H2D-D2D-D2H", ""
* - unit/memory/hipMemcpy2D.cc
* Test requirements
* ------------------------
* - HIP_VERSION >= 6.0
* - HIP_VERSION >= 6.1
*/
TEMPLATE_TEST_CASE("Unit_hipMemcpy2D_H2D-D2D-D2H_WithOffset", ""
@@ -159,6 +166,7 @@ TEMPLATE_TEST_CASE("Unit_hipMemcpy2D_H2D-D2D-D2H_WithOffset", ""
CHECK_IMAGE_SUPPORT
// 1 refers to pinned host memory
auto mem_type = GENERATE(0, 1);
auto memcpy_d2d_type = GENERATE(0, 1);
HIP_CHECK(hipSetDevice(0));
TestType *A_h{nullptr}, *B_h{nullptr}, *C_h{nullptr}, *A_d{nullptr},
*B_d{nullptr};
@@ -173,6 +181,12 @@ TEMPLATE_TEST_CASE("Unit_hipMemcpy2D_H2D-D2D-D2H_WithOffset", ""
HipTest::initArrays<TestType>(nullptr, nullptr, nullptr,
&A_h, &B_h, &C_h, NUM_W*NUM_H, false);
}
hipMemcpyKind d2d_type;
if (memcpy_d2d_type) {
d2d_type = hipMemcpyDeviceToDevice;
} else {
d2d_type = hipMemcpyDeviceToDeviceNoCU;
}
HIP_CHECK(hipMallocPitch(reinterpret_cast<void**>(&A_d),
&pitch_A, width, NUM_H));
HIP_CHECK(hipMallocPitch(reinterpret_cast<void**>(&B_d),
@@ -190,7 +204,7 @@ TEMPLATE_TEST_CASE("Unit_hipMemcpy2D_H2D-D2D-D2H_WithOffset", ""
HIP_CHECK(hipMemcpy2D(B_d+COLUMNS*sizeof(TestType), pitch_B,
A_d+COLUMNS*sizeof(TestType),
pitch_A, COLUMNS*sizeof(TestType),
ROWS, hipMemcpyDeviceToDevice));
ROWS, d2d_type));
// hipMemcpy2D Device to Host
HIP_CHECK(hipMemcpy2D(B_h, COLUMNS*sizeof(TestType),
projects/hip-tests/catch/unit/memory/hipMemcpyDeviceToDeviceNoCU.cc
+294
Bestand weergeven
@@ -0,0 +1,294 @@
/*
Copyright (c) 2024 Advanced Micro Devices, Inc. All rights reserved.
Permission is hereby granted, free of charge, to any person obtaining a copy
of this software and associated documentation files (the "Software"), to deal
in the Software without restriction, including without limitation the rights
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
copies of the Software, and to permit persons to whom the Software is
furnished to do so, subject to the following conditions:
The above copyright notice and this permission notice shall be included in
all copies or substantial portions of the Software.
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE.
*/
#include <hip_test_common.hh>
#include <hip_test_kernels.hh>
/**
* @addtogroup hipMemcpy hipMemcpy
* @{
* @ingroup MemoryTest
* `hipError_t hipMemcpy (void *dst, const void *src,
* size_t sizeBytes, hipMemcpyKind kind)` -
* Copy data from src to dst.
*/
static void fillDataTransfer2Dev(int *hostBuf, size_t len) {
unsigned int seed = time(nullptr);
for (size_t i = 0; i < len; i++) {
hostBuf[i] = (HipTest::RAND_R(&seed) & 0xFF);
}
}
/**
* Test Description
* ------------------------
* - Create 2 device memory chunks (Ad, Bd) and 2 host memory chunk (Ah, Bh).
* On default/user stream, perform the following memcpies in queue => Ah->Ad, Ad->Bd
* (hipMemcpyDeviceToDeviceNoCU), Kernel Launch (Bd*Bd), Bd->Bh. Verify data in Bh.
* ------------------------
* - catch\unit\memory\hipMemcpyDeviceToDeviceNoCU.cc
* Test requirements
* ------------------------
* - HIP_VERSION >= 6.1
*/
TEST_CASE("Unit_hipMemcpyDeviceToDeviceNoCU_SingleStream") {
auto testAsync = GENERATE(0, 1);
auto isDefaultStrm = GENERATE(0, 1);
constexpr int N = 1 << 18;
size_t buffer_size = N * sizeof(int);
constexpr unsigned threadsPerBlock = 128;
constexpr unsigned blocks = 64;
// Allocate device resources
int *Ad, *Bd;
HIP_CHECK(hipMalloc(&Ad, buffer_size));
HIP_CHECK(hipMalloc(&Bd, buffer_size));
// Allocate host resources
int *Ah = new int[N];
REQUIRE(Ah != nullptr);
int *Bh = new int[N];
REQUIRE(Bh != nullptr);
// Check whether to execute on default stream or user stream
hipStream_t strm = 0;
if (isDefaultStrm == 1) {
HIP_CHECK(hipStreamCreate(&strm));
}
// fill Ah with random data
fillDataTransfer2Dev(Ah, N);
if (0 == testAsync) {
HIP_CHECK(hipMemcpy(Ad, Ah, N*sizeof(int), hipMemcpyHostToDevice));
HIP_CHECK(hipMemcpy(Bd, Ad, N*sizeof(int), hipMemcpyDeviceToDeviceNoCU));
hipLaunchKernelGGL(HipTest::vector_square, dim3(blocks),
dim3(threadsPerBlock), 0, strm, Bd, Bd, N);
HIP_CHECK(hipMemcpy(Bh, Bd, N*sizeof(int), hipMemcpyDeviceToHost));
} else {
HIP_CHECK(hipMemcpyAsync(Ad, Ah, N*sizeof(int),
hipMemcpyHostToDevice, strm));
HIP_CHECK(hipMemcpyAsync(Bd, Ad, N*sizeof(int),
hipMemcpyDeviceToDeviceNoCU, strm));
hipLaunchKernelGGL(HipTest::vector_square, dim3(blocks),
dim3(threadsPerBlock), 0, strm, Bd, Bd, N);
HIP_CHECK(hipMemcpyAsync(Bh, Bd, N*sizeof(int),
hipMemcpyDeviceToHost, strm));
}
HIP_CHECK(hipDeviceSynchronize());
for (int i = 0; i < N; i++) {
REQUIRE(Bh[i] == (Ah[i]*Ah[i]));
}
if (isDefaultStrm == 1) {
HIP_CHECK(hipStreamDestroy(strm));
}
// Delete resources
delete[] Ah;
delete[] Bh;
HIP_CHECK(hipFree(Ad));
HIP_CHECK(hipFree(Bd));
}
/**
* Test Description
* ------------------------
* - Create 3 device memory chunks (Ad, Bd, Cd) and 2 host memory chunk (Ah, Bh).
* On default/user stream, perform the following memcpies in queue => Ah->Ad, Ad->Bd
* (hipMemcpyAsync with flag = hipMemcpyDeviceToDevice), Bd->Cd (with flag =
* hipMemcpyDeviceToDeviceNoCU, Cd->Bh. Verify data in Bh.
* ------------------------
* - catch\unit\memory\hipMemcpyDeviceToDeviceNoCU.cc
* Test requirements
* ------------------------
* - HIP_VERSION >= 6.1
*/
TEST_CASE("Unit_hipMemcpyDeviceToDeviceNoCU_WithCU_NoCU_Comb_SingleStrm") {
constexpr int N = 1 << 18;
size_t buffer_size = N * sizeof(int);
// Allocate device resources
int *Ad, *Bd, *Cd;
HIP_CHECK(hipMalloc(&Ad, buffer_size));
HIP_CHECK(hipMalloc(&Bd, buffer_size));
HIP_CHECK(hipMalloc(&Cd, buffer_size));
// Allocate host resources
int *Ah = new int[N];
REQUIRE(Ah != nullptr);
int *Bh = new int[N];
REQUIRE(Bh != nullptr);
// Check whether to execute on default stream or user stream
hipStream_t strm = 0;
HIP_CHECK(hipStreamCreate(&strm));
// fill Ah with random data
fillDataTransfer2Dev(Ah, N);
HIP_CHECK(hipMemcpyAsync(Ad, Ah, N*sizeof(int),
hipMemcpyHostToDevice, strm));
SECTION("Memcpy withCU first then without CU") {
HIP_CHECK(hipMemcpyAsync(Bd, Ad, N*sizeof(int),
hipMemcpyDeviceToDevice, strm));
HIP_CHECK(hipMemcpyAsync(Cd, Bd, N*sizeof(int),
hipMemcpyDeviceToDeviceNoCU, strm));
}
SECTION("Memcpy without CU first then with CU") {
HIP_CHECK(hipMemcpyAsync(Bd, Ad, N*sizeof(int),
hipMemcpyDeviceToDeviceNoCU, strm));
HIP_CHECK(hipMemcpyAsync(Cd, Bd, N*sizeof(int),
hipMemcpyDeviceToDevice, strm));
}
SECTION("Memcpy without CU twice") {
HIP_CHECK(hipMemcpyAsync(Bd, Ad, N*sizeof(int),
hipMemcpyDeviceToDeviceNoCU, strm));
HIP_CHECK(hipMemcpyAsync(Cd, Bd, N*sizeof(int),
hipMemcpyDeviceToDeviceNoCU, strm));
}
HIP_CHECK(hipMemcpyAsync(Bh, Cd, N*sizeof(int),
hipMemcpyDeviceToHost, strm));
HIP_CHECK(hipStreamSynchronize(strm));
for (int i = 0; i < N; i++) {
REQUIRE(Bh[i] == Ah[i]);
}
HIP_CHECK(hipStreamDestroy(strm));
// Delete resources
delete[] Ah;
delete[] Bh;
HIP_CHECK(hipFree(Ad));
HIP_CHECK(hipFree(Bd));
HIP_CHECK(hipFree(Cd));
}
/**
* Test Description
* ------------------------
* - Create 3 device memory chunks (Ad, Bd, Cd) and 3 host memory chunk
* (Ah, Bh, Ch). On user stream1, perform the following memcpies in queue
* => On default stream, perform Ah->Ad memcpy. On stream1, perform Ad->Bd
* (hipMemcpyAsync with flag = hipMemcpyDeviceToDeviceNoCU) and Bd->Bh.
* On stream2, perform Ad->Cd (hipMemcpyAsync with flag =
* hipMemcpyDeviceToDeviceNoCU) and Cd -> Ch. Wait for both stream1
* and stream2. Verify output in Bh, Ch.
* ------------------------
* - catch\unit\memory\hipMemcpyDeviceToDeviceNoCU.cc
* Test requirements
* ------------------------
* - HIP_VERSION >= 6.1
*/
TEST_CASE("Unit_hipMemcpyDeviceToDeviceNoCU_NoCU_MulStrm") {
constexpr int N = 1 << 18;
size_t buffer_size = N * sizeof(int);
// Allocate device resources
int *Ad, *Bd, *Cd;
HIP_CHECK(hipMalloc(&Ad, buffer_size));
HIP_CHECK(hipMalloc(&Bd, buffer_size));
HIP_CHECK(hipMalloc(&Cd, buffer_size));
// Allocate host resources
int *Ah = new int[N];
REQUIRE(Ah != nullptr);
int *Bh = new int[N];
REQUIRE(Bh != nullptr);
int *Ch = new int[N];
REQUIRE(Ch != nullptr);
// fill Ah with random data
fillDataTransfer2Dev(Ah, N);
HIP_CHECK(hipMemcpyAsync(Ad, Ah, N*sizeof(int),
hipMemcpyHostToDevice, 0));
hipStream_t strm1, strm2;
HIP_CHECK(hipStreamCreate(&strm1));
HIP_CHECK(hipStreamCreate(&strm2));
HIP_CHECK(hipMemcpyAsync(Bd, Ad, N*sizeof(int),
hipMemcpyDeviceToDeviceNoCU, strm1));
HIP_CHECK(hipMemcpyAsync(Cd, Ad, N*sizeof(int),
hipMemcpyDeviceToDeviceNoCU, strm2));
HIP_CHECK(hipMemcpyAsync(Bh, Bd, N*sizeof(int),
hipMemcpyDeviceToHost, strm1));
HIP_CHECK(hipMemcpyAsync(Ch, Cd, N*sizeof(int),
hipMemcpyDeviceToHost, strm2));
HIP_CHECK(hipStreamSynchronize(strm1));
HIP_CHECK(hipStreamSynchronize(strm2));
for (int i = 0; i < N; i++) {
REQUIRE(Bh[i] == Ah[i]);
REQUIRE(Ch[i] == Ah[i]);
}
HIP_CHECK(hipStreamDestroy(strm2));
HIP_CHECK(hipStreamDestroy(strm1));
// Delete resources
delete[] Ah;
delete[] Bh;
delete[] Ch;
HIP_CHECK(hipFree(Ad));
HIP_CHECK(hipFree(Bd));
HIP_CHECK(hipFree(Cd));
}
/**
* Test Description
* ------------------------
* - Create 3 device memory chunks (Ad, Bd, Cd) and 1 host memory
* chunk (Ah) of size 1024 MB. Initialize the host memory chunk (Ah)
* with random data. On default stream, copy from host to device
* (Ah -> Ad) the random data. Create 2 streams stream1 and stream2.
* Once completed, on stream1 copy from device to device (Ad -> Bd)
* using hipMemcpyDeviceToDeviceNoCU (Async Copy) and on stream2,
* launch a kernel to copy from (Ad -> Cd). Validate data in Bd and
* Cd are same.
* ------------------------
* - catch\unit\memory\hipMemcpyDeviceToDeviceNoCU.cc
* Test requirements
* ------------------------
* - HIP_VERSION >= 6.1
*/
TEST_CASE("Unit_hipMemcpyDeviceToDeviceNoCU_Memcpy_Kernel_InParallel") {
constexpr int N = 1 << 26;
size_t buffer_size = N * sizeof(int);
constexpr unsigned threadsPerBlock = 128;
constexpr unsigned blocks = 64;
// Allocate device resources
int *Ad, *Bd, *Cd;
HIP_CHECK(hipMalloc(&Ad, buffer_size));
HIP_CHECK(hipMalloc(&Bd, buffer_size));
HIP_CHECK(hipMalloc(&Cd, buffer_size));
// Allocate host resources
int *Ah = new int[N];
REQUIRE(Ah != nullptr);
int *Bh = new int[N];
REQUIRE(Bh != nullptr);
int *Ch = new int[N];
REQUIRE(Ch != nullptr);
// fill Ah with random data
fillDataTransfer2Dev(Ah, N);
HIP_CHECK(hipMemcpyAsync(Ad, Ah, N*sizeof(int),
hipMemcpyHostToDevice, 0));
hipStream_t strm1, strm2;
HIP_CHECK(hipStreamCreate(&strm1));
HIP_CHECK(hipStreamCreate(&strm2));
HIP_CHECK(hipMemcpyAsync(Bd, Ad, N*sizeof(int),
hipMemcpyDeviceToDeviceNoCU, strm1));
hipLaunchKernelGGL(HipTest::vector_square, dim3(blocks),
dim3(threadsPerBlock), 0, strm2, Ad, Cd, N);
HIP_CHECK(hipMemcpyAsync(Bh, Bd, N*sizeof(int),
hipMemcpyDeviceToHost, strm1));
HIP_CHECK(hipMemcpyAsync(Ch, Cd, N*sizeof(int),
hipMemcpyDeviceToHost, strm2));
HIP_CHECK(hipStreamSynchronize(strm1));
HIP_CHECK(hipStreamSynchronize(strm2));
for (int i = 0; i < N; i++) {
REQUIRE(Bh[i] == Ah[i]);
REQUIRE(Ch[i] == (Ah[i]*Ah[i]));
}
HIP_CHECK(hipStreamDestroy(strm2));
HIP_CHECK(hipStreamDestroy(strm1));
// Delete resources
delete[] Ah;
delete[] Bh;
delete[] Ch;
HIP_CHECK(hipFree(Ad));
HIP_CHECK(hipFree(Bd));
HIP_CHECK(hipFree(Cd));
}