diff --git a/projects/hip-tests/catch/unit/graph/hipGraphAddMemcpyNode1D_old.cc b/projects/hip-tests/catch/unit/graph/hipGraphAddMemcpyNode1D_old.cc index 40c8ef1847..529af8a4ee 100644 --- a/projects/hip-tests/catch/unit/graph/hipGraphAddMemcpyNode1D_old.cc +++ b/projects/hip-tests/catch/unit/graph/hipGraphAddMemcpyNode1D_old.cc @@ -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 #include -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)); diff --git a/projects/hip-tests/catch/unit/memory/CMakeLists.txt b/projects/hip-tests/catch/unit/memory/CMakeLists.txt index 5dd1554e7f..29f966d9a4 100644 --- a/projects/hip-tests/catch/unit/memory/CMakeLists.txt +++ b/projects/hip-tests/catch/unit/memory/CMakeLists.txt @@ -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 diff --git a/projects/hip-tests/catch/unit/memory/hipMemcpy2DAsync_old.cc b/projects/hip-tests/catch/unit/memory/hipMemcpy2DAsync_old.cc index 1ca39bd6c9..867064fe4c 100644 --- a/projects/hip-tests/catch/unit/memory/hipMemcpy2DAsync_old.cc +++ b/projects/hip-tests/catch/unit/memory/hipMemcpy2DAsync_old.cc @@ -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(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(&A_d), &pitch_A, width, NUM_H)); HIP_CHECK(hipMallocPitch(reinterpret_cast(&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, diff --git a/projects/hip-tests/catch/unit/memory/hipMemcpy2D_old.cc b/projects/hip-tests/catch/unit/memory/hipMemcpy2D_old.cc index d7e21e42f2..b3cd9eb865 100644 --- a/projects/hip-tests/catch/unit/memory/hipMemcpy2D_old.cc +++ b/projects/hip-tests/catch/unit/memory/hipMemcpy2D_old.cc @@ -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(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(&A_d), &pitch_A, width, NUM_H)); HIP_CHECK(hipMallocPitch(reinterpret_cast(&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(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(&A_d), &pitch_A, width, NUM_H)); HIP_CHECK(hipMallocPitch(reinterpret_cast(&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), diff --git a/projects/hip-tests/catch/unit/memory/hipMemcpyDeviceToDeviceNoCU.cc b/projects/hip-tests/catch/unit/memory/hipMemcpyDeviceToDeviceNoCU.cc new file mode 100644 index 0000000000..7385dc9fcd --- /dev/null +++ b/projects/hip-tests/catch/unit/memory/hipMemcpyDeviceToDeviceNoCU.cc @@ -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 +#include +/** + * @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)); +}