From 18591bc68f0320526f41aead5a982d1970e31ff2 Mon Sep 17 00:00:00 2001 From: DURGESH KROTTAPALLI Date: Tue, 6 Jul 2021 23:48:24 +0530 Subject: [PATCH] SWDEV-292393 - [catch2][dtest] Tests for hipMemcpy related apis. Migrated all hipMemcpy related APIs to CATCH2 framework by optmizing the code and moving the stress related tests to stress folder. Change-Id: Id47669b49304c35d1a68fabdaaf3f6e3ab0428a5 --- catch/hipTestMain/CMakeLists.txt | 15 + catch/include/hip_test_checkers.hh | 78 +- catch/include/hip_test_common.hh | 18 + catch/include/hip_test_kernels.hh | 10 +- catch/multiproc/hipMallocConcurrency.cc | 8 +- catch/stress/CMakeLists.txt | 8 - catch/stress/memory/CMakeLists.txt | 1 + catch/stress/memory/hipMemcpyMThreadMSize.cc | 275 +++++++ catch/unit/memory/CMakeLists.txt | 8 + catch/unit/memory/hipMemcpy.cc | 645 +++++++++++++++++ catch/unit/memory/hipMemcpy2DFromArray.cc | 13 +- .../unit/memory/hipMemcpy2DFromArrayAsync.cc | 11 +- catch/unit/memory/hipMemcpy3DAsync.cc | 12 +- catch/unit/memory/hipMemcpyAllApiNegative.cc | 326 +++++++++ catch/unit/memory/hipMemcpyAsync.cc | 403 +++++++++++ catch/unit/memory/hipMemcpyDtoD.cc | 95 +++ catch/unit/memory/hipMemcpyDtoDAsync.cc | 100 +++ catch/unit/memory/hipMemcpyWithStream.cc | 628 ++++++++++++++++ .../memory/hipMemcpyWithStreamMultiThread.cc | 679 ++++++++++++++++++ catch/unit/memory/hipMemcpy_MultiThread.cc | 322 +++++++++ 20 files changed, 3593 insertions(+), 62 deletions(-) create mode 100644 catch/stress/memory/hipMemcpyMThreadMSize.cc create mode 100644 catch/unit/memory/hipMemcpy.cc create mode 100644 catch/unit/memory/hipMemcpyAllApiNegative.cc create mode 100644 catch/unit/memory/hipMemcpyAsync.cc create mode 100644 catch/unit/memory/hipMemcpyDtoD.cc create mode 100644 catch/unit/memory/hipMemcpyDtoDAsync.cc create mode 100644 catch/unit/memory/hipMemcpyWithStream.cc create mode 100644 catch/unit/memory/hipMemcpyWithStreamMultiThread.cc create mode 100644 catch/unit/memory/hipMemcpy_MultiThread.cc diff --git a/catch/hipTestMain/CMakeLists.txt b/catch/hipTestMain/CMakeLists.txt index c80bbb585e..2b3cb3f60c 100644 --- a/catch/hipTestMain/CMakeLists.txt +++ b/catch/hipTestMain/CMakeLists.txt @@ -54,3 +54,18 @@ if(UNIX) catch_discover_tests(MultiProcTests PROPERTIES SKIP_REGULAR_EXPRESSION "HIP_SKIP_THIS_TEST") add_dependencies(build_tests MultiProcTests) endif() + +add_executable(StressTest EXCLUDE_FROM_ALL main.cc hip_test_context.cc) +add_custom_target(build_stress_test) +if(HIP_PLATFORM MATCHES "amd") + set_property(TARGET StressTest PROPERTY CXX_STANDARD 17) +else() + target_compile_options(StressTest PUBLIC -std=c++17) +endif() +if(HIP_PLATFORM MATCHES "amd") +target_link_libraries(StressTest PRIVATE printf stream) +endif() +target_link_libraries(StressTest PRIVATE memory stdc++fs) +add_dependencies(build_stress_test StressTest) +add_custom_target(stress_test COMMAND StressTest) + diff --git a/catch/include/hip_test_checkers.hh b/catch/include/hip_test_checkers.hh index e3941d17c6..0ec7b61522 100644 --- a/catch/include/hip_test_checkers.hh +++ b/catch/include/hip_test_checkers.hh @@ -1,14 +1,13 @@ #pragma once #include "hip_test_common.hh" #include -using namespace std; -#define guarantee(cond, str) \ - { \ - if (!(cond)) { \ - std::cout << str << std::endl; \ - abort(); \ - } \ - } +#define guarantee(cond, str) \ + { \ + if (!(cond)) { \ + INFO("guarantee failed: " << str); \ + abort(); \ + } \ + } namespace HipTest { @@ -49,31 +48,20 @@ size_t checkVectors(T* A, T* B, T* Out, size_t N, T (*F)(T a, T b), bool expectM return mismatchCount; } template // pointer type -void checkArray(T hData, T hOutputData, size_t width, size_t height,size_t depth) { +bool checkArray(T* hData, T* hOutputData, size_t width, size_t height,size_t depth = 1) { for (size_t i = 0; i < depth; i++) { for (size_t j = 0; j < height; j++) { for (size_t k = 0; k < width; k++) { int offset = i*width*height + j*width + k; if (hData[offset] != hOutputData[offset]) { - cerr << '[' << i << ',' << j << ',' << k << "]:" << hData[offset] << "----" - << hOutputData[offset]<<" "; - cout << "mistmatch at: " << i<< j< // pointer type -bool checkArray(T *result, T *compare, size_t width, size_t height) { - for (size_t i = 0; i < height; i++) { - for (size_t j = 0; j < width; j++) { - if (result[(i*width) + j] != compare[(i*width) + j]) { - std::cout << result[(i*width) + j] << "\t" << compare[(i*width) + j] << std::endl; - return false; - } - } - } return true; } @@ -103,17 +91,17 @@ template void setDefaultData(size_t numElements, T* A_h, T* B_h, T* for (size_t i = 0; i < numElements; i++) { if (std::is_same::value || std::is_same::value) { - if (A_h) (A_h)[i] = 3; - if (B_h) (B_h)[i] = 4; - if (C_h) (C_h)[i] = 5; + if (A_h) A_h[i] = 3; + if (B_h) B_h[i] = 4; + if (C_h) C_h[i] = 5; } else if(std::is_same::value || std::is_same::value) { - if (A_h) (A_h)[i] = 'a'; - if (B_h) (B_h)[i] = 'b'; - if (C_h) (C_h)[i] = 'c'; + if (A_h) A_h[i] = 'a'; + if (B_h) B_h[i] = 'b'; + if (C_h) C_h[i] = 'c'; } else { - if (A_h) (A_h)[i] = 3.146f + i; - if (B_h) (B_h)[i] = 1.618f + i; - if (C_h) (C_h)[i] = 1.4f + i; + if (A_h) A_h[i] = 3.146f + i; + if (B_h) B_h[i] = 1.618f + i; + if (C_h) C_h[i] = 1.4f + i; } } } @@ -135,21 +123,21 @@ bool initArraysForHost(T** A_h, T** B_h, T** C_h, size_t N, bool usePinnedHost = } else { if (A_h) { *A_h = (T*)malloc(Nbytes); - REQUIRE(*A_h != NULL); + REQUIRE(*A_h != nullptr); } if (B_h) { *B_h = (T*)malloc(Nbytes); - REQUIRE(*B_h != NULL); + REQUIRE(*B_h != nullptr); } if (C_h) { *C_h = (T*)malloc(Nbytes); - REQUIRE(*C_h != NULL); + REQUIRE(*C_h != nullptr); } } - setDefaultData(N, A_h ? *A_h : NULL, B_h ? *B_h : NULL, C_h ? *C_h : NULL); + setDefaultData(N, A_h ? *A_h : nullptr, B_h ? *B_h : nullptr, C_h ? *C_h : nullptr); return true; } @@ -210,4 +198,20 @@ bool freeArrays(T* A_d, T* B_d, T* C_d, T* A_h, T* B_h, T* C_h, bool usePinnedHo return freeArraysForHost(A_h, B_h, C_h, usePinnedHost); } + +template +unsigned setNumBlocks(T blocksPerCU, T threadsPerBlock, + size_t N) { + int device; + HIP_CHECK(hipGetDevice(&device)); + hipDeviceProp_t props; + HIP_CHECK(hipGetDeviceProperties(&props, device)); + + unsigned blocks = props.multiProcessorCount * blocksPerCU; + if (blocks * threadsPerBlock > N) { + blocks = (N + threadsPerBlock - 1) / threadsPerBlock; + } + return blocks; +} + } // namespace HipTest diff --git a/catch/include/hip_test_common.hh b/catch/include/hip_test_common.hh index 59a0971f0b..9107740333 100644 --- a/catch/include/hip_test_common.hh +++ b/catch/include/hip_test_common.hh @@ -55,6 +55,24 @@ THE SOFTWARE. #include #endif +#define HIPCHECK(error) \ + { \ + hipError_t localError = error; \ + if ((localError != hipSuccess) && (localError != hipErrorPeerAccessAlreadyEnabled)) { \ + printf("error: '%s'(%d) from %s at %s:%d\n", hipGetErrorString(localError), \ + localError, #error, __FILE__, __LINE__); \ + abort(); \ + } \ + } + +#define HIPASSERT(condition) \ + if (!(condition)) { \ + printf("assertion %s at %s:%d \n", #condition, __FILE__, __LINE__); \ + abort(); \ + } + + + // Utility Functions namespace HipTest { static inline int getDeviceCount() { diff --git a/catch/include/hip_test_kernels.hh b/catch/include/hip_test_kernels.hh index d871177f22..04b00b5ad3 100644 --- a/catch/include/hip_test_kernels.hh +++ b/catch/include/hip_test_kernels.hh @@ -72,7 +72,6 @@ __global__ void addCountReverse(const T* A_d, T* C_d, int64_t NELEM, int count) } } - template __global__ void memsetReverse(T* C_d, T val, int64_t NELEM) { size_t offset = (blockIdx.x * blockDim.x + threadIdx.x); size_t stride = blockDim.x * gridDim.x; @@ -81,4 +80,13 @@ template __global__ void memsetReverse(T* C_d, T val, int64_t NELEM C_d[i] = val; } } + +template __global__ void vector_square(const T* A_d, T* C_d, size_t N_ELMTS) { + size_t gputhread = (blockIdx.x * blockDim.x + threadIdx.x); + size_t stride = blockDim.x * gridDim.x; + for (size_t i = gputhread; i < N_ELMTS; i += stride) { + C_d[i] = A_d[i] * A_d[i]; + } +} + } // namespace HipTest diff --git a/catch/multiproc/hipMallocConcurrency.cc b/catch/multiproc/hipMallocConcurrency.cc index f2a01235e8..f6520403d5 100644 --- a/catch/multiproc/hipMallocConcurrency.cc +++ b/catch/multiproc/hipMallocConcurrency.cc @@ -12,10 +12,10 @@ #include -size_t N = 4 * 1024 * 1024; -unsigned blocksPerCU = 6; // to hide latency -unsigned threadsPerBlock = 256; +static constexpr size_t N = 4 * 1024 * 1024; +static constexpr unsigned blocksPerCU = 6; // to hide latency +static constexpr unsigned threadsPerBlock = 256; /** * Validates data consitency on supplied gpu */ @@ -31,7 +31,7 @@ bool validateMemoryOnGPU(int gpu, bool concurOnOneGPU = false) { printf("tgs allocating..\n"); HipTest::initArrays(&A_d, &B_d, &C_d, &A_h, &B_h, &C_h, N, false); - unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N); + unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N); HIP_CHECK(hipMemcpy(A_d, A_h, Nbytes, hipMemcpyHostToDevice)); HIP_CHECK(hipMemcpy(B_d, B_h, Nbytes, hipMemcpyHostToDevice)); diff --git a/catch/stress/CMakeLists.txt b/catch/stress/CMakeLists.txt index 8aea253b33..00e12f2c2c 100644 --- a/catch/stress/CMakeLists.txt +++ b/catch/stress/CMakeLists.txt @@ -1,13 +1,5 @@ -add_custom_target(build_stress_test) -add_executable(StressTest EXCLUDE_FROM_ALL ../hipTestMain/main.cc ../hipTestMain/hip_test_context.cc) -set_property(TARGET StressTest PROPERTY CXX_STANDARD 17) -target_link_libraries(StressTest PRIVATE stdc++fs) -add_dependencies(build_stress_test StressTest) -add_custom_target(stress_test COMMAND StressTest) add_subdirectory(memory) if(HIP_PLATFORM MATCHES "amd") add_subdirectory(printf) add_subdirectory(stream) -target_link_libraries(StressTest PRIVATE printf stream) endif() -target_link_libraries(StressTest PRIVATE memory) diff --git a/catch/stress/memory/CMakeLists.txt b/catch/stress/memory/CMakeLists.txt index faf6fafa18..5afd69ee11 100644 --- a/catch/stress/memory/CMakeLists.txt +++ b/catch/stress/memory/CMakeLists.txt @@ -1,6 +1,7 @@ # Common Tests - Test independent of all platforms set(TEST_SRC memcpy.cc + hipMemcpyMThreadMSize.cc ) # Create shared lib of all tests diff --git a/catch/stress/memory/hipMemcpyMThreadMSize.cc b/catch/stress/memory/hipMemcpyMThreadMSize.cc new file mode 100644 index 0000000000..791f15154e --- /dev/null +++ b/catch/stress/memory/hipMemcpyMThreadMSize.cc @@ -0,0 +1,275 @@ +/* +Copyright (c) 2021 - 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. +*/ + +#include +#include +#include +#include +#include +/* +This testfile verifies the following scenarios of all hipMemcpy API +1. Multi thread +2. Multi size +*/ + +static auto Available_Gpus{0}; +static constexpr auto MAX_GPU{256}; + +enum apiToTest {TEST_MEMCPY, TEST_MEMCPYH2D, TEST_MEMCPYD2H, TEST_MEMCPYD2D, + TEST_MEMCPYASYNC, TEST_MEMCPYH2DASYNC, TEST_MEMCPYD2HASYNC, + TEST_MEMCPYD2DASYNC}; + +template +void Memcpy_And_verify(int NUM_ELM) { + TestType *A_h, *B_h; + for (apiToTest api = TEST_MEMCPY; api <= TEST_MEMCPYD2DASYNC; + api = apiToTest(api + 1)) { + HipTest::initArrays(nullptr, nullptr, nullptr, + &A_h, &B_h, nullptr, + NUM_ELM); + HIP_CHECK(hipGetDeviceCount(&Available_Gpus)); + TestType *A_d[MAX_GPU]; + hipStream_t stream[MAX_GPU]; + for (int i = 0; i < Available_Gpus; ++i) { + HIP_CHECK(hipSetDevice(i)); + HIP_CHECK(hipMalloc(&A_d[i], NUM_ELM * sizeof(TestType))); + if (api >= TEST_MEMCPYD2D) { + HIP_CHECK(hipStreamCreate(&stream[i])); + } + } + HIP_CHECK(hipSetDevice(0)); + int canAccessPeer = 0; + switch (api) { + case TEST_MEMCPY: + { + // To test hipMemcpy() + // Copying data from host to individual devices followed by copying + // back to host and verifying the data consistency. + for (int i = 0; i < Available_Gpus; ++i) { + HIP_CHECK(hipMemcpy(A_d[i], A_h, NUM_ELM * sizeof(TestType), + hipMemcpyHostToDevice)); + HIP_CHECK(hipMemcpy(B_h, A_d[i], NUM_ELM * sizeof(TestType), + hipMemcpyDeviceToHost)); + HipTest::checkTest(A_h, B_h, NUM_ELM); + } + // Device to Device copying for all combinations + for (int i = 0; i < Available_Gpus; ++i) { + for (int j = i+1; j < Available_Gpus; ++j) { + canAccessPeer = 0; + hipDeviceCanAccessPeer(&canAccessPeer, i, j); + if (canAccessPeer) { + HIP_CHECK(hipMemcpy(A_d[j], A_d[i], NUM_ELM * sizeof(TestType), + hipMemcpyDefault)); + // Copying in reverse dir of above to check if bidirectional + // access is happening without any error + HIP_CHECK(hipMemcpy(A_d[i], A_d[j], NUM_ELM * sizeof(TestType), + hipMemcpyDefault)); + // Copying data to host to verify the content + HIP_CHECK(hipMemcpy(B_h, A_d[j], NUM_ELM * sizeof(TestType), + hipMemcpyDefault)); + HipTest::checkTest(A_h, B_h, NUM_ELM); + } + } + } + break; + } + case TEST_MEMCPYH2D: // To test hipMemcpyHtoD() + { + for (int i = 0; i < Available_Gpus; ++i) { + HIP_CHECK(hipMemcpyHtoD(hipDeviceptr_t(A_d[i]), + A_h, NUM_ELM * sizeof(TestType))); + // Copying data from device to host to check data consistency + HIP_CHECK(hipMemcpy(B_h, A_d[i], NUM_ELM * sizeof(TestType), + hipMemcpyDeviceToHost)); + HipTest::checkTest(A_h, B_h, NUM_ELM); + } + break; + } + case TEST_MEMCPYD2H: // To test hipMemcpyDtoH()--done + { + for (int i = 0; i < Available_Gpus; ++i) { + HIP_CHECK(hipMemcpy(A_d[i], A_h, NUM_ELM * sizeof(TestType), + hipMemcpyHostToDevice)); + HIP_CHECK(hipMemcpyDtoH(B_h, hipDeviceptr_t(A_d[i]), + NUM_ELM * sizeof(TestType))); + HipTest::checkTest(A_h, B_h, NUM_ELM); + } + break; + } + case TEST_MEMCPYD2D: // To test hipMemcpyDtoD() + { + if (Available_Gpus > 1) { + // First copy data from H to D and then + // from D to D followed by D to H + // HIP_CHECK(hipMemcpyHtoD(A_d[0], A_h, + // NUM_ELM * sizeof(TestType))); + int canAccessPeer = 0; + for (int i = 0; i < Available_Gpus; ++i) { + for (int j = i+1; j < Available_Gpus; ++j) { + hipDeviceCanAccessPeer(&canAccessPeer, i, j); + if (canAccessPeer) { + HIP_CHECK(hipMemcpyHtoD(hipDeviceptr_t(A_d[i]), + A_h, NUM_ELM * sizeof(TestType))); + HIP_CHECK(hipMemcpyDtoD(hipDeviceptr_t(A_d[j]), + hipDeviceptr_t(A_d[i]), NUM_ELM * sizeof(TestType))); + // Copying in direction reverse of above to check if + // bidirectional + // access is happening without any error + HIP_CHECK(hipMemcpyDtoD(hipDeviceptr_t(A_d[i]), + hipDeviceptr_t(A_d[j]), NUM_ELM * sizeof(TestType))); + HIP_CHECK(hipMemcpy(B_h, A_d[i], NUM_ELM * sizeof(TestType), + hipMemcpyDeviceToHost)); + HipTest::checkTest(A_h, B_h, NUM_ELM); + } + } + } + } else { + // As DtoD is not possible transfer data from HtH(A_h to B_h) + // so as to get through verification step + HIP_CHECK(hipMemcpy(B_h, A_h, NUM_ELM * sizeof(TestType), + hipMemcpyHostToHost)); + HipTest::checkTest(A_h, B_h, NUM_ELM); + } + break; + } + case TEST_MEMCPYASYNC: + { + // To test hipMemcpyAsync() + // Copying data from host to individual devices followed by copying + // back to host and verifying the data consistency. + for (int i = 0; i < Available_Gpus; ++i) { + HIP_CHECK(hipMemcpyAsync(A_d[i], A_h, NUM_ELM * sizeof(TestType), + hipMemcpyHostToDevice, stream[i])); + HIP_CHECK(hipMemcpyAsync(B_h, A_d[i], NUM_ELM * sizeof(TestType), + hipMemcpyDeviceToHost, stream[i])); + HIP_CHECK(hipStreamSynchronize(stream[i])); + HipTest::checkTest(A_h, B_h, NUM_ELM); + } + // Device to Device copying for all combinations + for (int i = 0; i < Available_Gpus; ++i) { + for (int j = i+1; j < Available_Gpus; ++j) { + canAccessPeer = 0; + hipDeviceCanAccessPeer(&canAccessPeer, i, j); + if (canAccessPeer) { + HIP_CHECK(hipMemcpyAsync(A_d[j], A_d[i], + NUM_ELM * sizeof(TestType), + hipMemcpyDefault, stream[i])); + // Copying in direction reverse of above to + // check if bidirectional + // access is happening without any error + HIP_CHECK(hipMemcpyAsync(A_d[i], A_d[j], + NUM_ELM * sizeof(TestType), + hipMemcpyDefault, stream[i])); + HIP_CHECK(hipStreamSynchronize(stream[i])); + HIP_CHECK(hipMemcpy(B_h, A_d[j], NUM_ELM * sizeof(TestType), + hipMemcpyDefault)); + HipTest::checkTest(A_h, B_h, NUM_ELM); + } + } + } + break; + } + case TEST_MEMCPYH2DASYNC: // To test hipMemcpyHtoDAsync() + { + for (int i = 0; i < Available_Gpus; ++i) { + HIP_CHECK(hipMemcpyHtoDAsync(hipDeviceptr_t(A_d[i]), A_h, + NUM_ELM * sizeof(TestType), stream[i])); + HIP_CHECK(hipStreamSynchronize(stream[i])); + // Copying data from device to host to check data consistency + HIP_CHECK(hipMemcpy(B_h, A_d[i], NUM_ELM * sizeof(TestType), + hipMemcpyDeviceToHost)); + HipTest::checkTest(A_h, B_h, NUM_ELM); + } + break; + } + case TEST_MEMCPYD2HASYNC: // To test hipMemcpyDtoHAsync() + { + for (int i = 0; i < Available_Gpus; ++i) { + HIP_CHECK(hipMemcpy(A_d[i], A_h, NUM_ELM * sizeof(TestType), + hipMemcpyHostToDevice)); + HIP_CHECK(hipMemcpyDtoHAsync(B_h, hipDeviceptr_t(A_d[i]), + NUM_ELM * sizeof(TestType), stream[i])); + HIP_CHECK(hipStreamSynchronize(stream[i])); + HipTest::checkTest(A_h, B_h, NUM_ELM); + } + break; + } + case TEST_MEMCPYD2DASYNC: // To test hipMemcpyDtoDAsync() + { + if (Available_Gpus > 1) { + // First copy data from H to D and then from D to D followed by D2H + HIP_CHECK(hipMemcpyHtoD(hipDeviceptr_t(A_d[0]), + A_h, NUM_ELM * sizeof(TestType))); + for (int i = 0; i < Available_Gpus; ++i) { + for (int j = i+1; j < Available_Gpus; ++j) { + canAccessPeer = 0; + hipDeviceCanAccessPeer(&canAccessPeer, i, j); + if (canAccessPeer) { + HIP_CHECK(hipSetDevice(j)); + HIP_CHECK(hipMemcpyDtoDAsync(hipDeviceptr_t(A_d[j]), + hipDeviceptr_t(A_d[i]), NUM_ELM * sizeof(TestType), + stream[i])); + // Copying in direction reverse of above to check if + // bidirectional + // access is happening without any error + HIP_CHECK(hipMemcpyDtoDAsync(hipDeviceptr_t(A_d[i]), + hipDeviceptr_t(A_d[j]), NUM_ELM * sizeof(TestType), + stream[i])); + HIP_CHECK(hipStreamSynchronize(stream[i])); + HIP_CHECK(hipMemcpy(B_h, A_d[i], NUM_ELM * sizeof(TestType), + hipMemcpyDeviceToHost)); + HipTest::checkTest(A_h, B_h, NUM_ELM); + } + } + } + } else { + // As DtoD is not possible we will transfer data + // from HtH(A_h to B_h) + // so as to get through verification step + HIP_CHECK(hipMemcpy(B_h, A_h, NUM_ELM * sizeof(TestType), + hipMemcpyHostToHost)); + HipTest::checkTest(A_h, B_h, NUM_ELM); + } + break; + } + } + for (int i = 0; i < Available_Gpus; ++i) { + HIP_CHECK(hipSetDevice(i)); + HIP_CHECK(hipFree((A_d[i]))); + if (api >= TEST_MEMCPYD2D) { + HIP_CHECK(hipStreamDestroy(stream[i])); + } + } + HipTest::freeArrays(nullptr, nullptr, nullptr, + A_h, B_h, nullptr, false); + } +} + +TEMPLATE_TEST_CASE("Unit_hipMemcpy_multiDevice-AllAPIs", "", + char, int, size_t, long double) { + auto diff_size = GENERATE(1, 5, 10, 100, 1024, 10*1024, 100*1024, + 1024*1024, 10*1024*1024, 100*1024*1024, + 1024*1024*1024); + size_t free = 0, total = 0; + HIP_CHECK(hipMemGetInfo(&free, &total)); + if ((diff_size * sizeof(TestType)) <= free) { + Memcpy_And_verify(diff_size); + HIP_CHECK(hipDeviceSynchronize()); + } +} diff --git a/catch/unit/memory/CMakeLists.txt b/catch/unit/memory/CMakeLists.txt index b82efb6b32..83186db858 100644 --- a/catch/unit/memory/CMakeLists.txt +++ b/catch/unit/memory/CMakeLists.txt @@ -16,6 +16,14 @@ set(TEST_SRC hipMemcpy2DFromArrayAsync.cc hipMemcpyAtoH.cc hipMemcpyHtoA.cc + hipMemcpyDtoD.cc + hipMemcpyDtoDAsync.cc + hipMemcpyAsync.cc + hipMemcpy.cc + hipMemcpyWithStream.cc + hipMemcpyAllApiNegative.cc + hipMemcpyWithStreamMultiThread.cc + hipMemcpy_MultiThread.cc ) # Create shared lib of all tests diff --git a/catch/unit/memory/hipMemcpy.cc b/catch/unit/memory/hipMemcpy.cc new file mode 100644 index 0000000000..71737538a0 --- /dev/null +++ b/catch/unit/memory/hipMemcpy.cc @@ -0,0 +1,645 @@ +/* +Copyright (c) 2021 - 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. +*/ + +/* +This testcase verifies following scenarios +1. hipMemcpy API along with kernel launch with different data types +2. H2D-D2D-D2H scenarios for unpinned and pinned memory +3. Boundary checks with different sizes +4. Multithread scenario +5. device offset scenario +*/ + +#include +#include +#include + +#ifdef _WIN32 +#define WIN32_LEAN_AND_MEAN +#include +#else +#include "sys/types.h" +#include "sys/sysinfo.h" +#endif + + +static constexpr auto NUM_ELM{4*1024 * 1024}; +static unsigned blocksPerCU{6}; // to hide latency +static unsigned threadsPerBlock{256}; + +template +class DeviceMemory { + public: + explicit DeviceMemory(size_t numElements); + DeviceMemory() = delete; + ~DeviceMemory(); + T* A_d() const { return _A_d + _offset; } + T* B_d() const { return _B_d + _offset; } + T* C_d() const { return _C_d + _offset; } + T* C_dd() const { return _C_dd + _offset; } + size_t maxNumElements() const { return _maxNumElements; } + void offset(int offset) { _offset = offset; } + int offset() const { return _offset; } + private: + T* _A_d; + T* _B_d; + T* _C_d; + T* _C_dd; + size_t _maxNumElements; + int _offset; +}; + +template +DeviceMemory::DeviceMemory(size_t numElements) : + _maxNumElements(numElements), _offset(0) { + T** np = nullptr; + HipTest::initArrays(&_A_d, &_B_d, &_C_d, np, np, np, numElements, 0); + size_t sizeElements = numElements * sizeof(T); + HIP_CHECK(hipMalloc(&_C_dd, sizeElements)); +} + + +template +DeviceMemory::~DeviceMemory() { + T* np = nullptr; + HipTest::freeArrays(_A_d, _B_d, _C_d, np, np, np, 0); + HIP_CHECK(hipFree(_C_dd)); + _C_dd = NULL; +} + +template +class HostMemory { + public: + HostMemory(size_t numElements, bool usePinnedHost); + HostMemory() = delete; + void reset(size_t numElements, bool full = false); + ~HostMemory(); + T* A_h() const { return _A_h + _offset; } + T* B_h() const { return _B_h + _offset; } + T* C_h() const { return _C_h + _offset; } + + size_t maxNumElements() const { return _maxNumElements; } + void offset(int offset) { _offset = offset; } + int offset() const { return _offset; } + + // Host arrays, secondary copy + T* A_hh; + T* B_hh; + bool _usePinnedHost; + + private: + size_t _maxNumElements; + int _offset; + + // Host arrays + T* _A_h; + T* _B_h; + T* _C_h; +}; + + template +HostMemory::HostMemory(size_t numElements, bool usePinnedHost) + : _usePinnedHost(usePinnedHost), _maxNumElements(numElements), _offset(0) { + T** np = nullptr; + HipTest::initArrays(np, np, np, &_A_h, &_B_h, &_C_h, + numElements, usePinnedHost); + + A_hh = NULL; + B_hh = NULL; + + + size_t sizeElements = numElements * sizeof(T); + + if (usePinnedHost) { + HIP_CHECK(hipHostMalloc(reinterpret_cast(&A_hh), sizeElements, + hipHostMallocDefault)); + HIP_CHECK(hipHostMalloc(reinterpret_cast(&B_hh), sizeElements, + hipHostMallocDefault)); + } else { + A_hh = reinterpret_cast(malloc(sizeElements)); + B_hh = reinterpret_cast(malloc(sizeElements)); + } + } + +template +void HostMemory::reset(size_t numElements, bool full) { + // Initialize the host data: + for (size_t i = 0; i < numElements; i++) { + (A_hh)[i] = 1097.0 + i; + (B_hh)[i] = 1492.0 + i; // Phi + + if (full) { + (_A_h)[i] = 3.146f + i; // Pi + (_B_h)[i] = 1.618f + i; // Phi + } + } +} + +template +HostMemory::~HostMemory() { + HipTest::freeArraysForHost(_A_h, _B_h, _C_h, _usePinnedHost); + + if (_usePinnedHost) { + HIP_CHECK(hipHostFree(A_hh)); + HIP_CHECK(hipHostFree(B_hh)); + + } else { + free(A_hh); + free(B_hh); + } +} + +#ifdef _WIN32 +void memcpytest2_get_host_memory(size_t *free, size_t *total) { + MEMORYSTATUSEX status; + status.dwLength = sizeof(status); + GlobalMemoryStatusEx(&status); + // Windows doesn't allow allocating more than half of system memory to the gpu + // Since the runtime also needs space for its internal allocations, + // we should not try to allocate more than 40% of reported system memory, + // otherwise we can run into OOM issues. + *free = static_cast(0.4 * status.ullAvailPhys); + *total = static_cast(0.4 * status.ullTotalPhys); +} +#else +struct sysinfo memInfo; +void memcpytest2_get_host_memory(size_t *free, size_t *total) { + sysinfo(&memInfo); + uint64_t freePhysMem = memInfo.freeram; + freePhysMem *= memInfo.mem_unit; + *free = freePhysMem; + uint64_t totalPhysMem = memInfo.totalram; + totalPhysMem *= memInfo.mem_unit; + *total = totalPhysMem; +} +#endif + +//--- +// Test many different kinds of memory copies. +// The subroutine allocates memory , copies to device, runs a vector +// add kernel, copies back, and +// checks the result. +// +// IN: numElements controls the number of elements used for allocations. +// IN: usePinnedHost : If true, allocate host with hipHostMalloc and is pinned +// else allocate host +// memory with malloc. IN: useHostToHost : If true, add an extra +// host-to-host copy. IN: +// useDeviceToDevice : If true, add an extra deviceto-device copy after +// result is produced. IN: +// useMemkindDefault : If true, use memkinddefault +// (runtime figures out direction). if false, use +// explicit memcpy direction. +// +template +void memcpytest2(DeviceMemory* dmem, HostMemory* hmem, + size_t numElements, bool useHostToHost, + bool useDeviceToDevice, bool useMemkindDefault) { + size_t sizeElements = numElements * sizeof(T); + + hmem->reset(numElements); + + assert(numElements <= dmem->maxNumElements()); + assert(numElements <= hmem->maxNumElements()); + + + if (useHostToHost) { + // Do some extra host-to-host copies here to mix things up: + HIP_CHECK(hipMemcpy(hmem->A_hh, hmem->A_h(), sizeElements, + useMemkindDefault ? hipMemcpyDefault : hipMemcpyHostToHost)); + HIP_CHECK(hipMemcpy(hmem->B_hh, hmem->B_h(), sizeElements, + useMemkindDefault ? hipMemcpyDefault : hipMemcpyHostToHost)); + + + HIP_CHECK(hipMemcpy(dmem->A_d(), hmem->A_hh, sizeElements, + useMemkindDefault ? hipMemcpyDefault : hipMemcpyHostToDevice)); + HIP_CHECK(hipMemcpy(dmem->B_d(), hmem->B_hh, sizeElements, + useMemkindDefault ? hipMemcpyDefault : hipMemcpyHostToDevice)); + } else { + HIP_CHECK(hipMemcpy(dmem->A_d(), hmem->A_h(), sizeElements, + useMemkindDefault ? hipMemcpyDefault : hipMemcpyHostToDevice)); + HIP_CHECK(hipMemcpy(dmem->B_d(), hmem->B_h(), sizeElements, + useMemkindDefault ? hipMemcpyDefault : hipMemcpyHostToDevice)); + } + + hipLaunchKernelGGL(HipTest::vectorADD, dim3(1), dim3(1), 0, 0, + static_cast(dmem->A_d()), static_cast(dmem->B_d()), + dmem->C_d(), numElements); + + if (useDeviceToDevice) { + // Do an extra device-to-device copy here to mix things up: + HIP_CHECK(hipMemcpy(dmem->C_dd(), dmem->C_d(), sizeElements, + useMemkindDefault ? hipMemcpyDefault : hipMemcpyDeviceToDevice)); + + // Destroy the original dmem->C_d(): + HIP_CHECK(hipMemset(dmem->C_d(), 0x5A, sizeElements)); + + HIP_CHECK(hipMemcpy(hmem->C_h(), dmem->C_dd(), sizeElements, + useMemkindDefault ? hipMemcpyDefault : hipMemcpyDeviceToHost)); + } else { + HIP_CHECK(hipMemcpy(hmem->C_h(), dmem->C_d(), sizeElements, + useMemkindDefault ? hipMemcpyDefault : hipMemcpyDeviceToHost)); + } + + HIP_CHECK(hipDeviceSynchronize()); + HipTest::checkVectorADD(hmem->A_h(), hmem->B_h(), hmem->C_h(), numElements); + + + printf(" %s success\n", __func__); +} + +// Try all the 16 possible combinations to memcpytest2 - usePinnedHost, +// useHostToHost, +// useDeviceToDevice, useMemkindDefault +template +void memcpytest2_for_type(size_t numElements) { + DeviceMemory memD(numElements); + HostMemory memU(numElements, 0 /*usePinnedHost*/); + HostMemory memP(numElements, 1 /*usePinnedHost*/); + + for (int usePinnedHost = 0; usePinnedHost <= 1; usePinnedHost++) { + for (int useHostToHost = 0; useHostToHost <= 1; useHostToHost++) { + for (int useDeviceToDevice = 0; useDeviceToDevice <= 1; + useDeviceToDevice++) { + for (int useMemkindDefault = 0; useMemkindDefault <= 1; + useMemkindDefault++) { + memcpytest2(&memD, usePinnedHost ? &memP : &memU, + numElements, useHostToHost, + useDeviceToDevice, useMemkindDefault); + } + } + } + } +} + +// Try many different sizes to memory copy. +template +void memcpytest2_sizes(size_t maxElem = 0) { + int deviceId; + HIP_CHECK(hipGetDevice(&deviceId)); + + size_t free, total, freeCPU, totalCPU; + HIP_CHECK(hipMemGetInfo(&free, &total)); + memcpytest2_get_host_memory(&freeCPU, &totalCPU); + + if (maxElem == 0) { + // Use lesser maxElem if not enough host memory available + size_t maxElemGPU = free / sizeof(T) / 8; + size_t maxElemCPU = freeCPU / sizeof(T) / 8; + maxElem = maxElemGPU < maxElemCPU ? maxElemGPU : maxElemCPU; + } + + HIP_CHECK(hipDeviceReset()); + DeviceMemory memD(maxElem); + HostMemory memU(maxElem, 0 /*usePinnedHost*/); + HostMemory memP(maxElem, 1 /*usePinnedHost*/); + + for (size_t elem = 1; elem <= maxElem; elem *= 2) { + memcpytest2(&memD, &memU, elem, 1, 1, 0); // unpinned host + memcpytest2(&memD, &memP, elem, 1, 1, 0); // pinned host + } +} + +// Try many different sizes to memory copy. +template +void memcpytest2_offsets(size_t maxElem, bool devOffsets, bool hostOffsets) { + int deviceId; + HIP_CHECK(hipGetDevice(&deviceId)); + + size_t free, total; + HIP_CHECK(hipMemGetInfo(&free, &total)); + + HIP_CHECK(hipDeviceReset()); + DeviceMemory memD(maxElem); + HostMemory memU(maxElem, 0 /*usePinnedHost*/); + HostMemory memP(maxElem, 1 /*usePinnedHost*/); + + size_t elem = maxElem / 2; + + for (size_t offset = 0; offset < 512; offset++) { + assert(elem + offset < maxElem); + if (devOffsets) { + memD.offset(offset); + } + if (hostOffsets) { + memU.offset(offset); + memP.offset(offset); + } + memcpytest2(&memD, &memU, elem, 1, 1, 0); // unpinned host + memcpytest2(&memD, &memP, elem, 1, 1, 0); // pinned host + } + + for (size_t offset = 512; offset < elem; offset *= 2) { + assert(elem + offset < maxElem); + if (devOffsets) { + memD.offset(offset); + } + if (hostOffsets) { + memU.offset(offset); + memP.offset(offset); + } + memcpytest2(&memD, &memU, elem, 1, 1, 0); // unpinned host + memcpytest2(&memD, &memP, elem, 1, 1, 0); // pinned host + } +} + +// Create multiple threads to stress multi-thread locking behavior in the +// allocation/deallocation/tracking logic: +template +void multiThread_1(bool serialize, bool usePinnedHost) { + DeviceMemory memD(NUM_ELM); + HostMemory mem1(NUM_ELM, usePinnedHost); + HostMemory mem2(NUM_ELM, usePinnedHost); + + std::thread t1(memcpytest2, &memD, &mem1, NUM_ELM, 0, 0, 0); + if (serialize) { + t1.join(); + } + + + std::thread t2(memcpytest2, &memD, &mem2, NUM_ELM, 0, 0, 0); + if (serialize) { + t2.join(); + } +} + + + +/* +This testcase verifies hipMemcpy API +Initializes device variables +Launches kernel and performs the sum of device variables +copies the result to host variable and validates the result. +*/ +TEMPLATE_TEST_CASE("Unit_hipMemcpy_KernelLaunch", "", int, float, + double) { + size_t Nbytes = NUM_ELM * sizeof(TestType); + + TestType *A_d{nullptr}, *B_d{nullptr}, *C_d{nullptr}; + TestType *A_h{nullptr}, *B_h{nullptr}, *C_h{nullptr}; + + HipTest::initArrays(&A_d, &B_d, &C_d, &A_h, &B_h, &C_h, NUM_ELM, false); + + HIP_CHECK(hipMemcpy(A_d, A_h, Nbytes, hipMemcpyHostToDevice)); + HIP_CHECK(hipMemcpy(B_d, B_h, Nbytes, hipMemcpyHostToDevice)); + + hipLaunchKernelGGL(HipTest::vectorADD, dim3(1), dim3(1), 0, 0, + static_cast(A_d), + static_cast(B_d), C_d, NUM_ELM); + + HIP_CHECK(hipMemcpy(C_h, C_d, Nbytes, hipMemcpyDeviceToHost)); + + HIP_CHECK(hipDeviceSynchronize()); + HipTest::checkVectorADD(A_h, B_h, C_h, NUM_ELM); + + HipTest::freeArrays(A_d, B_d, C_d, A_h, B_h, C_h, false); +} + +/* +This testcase verifies the following scenarios +1. H2H,H2PinMem and PinnedMem2Host +2. H2D-D2D-D2H in same GPU +3. Pinned Host Memory to device variables in same GPU +4. Device context change +5. H2D-D2D-D2H peer GPU +*/ +TEMPLATE_TEST_CASE("Unit_hipMemcpy_H2H-H2D-D2H-H2PinMem", "", int, + float, double) { + TestType *A_d{nullptr}, *B_d{nullptr}; + TestType *A_h{nullptr}, *B_h{nullptr}; + TestType *A_Ph{nullptr}, *B_Ph{nullptr}; + HIP_CHECK(hipSetDevice(0)); + HipTest::initArrays(&A_d, &B_d, nullptr, + &A_h, &B_h, nullptr, + NUM_ELM*sizeof(TestType)); + HipTest::initArrays(nullptr, nullptr, nullptr, + &A_Ph, &B_Ph, nullptr, + NUM_ELM*sizeof(TestType), true); + + SECTION("H2H, H2PinMem and PinMem2H") { + HIP_CHECK(hipMemcpy(B_h, A_h, NUM_ELM*sizeof(TestType), + hipMemcpyDefault)); + HIP_CHECK(hipMemcpy(A_Ph, B_h, NUM_ELM*sizeof(TestType), + hipMemcpyDefault)); + HIP_CHECK(hipMemcpy(B_Ph, A_Ph, NUM_ELM*sizeof(TestType), + hipMemcpyDefault)); + HipTest::checkTest(A_h, B_Ph, NUM_ELM); + } + + SECTION("H2D-D2D-D2H-SameGPU") { + HIP_CHECK(hipMemcpy(A_d, A_h, NUM_ELM*sizeof(TestType), hipMemcpyDefault)); + HIP_CHECK(hipMemcpy(B_d, A_d, NUM_ELM*sizeof(TestType), hipMemcpyDefault)); + HIP_CHECK(hipMemcpy(B_h, B_d, NUM_ELM*sizeof(TestType), hipMemcpyDefault)); + HipTest::checkTest(A_h, B_h, NUM_ELM); + } + + SECTION("pH2D-D2D-D2pH-SameGPU") { + HIP_CHECK(hipMemcpy(A_d, A_Ph, NUM_ELM*sizeof(TestType), + hipMemcpyDefault)); + HIP_CHECK(hipMemcpy(B_d, A_d, NUM_ELM*sizeof(TestType), hipMemcpyDefault)); + HIP_CHECK(hipMemcpy(B_Ph, B_d, NUM_ELM*sizeof(TestType), + hipMemcpyDefault)); + HipTest::checkTest(A_Ph, B_Ph, NUM_ELM); + } + SECTION("H2D-D2D-D2H-DeviceContextChange") { + int deviceCount = 0; + HIP_CHECK(hipGetDeviceCount(&deviceCount)); + if (deviceCount < 2) { + SUCCEED("deviceCount less then 2"); + } else { + int canAccessPeer = 0; + HIP_CHECK(hipDeviceCanAccessPeer(&canAccessPeer, 0, 1)); + if (canAccessPeer) { + HIP_CHECK(hipSetDevice(1)); + HIP_CHECK(hipMemcpy(A_d, A_h, NUM_ELM*sizeof(TestType), + hipMemcpyDefault)); + HIP_CHECK(hipMemcpy(B_d, A_d, NUM_ELM*sizeof(TestType), + hipMemcpyDefault)); + HIP_CHECK(hipMemcpy(B_h, B_d, NUM_ELM*sizeof(TestType), + hipMemcpyDefault)); + HipTest::checkTest(A_h, B_h, NUM_ELM); + } else { + SUCCEED("P2P capability is not present"); + } + } + } + + SECTION("H2D-D2D-D2H-PeerGPU") { + int deviceCount = 0; + HIP_CHECK(hipGetDeviceCount(&deviceCount)); + if (deviceCount < 2) { + SUCCEED("deviceCount less then 2"); + } else { + int canAccessPeer = 0; + HIP_CHECK(hipDeviceCanAccessPeer(&canAccessPeer, 0, 1)); + if (canAccessPeer) { + HIP_CHECK(hipSetDevice(1)); + TestType *C_d{nullptr}; + HipTest::initArrays(nullptr, nullptr, &C_d, + nullptr, nullptr, nullptr, + NUM_ELM*sizeof(TestType)); + HIP_CHECK(hipMemcpy(A_d, A_h, NUM_ELM*sizeof(TestType), + hipMemcpyDefault)); + HIP_CHECK(hipMemcpy(C_d, A_d, NUM_ELM*sizeof(TestType), + hipMemcpyDefault)); + HIP_CHECK(hipMemcpy(B_h, C_d, NUM_ELM*sizeof(TestType), + hipMemcpyDefault)); + HipTest::checkTest(A_h, B_h, NUM_ELM); + HIP_CHECK(hipFree(C_d)); + } else { + SUCCEED("P2P capability is not present"); + } + } + } + + HipTest::freeArrays(A_d, B_d, nullptr, A_h, B_h, nullptr, false); + HipTest::freeArrays(nullptr, nullptr, nullptr, A_Ph, + B_Ph, nullptr, true); +} +/* +This testcase verfies the boundary checks of hipMemcpy API for different sizes +*/ +TEST_CASE("Unit_hipMemcpy_BoundaryCheck") { + size_t maxElem = 32 * 1024 * 1024; + DeviceMemory memD(maxElem); + HostMemory memU(maxElem, 0 /*usePinnedHost*/); + HostMemory memP(maxElem, 0 /*usePinnedHost*/); + memcpytest2(&memD, &memU, 32 * 1024 * 1024, 0, 0, 0); + auto sizes = GENERATE(15 * 1024 * 1024, 16 * 1024 * 1024, + 16 * 1024 * 1024 + 16 * 1024, + 16 * 1024 * 1024 + 512 * 1024, + 17 * 1024 * 1024 + 1024, + 32 * 1024 * 1024); + memcpytest2(&memD, &memP, sizes, 0, 0, 0); +} + +/* +This testcase verifies the multi thread scenario +*/ +TEST_CASE("Unit_hipMemcpy_MultiThreadWithSerialization") { + HIP_CHECK(hipDeviceReset()); + + // Simplest cases: serialize the threads, and also used pinned memory: + // This verifies that the sub-calls to memcpytest2 are correct. + multiThread_1(true, true); + + // Serialize, but use unpinned memory to stress the unpinned memory xfer path. + multiThread_1(true, false); +} + +/* +This testcase verifies the device offsets +*/ +TEMPLATE_TEST_CASE("Unit_hipMemcpy_DeviceOffsets", "", float, double) { + HIP_CHECK(hipDeviceReset()); + size_t maxSize = 256 * 1024; + memcpytest2_offsets(maxSize, true, false); + memcpytest2_offsets(maxSize, false, true); +} + +/* +This testcase verifies hipMemcpy API with pinnedMemory and hostRegister +along with kernel launches +*/ + +TEMPLATE_TEST_CASE("Unit_hipMemcpy_PinnedRegMemWithKernelLaunch", + "", int, float, double) { + int numDevices = 0; + HIP_CHECK(hipGetDeviceCount(&numDevices)); + if (numDevices < 2) { + SUCCEED("No of devices are less than 2"); + } else { + // 1 refers to pinned Memory + // 2 refers to register Memory + int MallocPinType = GENERATE(0, 1); + size_t Nbytes = NUM_ELM * sizeof(TestType); + unsigned blocks = HipTest::setNumBlocks(blocksPerCU, + threadsPerBlock, NUM_ELM); + + TestType *A_d{nullptr}, *B_d{nullptr}, *C_d{nullptr}; + TestType *X_d{nullptr}, *Y_d{nullptr}, *Z_d{nullptr}; + TestType *A_h{nullptr}, *B_h{nullptr}, *C_h{nullptr}; + if (MallocPinType) { + HipTest::initArrays(&A_d, &B_d, &C_d, &A_h, &B_h, &C_h, NUM_ELM, true); + } else { + A_h = reinterpret_cast(malloc(Nbytes)); + HIP_CHECK(hipHostRegister(A_h, Nbytes, hipHostRegisterDefault)); + B_h = reinterpret_cast(malloc(Nbytes)); + HIP_CHECK(hipHostRegister(B_h, Nbytes, hipHostRegisterDefault)); + C_h = reinterpret_cast(malloc(Nbytes)); + HIP_CHECK(hipHostRegister(C_h, Nbytes, hipHostRegisterDefault)); + HipTest::initArrays(&A_d, &B_d, &C_d, nullptr, nullptr, + nullptr, NUM_ELM, false); + HipTest::setDefaultData(NUM_ELM, A_h, B_h, C_h); + } + HIP_CHECK(hipMemcpy(A_d, A_h, Nbytes, hipMemcpyHostToDevice)); + HIP_CHECK(hipMemcpy(B_d, B_h, Nbytes, hipMemcpyHostToDevice)); + + hipLaunchKernelGGL(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock), + 0, 0, static_cast(A_d), + static_cast(B_d), C_d, NUM_ELM); + + HIP_CHECK(hipMemcpy(C_h, C_d, Nbytes, hipMemcpyDeviceToHost)); + HipTest::checkVectorADD(A_h, B_h, C_h, NUM_ELM); + + unsigned int seed = time(0); + HIP_CHECK(hipSetDevice(rand_r(&seed) % (numDevices-1)+1)); + + int device; + hipGetDevice(&device); + std::cout <<"hipMemcpy is set to happen between device 0 and device " + <(&X_d, &Y_d, &Z_d, nullptr, + nullptr, nullptr, NUM_ELM, false); + + for (int j = 0; j < NUM_ELM; j++) { + A_h[j] = 0; + B_h[j] = 0; + C_h[j] = 0; + } + + hipMemcpy(A_h, A_d, Nbytes, hipMemcpyDeviceToHost); + hipMemcpy(X_d, A_h, Nbytes, hipMemcpyHostToDevice); + hipMemcpy(B_h, B_d, Nbytes, hipMemcpyDeviceToHost); + hipMemcpy(Y_d, B_h, Nbytes, hipMemcpyHostToDevice); + + hipLaunchKernelGGL(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock), + 0, 0, static_cast(X_d), + static_cast(Y_d), Z_d, NUM_ELM); + + HIP_CHECK(hipMemcpy(C_h, Z_d, Nbytes, hipMemcpyDeviceToHost)); + + HipTest::checkVectorADD(A_h, B_h, C_h, NUM_ELM); + + if (MallocPinType) { + HipTest::freeArrays(A_d, B_d, C_d, A_h, B_h, C_h, true); + } else { + HIP_CHECK(hipHostUnregister(A_h)); + free(A_h); + HIP_CHECK(hipHostUnregister(B_h)); + free(B_h); + HIP_CHECK(hipHostUnregister(C_h)); + free(C_h); + HipTest::freeArrays(A_d, B_d, C_d, nullptr, + nullptr, nullptr, false); + } + HipTest::freeArrays(X_d, Y_d, Z_d, nullptr, + nullptr, nullptr, false); + } +} diff --git a/catch/unit/memory/hipMemcpy2DFromArray.cc b/catch/unit/memory/hipMemcpy2DFromArray.cc index bc5fcc2e77..0e6a7aacde 100644 --- a/catch/unit/memory/hipMemcpy2DFromArray.cc +++ b/catch/unit/memory/hipMemcpy2DFromArray.cc @@ -76,11 +76,14 @@ TEST_CASE("Unit_hipMemcpy2DFromArray_ExtentValidation") { HIP_CHECK(hipSetDevice(0)); hipArray *A_d{nullptr}; size_t width{sizeof(float)*NUM_W}; - float *A_h{nullptr}, *hData{nullptr}; + float *A_h{nullptr}, *hData{nullptr}, *valData{nullptr}; // Initialization of variables HipTest::initArrays(nullptr, nullptr, nullptr, &A_h, &hData, nullptr, width*NUM_H, false); + HipTest::initArrays(nullptr, nullptr, nullptr, + nullptr, &valData, nullptr, + width*NUM_H, false); hipChannelFormatDesc desc = hipCreateChannelDesc(); HIP_CHECK(hipMallocArray(&A_d, &desc, NUM_W, NUM_H, hipArrayDefault)); @@ -101,9 +104,9 @@ TEST_CASE("Unit_hipMemcpy2DFromArray_ExtentValidation") { A_h, width, width, NUM_H, hipMemcpyHostToDevice)); HIP_CHECK(hipMemcpy2DFromArray(hData, width, A_d, - 0, 0, NUM_W*sizeof(float), + 0, 0, width, 0, hipMemcpyDeviceToHost)); - REQUIRE(HipTest::checkArray(hData, A_h, NUM_W, NUM_H) != true); + REQUIRE(HipTest::checkArray(hData, valData, NUM_W, NUM_H) == true); } // hipMemcpy2DFromArray API would return success for width and height as 0 // and does not perform any copy @@ -120,13 +123,15 @@ TEST_CASE("Unit_hipMemcpy2DFromArray_ExtentValidation") { HIP_CHECK(hipMemcpy2DFromArray(hData, width, A_d, 0, 0, 0, NUM_H, hipMemcpyDeviceToHost)); - REQUIRE(HipTest::checkArray(hData, A_h, NUM_W, NUM_H) != true); + REQUIRE(HipTest::checkArray(hData, valData, NUM_W, NUM_H) == true); } // Cleaning the memory HIP_CHECK(hipFreeArray(A_d)); HipTest::freeArrays(nullptr, nullptr, nullptr, A_h, hData, nullptr, false); + HipTest::freeArrays(nullptr, nullptr, nullptr, + nullptr, valData, nullptr, false); } /* * This Scenario Verifies hipMemcpy2DFromArray API by copying the diff --git a/catch/unit/memory/hipMemcpy2DFromArrayAsync.cc b/catch/unit/memory/hipMemcpy2DFromArrayAsync.cc index bf92878ca5..a8c9ac2944 100644 --- a/catch/unit/memory/hipMemcpy2DFromArrayAsync.cc +++ b/catch/unit/memory/hipMemcpy2DFromArrayAsync.cc @@ -83,13 +83,16 @@ TEST_CASE("Unit_hipMemcpy2DFromArrayAsync_ExtentValidation") { HIP_CHECK(hipSetDevice(0)); hipArray *A_d{nullptr}; size_t width{sizeof(float)*NUM_W}; - float *A_h{nullptr}, *hData{nullptr}; + float *A_h{nullptr}, *hData{nullptr}, *valData{nullptr}; hipStream_t stream; // Initialization of variables HipTest::initArrays(nullptr, nullptr, nullptr, &A_h, &hData, nullptr, width*NUM_H, false); + HipTest::initArrays(nullptr, nullptr, nullptr, + nullptr, &valData, nullptr, + width*NUM_H, false); hipChannelFormatDesc desc = hipCreateChannelDesc(); HIP_CHECK(hipMallocArray(&A_d, &desc, NUM_W, NUM_H, hipArrayDefault)); HIP_CHECK(hipStreamCreate(&stream)); @@ -116,7 +119,7 @@ TEST_CASE("Unit_hipMemcpy2DFromArrayAsync_ExtentValidation") { 0, 0, NUM_W*sizeof(float), 0, hipMemcpyDeviceToHost, stream)); HIP_CHECK(hipStreamSynchronize(stream)); - REQUIRE(HipTest::checkArray(hData, A_h, NUM_W, NUM_H) != true); + REQUIRE(HipTest::checkArray(hData, valData, NUM_W, NUM_H) == true); } // hipMemcpy2DFromArrayAsync API would return success for // width and height as 0 @@ -135,7 +138,7 @@ TEST_CASE("Unit_hipMemcpy2DFromArrayAsync_ExtentValidation") { 0, 0, 0, NUM_H, hipMemcpyDeviceToHost, stream)); HIP_CHECK(hipStreamSynchronize(stream)); - REQUIRE(HipTest::checkArray(hData, A_h, NUM_W, NUM_H) != true); + REQUIRE(HipTest::checkArray(hData, valData, NUM_W, NUM_H) == true); } // Cleaning the memory @@ -143,6 +146,8 @@ TEST_CASE("Unit_hipMemcpy2DFromArrayAsync_ExtentValidation") { HIP_CHECK(hipStreamDestroy(stream)); HipTest::freeArrays(nullptr, nullptr, nullptr, A_h, hData, nullptr, false); + HipTest::freeArrays(nullptr, nullptr, nullptr, + nullptr, valData, nullptr, false); } /* * This Scenario Verifies hipMemcpy2DFromArrayAsync API by copying the diff --git a/catch/unit/memory/hipMemcpy3DAsync.cc b/catch/unit/memory/hipMemcpy3DAsync.cc index c231a29bf1..017556967c 100644 --- a/catch/unit/memory/hipMemcpy3DAsync.cc +++ b/catch/unit/memory/hipMemcpy3DAsync.cc @@ -479,11 +479,12 @@ void Memcpy3DAsync::NegativeTests() { template void Memcpy3DAsync::D2D_SameDeviceMem_StreamDiffDevice() { HIP_CHECK(hipSetDevice(0)); + // Allocating the Memory AllocateMemory(); HIP_CHECK(hipSetDevice(1)); HIP_CHECK(hipStreamCreate(&stream)); - SetDefaultData(); memset(&myparms, 0x0, sizeof(hipMemcpy3DParms)); + SetDefaultData(); // Host to Device myparms.srcPtr = make_hipPitchedPtr(hData, width * sizeof(T), width, height); @@ -507,12 +508,13 @@ void Memcpy3DAsync::D2D_SameDeviceMem_StreamDiffDevice() { myparms.kind = hipMemcpyDeviceToDevice; #endif REQUIRE(hipMemcpy3DAsync(&myparms, stream) == hipSuccess); - memset(&myparms, 0x0, sizeof(hipMemcpy3DParms)); + HIP_CHECK(hipStreamSynchronize(stream)); T *hOutputData = reinterpret_cast(malloc(size)); memset(hOutputData, 0, size); - SetDefaultData(); // Device to host + memset(&myparms, 0x0, sizeof(hipMemcpy3DParms)); + SetDefaultData(); myparms.dstPtr = make_hipPitchedPtr(hOutputData, width * sizeof(T), width, height); myparms.srcArray = arr1; @@ -727,8 +729,8 @@ TEST_CASE("Unit_hipMemcpy3DAsync_multiDevice-DiffStream") { int numDevices = 0; HIP_CHECK(hipGetDeviceCount(&numDevices)); if (numDevices > 1) { - Memcpy3DAsync memcpy3dAsync(width, height, depth, - hipChannelFormatKindSigned); + Memcpy3DAsync memcpy3dAsync(width, height, depth, + hipChannelFormatKindFloat); memcpy3dAsync.D2D_SameDeviceMem_StreamDiffDevice(); } else { SUCCEED("skipping the testcases as numDevices < 2"); diff --git a/catch/unit/memory/hipMemcpyAllApiNegative.cc b/catch/unit/memory/hipMemcpyAllApiNegative.cc new file mode 100644 index 0000000000..76a5da166a --- /dev/null +++ b/catch/unit/memory/hipMemcpyAllApiNegative.cc @@ -0,0 +1,326 @@ +/* +Copyright (c) 2021 - 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. +*/ + +#include +#include +#include +#include +#include +/* +This testfile verifies the following scenarios of all hipMemcpy API +1. Negative Scenarios +2. Half Memory copy scenarios +3. Null check scenario +*/ + +static constexpr auto NUM_ELM{1024*1024}; + + +/*This testcase verifies the negative scenarios of hipMemcpy APIs +*/ +TEST_CASE("Unit_hipMemcpy_Negative") { + // Initialization of variables + float *A_d{nullptr}, *B_d{nullptr}, *C_d{nullptr}; + float *A_h{nullptr}, *B_h{nullptr}, *C_h{nullptr}; + HIP_CHECK(hipSetDevice(0)); + HipTest::initArrays(&A_d, &B_d, &C_d, + &A_h, &B_h, &C_h, + NUM_ELM*sizeof(float)); + hipStream_t stream; + HIP_CHECK(hipStreamCreate(&stream)); + SECTION("Pass nullptr to destination pointer for all Memcpy APIs") { + REQUIRE(hipMemcpy(nullptr, A_d, NUM_ELM * sizeof(float), + hipMemcpyDefault) != hipSuccess); + REQUIRE(hipMemcpyAsync(nullptr, A_h, NUM_ELM * sizeof(float), + hipMemcpyDefault, stream) != hipSuccess); + REQUIRE(hipMemcpyHtoD(hipDeviceptr_t(nullptr), A_h, + NUM_ELM * sizeof(float)) != hipSuccess); + REQUIRE(hipMemcpyHtoDAsync(hipDeviceptr_t(nullptr), A_h, + NUM_ELM * sizeof(float), + stream) != hipSuccess); + REQUIRE(hipMemcpyDtoH(nullptr, hipDeviceptr_t(A_d), + NUM_ELM * sizeof(float)) != hipSuccess); + REQUIRE(hipMemcpyDtoHAsync(nullptr, hipDeviceptr_t(A_d), + NUM_ELM * sizeof(float), + stream) != hipSuccess); + REQUIRE(hipMemcpyDtoD(hipDeviceptr_t(nullptr), + hipDeviceptr_t(A_d), NUM_ELM * sizeof(float)) + != hipSuccess); + REQUIRE(hipMemcpyDtoDAsync(hipDeviceptr_t(nullptr), + hipDeviceptr_t(A_d), + NUM_ELM * sizeof(float), stream) + != hipSuccess); + } + + SECTION("Passing nullptr to source pointer") { + REQUIRE(hipMemcpy(A_h, nullptr, NUM_ELM * sizeof(float), + hipMemcpyDefault) != hipSuccess); + REQUIRE(hipMemcpyAsync(A_d, nullptr, + NUM_ELM * sizeof(float), + hipMemcpyDefault, stream) != hipSuccess); + REQUIRE(hipMemcpyHtoD(hipDeviceptr_t(A_d), nullptr, + NUM_ELM * sizeof(float)) != hipSuccess); + REQUIRE(hipMemcpyHtoDAsync(hipDeviceptr_t(A_d), nullptr, + NUM_ELM * sizeof(float), + stream) != hipSuccess); + REQUIRE(hipMemcpyDtoH(A_h, hipDeviceptr_t(nullptr), + NUM_ELM * sizeof(float)) != hipSuccess); + REQUIRE(hipMemcpyDtoHAsync(A_h, hipDeviceptr_t(nullptr), + NUM_ELM * sizeof(float), + stream) != hipSuccess); + REQUIRE(hipMemcpyDtoD(hipDeviceptr_t(A_d), + hipDeviceptr_t(nullptr), NUM_ELM * sizeof(float)) + != hipSuccess); + REQUIRE(hipMemcpyDtoDAsync(hipDeviceptr_t(A_d), + hipDeviceptr_t(nullptr), + NUM_ELM * sizeof(float), stream) + != hipSuccess); + } + + SECTION("Passing nullptr to both source and dest pointer") { + REQUIRE(hipMemcpy(nullptr, nullptr, NUM_ELM * sizeof(float), + hipMemcpyDefault) != hipSuccess); + REQUIRE(hipMemcpyAsync(nullptr, nullptr, NUM_ELM * sizeof(float), + hipMemcpyDefault, stream) != hipSuccess); + REQUIRE(hipMemcpyHtoD(hipDeviceptr_t(nullptr), nullptr, + NUM_ELM * sizeof(float)) != hipSuccess); + REQUIRE(hipMemcpyHtoDAsync(hipDeviceptr_t(nullptr), nullptr, + NUM_ELM * sizeof(float), + stream) != hipSuccess); + REQUIRE(hipMemcpyDtoH(nullptr, hipDeviceptr_t(nullptr), + NUM_ELM * sizeof(float)) != hipSuccess); + REQUIRE(hipMemcpyDtoHAsync(nullptr, hipDeviceptr_t(nullptr), + NUM_ELM * sizeof(float), + stream) != hipSuccess); + REQUIRE(hipMemcpyDtoD(hipDeviceptr_t(nullptr), + hipDeviceptr_t(nullptr), NUM_ELM * sizeof(float)) + != hipSuccess); + REQUIRE(hipMemcpyDtoDAsync(hipDeviceptr_t(nullptr), + hipDeviceptr_t(nullptr), + NUM_ELM * sizeof(float), stream) + != hipSuccess); + } + + SECTION("Passing same pointers") { + HIP_CHECK(hipMemcpy(A_d, A_d, (NUM_ELM/2) * sizeof(float), + hipMemcpyDefault)); + HIP_CHECK(hipMemcpy(A_h, A_h, (NUM_ELM/2) * sizeof(float), + hipMemcpyDefault)); + HIP_CHECK(hipMemcpyAsync(A_d, A_d, (NUM_ELM/2) * sizeof(float), + hipMemcpyDefault, stream)); + HIP_CHECK(hipMemcpyAsync(A_h, A_h, (NUM_ELM/2) * sizeof(float), + hipMemcpyDefault, stream)); + HIP_CHECK(hipMemcpyDtoD(hipDeviceptr_t(A_d), + hipDeviceptr_t(A_d), + NUM_ELM * sizeof(float))); + HIP_CHECK(hipMemcpyDtoDAsync(hipDeviceptr_t(A_d), + hipDeviceptr_t(A_d), + NUM_ELM * sizeof(float), stream)); + } + + HipTest::freeArrays(A_d, B_d, C_d, A_h, B_h, C_h, false); + HIP_CHECK(hipStreamDestroy(stream)); +} + +/* +This testcase verifies the Nullcheck for all the 8 Memcpy APIs +*/ +TEST_CASE("Unit_hipMemcpy_NullCheck") { + // Initialization of variables + float *A_d{nullptr}, *B_d{nullptr}, *C_d{nullptr}; + float *A_h{nullptr}, *B_h{nullptr}, *C_h{nullptr}; + HIP_CHECK(hipSetDevice(0)); + HipTest::initArrays(&A_d, &B_d, &C_d, + &A_h, &B_h, &C_h, + NUM_ELM*sizeof(float)); + hipStream_t stream; + hipStreamCreate(&stream); + HIP_CHECK(hipMemcpy(A_d, C_h, + NUM_ELM*sizeof(float), + hipMemcpyHostToDevice)); + + SECTION("hipMemcpyHtoD API null size check") { + REQUIRE(hipMemcpyHtoD(hipDeviceptr_t(A_d), A_h, + 0) == hipSuccess); + HIP_CHECK(hipMemcpy(B_h, A_d, + NUM_ELM*sizeof(float), + hipMemcpyDeviceToHost)); + HipTest::checkTest(C_h, B_h, NUM_ELM); + } + + SECTION("hipMemcpyHtoDAsync API null size check") { + HIP_CHECK(hipMemcpyHtoDAsync(hipDeviceptr_t(A_d), A_h, + 0, stream)); + HIP_CHECK(hipStreamSynchronize(stream)); + HIP_CHECK(hipMemcpy(B_h, A_d, + NUM_ELM*sizeof(float), + hipMemcpyDeviceToHost)); + HipTest::checkTest(C_h, B_h, NUM_ELM); + } + SECTION("hipMemcpy API null size check") { + HIP_CHECK(hipMemcpy(A_d, B_h, 0, hipMemcpyDefault)); + HIP_CHECK(hipMemcpy(B_h, A_d, + NUM_ELM*sizeof(float), + hipMemcpyDeviceToHost)); + HipTest::checkTest(C_h, B_h, NUM_ELM); + } + SECTION("hipMemcpyAsync API null size check") { + HIP_CHECK(hipMemcpyAsync(A_d, B_h, 0, hipMemcpyDefault, stream)); + HIP_CHECK(hipMemcpy(B_h, A_d, + NUM_ELM*sizeof(float), + hipMemcpyDeviceToHost)); + HipTest::checkTest(C_h, B_h, NUM_ELM); + } + SECTION("hipMemcpyDtoH API null size check") { + HIP_CHECK(hipMemcpyDtoH(C_h, hipDeviceptr_t(A_d), 0)); + HIP_CHECK(hipMemcpy(B_h, A_d, + NUM_ELM*sizeof(float), + hipMemcpyDeviceToHost)); + HipTest::checkTest(C_h, B_h, NUM_ELM); + } + SECTION("hipMemcpyDtoHAsync API null size check") { + HIP_CHECK(hipMemcpyDtoHAsync(C_h, hipDeviceptr_t(A_d), 0, stream)); + HIP_CHECK(hipMemcpy(B_h, A_d, + NUM_ELM*sizeof(float), + hipMemcpyDeviceToHost)); + HipTest::checkTest(C_h, B_h, NUM_ELM); + } + SECTION("hipMemcpyDtoD API null size check") { + HIP_CHECK(hipMemcpy(C_d, A_h, + NUM_ELM*sizeof(float), + hipMemcpyHostToDevice)); + HIP_CHECK(hipMemcpyDtoD(hipDeviceptr_t(C_d), hipDeviceptr_t(A_d), 0)); + HIP_CHECK(hipMemcpy(B_h, C_d, + NUM_ELM*sizeof(float), + hipMemcpyDeviceToHost)); + HipTest::checkTest(A_h, B_h, NUM_ELM); + } + SECTION("hipMemcpyDtoDAsync API null size check") { + HIP_CHECK(hipMemcpy(C_d, A_h, + NUM_ELM*sizeof(float), + hipMemcpyHostToDevice)); + HIP_CHECK(hipMemcpyDtoDAsync(hipDeviceptr_t(C_d), hipDeviceptr_t(A_d), + 0, stream)); + HIP_CHECK(hipMemcpy(B_h, C_d, + NUM_ELM*sizeof(float), + hipMemcpyDeviceToHost)); + HipTest::checkTest(A_h, B_h, NUM_ELM); + } + + HipTest::freeArrays(A_d, B_d, C_d, A_h, B_h, C_h, false); + HIP_CHECK(hipStreamDestroy(stream)); +} + +/* +This testcase verifies all the hipMemcpy APIs by +copying half the memory. +*/ +TEST_CASE("Unit_hipMemcpy_HalfMemCopy") { + // Initialization of variables + float *A_d{nullptr}, *B_d{nullptr}, *C_d{nullptr}; + float *A_h{nullptr}, *B_h{nullptr}, *C_h{nullptr}; + HIP_CHECK(hipSetDevice(0)); + HipTest::initArrays(&A_d, &B_d, &C_d, + &A_h, &B_h, &C_h, + NUM_ELM*sizeof(float)); + hipStream_t stream; + hipStreamCreate(&stream); + + SECTION("hipMemcpyHtoD half memory copy") { + HIP_CHECK(hipMemcpyHtoD(hipDeviceptr_t(A_d), A_h, + (NUM_ELM * sizeof(float))/2)); + HIP_CHECK(hipMemcpy(B_h, A_d, + (NUM_ELM*sizeof(float))/2, + hipMemcpyDeviceToHost)); + HipTest::checkTest(A_h, B_h, NUM_ELM/2); + } + + SECTION("hipMemcpyHtoDAsync half memory copy") { + HIP_CHECK(hipMemcpyHtoDAsync(hipDeviceptr_t(A_d), A_h, + (NUM_ELM * sizeof(float))/2, stream)); + HIP_CHECK(hipStreamSynchronize(stream)); + HIP_CHECK(hipMemcpy(B_h, A_d, + (NUM_ELM*sizeof(float))/2, + hipMemcpyDeviceToHost)); + HipTest::checkTest(A_h, B_h, NUM_ELM/2); + } + + SECTION("hipMemcpyDtoH half memory copy") { + HIP_CHECK(hipMemcpyHtoD(hipDeviceptr_t(A_d), A_h, + (NUM_ELM * sizeof(float)))); + HIP_CHECK(hipMemcpyDtoH(B_h, hipDeviceptr_t(A_d), + (NUM_ELM * sizeof(float))/2)); + HipTest::checkTest(A_h, B_h, NUM_ELM/2); + } + + SECTION("hipMemcpyDtoHAsync half memory copy") { + HIP_CHECK(hipMemcpyHtoD(hipDeviceptr_t(A_d), A_h, + (NUM_ELM * sizeof(float)))); + HIP_CHECK(hipMemcpyDtoHAsync(B_h, hipDeviceptr_t(A_d), + (NUM_ELM * sizeof(float))/2, + stream)); + HIP_CHECK(hipStreamSynchronize(stream)); + HipTest::checkTest(A_h, B_h, NUM_ELM/2); + } + + SECTION("hipMemcpyDtoD half memory copy") { + HIP_CHECK(hipMemcpyHtoD(hipDeviceptr_t(A_d), A_h, + (NUM_ELM * sizeof(float))/2)); + HIP_CHECK(hipMemcpyDtoD(hipDeviceptr_t(B_d), hipDeviceptr_t(A_d), + (NUM_ELM*sizeof(float))/2)); + HIP_CHECK(hipMemcpy(B_h, B_d, + (NUM_ELM*sizeof(float))/2, + hipMemcpyDeviceToHost)); + HipTest::checkTest(A_h, B_h, NUM_ELM/2); + } + + SECTION("hipMemcpyDtoDAsync half memory copy") { + HIP_CHECK(hipMemcpyHtoD(hipDeviceptr_t(A_d), A_h, + (NUM_ELM * sizeof(float))/2)); + HIP_CHECK(hipMemcpyDtoDAsync(hipDeviceptr_t(B_d), hipDeviceptr_t(A_d), + (NUM_ELM*sizeof(float))/2, + stream)); + HIP_CHECK(hipMemcpy(B_h, B_d, + (NUM_ELM*sizeof(float))/2, + hipMemcpyDeviceToHost)); + HipTest::checkTest(A_h, B_h, NUM_ELM/2); + } + + SECTION("hipMemcpy half memory copy") { + HIP_CHECK(hipMemcpy(A_d, A_h + , (NUM_ELM*sizeof(float)), + hipMemcpyDefault)); + HIP_CHECK(hipMemcpy(B_h, A_d, + (NUM_ELM/2)*sizeof(float), + hipMemcpyDeviceToHost)); + HipTest::checkTest(A_h, B_h, NUM_ELM/2); + } + + SECTION("hipMemcpyAsync half memory copy") { + HIP_CHECK(hipMemcpy(A_d, A_h, + (NUM_ELM*sizeof(float)), + hipMemcpyDefault)); + HIP_CHECK(hipMemcpyAsync(B_h, A_d, + (NUM_ELM/2)*sizeof(float), + hipMemcpyDeviceToHost, stream)); + HipTest::checkTest(A_h, B_h, NUM_ELM/2); + } + HipTest::freeArrays(A_d, B_d, C_d, A_h, B_h, C_h, false); + HIP_CHECK(hipStreamDestroy(stream)); +} diff --git a/catch/unit/memory/hipMemcpyAsync.cc b/catch/unit/memory/hipMemcpyAsync.cc new file mode 100644 index 0000000000..a7e327354d --- /dev/null +++ b/catch/unit/memory/hipMemcpyAsync.cc @@ -0,0 +1,403 @@ +/* +Copyright (c) 2021 - 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. +*/ + +/* +This testcase verifies the following scenarios +1. hipMemcpyAsync with kernel launch +2. H2D-D2D-D2H-H2PinnMem and device context change scenarios +3. This test launches multiple threads which uses same stream to deploy kernel + and also launch hipMemcpyAsync() api. This test case is simulate the scenario + reported in SWDEV-181598. +*/ + +#include +#include +#include +#include + +#define NUM_THREADS 16 + +static constexpr auto NUM_ELM{1024 * 1024}; + + + +static constexpr size_t N_ELMTS{32 * 1024}; +std::atomic Thread_count { 0 }; +static unsigned blocksPerCU{6}; // to hide latency +static unsigned threadsPerBlock{256}; + +template +void Thread_func(T *A_d, T *B_d, T* C_d, T* C_h, size_t Nbytes, + hipStream_t mystream) { + unsigned blocks = HipTest::setNumBlocks(blocksPerCU, + threadsPerBlock, N_ELMTS); + hipLaunchKernelGGL(HipTest::vector_square, dim3(blocks), + dim3(threadsPerBlock), 0, + mystream, A_d, C_d, N_ELMTS); + HIP_CHECK(hipMemcpyAsync(C_h, C_d, Nbytes, hipMemcpyDeviceToHost, mystream)); + // The following two MemcpyAsync calls are for sole + // purpose of loading stream with multiple async calls + HIP_CHECK(hipMemcpyAsync(B_d, A_d, Nbytes, + hipMemcpyDeviceToDevice, mystream)); + HIP_CHECK(hipMemcpyAsync(B_d, A_d, Nbytes, + hipMemcpyDeviceToDevice, mystream)); + Thread_count++; +} + +template +void Thread_func_MultiStream() { + int Data_mismatch = 0; + T *A_d{nullptr}, *B_d{nullptr}, *C_d{nullptr}; + T *A_h{nullptr}, *B_h{nullptr}, *C_h{nullptr}; + size_t Nbytes = N_ELMTS * sizeof(T); + unsigned blocks = HipTest::setNumBlocks(blocksPerCU, + threadsPerBlock, N_ELMTS); + + HipTest::initArrays(&A_d, &B_d, &C_d, &A_h, &B_h, &C_h, N_ELMTS, false); + hipStream_t mystream; + HIP_CHECK(hipStreamCreateWithFlags(&mystream, hipStreamNonBlocking)); + HIP_CHECK(hipMemcpyAsync(A_d, A_h, Nbytes, hipMemcpyHostToDevice, mystream)); + hipLaunchKernelGGL((HipTest::vector_square), dim3(blocks), + dim3(threadsPerBlock), 0, + mystream, A_d, C_d, N_ELMTS); + HIP_CHECK(hipMemcpyAsync(C_h, C_d, Nbytes, hipMemcpyDeviceToHost, mystream)); + // The following hipMemcpyAsync() is called only to + // load stream with multiple Async calls + HIP_CHECK(hipMemcpyAsync(B_d, A_d, Nbytes, + hipMemcpyDeviceToDevice, mystream)); + Thread_count++; + + HIP_CHECK(hipStreamSynchronize(mystream)); + HIP_CHECK(hipStreamDestroy(mystream)); + // Verifying result of the kernel computation + for (size_t i = 0; i < N_ELMTS; i++) { + if (C_h[i] != A_h[i] * A_h[i]) { + Data_mismatch++; + } + } + // Releasing resources + HipTest::freeArrays(A_d, B_d, C_d, A_h, B_h, C_h, false); + REQUIRE(Data_mismatch == 0); +} + +/* +This testcase verifies hipMemcpyAsync API +Initializes device variables +Launches kernel and performs the sum of device variables +copies the result to host variable and validates the result. +*/ +TEMPLATE_TEST_CASE("Unit_hipMemcpyAsync_KernelLaunch", "", int, float, + double) { + size_t Nbytes = NUM_ELM * sizeof(TestType); + + TestType *A_d{nullptr}, *B_d{nullptr}, *C_d{nullptr}; + TestType *A_h{nullptr}, *B_h{nullptr}, *C_h{nullptr}; + HIP_CHECK(hipSetDevice(0)); + hipStream_t stream; + hipStreamCreate(&stream); + + HipTest::initArrays(&A_d, &B_d, &C_d, &A_h, &B_h, &C_h, NUM_ELM, false); + + HIP_CHECK(hipMemcpyAsync(A_d, A_h, Nbytes, hipMemcpyHostToDevice, stream)); + HIP_CHECK(hipMemcpyAsync(B_d, B_h, Nbytes, hipMemcpyHostToDevice, stream)); + HIP_CHECK(hipStreamSynchronize(stream)); + + hipLaunchKernelGGL(HipTest::vectorADD, dim3(1), dim3(1), 0, 0, + static_cast(A_d), + static_cast(B_d), C_d, NUM_ELM); + + HIP_CHECK(hipMemcpyAsync(C_h, C_d, Nbytes, hipMemcpyDeviceToHost, stream)); + HIP_CHECK(hipStreamSynchronize(stream)); + + HipTest::checkVectorADD(A_h, B_h, C_h, NUM_ELM); + + HipTest::freeArrays(A_d, B_d, C_d, A_h, B_h, C_h, false); +} +/* +This testcase verifies the following scenarios +1. H2H,H2PinMem and PinnedMem2Host +2. H2D-D2D-D2H in same GPU +3. Pinned Host Memory to device variables in same GPU +4. Device context change +5. H2D-D2D-D2H peer GPU +*/ +TEMPLATE_TEST_CASE("Unit_hipMemcpyAsync_H2H-H2D-D2H-H2PinMem", "", char, int, + float, double) { + TestType *A_d{nullptr}, *B_d{nullptr}; + TestType *A_h{nullptr}, *B_h{nullptr}; + TestType *A_Ph{nullptr}, *B_Ph{nullptr}; + HIP_CHECK(hipSetDevice(0)); + hipStream_t stream; + hipStreamCreate(&stream); + HipTest::initArrays(&A_d, &B_d, nullptr, + &A_h, &B_h, nullptr, + NUM_ELM*sizeof(TestType)); + HipTest::initArrays(nullptr, nullptr, nullptr, + &A_Ph, &B_Ph, nullptr, + NUM_ELM*sizeof(TestType), true); + + SECTION("H2H, H2PinMem and PinMem2H") { + HIP_CHECK(hipMemcpyAsync(B_h, A_h, NUM_ELM*sizeof(TestType), + hipMemcpyHostToHost, stream)); + HIP_CHECK(hipMemcpyAsync(A_Ph, B_h, NUM_ELM*sizeof(TestType), + hipMemcpyHostToHost, stream)); + HIP_CHECK(hipMemcpyAsync(B_Ph, A_Ph, NUM_ELM*sizeof(TestType), + hipMemcpyHostToHost, stream)); + HIP_CHECK(hipStreamSynchronize(stream)); + HipTest::checkTest(A_h, B_Ph, NUM_ELM); + } + + SECTION("H2D-D2D-D2H-SameGPU") { + HIP_CHECK(hipMemcpyAsync(A_d, A_h, NUM_ELM*sizeof(TestType), + hipMemcpyHostToDevice, stream)); + HIP_CHECK(hipMemcpyAsync(B_d, A_d, NUM_ELM*sizeof(TestType), + hipMemcpyDeviceToDevice, stream)); + HIP_CHECK(hipMemcpyAsync(B_h, B_d, NUM_ELM*sizeof(TestType), + hipMemcpyDeviceToHost, stream)); + HIP_CHECK(hipStreamSynchronize(stream)); + HipTest::checkTest(A_h, B_h, NUM_ELM); + } + + SECTION("pH2D-D2D-D2pH-SameGPU") { + HIP_CHECK(hipMemcpyAsync(A_d, A_Ph, NUM_ELM*sizeof(TestType), + hipMemcpyHostToDevice, stream)); + HIP_CHECK(hipMemcpyAsync(B_d, A_d, NUM_ELM*sizeof(TestType), + hipMemcpyDeviceToDevice, stream)); + HIP_CHECK(hipMemcpyAsync(B_Ph, B_d, NUM_ELM*sizeof(TestType), + hipMemcpyDeviceToHost, stream)); + HIP_CHECK(hipStreamSynchronize(stream)); + HipTest::checkTest(A_Ph, B_Ph, NUM_ELM); + } + SECTION("H2D-D2D-D2H-DeviceContextChange") { + int deviceCount = 0; + HIP_CHECK(hipGetDeviceCount(&deviceCount)); + if (deviceCount < 2) { + SUCCEED("deviceCount less then 2"); + } else { + int canAccessPeer = 0; + HIP_CHECK(hipDeviceCanAccessPeer(&canAccessPeer, 0, 1)); + if (canAccessPeer) { + HIP_CHECK(hipSetDevice(1)); + HIP_CHECK(hipMemcpyAsync(A_d, A_h, NUM_ELM*sizeof(TestType), + hipMemcpyHostToDevice, stream)); + HIP_CHECK(hipMemcpyAsync(B_d, A_d, NUM_ELM*sizeof(TestType), + hipMemcpyDeviceToDevice, stream)); + HIP_CHECK(hipMemcpyAsync(B_h, B_d, NUM_ELM*sizeof(TestType), + hipMemcpyDeviceToHost, stream)); + HIP_CHECK(hipStreamSynchronize(stream)); + HipTest::checkTest(A_h, B_h, NUM_ELM); + + } else { + SUCCEED("P2P capability is not present"); + } + } + } + + SECTION("H2D-D2D-D2H-PeerGPU") { + int deviceCount = 0; + HIP_CHECK(hipGetDeviceCount(&deviceCount)); + if (deviceCount < 2) { + SUCCEED("deviceCount less then 2"); + } else { + int canAccessPeer = 0; + HIP_CHECK(hipDeviceCanAccessPeer(&canAccessPeer, 0, 1)); + if (canAccessPeer) { + HIP_CHECK(hipSetDevice(1)); + TestType *C_d{nullptr}; + HipTest::initArrays(nullptr, nullptr, &C_d, + nullptr, nullptr, nullptr, + NUM_ELM*sizeof(TestType)); + HIP_CHECK(hipMemcpyAsync(A_d, A_h, NUM_ELM*sizeof(TestType), + hipMemcpyHostToDevice, stream)); + HIP_CHECK(hipMemcpyAsync(C_d, A_d, NUM_ELM*sizeof(TestType), + hipMemcpyDeviceToDevice, stream)); + HIP_CHECK(hipMemcpyAsync(B_h, C_d, NUM_ELM*sizeof(TestType), + hipMemcpyDeviceToHost, stream)); + HIP_CHECK(hipStreamSynchronize(stream)); + HipTest::checkTest(A_h, B_h, NUM_ELM); + HIP_CHECK(hipFree(C_d)); + + } else { + SUCCEED("P2P capability is not present"); + } + } + } + + HipTest::freeArrays(A_d, B_d, nullptr, A_h, B_h, nullptr, false); + HipTest::freeArrays(nullptr, nullptr, nullptr, A_Ph, + B_Ph, nullptr, true); +} + +// This test launches multiple threads which uses same stream to deploy kernel +// and also launch hipMemcpyAsync() api. This test case is simulate the scenario +// reported in SWDEV-181598 + +TEMPLATE_TEST_CASE("Unit_hipMemcpyAsync_hipMultiMemcpyMultiThread", "", + int, float, double) { + size_t Nbytes = N_ELMTS * sizeof(TestType); + + int Data_mismatch = 0; + hipStream_t mystream; + TestType *A_d{nullptr}, *B_d{nullptr}, *C_d{nullptr}; + TestType *A_h{nullptr}, *B_h{nullptr}, *C_h{nullptr}; + + HipTest::initArrays(&A_d, &B_d, &C_d, &A_h, &B_h, &C_h, N_ELMTS, false); + + HIP_CHECK(hipStreamCreateWithFlags(&mystream, hipStreamNonBlocking)); + HIP_CHECK(hipMemcpyAsync(A_d, A_h, Nbytes, hipMemcpyHostToDevice, mystream)); + + std::thread T[NUM_THREADS]; + for (int i = 0; i < NUM_THREADS; i++) { + T[i] = std::thread(Thread_func, A_d, B_d, C_d, + C_h, Nbytes, mystream); + } + + // Wait until all the threads finish their execution + for (int i = 0; i < NUM_THREADS; i++) { + T[i].join(); + } + + HIP_CHECK(hipStreamSynchronize(mystream)); + HIP_CHECK(hipStreamDestroy(mystream)); + + // Verifying the result of the kernel computation + for (size_t i = 0; i < N_ELMTS; i++) { + if (C_h[i] != A_h[i] * A_h[i]) { + Data_mismatch++; + } + } + REQUIRE(Thread_count.load() == NUM_THREADS); + REQUIRE(Data_mismatch == 0); + HipTest::freeArrays(A_d, B_d, C_d, A_h, B_h, C_h, false); + Thread_count.exchange(0); +} + +TEMPLATE_TEST_CASE("Unit_hipMemcpyAsync_hipMultiMemcpyMultiThreadMultiStream", + "", int, float, double) { + std::thread T[NUM_THREADS]; + for (int i = 0; i < NUM_THREADS; i++) { + T[i] = std::thread(Thread_func_MultiStream); + } + + // Wait until all the threads finish their execution + for (int i = 0; i < NUM_THREADS; i++) { + T[i].join(); + } + + REQUIRE(Thread_count.load() == NUM_THREADS); + Thread_count.exchange(0); +} + +/* +This testcase verifies hipMemcpy API with pinnedMemory and hostRegister +along with kernel launches +*/ + +TEMPLATE_TEST_CASE("Unit_hipMemcpyAsync_PinnedRegMemWithKernelLaunch", + "", int, float, double) { + int numDevices = 0; + HIP_CHECK(hipGetDeviceCount(&numDevices)); + if (numDevices < 2) { + SUCCEED("No of devices are less than 2"); + } else { + // 1 refers to pinned Memory + // 2 refers to register Memory + int MallocPinType = GENERATE(0, 1); + size_t Nbytes = NUM_ELM * sizeof(TestType); + unsigned blocks = HipTest::setNumBlocks(blocksPerCU, + threadsPerBlock, NUM_ELM); + + TestType *A_d{nullptr}, *B_d{nullptr}, *C_d{nullptr}; + TestType *X_d{nullptr}, *Y_d{nullptr}, *Z_d{nullptr}; + TestType *A_h{nullptr}, *B_h{nullptr}, *C_h{nullptr}; + if (MallocPinType) { + HipTest::initArrays(&A_d, &B_d, &C_d, &A_h, &B_h, &C_h, NUM_ELM, true); + } else { + A_h = reinterpret_cast(malloc(Nbytes)); + HIP_CHECK(hipHostRegister(A_h, Nbytes, hipHostRegisterDefault)); + B_h = reinterpret_cast(malloc(Nbytes)); + HIP_CHECK(hipHostRegister(B_h, Nbytes, hipHostRegisterDefault)); + C_h = reinterpret_cast(malloc(Nbytes)); + HIP_CHECK(hipHostRegister(C_h, Nbytes, hipHostRegisterDefault)); + HipTest::initArrays(&A_d, &B_d, &C_d, nullptr, nullptr, + nullptr, NUM_ELM, false); + HipTest::setDefaultData(NUM_ELM, A_h, B_h, C_h); + } + HIP_CHECK(hipMemcpy(A_d, A_h, Nbytes, hipMemcpyHostToDevice)); + HIP_CHECK(hipMemcpy(B_d, B_h, Nbytes, hipMemcpyHostToDevice)); + + hipLaunchKernelGGL(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock), + 0, 0, static_cast(A_d), + static_cast(B_d), C_d, NUM_ELM); + + HIP_CHECK(hipMemcpy(C_h, C_d, Nbytes, hipMemcpyDeviceToHost)); + HipTest::checkVectorADD(A_h, B_h, C_h, NUM_ELM); + + unsigned int seed = time(0); + HIP_CHECK(hipSetDevice(rand_r(&seed) % (numDevices-1)+1)); + + int device; + HIP_CHECK(hipGetDevice(&device)); + INFO("hipMemcpy is set to happen between device 0 and device " + << device); + HipTest::initArrays(&X_d, &Y_d, &Z_d, nullptr, + nullptr, nullptr, NUM_ELM, false); + + hipStream_t gpu1Stream; + HIP_CHECK(hipStreamCreate(&gpu1Stream)); + + for (int j = 0; j < NUM_ELM; j++) { + A_h[j] = 0; + B_h[j] = 0; + C_h[j] = 0; + } + + hipMemcpy(A_h, A_d, Nbytes, hipMemcpyDeviceToHost); + hipMemcpyAsync(X_d, A_h, Nbytes, hipMemcpyHostToDevice, gpu1Stream); + hipMemcpy(B_h, B_d, Nbytes, hipMemcpyDeviceToHost); + hipMemcpyAsync(Y_d, B_h, Nbytes, hipMemcpyHostToDevice, gpu1Stream); + + hipLaunchKernelGGL(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock), + 0, 0, static_cast(X_d), + static_cast(Y_d), Z_d, NUM_ELM); + + HIP_CHECK(hipMemcpyAsync(C_h, Z_d, Nbytes, + hipMemcpyDeviceToHost, gpu1Stream)); + HIP_CHECK(hipStreamSynchronize(gpu1Stream)); + + HipTest::checkVectorADD(A_h, B_h, C_h, NUM_ELM); + + if (MallocPinType) { + HipTest::freeArrays(A_d, B_d, C_d, A_h, B_h, C_h, true); + } else { + HIP_CHECK(hipHostUnregister(A_h)); + free(A_h); + HIP_CHECK(hipHostUnregister(B_h)); + free(B_h); + HIP_CHECK(hipHostUnregister(C_h)); + free(C_h); + HipTest::freeArrays(A_d, B_d, C_d, nullptr, + nullptr, nullptr, false); + } + HipTest::freeArrays(X_d, Y_d, Z_d, nullptr, + nullptr, nullptr, false); + HIP_CHECK(hipStreamDestroy(gpu1Stream)); + } +} + diff --git a/catch/unit/memory/hipMemcpyDtoD.cc b/catch/unit/memory/hipMemcpyDtoD.cc new file mode 100644 index 0000000000..ab6dc99ebe --- /dev/null +++ b/catch/unit/memory/hipMemcpyDtoD.cc @@ -0,0 +1,95 @@ +/* +Copyright (c) 2021 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 WARRANNTY OF ANY KIND, EXPRESS OR +IMPLIED, INNCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANNY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +/* +This testcase verifies the hipMemcpyDtoD basic scenario +1. H2D-KernelLaunch-D2H scenario +*/ + +#include +#include +#include + +static constexpr auto NUM_ELM{1024}; + +/* +This testcase verifies hipMemcpyDtoD API +1.Initializes device variables +2.Launches kernel and performs the sum of device variables +3.Copies the result to host variable and validates the result. +4.Sets the peer device +5.D2D copy from GPU-0 to GPU-1 +6.Kernel Launch +7.DtoH copy and validating the result +*/ +TEMPLATE_TEST_CASE("Unit_hipMemcpyDtoD_Basic", "", + int, float, double) { + size_t Nbytes = NUM_ELM * sizeof(TestType); + int numDevices = 0; + TestType *A_d{nullptr}, *B_d{nullptr}, *C_d{nullptr}, + *X_d{nullptr}, *Y_d{nullptr}, *Z_d{nullptr}; + TestType *A_h{nullptr}, *B_h{nullptr}, *C_h{nullptr}; + + HIP_CHECK(hipGetDeviceCount(&numDevices)); + if (numDevices > 1) { + int canAccessPeer = 0; + HIP_CHECK(hipDeviceCanAccessPeer(&canAccessPeer, 0, 1)); + if (canAccessPeer) { + HIP_CHECK(hipSetDevice(0)); + HipTest::initArrays(&A_d, &B_d, &C_d, &A_h, &B_h, &C_h, + NUM_ELM, false); + HIP_CHECK(hipSetDevice(1)); + HIP_CHECK(hipMalloc(&X_d, Nbytes)); + HIP_CHECK(hipMalloc(&Y_d, Nbytes)); + HIP_CHECK(hipMalloc(&Z_d, Nbytes)); + + HIP_CHECK(hipSetDevice(0)); + HIP_CHECK(hipMemcpy(A_d, A_h, Nbytes, hipMemcpyHostToDevice)); + HIP_CHECK(hipMemcpy(B_d, B_h, Nbytes, hipMemcpyHostToDevice)); + hipLaunchKernelGGL(HipTest::vectorADD, dim3(1), + dim3(1), 0, 0, + static_cast(A_d), + static_cast(B_d), C_d, NUM_ELM); + HIP_CHECK(hipMemcpy(C_h, C_d, Nbytes, hipMemcpyDeviceToHost)); + HIP_CHECK(hipDeviceSynchronize()); + HipTest::checkVectorADD(A_h, B_h, C_h, NUM_ELM); + + HIP_CHECK(hipSetDevice(1)); + HIP_CHECK(hipMemcpyDtoD((hipDeviceptr_t)X_d, (hipDeviceptr_t)A_d, + Nbytes)); + HIP_CHECK(hipMemcpyDtoD((hipDeviceptr_t)Y_d, (hipDeviceptr_t)B_d, + Nbytes)); + + hipLaunchKernelGGL(HipTest::vectorADD, dim3(1), + dim3(1), 0, 0, + static_cast(X_d), + static_cast(Y_d), Z_d, NUM_ELM); + HIP_CHECK(hipMemcpyDtoH(C_h, (hipDeviceptr_t)Z_d, Nbytes)); + HIP_CHECK(hipDeviceSynchronize()); + HipTest::checkVectorADD(A_h, B_h, C_h, NUM_ELM); + + HipTest::freeArrays(A_d, B_d, C_d, A_h, B_h, C_h, false); + HIP_CHECK(hipFree(X_d)); + HIP_CHECK(hipFree(Y_d)); + HIP_CHECK(hipFree(Z_d)); + } else { + SUCCEED("Machine does not seem to have P2P Capabilities"); + } + } +} diff --git a/catch/unit/memory/hipMemcpyDtoDAsync.cc b/catch/unit/memory/hipMemcpyDtoDAsync.cc new file mode 100644 index 0000000000..af5be6d284 --- /dev/null +++ b/catch/unit/memory/hipMemcpyDtoDAsync.cc @@ -0,0 +1,100 @@ +/* +Copyright (c) 2021 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 WARRANNTY OF ANY KIND, EXPRESS OR +IMPLIED, INNCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANNY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +/* +This testcase verifies the Basic scenario +1. H2D-KernelLaunch-D2H then D2D-KernelLaunch-D2H in peer GPU +*/ + +#include +#include +#include + +static constexpr auto NUM_ELM{1024}; + +/* +This testcase verifies hipMemcpyDtoDAsync API +1.Initializes device variables +2.Launches kernel and performs the sum of device variables +3.Copies the result to host variable and validates the result. +4.Sets the peer device +5.D2D copy from GPU-0 to GPU-1 +6.Kernel Launch +7.DtoH copy and validating the result +*/ + +TEMPLATE_TEST_CASE("Unit_hipMemcpyDtoDAsync_Basic", "", + int, float, double) { + size_t Nbytes = NUM_ELM * sizeof(TestType); + int numDevices = 0; + TestType *A_d{nullptr}, *B_d{nullptr}, *C_d{nullptr}, + *X_d{nullptr}, *Y_d{nullptr}, *Z_d{nullptr}; + TestType *A_h{nullptr}, *B_h{nullptr}, *C_h{nullptr}; + hipStream_t stream; + + HIP_CHECK(hipGetDeviceCount(&numDevices)); + if (numDevices > 1) { + int canAccessPeer = 0; + HIP_CHECK(hipDeviceCanAccessPeer(&canAccessPeer, 0, 1)); + if (canAccessPeer) { + HIP_CHECK(hipSetDevice(0)); + HipTest::initArrays(&A_d, &B_d, &C_d, &A_h, &B_h, &C_h, + NUM_ELM, false); + HIP_CHECK(hipSetDevice(1)); + HIP_CHECK(hipMalloc(&X_d, Nbytes)); + HIP_CHECK(hipMalloc(&Y_d, Nbytes)); + HIP_CHECK(hipMalloc(&Z_d, Nbytes)); + + HIP_CHECK(hipSetDevice(0)); + HIP_CHECK(hipMemcpy(A_d, A_h, Nbytes, hipMemcpyHostToDevice)); + HIP_CHECK(hipMemcpy(B_d, B_h, Nbytes, hipMemcpyHostToDevice)); + hipLaunchKernelGGL(HipTest::vectorADD, dim3(1), + dim3(1), 0, 0, + static_cast(A_d), + static_cast(B_d), C_d, NUM_ELM); + HIP_CHECK(hipMemcpy(C_h, C_d, Nbytes, hipMemcpyDeviceToHost)); + HIP_CHECK(hipDeviceSynchronize()); + HipTest::checkVectorADD(A_h, B_h, C_h, NUM_ELM); + + HIP_CHECK(hipSetDevice(1)); + HIP_CHECK(hipStreamCreate(&stream)); + HIP_CHECK(hipMemcpyDtoDAsync((hipDeviceptr_t)X_d, (hipDeviceptr_t)A_d, + Nbytes, stream)); + HIP_CHECK(hipMemcpyDtoDAsync((hipDeviceptr_t)Y_d, (hipDeviceptr_t)B_d, + Nbytes, stream)); + HIP_CHECK(hipStreamSynchronize(stream)); + + hipLaunchKernelGGL(HipTest::vectorADD, dim3(1), + dim3(1), 0, 0, + static_cast(X_d), + static_cast(Y_d), Z_d, NUM_ELM); + HIP_CHECK(hipMemcpyDtoHAsync(C_h, (hipDeviceptr_t)Z_d, Nbytes, stream)); + HIP_CHECK(hipStreamSynchronize(stream)); + HIP_CHECK(hipDeviceSynchronize()); + HipTest::checkVectorADD(A_h, B_h, C_h, NUM_ELM); + + HipTest::freeArrays(A_d, B_d, C_d, A_h, B_h, C_h, false); + HIP_CHECK(hipFree(X_d)); + HIP_CHECK(hipFree(Y_d)); + HIP_CHECK(hipFree(Z_d)); + } else { + SUCCEED("Machine does not seem to have P2P Capabilities"); + } + } +} diff --git a/catch/unit/memory/hipMemcpyWithStream.cc b/catch/unit/memory/hipMemcpyWithStream.cc new file mode 100644 index 0000000000..275d98ff19 --- /dev/null +++ b/catch/unit/memory/hipMemcpyWithStream.cc @@ -0,0 +1,628 @@ +/* +Copyright (c) 2021-22-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); + */ +/* +This testfile verifies the following scenarios +1. hipMemcpyWithStream with one stream +2. hipMemcpyWithStream with two streams +3. Multi GPU and single stream +4. hipMemcpyWithStream API with testkind DtoH +5. hipMemcpyWithStream API with testkind DtoD +6. hipMemcpyWithStream API with testkind HtoH +7. hipMemcpyWithStream API with testkind TestkindDefault +8. hipMemcpyWithStream API with testkind TestkindDefaultForDtoD +9. hipMemcpyWithStream API DtoD on same device +10.Multi threaded scenario +*/ + + +#include +#include +#include + +#include +#include +#include + +#define LEN 64 +#define SIZE LEN << 2 +#define THREADS 2 +#define MAX_THREADS 16 + +static constexpr size_t N{4 * 1024 * 1024}; +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 +}; + +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 TestwithOnestream(void) { + size_t Nbytes = N * sizeof(int); + 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; + HIP_CHECK(hipStreamCreate(&stream)); + + HIP_CHECK(hipMemcpyWithStream(A_d, A_h, Nbytes, + hipMemcpyHostToDevice, stream)); + HIP_CHECK(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(hipStreamSynchronize(stream)); + HIP_CHECK(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); + HIP_CHECK(hipStreamDestroy(stream)); +} + +void TestwithTwoStream(void) { + size_t Nbytes = N * sizeof(int); + const int NUM_STREAMS = 2; + int *A_d[NUM_STREAMS], *B_d[NUM_STREAMS], *C_d[NUM_STREAMS]; + int *A_h[NUM_STREAMS], *B_h[NUM_STREAMS], *C_h[NUM_STREAMS]; + + unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N); + + for (int i=0; i < NUM_STREAMS; ++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[NUM_STREAMS]; + for (int i=0; i < NUM_STREAMS; ++i) { + HIP_CHECK(hipStreamCreate(&stream[i])); + } + + for (int i=0; i < NUM_STREAMS; ++i) { + HIP_CHECK(hipMemcpyWithStream(A_d[i], A_h[i], Nbytes, + hipMemcpyHostToDevice, stream[i])); + HIP_CHECK(hipMemcpyWithStream(B_d[i], B_h[i], Nbytes, + hipMemcpyHostToDevice, stream[i])); + } + + for (int i=0; i < NUM_STREAMS; ++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 < NUM_STREAMS; ++i) { + HIP_CHECK(hipStreamSynchronize(stream[i])); + HIP_CHECK(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 < NUM_STREAMS; ++i) { + HipTest::freeArrays(A_d[i], B_d[i], C_d[i], A_h[i], B_h[i], C_h[i], false); + HIP_CHECK(hipStreamDestroy(stream[i])); + } +} + +void TestDtoDonSameDevice(void) { + size_t Nbytes = N * sizeof(int); + const int NUM_STREAMS = 2; + int *A_d[NUM_STREAMS], *B_d[NUM_STREAMS], *C_d[NUM_STREAMS]; + int *A_h[NUM_STREAMS], *B_h[NUM_STREAMS], *C_h[NUM_STREAMS]; + + 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[NUM_STREAMS]; + for (int i=0; i < NUM_STREAMS; ++i) { + HIP_CHECK(hipSetDevice(0)); + HIP_CHECK(hipStreamCreate(&stream[i])); + } + + HIP_CHECK(hipSetDevice(0)); + HIP_CHECK(hipMalloc(&A_d[1], Nbytes)); + HIP_CHECK(hipMalloc(&B_d[1], Nbytes)); + HIP_CHECK(hipMalloc(&C_d[1], Nbytes)); + C_h[1] = reinterpret_cast(malloc(Nbytes)); + HIP_ASSERT(C_h[1] != NULL); + + HIP_CHECK(hipMemcpyWithStream(A_d[0], A_h[0], Nbytes, + hipMemcpyHostToDevice, stream[0])); + HIP_CHECK(hipMemcpyWithStream(B_d[0], B_h[0], Nbytes, + hipMemcpyHostToDevice, stream[0])); + + HIP_CHECK(hipMemcpyWithStream(A_d[1], A_d[0], Nbytes, + hipMemcpyDeviceToDevice, stream[1])); + HIP_CHECK(hipMemcpyWithStream(B_d[1], B_d[0], Nbytes, + hipMemcpyDeviceToDevice, stream[1])); + + + for (int i=0; i < NUM_STREAMS; ++i) { + HIP_CHECK(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 < NUM_STREAMS; ++i) { + HIP_CHECK(hipSetDevice(0)); + HIP_CHECK(hipStreamSynchronize(stream[i])); + HIP_CHECK(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]) { + HIP_CHECK(hipFree(A_d[1])); + } + if (B_d[1]) { + HIP_CHECK(hipFree(B_d[1])); + } + if (C_d[1]) { + HIP_CHECK(hipFree(C_d[1])); + } + if (C_h[1]) { + free(C_h[1]); + } + + + for (int i=0; i < NUM_STREAMS; ++i) { + HIP_CHECK(hipStreamDestroy(stream[i])); + } +} + +void TestOnMultiGPUwithOneStream(void) { + size_t Nbytes = N * sizeof(int); + int NumDevices = 0; + + HIP_CHECK(hipGetDeviceCount(&NumDevices)); + // If you have single GPU machine the return + if (NumDevices <= 1) { + SUCCEED("NumDevices <2"); + } else { + unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N); + int *A_d[MaxGPUDevices], *B_d[MaxGPUDevices], *C_d[MaxGPUDevices]; + int *A_h[MaxGPUDevices], *B_h[MaxGPUDevices], *C_h[MaxGPUDevices]; + + hipStream_t stream[MaxGPUDevices]; + for (int i=0; i < NumDevices; ++i) { + HIP_CHECK(hipSetDevice(i)); + HIP_CHECK(hipStreamCreate(&stream[i])); + } + + for (int i=0; i < NumDevices; ++i) { + HIP_CHECK(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) { + HIP_CHECK(hipSetDevice(i)); + HIP_CHECK(hipMemcpyWithStream(A_d[i], A_h[i], Nbytes, + hipMemcpyHostToDevice, stream[i])); + HIP_CHECK(hipMemcpyWithStream(B_d[i], B_h[i], Nbytes, + hipMemcpyHostToDevice, stream[i])); + } + + for (int i=0; i < NumDevices; ++i) { + HIP_CHECK(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) { + HIP_CHECK(hipSetDevice(i)); + HIP_CHECK(hipStreamSynchronize(stream[i])); + HIP_CHECK(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) { + HIP_CHECK(hipSetDevice(i)); + HipTest::freeArrays(A_d[i], B_d[i], C_d[i], + A_h[i], B_h[i], C_h[i], false); + HIP_CHECK(hipStreamDestroy(stream[i])); + } + } +} + +void TestkindDtoH(void) { + size_t Nbytes = N * sizeof(int); + 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; + HIP_CHECK(hipStreamCreate(&stream)); + + HIP_CHECK(hipMemcpyWithStream(A_d, A_h, Nbytes, + hipMemcpyHostToDevice, stream)); + HIP_CHECK(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(hipStreamSynchronize(stream)); + HIP_CHECK(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); + HIP_CHECK(hipStreamDestroy(stream)); +} + +void TestkindDtoD(void) { + size_t Nbytes = N * sizeof(int); + int NumDevices = 0; + + unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N); + HIP_CHECK(hipGetDeviceCount(&NumDevices)); + // If you have single GPU machine the return + if (NumDevices <= 1) { + SUCCEED("NumDevices are less than 2"); + } else { + int *A_d[MaxGPUDevices], *B_d[MaxGPUDevices], *C_d[MaxGPUDevices]; + int *A_h[MaxGPUDevices], *B_h[MaxGPUDevices], *C_h[MaxGPUDevices]; + + hipStream_t stream[MaxGPUDevices]; + for (int i=0; i < NumDevices; ++i) { + HIP_CHECK(hipSetDevice(i)); + HIP_CHECK(hipStreamCreate(&stream[i])); + } + + // Initialize and create the host and device elements for first device + HIP_CHECK(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) { + HIP_CHECK(hipSetDevice(i)) + HIP_CHECK(hipMalloc(&A_d[i], Nbytes)); + HIP_CHECK(hipMalloc(&B_d[i], Nbytes)); + HIP_CHECK(hipMalloc(&C_d[i], Nbytes)); + C_h[i] = reinterpret_cast(malloc(Nbytes)); + HIP_ASSERT(C_h[i] != NULL); + } + + HIP_CHECK(hipSetDevice(0)); + HIP_CHECK(hipMemcpyWithStream(A_d[0], A_h[0], Nbytes, + hipMemcpyHostToDevice, stream[0])); + HIP_CHECK(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) { + HIP_CHECK(hipSetDevice(i)); + HIP_CHECK(hipMemcpyWithStream(A_d[i], A_d[0], Nbytes, + hipMemcpyDeviceToDevice, stream[i])); + HIP_CHECK(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) { + HIP_CHECK(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) { + HIP_CHECK(hipSetDevice(i)); + HIP_CHECK(hipStreamSynchronize(stream[i])); + HIP_CHECK(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); + HIP_CHECK(hipStreamDestroy(stream[0])); + + for (int i=1; i < NumDevices; ++i) { + if (A_d[i]) { + HIP_CHECK(hipFree(A_d[i])); + } + if (B_d[i]) { + HIP_CHECK(hipFree(B_d[i])); + } + if (C_d[i]) { + HIP_CHECK(hipFree(C_d[i])); + } + if (C_h[i]) { + free(C_h[i]); + } + HIP_CHECK(hipStreamDestroy(stream[i])); + } + } +} + +void TestkindDefault(void) { + size_t Nbytes = N * sizeof(int); + 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; + HIP_CHECK(hipStreamCreate(&stream)); + + HIP_CHECK(hipMemcpyWithStream(A_d, A_h, Nbytes, hipMemcpyDefault, stream)); + HIP_CHECK(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(hipStreamSynchronize(stream)); + HIP_CHECK(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); + HIP_CHECK(hipStreamDestroy(stream)); +} + +void TestkindDefaultForDtoD(void) { + size_t Nbytes = N * sizeof(int); + int NumDevices = 0; + + unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N); + HIP_CHECK(hipGetDeviceCount(&NumDevices)); + // Test case will not run on single GPU setup. + if (NumDevices <= 1) { + SUCCEED("No of Devices < 2"); + } else { + int *A_d[MaxGPUDevices], *B_d[MaxGPUDevices], *C_d[MaxGPUDevices]; + int *A_h[MaxGPUDevices], *B_h[MaxGPUDevices], *C_h[MaxGPUDevices]; + + // Initialize and create the host and device elements for first device + HIP_CHECK(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) { + HIP_CHECK(hipSetDevice(i)); + HIP_CHECK(hipMalloc(&A_d[i], Nbytes)); + HIP_CHECK(hipMalloc(&B_d[i], Nbytes)); + HIP_CHECK(hipMalloc(&C_d[i], Nbytes)); + C_h[i] = reinterpret_cast(malloc(Nbytes)); + HIP_ASSERT(C_h[i] != NULL); + } + + hipStream_t stream[MaxGPUDevices]; + for (int i=0; i < NumDevices; ++i) { + HIP_CHECK(hipSetDevice(i)); + HIP_CHECK(hipStreamCreate(&stream[i])); + } + + HIP_CHECK(hipMemcpyWithStream(A_d[0], A_h[0], Nbytes, + hipMemcpyHostToDevice, stream[0])); + HIP_CHECK(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) { + HIP_CHECK(hipMemcpyWithStream(A_d[i], A_d[0], Nbytes, + hipMemcpyDefault, stream[i])); + HIP_CHECK(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); + } + + for (int i=0; i < NumDevices; ++i) { + HIP_CHECK(hipSetDevice(i)); // hipMemcpy will be on this device + HIP_CHECK(hipStreamSynchronize(stream[i])); + HIP_CHECK(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); + HIP_CHECK(hipStreamDestroy(stream[0])); + + for (int i=1; i < NumDevices; ++i) { + if (A_d[i]) { + HIP_CHECK(hipFree(A_d[i])); + } + if (B_d[i]) { + HIP_CHECK(hipFree(B_d[i])); + } + if (C_d[i]) { + HIP_CHECK(hipFree(C_d[i])); + } + if (C_h[i]) { + free(C_h[i]); + } + HIP_CHECK(hipStreamDestroy(stream[i])); + } + } +} + +void TestkindHtoH(void) { + size_t Nbytes = N * sizeof(int); + int *A_h, *B_h; + + + // Allocate memory to A_h and B_h + A_h = static_cast(malloc(Nbytes)); + HIP_ASSERT(A_h != NULL); + B_h = static_cast(malloc(Nbytes)); + HIP_ASSERT(B_h != NULL); + + for (size_t i = 0; i < N; ++i) { + if (A_h) { + (A_h)[i] = 3.146f + i; // Pi + } + } + + hipStream_t stream; + HIP_CHECK(hipStreamCreate(&stream)); + + HIP_CHECK(hipMemcpyWithStream(B_h, A_h, Nbytes, hipMemcpyHostToHost, stream)); + HIP_CHECK(hipStreamSynchronize(stream)); + + for (size_t i = 0; i < N; i++) { + HIP_ASSERT(A_h[i] == B_h[i]); + } + + if (A_h) { + free(A_h); + } + if (B_h) { + free(B_h); + } + HIP_CHECK(hipStreamDestroy(stream)); +} + +TEST_CASE("Unit_hipMemcpyWithStream_MultiThread") { + size_t thread_count = 10; + std::vector threads; + int deviceCount = 0; + HIP_CHECK(hipGetDeviceCount(&deviceCount)); + if (deviceCount < 2) { + SUCCEED("deviceCount < 2"); + } else { + for (int op = static_cast(ops::TestwithOnestream); + op < static_cast(ops::END_OF_LIST); ++op) { + for (uint32_t i = 0; i < thread_count; i++) { + threads.emplace_back(std::thread{[&] { + switch ( op ) { + case static_cast(ops::TestwithOnestream): + TestwithOnestream(); + break; + case static_cast(ops::TestwithTwoStream): + TestwithTwoStream(); + break; + case static_cast(ops::TestkindDtoH): + TestkindDtoH(); + break; + case static_cast(ops::TestkindHtoH): + TestkindHtoH(); + break; + case static_cast(ops::TestkindDtoD): + TestkindDtoD(); + break; + case static_cast(ops::TestOnMultiGPUwithOneStream): + TestOnMultiGPUwithOneStream(); + break; + case static_cast(ops::TestkindDefault): + TestkindDefault(); + break; +#ifndef __HIP_PLATFORM_NVCC__ + case static_cast(ops::TestkindDefaultForDtoD): + TestkindDefaultForDtoD(); + break; +#endif + case static_cast(ops::TestDtoDonSameDevice): + TestDtoDonSameDevice(); + break; + default:{} + } + }}); + } + } + } +} + +TEST_CASE("Unit_hipMemcpyWithStream_TestWithOneStream") { + TestwithOnestream(); +} + +TEST_CASE("Unit_hipMemcpyWithStream_TestwithTwoStream") { + TestwithTwoStream(); +} + +TEST_CASE("Unit_hipMemcpyWithStream_TestkindDtoH") { + TestkindDtoH(); +} + +TEST_CASE("Unit_hipMemcpyWithStream_TestkindHtoH") { + TestkindHtoH(); +} + +TEST_CASE("Unit_hipMemcpyWithStream_TestkindDtoD") { + TestkindDtoD(); +} + +TEST_CASE("Unit_hipMemcpyWithStream_TestOnMultiGPUwithOneStream") { + TestOnMultiGPUwithOneStream(); +} + +TEST_CASE("Unit_hipMemcpyWithStream_TestkindDefault") { + TestkindDefault(); +} +#ifndef __HIP_PLATFORM_NVCC__ +TEST_CASE("Unit_hipMemcpyWithStream_TestkindDefaultForDtoD") { + TestkindDefaultForDtoD(); +} +#endif + +TEST_CASE("Unit_hipMemcpyWithStream_TestDtoDonSameDevice") { + TestDtoDonSameDevice(); +} diff --git a/catch/unit/memory/hipMemcpyWithStreamMultiThread.cc b/catch/unit/memory/hipMemcpyWithStreamMultiThread.cc new file mode 100644 index 0000000000..3f0c28239b --- /dev/null +++ b/catch/unit/memory/hipMemcpyWithStreamMultiThread.cc @@ -0,0 +1,679 @@ +/* +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); + */ + + +#include +#include +#include + +#define LEN 64 +#define SIZE LEN << 2 +#define THREADS 2 +#define MAX_THREADS 16 + +static constexpr auto N{1024}; +static constexpr auto Nbytes{N * sizeof(int)}; +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 +}; + +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)); + + unsigned blocks = props.multiProcessorCount * blocksPerCU; + if (blocks * threadsPerBlock > N) { + blocks = (N + threadsPerBlock - 1) / threadsPerBlock; + } + return blocks; + } +} // namespace MemcpyStream + + +class HipMemcpyWithStreamMultiThreadtests { + public: + // Test hipMemcpyWithStream with one streams and launch kernel in + // that stream, verify the data. + void TestwithOnestream(bool &val_res); + // Test hipMemcpyWithStream with two streams and launch kernels in + // two streams, verify the data. + 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); + // Test hipMemcpyWithStream to copy data from + // device to host (hipMemcpyDeviceToHost). + void TestkindDtoH(bool &val_res); + // Test hipMemcpyWithStream with hipMemcpyDeviceToDevice on MultiGPU. + void TestkindDtoD(bool &val_res); + // Test hipMemcpyWithStream with hipMemcpyHostToHost. + void TestkindHtoH(bool &val_res); + // Test hipMemcpyWithStream with hipMemcpyDefault. + void TestkindDefault(bool &val_res); + // Test hipMemcpyWithStream with hipMemcpyDefault for + // device to device transfer case. + void TestkindDefaultForDtoD(bool &val_res); + // Test hipMemcpyWithStream with hipMemcpyDeviceToDevice on same device. + 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); + // DeAllocate Memory + 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); +}; + +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)); + *A_h = reinterpret_cast(malloc(Nbytes)); + *B_h = reinterpret_cast(malloc(Nbytes)); + *C_h = reinterpret_cast(malloc(Nbytes)); + + for (size_t i = 0; i < N; i++) { + if (*A_h) (*A_h)[i] = 3; + if (*B_h) (*B_h)[i] = 4; + if (*C_h) (*C_h)[i] = 5; + } +} + +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)); + free(A_h); + free(B_h); + free(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]) { + TestPassed = false; + break; + } + } + return TestPassed; +} + + +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); + 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)); + 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)); +} + +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); + + 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) { + 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)); + val_res = ValidateResult(A_h[i], B_h[i], C_h[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])); + } +} + +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); + + 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])); + } + + 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)); + 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]); + + 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(bool &val_res) { + size_t Nbytes = N * sizeof(int); + int numDevices = 0; + + unsigned blocks = MemcpyStream::setNumBlocks(blocksPerCU, threadsPerBlock, N); + HIPCHECK(hipGetDeviceCount(&numDevices)); + // If you have single GPU machine the return + if (numDevices <= 1) { + return; + } + int *A_d[MaxGPUDevices], *B_d[MaxGPUDevices], *C_d[MaxGPUDevices]; + 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) { + 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) { + 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)); + val_res = ValidateResult(A_h[i], B_h[i], C_h[i]); + } + + for (int i=0; i < numDevices; ++i) { + HIPCHECK(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])); + } +} + +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); + AllocateMemory(&A_d, &B_d, &C_d, &A_h, &B_h, &C_h); + + 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)); + 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)); +} + + +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)); + // If you have single GPU machine the return + if (numDevices <= 1) { + return; + } + + int *A_d[MaxGPUDevices], *B_d[MaxGPUDevices], *C_d[MaxGPUDevices]; + 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])); + } + + // 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]); + + 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)); + 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])); + + 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(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); + AllocateMemory(&A_d, &B_d, &C_d, &A_h, &B_h, &C_h); + + 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)); + 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)); +} + +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)); + // Test case will not run on single GPU setup. + if (numDevices <= 1) { + return; + } + + int *A_d[MaxGPUDevices], *B_d[MaxGPUDevices], *C_d[MaxGPUDevices]; + 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]); + + 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); + } + + hipStream_t stream[MaxGPUDevices]; + for (int i=0; i < numDevices; ++i) { + HIPCHECK(hipSetDevice(i)); + HIPCHECK(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])); + + // 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=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); + } + + 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)); + // 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])); + + 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(bool &val_res) { + size_t Nbytes = N * sizeof(int); + int *A_h, *B_h; + + // Allocate memory to A_h and B_h + A_h = static_cast(malloc(Nbytes)); + B_h = static_cast(malloc(Nbytes)); + + 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++) { + if ((A_h[i] != B_h[i])) { + val_res = false; + break; + } + } + + if (A_h) { + free(A_h); + } + if (B_h) { + free(B_h); + } + HIPCHECK(hipStreamDestroy(stream)); +} + +TEST_CASE("Unit_hipMemcpyWithStream_NewMultiThread") { + const auto Threadcount{100}; + bool ret_val[Threadcount]; + std::thread th[Threadcount]; + 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) ) { + case ops::TestwithOnestream: + 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])); + break; + case ops::TestkindDtoH: + 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])); + break; + case ops::TestkindDtoD: + th[i] = std::thread(&HipMemcpyWithStreamMultiThreadtests:: + TestkindDtoD, + &tests, std::ref(ret_val[i])); + break; + case ops::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])); + break; + case ops::TestkindDefaultForDtoD: + 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])); + break; + default: {} + } + } + + for (uint32_t i = 0; i < Threadcount; i++) { + th[i].join(); + } + + for (uint32_t i = 0; i < Threadcount; i++) { + REQUIRE(ret_val[i] == true); + } + } +} diff --git a/catch/unit/memory/hipMemcpy_MultiThread.cc b/catch/unit/memory/hipMemcpy_MultiThread.cc new file mode 100644 index 0000000000..1e8f7e46a8 --- /dev/null +++ b/catch/unit/memory/hipMemcpy_MultiThread.cc @@ -0,0 +1,322 @@ +/* +Copyright (c) 2021 - 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. +*/ + +#include +#include +#include +#include +#include +/* +This testfile verifies the following scenarios of all hipMemcpy API +1. Multi thread +*/ +static constexpr auto NUM_ELM{1024}; +static constexpr auto NUM_THREADS{10}; +static auto Available_Gpus{0}; +static constexpr auto MAX_GPU{256}; + +enum apiToTest {TEST_MEMCPY, TEST_MEMCPYH2D, TEST_MEMCPYD2H, TEST_MEMCPYD2D, + TEST_MEMCPYASYNC, TEST_MEMCPYH2DASYNC, TEST_MEMCPYD2HASYNC, + TEST_MEMCPYD2DASYNC}; + + +template +class memcpyTests { + public: + T *A_h, *B_h; + apiToTest api; + explicit memcpyTests(apiToTest val); + memcpyTests() = delete; + void Memcpy_And_verify(bool *ret_val); + bool CheckTests(T* A_h, T* B_h, int NUM_ELEMENTS); + ~memcpyTests(); +}; + +template +bool memcpyTests::CheckTests(T *A_h, T *B_h, int NUM_ELEMENTS) { + for (auto i =0; i < NUM_ELEMENTS; i++) { + if (A_h[i] != B_h[i]) { + return false; + } + } + return true; +} + +template +memcpyTests::memcpyTests(apiToTest val) { + api = val; + A_h = reinterpret_cast(malloc(NUM_ELM * sizeof(T))); + B_h = reinterpret_cast(malloc(NUM_ELM * sizeof(T))); + if ((A_h == nullptr) || (B_h == nullptr)) { + exit(1); + } + + for (size_t i = 0; i < NUM_ELM; ++i) { + A_h[i] = 123; + B_h[i] = 0; + } +} + + +template +void memcpyTests::Memcpy_And_verify(bool *ret_val) { + HIPCHECK(hipGetDeviceCount(&Available_Gpus)); + T *A_d[MAX_GPU]; + hipStream_t stream[MAX_GPU]; + for (int i = 0; i < Available_Gpus; ++i) { + HIPCHECK(hipSetDevice(i)); + HIPCHECK(hipMalloc(&A_d[i], NUM_ELM * sizeof(T))); + if (api >= TEST_MEMCPYD2D) { + HIPCHECK(hipStreamCreate(&stream[i])); + } + } + HIPCHECK(hipSetDevice(0)); + int canAccessPeer = 0; + switch (api) { + case TEST_MEMCPY: + { + // To test hipMemcpy() + // Copying data from host to individual devices followed by copying + // back to host and verifying the data consistency. + for (int i = 0; i < Available_Gpus; ++i) { + HIPCHECK(hipMemcpy(A_d[i], A_h, NUM_ELM * sizeof(T), + hipMemcpyHostToDevice)); + HIPCHECK(hipMemcpy(B_h, A_d[i], NUM_ELM * sizeof(T), + hipMemcpyDeviceToHost)); + *ret_val = CheckTests(A_h, B_h, NUM_ELM); + } + // Device to Device copying for all combinations + for (int i = 0; i < Available_Gpus; ++i) { + for (int j = i+1; j < Available_Gpus; ++j) { + canAccessPeer = 0; + hipDeviceCanAccessPeer(&canAccessPeer, i, j); + if (canAccessPeer) { + HIPCHECK(hipMemcpy(A_d[j], A_d[i], NUM_ELM * sizeof(T), + hipMemcpyDefault)); + // Copying in reverse dir of above to check if bidirectional + // access is happening without any error + HIPCHECK(hipMemcpy(A_d[i], A_d[j], NUM_ELM * sizeof(T), + hipMemcpyDefault)); + // Copying data to host to verify the content + HIPCHECK(hipMemcpy(B_h, A_d[j], NUM_ELM * sizeof(T), + hipMemcpyDefault)); + *ret_val &= CheckTests(A_h, B_h, NUM_ELM); + } + } + } + break; + } + case TEST_MEMCPYH2D: // To test hipMemcpyHtoD() + { + for (int i = 0; i < Available_Gpus; ++i) { + HIPCHECK(hipMemcpyHtoD(hipDeviceptr_t(A_d[i]), + A_h, NUM_ELM * sizeof(T))); + // Copying data from device to host to check data consistency + HIPCHECK(hipMemcpy(B_h, A_d[i], NUM_ELM * sizeof(T), + hipMemcpyDeviceToHost)); + *ret_val &= CheckTests(A_h, B_h, NUM_ELM); + } + break; + } + case TEST_MEMCPYD2H: // To test hipMemcpyDtoH()--done + { + for (int i = 0; i < Available_Gpus; ++i) { + HIPCHECK(hipMemcpy(A_d[i], A_h, NUM_ELM * sizeof(T), + hipMemcpyHostToDevice)); + HIPCHECK(hipMemcpyDtoH(B_h, hipDeviceptr_t(A_d[i]), + NUM_ELM * sizeof(T))); + *ret_val &= CheckTests(A_h, B_h, NUM_ELM); + } + break; + } + case TEST_MEMCPYD2D: // To test hipMemcpyDtoD() + { + if (Available_Gpus > 1) { + // First copy data from H to D and then + // from D to D followed by D to H + // HIPCHECK(hipMemcpyHtoD(A_d[0], A_h, + // NUM_ELM * sizeof(T))); + int canAccessPeer = 0; + for (int i = 0; i < Available_Gpus; ++i) { + for (int j = i+1; j < Available_Gpus; ++j) { + hipDeviceCanAccessPeer(&canAccessPeer, i, j); + if (canAccessPeer) { + HIPCHECK(hipMemcpyHtoD(hipDeviceptr_t(A_d[i]), + A_h, NUM_ELM * sizeof(T))); + HIPCHECK(hipMemcpyDtoD(hipDeviceptr_t(A_d[j]), + hipDeviceptr_t(A_d[i]), NUM_ELM * sizeof(T))); + // Copying in direction reverse of above to check if + // bidirectional + // access is happening without any error + HIPCHECK(hipMemcpyDtoD(hipDeviceptr_t(A_d[i]), + hipDeviceptr_t(A_d[j]), NUM_ELM * sizeof(T))); + HIPCHECK(hipMemcpy(B_h, A_d[i], NUM_ELM * sizeof(T), + hipMemcpyDeviceToHost)); + *ret_val &= CheckTests(A_h, B_h, NUM_ELM); + } + } + } + } else { + // As DtoD is not possible transfer data from HtH(A_h to B_h) + // so as to get through verification step + HIPCHECK(hipMemcpy(B_h, A_h, NUM_ELM * sizeof(T), + hipMemcpyHostToHost)); + *ret_val &= CheckTests(A_h, B_h, NUM_ELM); + } + break; + } + case TEST_MEMCPYASYNC: + { + // To test hipMemcpyAsync() + // Copying data from host to individual devices followed by copying + // back to host and verifying the data consistency. + for (int i = 0; i < Available_Gpus; ++i) { + HIPCHECK(hipMemcpyAsync(A_d[i], A_h, NUM_ELM * sizeof(T), + hipMemcpyHostToDevice, stream[i])); + HIPCHECK(hipMemcpyAsync(B_h, A_d[i], NUM_ELM * sizeof(T), + hipMemcpyDeviceToHost, stream[i])); + HIPCHECK(hipStreamSynchronize(stream[i])); + *ret_val &= CheckTests(A_h, B_h, NUM_ELM); + } + // Device to Device copying for all combinations + for (int i = 0; i < Available_Gpus; ++i) { + for (int j = i+1; j < Available_Gpus; ++j) { + canAccessPeer = 0; + hipDeviceCanAccessPeer(&canAccessPeer, i, j); + if (canAccessPeer) { + HIPCHECK(hipMemcpyAsync(A_d[j], A_d[i], + NUM_ELM * sizeof(T), + hipMemcpyDefault, stream[i])); + // Copying in direction reverse of above to + // check if bidirectional + // access is happening without any error + HIPCHECK(hipMemcpyAsync(A_d[i], A_d[j], + NUM_ELM * sizeof(T), + hipMemcpyDefault, stream[i])); + HIPCHECK(hipStreamSynchronize(stream[i])); + HIPCHECK(hipMemcpy(B_h, A_d[j], NUM_ELM * sizeof(T), + hipMemcpyDefault)); + *ret_val &= CheckTests(A_h, B_h, NUM_ELM); + } + } + } + break; + } + case TEST_MEMCPYH2DASYNC: // To test hipMemcpyHtoDAsync() + { + for (int i = 0; i < Available_Gpus; ++i) { + HIPCHECK(hipMemcpyHtoDAsync(hipDeviceptr_t(A_d[i]), A_h, + NUM_ELM * sizeof(T), stream[i])); + HIPCHECK(hipStreamSynchronize(stream[i])); + // Copying data from device to host to check data consistency + HIPCHECK(hipMemcpy(B_h, A_d[i], NUM_ELM * sizeof(T), + hipMemcpyDeviceToHost)); + *ret_val &= CheckTests(A_h, B_h, NUM_ELM); + } + break; + } + case TEST_MEMCPYD2HASYNC: // To test hipMemcpyDtoHAsync() + { + for (int i = 0; i < Available_Gpus; ++i) { + HIPCHECK(hipMemcpy(A_d[i], A_h, NUM_ELM * sizeof(T), + hipMemcpyHostToDevice)); + HIPCHECK(hipMemcpyDtoHAsync(B_h, hipDeviceptr_t(A_d[i]), + NUM_ELM * sizeof(T), stream[i])); + HIPCHECK(hipStreamSynchronize(stream[i])); + *ret_val &= CheckTests(A_h, B_h, NUM_ELM); + } + break; + } + case TEST_MEMCPYD2DASYNC: // To test hipMemcpyDtoDAsync() + { + if (Available_Gpus > 1) { + // First copy data from H to D and then from D to D followed by D2H + HIPCHECK(hipMemcpyHtoD(hipDeviceptr_t(A_d[0]), + A_h, NUM_ELM * sizeof(T))); + for (int i = 0; i < Available_Gpus; ++i) { + for (int j = i+1; j < Available_Gpus; ++j) { + canAccessPeer = 0; + hipDeviceCanAccessPeer(&canAccessPeer, i, j); + if (canAccessPeer) { + HIPCHECK(hipSetDevice(j)); + HIPCHECK(hipMemcpyDtoDAsync(hipDeviceptr_t(A_d[j]), + hipDeviceptr_t(A_d[i]), NUM_ELM * sizeof(T), + stream[i])); + // Copying in direction reverse of above to check if + // bidirectional + // access is happening without any error + HIPCHECK(hipMemcpyDtoDAsync(hipDeviceptr_t(A_d[i]), + hipDeviceptr_t(A_d[j]), NUM_ELM * sizeof(T), + stream[i])); + HIPCHECK(hipStreamSynchronize(stream[i])); + HIPCHECK(hipMemcpy(B_h, A_d[i], NUM_ELM * sizeof(T), + hipMemcpyDeviceToHost)); + *ret_val &= CheckTests(A_h, B_h, NUM_ELM); + } + } + } + } else { + // As DtoD is not possible we will transfer data + // from HtH(A_h to B_h) + // so as to get through verification step + HIPCHECK(hipMemcpy(B_h, A_h, NUM_ELM * sizeof(T), + hipMemcpyHostToHost)); + *ret_val &= CheckTests(A_h, B_h, NUM_ELM); + } + break; + } + } + for (int i = 0; i < Available_Gpus; ++i) { + HIPCHECK(hipSetDevice(i)); + HIPCHECK(hipFree((A_d[i]))); + if (api >= TEST_MEMCPYD2D) { + HIPCHECK(hipStreamDestroy(stream[i])); + } + } +} +template +memcpyTests::~memcpyTests() { + free(A_h); + free(B_h); +} + +void Thread_func(bool &ret_val) { + for (apiToTest api = TEST_MEMCPY; api <= TEST_MEMCPYD2DASYNC; + api = apiToTest(api + 1)) { + memcpyTests obj(api); + obj.Memcpy_And_verify(&ret_val); + } +} + + + +TEST_CASE("Unit_hipMemcpy_MultiThread-AllAPIs") { + std::thread Thrd[NUM_THREADS]; + bool ret_val[NUM_THREADS]; + for (int i = 0; i < NUM_THREADS; i++) + Thrd[i] = std::thread(Thread_func, std::ref(ret_val[i])); + + // Thread join is being called separately so as to allow the + // threads run parallely + for (int i = 0; i < NUM_THREADS; i++) + Thrd[i].join(); + + for (int i = 0; i < NUM_THREADS; i++) + REQUIRE(ret_val[i] == true); +}