From f09b9af96993b0dbaa9cc3b317b57d198f306b10 Mon Sep 17 00:00:00 2001 From: Jatin Chaudhary Date: Wed, 6 Dec 2023 17:57:37 +0000 Subject: [PATCH] SWDEV-398975 - re-enable Unit_hipMemcpyWithStream_MultiThread test use correct macro inside the test Change-Id: I7034a97efe91d6c9e2d9b4e6026ec770714a01ac [ROCm/hip-tests commit: 9704c269d0e80a1f1c39edb6189b822bb19ad2b3] --- .../catch/hipTestMain/config/config_amd_linux | 2 - .../memory/hipMemcpyWithStreamMultiThread.cc | 555 ++++++++---------- 2 files changed, 255 insertions(+), 302 deletions(-) diff --git a/projects/hip-tests/catch/hipTestMain/config/config_amd_linux b/projects/hip-tests/catch/hipTestMain/config/config_amd_linux index 7a9025d125..4d065782ce 100644 --- a/projects/hip-tests/catch/hipTestMain/config/config_amd_linux +++ b/projects/hip-tests/catch/hipTestMain/config/config_amd_linux @@ -92,8 +92,6 @@ "Unit_hipMemcpyPeer_Positive_ZeroSize", "Unit_hipMemcpyPeerAsync_Positive_ZeroSize", "Fails in Stress test SWDEV-398971", - "SWDEV-398975 Seg faults in stress test", - "Unit_hipMemcpyWithStream_MultiThread", "SWDEV-398977 fails in stress tests", "Unit_hipMemset2DSync", "SWDEV-398981 fails in stress test", diff --git a/projects/hip-tests/catch/unit/memory/hipMemcpyWithStreamMultiThread.cc b/projects/hip-tests/catch/unit/memory/hipMemcpyWithStreamMultiThread.cc index 187654ac31..6e906d6ab4 100644 --- a/projects/hip-tests/catch/unit/memory/hipMemcpyWithStreamMultiThread.cc +++ b/projects/hip-tests/catch/unit/memory/hipMemcpyWithStreamMultiThread.cc @@ -39,33 +39,31 @@ static const auto MaxGPUDevices{256}; static constexpr unsigned blocksPerCU{6}; // to hide latency static constexpr unsigned threadsPerBlock{256}; -enum class ops -{ TestwithOnestream, - TestwithTwoStream, - TestOnMultiGPUwithOneStream, - TestkindDtoH, - TestkindDtoD, - TestkindHtoH, - TestkindDefault, - TestkindDefaultForDtoD, - TestDtoDonSameDevice, - END_OF_LIST +enum class ops { + TestwithOnestream, + TestwithTwoStream, + TestOnMultiGPUwithOneStream, + TestkindDtoH, + TestkindDtoD, + TestkindHtoH, + TestkindDefault, + TestkindDefaultForDtoD, + TestDtoDonSameDevice, + END_OF_LIST }; namespace MemcpyStream { - unsigned setNumBlocks(int blocksPerCU, int threadsPerBlock, - size_t N) { - int device; - HIP_CHECK(hipGetDevice(&device)); - hipDeviceProp_t props; - HIP_CHECK(hipGetDeviceProperties(&props, device)); +void setNumBlocks(int blocksPerCU, int threadsPerBlock, size_t N, unsigned& blocks) { + int device; + HIP_CHECK_THREAD(hipGetDevice(&device)); + hipDeviceProp_t props; + HIP_CHECK_THREAD(hipGetDeviceProperties(&props, device)); - unsigned blocks = props.multiProcessorCount * blocksPerCU; - if (blocks * threadsPerBlock > N) { - blocks = (N + threadsPerBlock - 1) / threadsPerBlock; - } - return blocks; + blocks = props.multiProcessorCount * blocksPerCU; + if (blocks * threadsPerBlock > N) { + blocks = (N + threadsPerBlock - 1) / threadsPerBlock; } +} } // namespace MemcpyStream @@ -73,47 +71,40 @@ class HipMemcpyWithStreamMultiThreadtests { public: // Test hipMemcpyWithStream with one streams and launch kernel in // that stream, verify the data. - void TestwithOnestream(bool &val_res); + void TestwithOnestream(bool& val_res); // Test hipMemcpyWithStream with two streams and launch kernels in // two streams, verify the data. - void TestwithTwoStream(bool &val_res); + void TestwithTwoStream(bool& val_res); // Test hipMemcpyWithStream with one stream for each gpu and launch // kernels in each, verify the data - void TestOnMultiGPUwithOneStream(bool &val_res); + void TestOnMultiGPUwithOneStream(bool& val_res); // Test hipMemcpyWithStream to copy data from // device to host (hipMemcpyDeviceToHost). - void TestkindDtoH(bool &val_res); + void TestkindDtoH(bool& val_res); // Test hipMemcpyWithStream with hipMemcpyDeviceToDevice on MultiGPU. - void TestkindDtoD(bool &val_res); + void TestkindDtoD(bool& val_res); // Test hipMemcpyWithStream with hipMemcpyHostToHost. - void TestkindHtoH(bool &val_res); + void TestkindHtoH(bool& val_res); // Test hipMemcpyWithStream with hipMemcpyDefault. - void TestkindDefault(bool &val_res); + void TestkindDefault(bool& val_res); // Test hipMemcpyWithStream with hipMemcpyDefault for // device to device transfer case. - void TestkindDefaultForDtoD(bool &val_res); + void TestkindDefaultForDtoD(bool& val_res); // Test hipMemcpyWithStream with hipMemcpyDeviceToDevice on same device. - void TestDtoDonSameDevice(bool &val_res); + void TestDtoDonSameDevice(bool& val_res); // Allocate Memory - void AllocateMemory(int** A_d, int** B_d, - int** C_d, int** A_h, - int** B_h, - int** C_h); + void AllocateMemory(int** A_d, int** B_d, int** C_d, int** A_h, int** B_h, int** C_h); // DeAllocate Memory - void DeAllocateMemory(int* A_d, int* B_d, - int* C_d, int* A_h, int* B_h, - int* C_h); + void DeAllocateMemory(int* A_d, int* B_d, int* C_d, int* A_h, int* B_h, int* C_h); // Validate Result - bool ValidateResult(int *A_h, int *B_h, int *C_h); + bool ValidateResult(int* A_h, int* B_h, int* C_h); }; -void HipMemcpyWithStreamMultiThreadtests::AllocateMemory(int** A_d, int** B_d, - int** C_d, int** A_h, - int** B_h, - int** C_h) { - HIPCHECK(hipMalloc(A_d, Nbytes)); - HIPCHECK(hipMalloc(B_d, Nbytes)); - HIPCHECK(hipMalloc(C_d, Nbytes)); +void HipMemcpyWithStreamMultiThreadtests::AllocateMemory(int** A_d, int** B_d, int** C_d, int** A_h, + int** B_h, int** C_h) { + HIP_CHECK_THREAD(hipMalloc(A_d, Nbytes)); + HIP_CHECK_THREAD(hipMalloc(B_d, Nbytes)); + HIP_CHECK_THREAD(hipMalloc(C_d, Nbytes)); *A_h = reinterpret_cast(malloc(Nbytes)); *B_h = reinterpret_cast(malloc(Nbytes)); *C_h = reinterpret_cast(malloc(Nbytes)); @@ -125,20 +116,18 @@ void HipMemcpyWithStreamMultiThreadtests::AllocateMemory(int** A_d, int** B_d, } } -void HipMemcpyWithStreamMultiThreadtests::DeAllocateMemory(int* A_d, int* B_d, - int* C_d, int* A_h, int* B_h, - int* C_h) { - HIP_CHECK(hipFree(A_d)); - HIP_CHECK(hipFree(B_d)); - HIP_CHECK(hipFree(C_d)); +void HipMemcpyWithStreamMultiThreadtests::DeAllocateMemory(int* A_d, int* B_d, int* C_d, int* A_h, + int* B_h, int* C_h) { + HIP_CHECK_THREAD(hipFree(A_d)); + HIP_CHECK_THREAD(hipFree(B_d)); + HIP_CHECK_THREAD(hipFree(C_d)); free(A_h); free(B_h); free(C_h); } -bool HipMemcpyWithStreamMultiThreadtests:: - ValidateResult(int *A_h, int *B_h, int *C_h) { +bool HipMemcpyWithStreamMultiThreadtests::ValidateResult(int* A_h, int* B_h, int* C_h) { bool TestPassed = true; for (size_t i = 0; i < N; i++) { if ((A_h[i] + B_h[i]) != C_h[i]) { @@ -150,123 +139,113 @@ bool HipMemcpyWithStreamMultiThreadtests:: } -void HipMemcpyWithStreamMultiThreadtests::TestwithOnestream(bool &val_res) { +void HipMemcpyWithStreamMultiThreadtests::TestwithOnestream(bool& val_res) { int *A_d, *B_d, *C_d; int *A_h, *B_h, *C_h; size_t Nbytes{N * sizeof(int)}; AllocateMemory(&A_d, &B_d, &C_d, &A_h, &B_h, &C_h); - unsigned blocks = MemcpyStream::setNumBlocks(blocksPerCU, threadsPerBlock, N); + unsigned blocks = 0; + MemcpyStream::setNumBlocks(blocksPerCU, threadsPerBlock, N, blocks); hipStream_t stream; - HIPCHECK(hipStreamCreate(&stream)); + HIP_CHECK_THREAD(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); - HIP_CHECK(hipGetLastError()); - HIPCHECK(hipStreamSynchronize(stream)); - HIPCHECK(hipMemcpy(C_h, C_d, Nbytes, hipMemcpyDeviceToHost)); + HIP_CHECK_THREAD(hipMemcpyWithStream(A_d, A_h, Nbytes, hipMemcpyHostToDevice, stream)); + HIP_CHECK_THREAD(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); + HIP_CHECK_THREAD(hipGetLastError()); + HIP_CHECK_THREAD(hipStreamSynchronize(stream)); + HIP_CHECK_THREAD(hipMemcpy(C_h, C_d, Nbytes, hipMemcpyDeviceToHost)); val_res = ValidateResult(A_h, B_h, C_h); DeAllocateMemory(A_d, B_d, C_d, A_h, B_h, C_h); - HIPCHECK(hipStreamDestroy(stream)); + HIP_CHECK_THREAD(hipStreamDestroy(stream)); } -void HipMemcpyWithStreamMultiThreadtests::TestwithTwoStream(bool &val_res) { +void HipMemcpyWithStreamMultiThreadtests::TestwithTwoStream(bool& val_res) { size_t Nbytes = N * sizeof(int); const 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 = MemcpyStream::setNumBlocks(blocksPerCU, threadsPerBlock, N); + unsigned blocks = 0; + MemcpyStream::setNumBlocks(blocksPerCU, threadsPerBlock, N, blocks); - for (int i=0; i < NoofStreams; ++i) { - AllocateMemory(&A_d[i], &B_d[i], &C_d[i], - &A_h[i], &B_h[i], &C_h[i]); + for (int i = 0; i < NoofStreams; ++i) { + AllocateMemory(&A_d[i], &B_d[i], &C_d[i], &A_h[i], &B_h[i], &C_h[i]); } hipStream_t stream[NoofStreams]; - for (int i=0; i < NoofStreams; ++i) { - HIPCHECK(hipStreamCreate(&stream[i])); + for (int i = 0; i < NoofStreams; ++i) { + HIP_CHECK_THREAD(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) { + HIP_CHECK_THREAD(hipMemcpyWithStream(A_d[i], A_h[i], Nbytes, hipMemcpyHostToDevice, stream[i])); + HIP_CHECK_THREAD(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); - HIP_CHECK(hipGetLastError()); + 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); + HIP_CHECK_THREAD(hipGetLastError()); } - for (int i=0; i < NoofStreams; ++i) { - HIPCHECK(hipStreamSynchronize(stream[i])); - HIPCHECK(hipMemcpy(C_h[i], C_d[i], Nbytes, hipMemcpyDeviceToHost)); + for (int i = 0; i < NoofStreams; ++i) { + HIP_CHECK_THREAD(hipStreamSynchronize(stream[i])); + HIP_CHECK_THREAD(hipMemcpy(C_h[i], C_d[i], Nbytes, hipMemcpyDeviceToHost)); val_res = ValidateResult(A_h[i], B_h[i], C_h[i]); } - for (int i=0; i < NoofStreams; ++i) { + for (int i = 0; i < NoofStreams; ++i) { DeAllocateMemory(A_d[i], B_d[i], C_d[i], A_h[i], B_h[i], C_h[i]); - HIPCHECK(hipStreamDestroy(stream[i])); + HIP_CHECK_THREAD(hipStreamDestroy(stream[i])); } } -void HipMemcpyWithStreamMultiThreadtests::TestDtoDonSameDevice(bool &val_res) { +void HipMemcpyWithStreamMultiThreadtests::TestDtoDonSameDevice(bool& val_res) { size_t Nbytes = N * sizeof(int); const 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 = MemcpyStream::setNumBlocks(blocksPerCU, threadsPerBlock, N); + unsigned blocks = 0; + MemcpyStream::setNumBlocks(blocksPerCU, threadsPerBlock, N, blocks); - AllocateMemory(&A_d[0], &B_d[0], &C_d[0], - &A_h[0], &B_h[0], &C_h[0]); + AllocateMemory(&A_d[0], &B_d[0], &C_d[0], &A_h[0], &B_h[0], &C_h[0]); hipStream_t stream[NoofStreams]; - for (int i=0; i < NoofStreams; ++i) { - HIPCHECK(hipSetDevice(0)); - HIPCHECK(hipStreamCreate(&stream[i])); + for (int i = 0; i < NoofStreams; ++i) { + HIP_CHECK_THREAD(hipSetDevice(0)); + HIP_CHECK_THREAD(hipStreamCreate(&stream[i])); } - HIPCHECK(hipSetDevice(0)); - HIPCHECK(hipMalloc(&A_d[1], Nbytes)); - HIPCHECK(hipMalloc(&B_d[1], Nbytes)); - HIPCHECK(hipMalloc(&C_d[1], Nbytes)); + HIP_CHECK_THREAD(hipSetDevice(0)); + HIP_CHECK_THREAD(hipMalloc(&A_d[1], Nbytes)); + HIP_CHECK_THREAD(hipMalloc(&B_d[1], Nbytes)); + HIP_CHECK_THREAD(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])); + HIP_CHECK_THREAD(hipMemcpyWithStream(A_d[0], A_h[0], Nbytes, hipMemcpyHostToDevice, stream[0])); + HIP_CHECK_THREAD(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])); + HIP_CHECK_THREAD(hipMemcpyWithStream(A_d[1], A_d[0], Nbytes, hipMemcpyDeviceToDevice, stream[1])); + HIP_CHECK_THREAD(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); - HIP_CHECK(hipGetLastError()); + for (int i = 0; i < NoofStreams; ++i) { + HIP_CHECK_THREAD(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); + HIP_CHECK_THREAD(hipGetLastError()); } - for (int i=0; i < NoofStreams; ++i) { - HIPCHECK(hipSetDevice(0)); - HIPCHECK(hipStreamSynchronize(stream[i])); - HIPCHECK(hipMemcpy(C_h[i], C_d[i], Nbytes, hipMemcpyDeviceToHost)); + for (int i = 0; i < NoofStreams; ++i) { + HIP_CHECK_THREAD(hipSetDevice(0)); + HIP_CHECK_THREAD(hipStreamSynchronize(stream[i])); + HIP_CHECK_THREAD(hipMemcpy(C_h[i], C_d[i], Nbytes, hipMemcpyDeviceToHost)); val_res = ValidateResult(A_h[0], B_h[0], C_h[i]); } @@ -274,31 +253,31 @@ void HipMemcpyWithStreamMultiThreadtests::TestDtoDonSameDevice(bool &val_res) { DeAllocateMemory(A_d[0], B_d[0], C_d[0], A_h[0], B_h[0], C_h[0]); if (A_d[1]) { - HIPCHECK(hipFree(A_d[1])); + HIP_CHECK_THREAD(hipFree(A_d[1])); } if (B_d[1]) { - HIPCHECK(hipFree(B_d[1])); + HIP_CHECK_THREAD(hipFree(B_d[1])); } if (C_d[1]) { - HIPCHECK(hipFree(C_d[1])); + HIP_CHECK_THREAD(hipFree(C_d[1])); } if (C_h[1]) { free(C_h[1]); } - for (int i=0; i < NoofStreams; ++i) { - HIPCHECK(hipStreamDestroy(stream[i])); + for (int i = 0; i < NoofStreams; ++i) { + HIP_CHECK_THREAD(hipStreamDestroy(stream[i])); } } -void HipMemcpyWithStreamMultiThreadtests:: - TestOnMultiGPUwithOneStream(bool &val_res) { +void HipMemcpyWithStreamMultiThreadtests::TestOnMultiGPUwithOneStream(bool& val_res) { size_t Nbytes = N * sizeof(int); int numDevices = 0; - unsigned blocks = MemcpyStream::setNumBlocks(blocksPerCU, threadsPerBlock, N); - HIPCHECK(hipGetDeviceCount(&numDevices)); + unsigned blocks = 0; + MemcpyStream::setNumBlocks(blocksPerCU, threadsPerBlock, N, blocks); + HIP_CHECK_THREAD(hipGetDeviceCount(&numDevices)); // If you have single GPU machine the return if (numDevices <= 1) { return; @@ -307,84 +286,78 @@ void HipMemcpyWithStreamMultiThreadtests:: int *A_h[MaxGPUDevices], *B_h[MaxGPUDevices], *C_h[MaxGPUDevices]; hipStream_t stream[MaxGPUDevices]; - for (int i=0; i < numDevices; ++i) { - HIPCHECK(hipSetDevice(i)); - HIPCHECK(hipStreamCreate(&stream[i])); + for (int i = 0; i < numDevices; ++i) { + HIP_CHECK_THREAD(hipSetDevice(i)); + HIP_CHECK_THREAD(hipStreamCreate(&stream[i])); } - for (int i=0; i < numDevices; ++i) { - HIPCHECK(hipSetDevice(i)); - AllocateMemory(&A_d[i], &B_d[i], &C_d[i], - &A_h[i], &B_h[i], &C_h[i]); + for (int i = 0; i < numDevices; ++i) { + HIP_CHECK_THREAD(hipSetDevice(i)); + AllocateMemory(&A_d[i], &B_d[i], &C_d[i], &A_h[i], &B_h[i], &C_h[i]); } - 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) { + HIP_CHECK_THREAD(hipSetDevice(i)); + HIP_CHECK_THREAD(hipMemcpyWithStream(A_d[i], A_h[i], Nbytes, hipMemcpyHostToDevice, stream[i])); + HIP_CHECK_THREAD(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); - HIP_CHECK(hipGetLastError()); + for (int i = 0; i < numDevices; ++i) { + HIP_CHECK_THREAD(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); + HIP_CHECK_THREAD(hipGetLastError()); } - for (int i=0; i < numDevices; ++i) { - HIPCHECK(hipSetDevice(i)); - HIPCHECK(hipStreamSynchronize(stream[i])); - HIPCHECK(hipMemcpy(C_h[i], C_d[i], Nbytes, hipMemcpyDeviceToHost)); + for (int i = 0; i < numDevices; ++i) { + HIP_CHECK_THREAD(hipSetDevice(i)); + HIP_CHECK_THREAD(hipStreamSynchronize(stream[i])); + HIP_CHECK_THREAD(hipMemcpy(C_h[i], C_d[i], Nbytes, hipMemcpyDeviceToHost)); val_res = ValidateResult(A_h[i], B_h[i], C_h[i]); } - for (int i=0; i < numDevices; ++i) { - HIPCHECK(hipSetDevice(i)); + for (int i = 0; i < numDevices; ++i) { + HIP_CHECK_THREAD(hipSetDevice(i)); DeAllocateMemory(A_d[i], B_d[i], C_d[i], A_h[i], B_h[i], C_h[i]); - HIPCHECK(hipStreamDestroy(stream[i])); + HIP_CHECK_THREAD(hipStreamDestroy(stream[i])); } } -void HipMemcpyWithStreamMultiThreadtests::TestkindDtoH(bool &val_res) { +void HipMemcpyWithStreamMultiThreadtests::TestkindDtoH(bool& val_res) { size_t Nbytes = N * sizeof(int); int *A_d, *B_d, *C_d; int *A_h, *B_h, *C_h; - unsigned blocks = MemcpyStream::setNumBlocks(blocksPerCU, threadsPerBlock, N); + unsigned blocks = 0; + MemcpyStream::setNumBlocks(blocksPerCU, threadsPerBlock, N, blocks); AllocateMemory(&A_d, &B_d, &C_d, &A_h, &B_h, &C_h); hipStream_t stream; - HIPCHECK(hipStreamCreate(&stream)); + HIP_CHECK_THREAD(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); - HIP_CHECK(hipGetLastError()); - HIPCHECK(hipStreamSynchronize(stream)); - HIPCHECK(hipMemcpyWithStream(C_h, C_d, Nbytes, - hipMemcpyDeviceToHost, stream)); + HIP_CHECK_THREAD(hipMemcpyWithStream(A_d, A_h, Nbytes, hipMemcpyHostToDevice, stream)); + HIP_CHECK_THREAD(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); + HIP_CHECK_THREAD(hipGetLastError()); + HIP_CHECK_THREAD(hipStreamSynchronize(stream)); + HIP_CHECK_THREAD(hipMemcpyWithStream(C_h, C_d, Nbytes, hipMemcpyDeviceToHost, stream)); val_res = ValidateResult(A_h, B_h, C_h); DeAllocateMemory(A_d, B_d, C_d, A_h, B_h, C_h); - HIPCHECK(hipStreamDestroy(stream)); + HIP_CHECK_THREAD(hipStreamDestroy(stream)); } -void HipMemcpyWithStreamMultiThreadtests::TestkindDtoD(bool &val_res) { +void HipMemcpyWithStreamMultiThreadtests::TestkindDtoD(bool& val_res) { size_t Nbytes = N * sizeof(int); int numDevices = 0; - unsigned blocks = MemcpyStream::setNumBlocks(blocksPerCU, threadsPerBlock, N); - HIPCHECK(hipGetDeviceCount(&numDevices)); + unsigned blocks = 0; + MemcpyStream::setNumBlocks(blocksPerCU, threadsPerBlock, N, blocks); + HIP_CHECK_THREAD(hipGetDeviceCount(&numDevices)); // If you have single GPU machine the return if (numDevices <= 1) { return; @@ -394,116 +367,110 @@ void HipMemcpyWithStreamMultiThreadtests::TestkindDtoD(bool &val_res) { int *A_h[MaxGPUDevices], *B_h[MaxGPUDevices], *C_h[MaxGPUDevices]; hipStream_t stream[MaxGPUDevices]; - for (int i=0; i < numDevices; ++i) { - HIPCHECK(hipSetDevice(i)); - HIPCHECK(hipStreamCreate(&stream[i])); + for (int i = 0; i < numDevices; ++i) { + HIP_CHECK_THREAD(hipSetDevice(i)); + HIP_CHECK_THREAD(hipStreamCreate(&stream[i])); } // Initialize and create the host and device elements for first device - HIPCHECK(hipSetDevice(0)); - AllocateMemory(&A_d[0], &B_d[0], &C_d[0], - &A_h[0], &B_h[0], &C_h[0]); + HIP_CHECK_THREAD(hipSetDevice(0)); + AllocateMemory(&A_d[0], &B_d[0], &C_d[0], &A_h[0], &B_h[0], &C_h[0]); - 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)); + for (int i = 1; i < numDevices; ++i) { + HIP_CHECK_THREAD(hipSetDevice(i)) + HIP_CHECK_THREAD(hipMalloc(&A_d[i], Nbytes)); + HIP_CHECK_THREAD(hipMalloc(&B_d[i], Nbytes)); + HIP_CHECK_THREAD(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])); + HIP_CHECK_THREAD(hipSetDevice(0)); + HIP_CHECK_THREAD(hipMemcpyWithStream(A_d[0], A_h[0], Nbytes, hipMemcpyHostToDevice, stream[0])); + HIP_CHECK_THREAD(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])); + for (int i = 1; i < numDevices; ++i) { + HIP_CHECK_THREAD(hipSetDevice(i)); + HIP_CHECK_THREAD( + hipMemcpyWithStream(A_d[i], A_d[0], Nbytes, hipMemcpyDeviceToDevice, stream[i])); + HIP_CHECK_THREAD( + 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); - HIP_CHECK(hipGetLastError()); + for (int i = 0; i < numDevices; ++i) { + HIP_CHECK_THREAD(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); + HIP_CHECK_THREAD(hipGetLastError()); } - for (int i=0; i < numDevices; ++i) { - HIPCHECK(hipSetDevice(i)); - HIPCHECK(hipStreamSynchronize(stream[i])); - HIPCHECK(hipMemcpy(C_h[i], C_d[i], Nbytes, hipMemcpyDeviceToHost)); + for (int i = 0; i < numDevices; ++i) { + HIP_CHECK_THREAD(hipSetDevice(i)); + HIP_CHECK_THREAD(hipStreamSynchronize(stream[i])); + HIP_CHECK_THREAD(hipMemcpy(C_h[i], C_d[i], Nbytes, hipMemcpyDeviceToHost)); val_res = ValidateResult(A_h[0], B_h[0], C_h[i]); } DeAllocateMemory(A_d[0], B_d[0], C_d[0], A_h[0], B_h[0], C_h[0]); - HIPCHECK(hipStreamDestroy(stream[0])); + HIP_CHECK_THREAD(hipStreamDestroy(stream[0])); - for (int i=1; i < numDevices; ++i) { + for (int i = 1; i < numDevices; ++i) { if (A_d[i]) { - HIPCHECK(hipFree(A_d[i])); + HIP_CHECK_THREAD(hipFree(A_d[i])); } if (B_d[i]) { - HIPCHECK(hipFree(B_d[i])); + HIP_CHECK_THREAD(hipFree(B_d[i])); } if (C_d[i]) { - HIPCHECK(hipFree(C_d[i])); + HIP_CHECK_THREAD(hipFree(C_d[i])); } if (C_h[i]) { free(C_h[i]); } - HIPCHECK(hipStreamDestroy(stream[i])); + HIP_CHECK_THREAD(hipStreamDestroy(stream[i])); } } -void HipMemcpyWithStreamMultiThreadtests:: - TestkindDefault(bool &val_res) { +void HipMemcpyWithStreamMultiThreadtests::TestkindDefault(bool& val_res) { size_t Nbytes = N * sizeof(int); int *A_d, *B_d, *C_d; int *A_h, *B_h, *C_h; - unsigned blocks = MemcpyStream::setNumBlocks(blocksPerCU, threadsPerBlock, N); + unsigned blocks = 0; + MemcpyStream::setNumBlocks(blocksPerCU, threadsPerBlock, N, blocks); AllocateMemory(&A_d, &B_d, &C_d, &A_h, &B_h, &C_h); hipStream_t stream; - HIPCHECK(hipStreamCreate(&stream)); + HIP_CHECK_THREAD(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); - HIP_CHECK(hipGetLastError()); - HIPCHECK(hipStreamSynchronize(stream)); - HIPCHECK(hipMemcpyWithStream(C_h, C_d, Nbytes, hipMemcpyDefault, stream)); + HIP_CHECK_THREAD(hipMemcpyWithStream(A_d, A_h, Nbytes, hipMemcpyDefault, stream)); + HIP_CHECK_THREAD(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); + HIP_CHECK_THREAD(hipGetLastError()); + HIP_CHECK_THREAD(hipStreamSynchronize(stream)); + HIP_CHECK_THREAD(hipMemcpyWithStream(C_h, C_d, Nbytes, hipMemcpyDefault, stream)); val_res = ValidateResult(A_h, B_h, C_h); DeAllocateMemory(A_d, B_d, C_d, A_h, B_h, C_h); - HIPCHECK(hipStreamDestroy(stream)); + HIP_CHECK_THREAD(hipStreamDestroy(stream)); } -void HipMemcpyWithStreamMultiThreadtests:: - TestkindDefaultForDtoD(bool &val_res) { +void HipMemcpyWithStreamMultiThreadtests::TestkindDefaultForDtoD(bool& val_res) { size_t Nbytes = N * sizeof(int); int numDevices = 0; - unsigned blocks = MemcpyStream::setNumBlocks(blocksPerCU, threadsPerBlock, N); - HIPCHECK(hipGetDeviceCount(&numDevices)); + unsigned blocks = 0; + MemcpyStream::setNumBlocks(blocksPerCU, threadsPerBlock, N, blocks); + HIP_CHECK_THREAD(hipGetDeviceCount(&numDevices)); // Test case will not run on single GPU setup. if (numDevices <= 1) { return; @@ -513,76 +480,70 @@ void HipMemcpyWithStreamMultiThreadtests:: int *A_h[MaxGPUDevices], *B_h[MaxGPUDevices], *C_h[MaxGPUDevices]; // Initialize and create the host and device elements for first device - HIPCHECK(hipSetDevice(0)); - AllocateMemory(&A_d[0], &B_d[0], &C_d[0], - &A_h[0], &B_h[0], &C_h[0]); + HIP_CHECK_THREAD(hipSetDevice(0)); + AllocateMemory(&A_d[0], &B_d[0], &C_d[0], &A_h[0], &B_h[0], &C_h[0]); - 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)); + for (int i = 1; i < numDevices; ++i) { + HIP_CHECK_THREAD(hipSetDevice(i)); + HIP_CHECK_THREAD(hipMalloc(&A_d[i], Nbytes)); + HIP_CHECK_THREAD(hipMalloc(&B_d[i], Nbytes)); + HIP_CHECK_THREAD(hipMalloc(&C_d[i], Nbytes)); C_h[i] = reinterpret_cast(malloc(Nbytes)); HIPASSERT(C_h[i] != NULL); } hipStream_t stream[MaxGPUDevices]; - for (int i=0; i < numDevices; ++i) { - HIPCHECK(hipSetDevice(i)); - HIPCHECK(hipStreamCreate(&stream[i])); + for (int i = 0; i < numDevices; ++i) { + HIP_CHECK_THREAD(hipSetDevice(i)); + HIP_CHECK_THREAD(hipStreamCreate(&stream[i])); } - HIPCHECK(hipMemcpyWithStream(A_d[0], A_h[0], Nbytes, - hipMemcpyHostToDevice, stream[0])); - HIPCHECK(hipMemcpyWithStream(B_d[0], B_h[0], Nbytes, - hipMemcpyHostToDevice, stream[0])); + HIP_CHECK_THREAD(hipMemcpyWithStream(A_d[0], A_h[0], Nbytes, hipMemcpyHostToDevice, stream[0])); + HIP_CHECK_THREAD(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(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 = 1; i < numDevices; ++i) { + HIP_CHECK_THREAD(hipMemcpyWithStream(A_d[i], A_d[0], Nbytes, hipMemcpyDefault, stream[i])); + HIP_CHECK_THREAD(hipMemcpyWithStream(B_d[i], B_d[0], Nbytes, hipMemcpyDefault, stream[i])); } - for (int i=0; i < numDevices; ++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); - HIP_CHECK(hipGetLastError()); + for (int i = 0; i < numDevices; ++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); + HIP_CHECK_THREAD(hipGetLastError()); } - for (int i=0; i < numDevices; ++i) { - HIPCHECK(hipSetDevice(i)); // hipMemcpy will be on this device - HIPCHECK(hipStreamSynchronize(stream[i])); - HIPCHECK(hipMemcpy(C_h[i], C_d[i], Nbytes, hipMemcpyDeviceToHost)); + for (int i = 0; i < numDevices; ++i) { + HIP_CHECK_THREAD(hipSetDevice(i)); // hipMemcpy will be on this device + HIP_CHECK_THREAD(hipStreamSynchronize(stream[i])); + HIP_CHECK_THREAD(hipMemcpy(C_h[i], C_d[i], Nbytes, hipMemcpyDeviceToHost)); // Output of each GPU is getting validated with input of 1st GPU. val_res = ValidateResult(A_h[0], B_h[0], C_h[i]); } DeAllocateMemory(A_d[0], B_d[0], C_d[0], A_h[0], B_h[0], C_h[0]); - HIPCHECK(hipStreamDestroy(stream[0])); + HIP_CHECK_THREAD(hipStreamDestroy(stream[0])); - for (int i=1; i < numDevices; ++i) { + for (int i = 1; i < numDevices; ++i) { if (A_d[i]) { - HIPCHECK(hipFree(A_d[i])); + HIP_CHECK_THREAD(hipFree(A_d[i])); } if (B_d[i]) { - HIPCHECK(hipFree(B_d[i])); + HIP_CHECK_THREAD(hipFree(B_d[i])); } if (C_d[i]) { - HIPCHECK(hipFree(C_d[i])); + HIP_CHECK_THREAD(hipFree(C_d[i])); } if (C_h[i]) { free(C_h[i]); } - HIPCHECK(hipStreamDestroy(stream[i])); + HIP_CHECK_THREAD(hipStreamDestroy(stream[i])); } } -void HipMemcpyWithStreamMultiThreadtests::TestkindHtoH(bool &val_res) { +void HipMemcpyWithStreamMultiThreadtests::TestkindHtoH(bool& val_res) { size_t Nbytes = N * sizeof(int); int *A_h, *B_h; @@ -597,15 +558,15 @@ void HipMemcpyWithStreamMultiThreadtests::TestkindHtoH(bool &val_res) { } hipStream_t stream; - HIPCHECK(hipStreamCreate(&stream)); + HIP_CHECK_THREAD(hipStreamCreate(&stream)); - HIPCHECK(hipMemcpyWithStream(B_h, A_h, Nbytes, hipMemcpyHostToHost, stream)); - HIPCHECK(hipStreamSynchronize(stream)); + HIP_CHECK_THREAD(hipMemcpyWithStream(B_h, A_h, Nbytes, hipMemcpyHostToHost, stream)); + HIP_CHECK_THREAD(hipStreamSynchronize(stream)); for (size_t i = 0; i < N; i++) { if ((A_h[i] != B_h[i])) { - val_res = false; - break; + val_res = false; + break; } } @@ -615,64 +576,56 @@ void HipMemcpyWithStreamMultiThreadtests::TestkindHtoH(bool &val_res) { if (B_h) { free(B_h); } - HIPCHECK(hipStreamDestroy(stream)); + HIP_CHECK_THREAD(hipStreamDestroy(stream)); } TEST_CASE("Unit_hipMemcpyWithStream_MultiThread") { const auto Threadcount{10}; bool ret_val[Threadcount]; std::thread th[Threadcount]; - for (int op = static_cast(ops::TestwithOnestream); - op < static_cast(ops::END_OF_LIST); ++op) { + for (int op = static_cast(ops::TestwithOnestream); op < static_cast(ops::END_OF_LIST); + ++op) { HipMemcpyWithStreamMultiThreadtests tests; for (uint32_t i = 0; i < Threadcount; i++) { - switch ( static_cast(op) ) { + switch (static_cast(op)) { case ops::TestwithOnestream: - th[i] = std::thread(&HipMemcpyWithStreamMultiThreadtests:: - TestwithOnestream, - &tests, std::ref(ret_val[i])); + th[i] = std::thread(&HipMemcpyWithStreamMultiThreadtests::TestwithOnestream, &tests, + std::ref(ret_val[i])); break; case ops::TestwithTwoStream: - th[i] = std::thread(&HipMemcpyWithStreamMultiThreadtests:: - TestwithTwoStream, - &tests, std::ref(ret_val[i])); + th[i] = std::thread(&HipMemcpyWithStreamMultiThreadtests::TestwithTwoStream, &tests, + std::ref(ret_val[i])); break; case ops::TestkindDtoH: - th[i] = std::thread(&HipMemcpyWithStreamMultiThreadtests:: - TestkindDtoH, - &tests, std::ref(ret_val[i])); + th[i] = std::thread(&HipMemcpyWithStreamMultiThreadtests::TestkindDtoH, &tests, + std::ref(ret_val[i])); break; case ops::TestkindHtoH: - th[i] = std::thread(&HipMemcpyWithStreamMultiThreadtests:: - TestkindHtoH, - &tests, std::ref(ret_val[i])); + th[i] = std::thread(&HipMemcpyWithStreamMultiThreadtests::TestkindHtoH, &tests, + std::ref(ret_val[i])); break; case ops::TestkindDtoD: - th[i] = std::thread(&HipMemcpyWithStreamMultiThreadtests:: - TestkindDtoD, - &tests, std::ref(ret_val[i])); + th[i] = std::thread(&HipMemcpyWithStreamMultiThreadtests::TestkindDtoD, &tests, + std::ref(ret_val[i])); break; case ops::TestOnMultiGPUwithOneStream: - th[i] = std::thread(&HipMemcpyWithStreamMultiThreadtests:: - TestOnMultiGPUwithOneStream, + th[i] = std::thread(&HipMemcpyWithStreamMultiThreadtests::TestOnMultiGPUwithOneStream, &tests, std::ref(ret_val[i])); break; case ops::TestkindDefault: - th[i] = std::thread(&HipMemcpyWithStreamMultiThreadtests:: - TestkindDefault, - &tests, std::ref(ret_val[i])); + th[i] = std::thread(&HipMemcpyWithStreamMultiThreadtests::TestkindDefault, &tests, + std::ref(ret_val[i])); break; case ops::TestkindDefaultForDtoD: - th[i] = std::thread(&HipMemcpyWithStreamMultiThreadtests:: - TestkindDefaultForDtoD, - &tests, std::ref(ret_val[i])); + th[i] = std::thread(&HipMemcpyWithStreamMultiThreadtests::TestkindDefaultForDtoD, &tests, + std::ref(ret_val[i])); break; case ops::TestDtoDonSameDevice: - th[i] = std::thread(&HipMemcpyWithStreamMultiThreadtests:: - TestDtoDonSameDevice, - &tests, std::ref(ret_val[i])); + th[i] = std::thread(&HipMemcpyWithStreamMultiThreadtests::TestDtoDonSameDevice, &tests, + std::ref(ret_val[i])); break; - default: {} + default: { + } } } @@ -680,6 +633,8 @@ TEST_CASE("Unit_hipMemcpyWithStream_MultiThread") { th[i].join(); } + HIP_CHECK_THREAD_FINALIZE(); + for (uint32_t i = 0; i < Threadcount; i++) { REQUIRE(ret_val[i] == true); }