From 1336a9de0ba26aacaa0760649eefb5e612397c22 Mon Sep 17 00:00:00 2001 From: rohit pathania Date: Tue, 2 Jun 2020 10:08:26 -0400 Subject: [PATCH] [dtest] hipMemcpyWithStream API multithreaded test scenarios 1.hipMemcpyWithStream with one stream 2.hipMemcpyWithStream with two streams 3.hipMemcpyWithStream multiGPU with one stream 4.hipMemcpyWithStream with kind DtoH 5.hipMemcpyWithStream with HtoH 6.hipMemcpyWithStream with DtoD 7.hipMemcpyWithStream with Default kind 8.hipMemcpyWithStream with Default kind on same device 9.hipMemcpyWithStream with DtoD on same device SWDEV-238517 for enhancing hip unit tests Change-Id: I5f55a12bdd7c8d28fcb06db94a491c2ad5ee3004 [ROCm/hip commit: 32a6e854b7e8f0b18abb9b0653bf0000dcd295dd] --- .../memory/hipMemcpyWithStreamMultiThread.cpp | 659 ++++++++++++++++++ 1 file changed, 659 insertions(+) create mode 100644 projects/hip/tests/src/runtimeApi/memory/hipMemcpyWithStreamMultiThread.cpp diff --git a/projects/hip/tests/src/runtimeApi/memory/hipMemcpyWithStreamMultiThread.cpp b/projects/hip/tests/src/runtimeApi/memory/hipMemcpyWithStreamMultiThread.cpp new file mode 100644 index 0000000000..aaaf8332bc --- /dev/null +++ b/projects/hip/tests/src/runtimeApi/memory/hipMemcpyWithStreamMultiThread.cpp @@ -0,0 +1,659 @@ +/* +Copyright (c) 2020-present 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. +*/ + +/* + * Different test for checking functionality of + * hipError_t hipMemcpyWithStream(void* dst, const void* src, size_t sizeBytes, + * hipMemcpyKind kind, hipStream_t stream); + */ + +/* HIT_START + * BUILD: %t %s ../../test_common.cpp + * TEST: %t + * HIT_END + */ + +#include +#include +#include +#include "test_common.h" + +#define LEN 64 +#define SIZE LEN << 2 +#define THREADS 2 +#define MAX_THREADS 16 + + +#define test_passed(test_name) printf("%s %s PASSED!%s\n", \ + KGRN, #test_name, KNRM); +#define test_failed(test_name) printf("%s %s FAILED!%s\n", \ + KRED, #test_name, KNRM); + +enum class ops +{ TestwithOnestream, + TestwithTwoStream, + TestOnMultiGPUwithOneStream, + TestkindDtoH, + TestkindDtoD, + TestkindHtoH, + TestkindDefault, + TestkindDefaultForDtoD, + TestDtoDonSameDevice, + END_OF_LIST +}; + + +class HipMemcpyWithStreamMultiThreadtests { + // Test hipMemcpyWithStream with one streams and launch kernel in + // that stream, verify the data. + void TestwithOnestream(void); + // Test hipMemcpyWithStream with two streams and launch kernels in + // two streams, verify the data. + void TestwithTwoStream(void); + // Test hipMemcpyWithStream with one stream for each gpu and launch + // kernels in each, verify the data + void TestOnMultiGPUwithOneStream(void); + // Test hipMemcpyWithStream to copy data from + // device to host (hipMemcpyDeviceToHost). + void TestkindDtoH(void); + // Test hipMemcpyWithStream with hipMemcpyDeviceToDevice on MultiGPU. + void TestkindDtoD(void); + // Test hipMemcpyWithStream with hipMemcpyHostToHost. + void TestkindHtoH(void); + // Test hipMemcpyWithStream with hipMemcpyDefault. + void TestkindDefault(void); + // Test hipMemcpyWithStream with hipMemcpyDefault for + // device to device transfer case. + void TestkindDefaultForDtoD(void); + // Test hipMemcpyWithStream with hipMemcpyDeviceToDevice on same device. + void TestDtoDonSameDevice(void); + + public: + // run all the tests on multithreaded. + void TestwithMultiThreaded(ops op); +}; + +struct joinable_thread : std::thread { + template + explicit joinable_thread(Xs&&... xs) : std::thread(std::forward(xs)...) + {} // NOLINT + + joinable_thread& operator=(joinable_thread&& other) = default; + joinable_thread(joinable_thread&& other) = default; + + ~joinable_thread() { + if (this->joinable()) + this->join(); + } +}; + +void HipMemcpyWithStreamMultiThreadtests::TestwithOnestream(void) { + size_t Nbytes = N * sizeof(int); + int numDevices = 0; + int *A_d, *B_d, *C_d; + int *A_h, *B_h, *C_h; + + unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N); + HipTest::initArrays(&A_d, &B_d, &C_d, &A_h, &B_h, &C_h, N, false); + + hipStream_t stream; + HIPCHECK(hipStreamCreate(&stream)); + + HIPCHECK(hipMemcpyWithStream(A_d, A_h, Nbytes, + hipMemcpyHostToDevice, stream)); + HIPCHECK(hipMemcpyWithStream(B_d, B_h, Nbytes, + hipMemcpyHostToDevice, stream)); + hipLaunchKernelGGL(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock), + 0, stream, static_cast(A_d), + static_cast(B_d), C_d, N); + HIPCHECK(hipStreamSynchronize(stream)); + HIPCHECK(hipMemcpy(C_h, C_d, Nbytes, hipMemcpyDeviceToHost)); + HipTest::checkVectorADD(A_h, B_h, C_h, N); + + HipTest::freeArrays(A_d, B_d, C_d, A_h, B_h, C_h, false); + HIPCHECK(hipStreamDestroy(stream)); +} + +void HipMemcpyWithStreamMultiThreadtests::TestwithTwoStream(void) { + size_t Nbytes = N * sizeof(int); + int numDevices = 0; + int noOfstreams = 2; + int *A_d[noOfstreams], *B_d[noOfstreams], *C_d[noOfstreams]; + int *A_h[noOfstreams], *B_h[noOfstreams], *C_h[noOfstreams]; + + unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N); + + for (int i=0; i < noOfstreams; ++i) { + HipTest::initArrays(&A_d[i], &B_d[i], &C_d[i], + &A_h[i], &B_h[i], &C_h[i], N, false); + } + + hipStream_t stream[noOfstreams]; + for (int i=0; i < noOfstreams; ++i) { + HIPCHECK(hipStreamCreate(&stream[i])); + } + + for (int i=0; i < noOfstreams; ++i) { + HIPCHECK(hipMemcpyWithStream(A_d[i], A_h[i], Nbytes, + hipMemcpyHostToDevice, stream[i])); + HIPCHECK(hipMemcpyWithStream(B_d[i], B_h[i], Nbytes, + hipMemcpyHostToDevice, stream[i])); + } + + for (int i=0; i < noOfstreams; ++i) { + hipLaunchKernelGGL(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock), + 0, stream[i], static_cast(A_d[i]), + static_cast(B_d[i]), C_d[i], N); + } + + for (int i=0; i < noOfstreams; ++i) { + HIPCHECK(hipStreamSynchronize(stream[i])); + HIPCHECK(hipMemcpy(C_h[i], C_d[i], Nbytes, hipMemcpyDeviceToHost)); + HipTest::checkVectorADD(A_h[i], B_h[i], C_h[i], N); + } + + for (int i=0; i < noOfstreams; ++i) { + HipTest::freeArrays(A_d[i], B_d[i], C_d[i], A_h[i], B_h[i], C_h[i], false); + HIPCHECK(hipStreamDestroy(stream[i])); + } +} + +void HipMemcpyWithStreamMultiThreadtests::TestDtoDonSameDevice(void) { + size_t Nbytes = N * sizeof(int); + int numDevices = 0; + int noOfstreams = 2; + int *A_d[noOfstreams], *B_d[noOfstreams], *C_d[noOfstreams]; + int *A_h[noOfstreams], *B_h[noOfstreams], *C_h[noOfstreams]; + + unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N); + + HipTest::initArrays(&A_d[0], &B_d[0], &C_d[0], + &A_h[0], &B_h[0], &C_h[0], N, false); + + + hipStream_t stream[noOfstreams]; + for (int i=0; i < noOfstreams; ++i) { + HIPCHECK(hipSetDevice(0)); + HIPCHECK(hipStreamCreate(&stream[i])); + } + + HIPCHECK(hipSetDevice(0)); + HIPCHECK(hipMalloc(&A_d[1], Nbytes)); + HIPCHECK(hipMalloc(&B_d[1], Nbytes)); + HIPCHECK(hipMalloc(&C_d[1], Nbytes)); + C_h[1] = reinterpret_cast(malloc(Nbytes)); + HIPASSERT(C_h[1] != NULL); + + HIPCHECK(hipMemcpyWithStream(A_d[0], A_h[0], Nbytes, + hipMemcpyHostToDevice, stream[0])); + HIPCHECK(hipMemcpyWithStream(B_d[0], B_h[0], Nbytes, + hipMemcpyHostToDevice, stream[0])); + + HIPCHECK(hipMemcpyWithStream(A_d[1], A_d[0], Nbytes, + hipMemcpyDeviceToDevice, stream[1])); + HIPCHECK(hipMemcpyWithStream(B_d[1], B_d[0], Nbytes, + hipMemcpyDeviceToDevice, stream[1])); + + + for (int i=0; i < noOfstreams; ++i) { + HIPCHECK(hipSetDevice(0)); + hipLaunchKernelGGL(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock), + 0, stream[i], static_cast(A_d[i]), + static_cast(B_d[i]), C_d[i], N); + } + + for (int i=0; i < noOfstreams; ++i) { + HIPCHECK(hipSetDevice(0)); + HIPCHECK(hipStreamSynchronize(stream[i])); + HIPCHECK(hipMemcpy(C_h[i], C_d[i], Nbytes, hipMemcpyDeviceToHost)); + HipTest::checkVectorADD(A_h[0], B_h[0], C_h[i], N); + } + + + HipTest::freeArrays(A_d[0], B_d[0], C_d[0], A_h[0], B_h[0], C_h[0], false); + + if (A_d[1]) { + HIPCHECK(hipFree(A_d[1])); + } + if (B_d[1]) { + HIPCHECK(hipFree(B_d[1])); + } + if (C_d[1]) { + HIPCHECK(hipFree(C_d[1])); + } + if (C_h[1]) { + free(C_h[1]); + } + + + for (int i=0; i < noOfstreams; ++i) { + HIPCHECK(hipStreamDestroy(stream[i])); + } +} + +void HipMemcpyWithStreamMultiThreadtests::TestOnMultiGPUwithOneStream(void) { + size_t Nbytes = N * sizeof(int); + int numDevices = 0; + + unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N); + HIPCHECK(hipGetDeviceCount(&numDevices)); + // If you have single GPU machine the return + if (numDevices <= 1) { + return; + } + int *A_d[numDevices], *B_d[numDevices], *C_d[numDevices]; + int *A_h[numDevices], *B_h[numDevices], *C_h[numDevices]; + + hipStream_t stream[numDevices]; + for (int i=0; i < numDevices; ++i) { + HIPCHECK(hipSetDevice(i)); + HIPCHECK(hipStreamCreate(&stream[i])); + } + + for (int i=0; i < numDevices; ++i) { + HIPCHECK(hipSetDevice(i)); + HipTest::initArrays(&A_d[i], &B_d[i], &C_d[i], + &A_h[i], &B_h[i], &C_h[i], N, false); + } + + + for (int i=0; i < numDevices; ++i) { + HIPCHECK(hipSetDevice(i)); + HIPCHECK(hipMemcpyWithStream(A_d[i], A_h[i], Nbytes, + hipMemcpyHostToDevice, stream[i])); + HIPCHECK(hipMemcpyWithStream(B_d[i], B_h[i], Nbytes, + hipMemcpyHostToDevice, stream[i])); + } + + for (int i=0; i < numDevices; ++i) { + HIPCHECK(hipSetDevice(i)); + hipLaunchKernelGGL(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock), + 0, stream[i], static_cast(A_d[i]), + static_cast(B_d[i]), C_d[i], N); + } + + for (int i=0; i < numDevices; ++i) { + HIPCHECK(hipSetDevice(i)); + HIPCHECK(hipStreamSynchronize(stream[i])); + HIPCHECK(hipMemcpy(C_h[i], C_d[i], Nbytes, hipMemcpyDeviceToHost)); + HipTest::checkVectorADD(A_h[i], B_h[i], C_h[i], N); + } + + for (int i=0; i < numDevices; ++i) { + HIPCHECK(hipSetDevice(i)); + HipTest::freeArrays(A_d[i], B_d[i], C_d[i], A_h[i], B_h[i], C_h[i], false); + HIPCHECK(hipStreamDestroy(stream[i])); + } +} + +void HipMemcpyWithStreamMultiThreadtests::TestkindDtoH(void) { + size_t Nbytes = N * sizeof(int); + int numDevices = 0; + int *A_d, *B_d, *C_d; + int *A_h, *B_h, *C_h; + + unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N); + HipTest::initArrays(&A_d, &B_d, &C_d, &A_h, &B_h, &C_h, N, false); + + hipStream_t stream; + HIPCHECK(hipStreamCreate(&stream)); + + HIPCHECK(hipMemcpyWithStream(A_d, A_h, Nbytes, + hipMemcpyHostToDevice, stream)); + HIPCHECK(hipMemcpyWithStream(B_d, B_h, Nbytes, + hipMemcpyHostToDevice, stream)); + hipLaunchKernelGGL(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock), + 0, stream, static_cast(A_d), + static_cast(B_d), C_d, N); + HIPCHECK(hipStreamSynchronize(stream)); + HIPCHECK(hipMemcpyWithStream(C_h, C_d, Nbytes, + hipMemcpyDeviceToHost, stream)); + HipTest::checkVectorADD(A_h, B_h, C_h, N); + + HipTest::freeArrays(A_d, B_d, C_d, A_h, B_h, C_h, false); + HIPCHECK(hipStreamDestroy(stream)); +} + + +void HipMemcpyWithStreamMultiThreadtests::TestkindDtoD(void) { + size_t Nbytes = N * sizeof(int); + int numDevices = 0; + + + unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N); + HIPCHECK(hipGetDeviceCount(&numDevices)); + // If you have single GPU machine the return + if (numDevices <= 1) { + return; + } + + int *A_d[numDevices], *B_d[numDevices], *C_d[numDevices]; + int *A_h[numDevices], *B_h[numDevices], *C_h[numDevices]; + + hipStream_t stream[numDevices]; + for (int i=0; i < numDevices; ++i) { + HIPCHECK(hipSetDevice(i)); + HIPCHECK(hipStreamCreate(&stream[i])); + } + + // Initialize and create the host and device elements for first device + HIPCHECK(hipSetDevice(0)); + HipTest::initArrays(&A_d[0], &B_d[0], &C_d[0], + &A_h[0], &B_h[0], &C_h[0], N, false); + + for (int i=1; i < numDevices; ++i) { + HIPCHECK(hipSetDevice(i)) + HIPCHECK(hipMalloc(&A_d[i], Nbytes)); + HIPCHECK(hipMalloc(&B_d[i], Nbytes)); + HIPCHECK(hipMalloc(&C_d[i], Nbytes)); + C_h[i] = reinterpret_cast(malloc(Nbytes)); + HIPASSERT(C_h[i] != NULL); + } + + + + HIPCHECK(hipSetDevice(0)); + HIPCHECK(hipMemcpyWithStream(A_d[0], A_h[0], Nbytes, + hipMemcpyHostToDevice, stream[0])); + HIPCHECK(hipMemcpyWithStream(B_d[0], B_h[0], Nbytes, + hipMemcpyHostToDevice, stream[0])); + + // Copying device data from 1st GPU to the rest of the the GPUs that is + // numDevices in the setup. 1st GPU start numbering from 0,1,2..n etc. + for (int i=1; i < numDevices; ++i) { + HIPCHECK(hipSetDevice(i)); + HIPCHECK(hipMemcpyWithStream(A_d[i], A_d[0], Nbytes, + hipMemcpyDeviceToDevice, stream[i])); + HIPCHECK(hipMemcpyWithStream(B_d[i], B_d[0], Nbytes, + hipMemcpyDeviceToDevice, stream[i])); + } + + + // Launching the kernel including the 1st GPU to the no of GPUs present + // in the setup. 1st GPU start numbering from 0,1,2..n etc. + for (int i=0; i < numDevices; ++i) { + HIPCHECK(hipSetDevice(i)); + hipLaunchKernelGGL(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock), + 0, stream[i], static_cast(A_d[i]), + static_cast(B_d[i]), C_d[i], N); + } + + for (int i=0; i < numDevices; ++i) { + HIPCHECK(hipSetDevice(i)); + HIPCHECK(hipStreamSynchronize(stream[i])); + HIPCHECK(hipMemcpy(C_h[i], C_d[i], Nbytes, hipMemcpyDeviceToHost)); + HipTest::checkVectorADD(A_h[0], B_h[0], C_h[i], N); + } + + HipTest::freeArrays(A_d[0], B_d[0], C_d[0], A_h[0], B_h[0], C_h[0], false); + HIPCHECK(hipStreamDestroy(stream[0])); + + for (int i=1; i < numDevices; ++i) { + if (A_d[i]) { + HIPCHECK(hipFree(A_d[i])); + } + if (B_d[i]) { + HIPCHECK(hipFree(B_d[i])); + } + if (C_d[i]) { + HIPCHECK(hipFree(C_d[i])); + } + if (C_h[i]) { + free(C_h[i]); + } + HIPCHECK(hipStreamDestroy(stream[i])); + } +} + +void HipMemcpyWithStreamMultiThreadtests::TestkindDefault(void) { + size_t Nbytes = N * sizeof(int); + int numDevices = 0; + int *A_d, *B_d, *C_d; + int *A_h, *B_h, *C_h; + + + unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N); + HipTest::initArrays(&A_d, &B_d, &C_d, &A_h, &B_h, &C_h, N, false); + + hipStream_t stream; + HIPCHECK(hipStreamCreate(&stream)); + + HIPCHECK(hipMemcpyWithStream(A_d, A_h, Nbytes, hipMemcpyDefault, stream)); + HIPCHECK(hipMemcpyWithStream(B_d, B_h, Nbytes, hipMemcpyDefault, stream)); + hipLaunchKernelGGL(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock), + 0, stream, static_cast(A_d), + static_cast(B_d), C_d, N); + HIPCHECK(hipStreamSynchronize(stream)); + HIPCHECK(hipMemcpyWithStream(C_h, C_d, Nbytes, hipMemcpyDefault, stream)); + HipTest::checkVectorADD(A_h, B_h, C_h, N); + + HipTest::freeArrays(A_d, B_d, C_d, A_h, B_h, C_h, false); + HIPCHECK(hipStreamDestroy(stream)); +} + +void HipMemcpyWithStreamMultiThreadtests::TestkindDefaultForDtoD(void) { + size_t Nbytes = N * sizeof(int); + int numDevices = 0; + + + unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N); + HIPCHECK(hipGetDeviceCount(&numDevices)); + // Test case will not run on single GPU setup. + if (numDevices <= 1) { + return; + } + + int *A_d[numDevices], *B_d[numDevices], *C_d[numDevices]; + int *A_h[numDevices], *B_h[numDevices], *C_h[numDevices]; + + // Initialize and create the host and device elements for first device + HipTest::initArrays(&A_d[0], &B_d[0], &C_d[0], + &A_h[0], &B_h[0], &C_h[0], N, false); + + for (int i=1; i < numDevices; ++i) { + HIPCHECK(hipMalloc(&A_d[i], Nbytes)); + HIPCHECK(hipMalloc(&B_d[i], Nbytes)); + HIPCHECK(hipMalloc(&C_d[i], Nbytes)); + C_h[i] = reinterpret_cast(malloc(Nbytes)); + HIPASSERT(C_h[i] != NULL); + } + + hipStream_t stream[numDevices]; + for (int i=0; i < numDevices; ++i) { + HIPCHECK(hipStreamCreate(&stream[i])); + } + + HIPCHECK(hipSetDevice(0)); + HIPCHECK(hipMemcpyWithStream(A_d[0], A_h[0], Nbytes, + hipMemcpyHostToDevice, stream[0])); + HIPCHECK(hipMemcpyWithStream(B_d[0], B_h[0], Nbytes, + hipMemcpyHostToDevice, stream[0])); + + // Copying device data from 1st GPU to the rest of the the GPUs + // using hipMemcpyDefault kind that is numDevices in the setup. + // 1st GPU start numbering from 0,1,2..n etc. + for (int i=1; i < numDevices; ++i) { + HIPCHECK(hipSetDevice(i)); + HIPCHECK(hipMemcpyWithStream(A_d[i], A_d[0], Nbytes, + hipMemcpyDefault, stream[i])); + HIPCHECK(hipMemcpyWithStream(B_d[i], B_d[0], Nbytes, + hipMemcpyDefault, stream[i])); + } + + for (int i=0; i < numDevices; ++i) { + HIPCHECK(hipSetDevice(i)); + hipLaunchKernelGGL(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock), + 0, stream[i], static_cast(A_d[i]), + static_cast(B_d[i]), C_d[i], N); + } + + for (int i=0; i < numDevices; ++i) { + HIPCHECK(hipSetDevice(i)); + HIPCHECK(hipStreamSynchronize(stream[i])); + HIPCHECK(hipMemcpy(C_h[i], C_d[i], Nbytes, hipMemcpyDeviceToHost)); + // Output of each GPU is getting validated with input of 1st GPU. + HipTest::checkVectorADD(A_h[0], B_h[0], C_h[i], N); + } + + HipTest::freeArrays(A_d[0], B_d[0], C_d[0], A_h[0], B_h[0], C_h[0], false); + HIPCHECK(hipStreamDestroy(stream[0])); + + for (int i=1; i < numDevices; ++i) { + if (A_d[i]) { + HIPCHECK(hipFree(A_d[i])); + } + if (B_d[i]) { + HIPCHECK(hipFree(B_d[i])); + } + if (C_d[i]) { + HIPCHECK(hipFree(C_d[i])); + } + if (C_h[i]) { + free(C_h[i]); + } + HIPCHECK(hipStreamDestroy(stream[i])); + } +} + +void HipMemcpyWithStreamMultiThreadtests::TestkindHtoH(void) { + size_t Nbytes = N * sizeof(int); + int numDevices = 0; + int *A_h, *B_h; + + unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N); + + // Allocate memory to A_h and B_h + A_h = static_cast(malloc(Nbytes)); + HIPASSERT(A_h != NULL); + B_h = static_cast(malloc(Nbytes)); + HIPASSERT(B_h != NULL); + + for (size_t i = 0; i < N; ++i) { + if (A_h) { + (A_h)[i] = 3.146f + i; // Pi + } + } + + hipStream_t stream; + HIPCHECK(hipStreamCreate(&stream)); + + HIPCHECK(hipMemcpyWithStream(B_h, A_h, Nbytes, hipMemcpyHostToHost, stream)); + HIPCHECK(hipStreamSynchronize(stream)); + + for (size_t i = 0; i < N; i++) { + HIPASSERT(A_h[i] == B_h[i]); + } + + if (A_h) { + free(A_h); + } + if (B_h) { + free(B_h); + } + HIPCHECK(hipStreamDestroy(stream)); +} + + +void HipMemcpyWithStreamMultiThreadtests::TestwithMultiThreaded(ops op) { + int n = min(THREADS * std::thread::hardware_concurrency(), MAX_THREADS); + std::vector threads; + + for (uint32_t i = 0; i < n; i++) { + threads.emplace_back(std::thread{[&] { + switch ( op ) { + case ops::TestwithOnestream: + TestwithOnestream(); + break; + case ops::TestwithTwoStream: + TestwithTwoStream(); + break; + case ops::TestkindDtoH: + TestkindDtoH(); + break; + case ops::TestkindHtoH: + TestkindHtoH(); + break; + case ops::TestkindDtoD: + TestkindDtoD(); + break; + case ops::TestOnMultiGPUwithOneStream: + TestOnMultiGPUwithOneStream(); + break; + case ops::TestkindDefault: + TestkindDefault(); + break; + case ops::TestkindDefaultForDtoD: + TestkindDefaultForDtoD(); + break; + case ops::TestDtoDonSameDevice: + TestDtoDonSameDevice(); + break; + default:{} + } + }}); + } +} + + +int main() { + HipMemcpyWithStreamMultiThreadtests tests; + for (int op = static_cast(ops::TestwithOnestream); + op < static_cast(ops::END_OF_LIST); ++op) { + tests.TestwithMultiThreaded(static_cast(op)); + switch ( static_cast(op) ) { + case ops::TestwithOnestream: + test_passed(HipMemcpyWithStreamMultiThreadtests + ::TestwithOnestream); + break; + case ops::TestwithTwoStream: + test_passed(HipMemcpyWithStreamMultiThreadtests + ::TestwithTwoStream); + break; + case ops::TestkindDtoH: + test_passed(HipMemcpyWithStreamMultiThreadtests + ::TestkindDtoH); + break; + case ops::TestkindHtoH: + test_passed(HipMemcpyWithStreamMultiThreadtests + ::TestkindHtoH); + break; + case ops::TestkindDtoD: + test_passed(HipMemcpyWithStreamMultiThreadtests + ::TestkindDtoD); + break; + case ops::TestOnMultiGPUwithOneStream: + test_passed(HipMemcpyWithStreamMultiThreadtests + ::TestOnMultiGPUwithOneStream); + break; + case ops::TestkindDefault: + test_passed(HipMemcpyWithStreamMultiThreadtests + ::TestkindDefault); + break; + case ops::TestkindDefaultForDtoD: + test_passed(HipMemcpyWithStreamMultiThreadtests + ::TestkindDefaultForDtoD); + break; + case ops::TestDtoDonSameDevice: + test_passed(HipMemcpyWithStreamMultiThreadtests + ::TestDtoDonSameDevice); + break; + default: { test_failed("No Operation to done with API"); } + } + } +}