diff --git a/projects/hip/include/hip/hcc_detail/hip_runtime.h b/projects/hip/include/hip/hcc_detail/hip_runtime.h index b277ba038c..7c5f7c7287 100644 --- a/projects/hip/include/hip/hcc_detail/hip_runtime.h +++ b/projects/hip/include/hip/hcc_detail/hip_runtime.h @@ -478,24 +478,36 @@ struct __hip_builtin_threadIdx_t { __HIP_DEVICE_BUILTIN(x,__hip_get_thread_idx_x()); __HIP_DEVICE_BUILTIN(y,__hip_get_thread_idx_y()); __HIP_DEVICE_BUILTIN(z,__hip_get_thread_idx_z()); +#ifdef __cplusplus + __device__ operator dim3() const { return dim3(x, y, z); } +#endif }; struct __hip_builtin_blockIdx_t { __HIP_DEVICE_BUILTIN(x,__hip_get_block_idx_x()); __HIP_DEVICE_BUILTIN(y,__hip_get_block_idx_y()); __HIP_DEVICE_BUILTIN(z,__hip_get_block_idx_z()); +#ifdef __cplusplus + __device__ operator dim3() const { return dim3(x, y, z); } +#endif }; struct __hip_builtin_blockDim_t { __HIP_DEVICE_BUILTIN(x,__hip_get_block_dim_x()); __HIP_DEVICE_BUILTIN(y,__hip_get_block_dim_y()); __HIP_DEVICE_BUILTIN(z,__hip_get_block_dim_z()); +#ifdef __cplusplus + __device__ operator dim3() const { return dim3(x, y, z); } +#endif }; struct __hip_builtin_gridDim_t { __HIP_DEVICE_BUILTIN(x,__hip_get_grid_dim_x()); __HIP_DEVICE_BUILTIN(y,__hip_get_grid_dim_y()); __HIP_DEVICE_BUILTIN(z,__hip_get_grid_dim_z()); +#ifdef __cplusplus + __device__ operator dim3() const { return dim3(x, y, z); } +#endif }; #undef __HIP_DEVICE_BUILTIN diff --git a/projects/hip/lpl_ca/CMakeLists.txt b/projects/hip/lpl_ca/CMakeLists.txt index 90e7143f66..c272273c09 100644 --- a/projects/hip/lpl_ca/CMakeLists.txt +++ b/projects/hip/lpl_ca/CMakeLists.txt @@ -1,4 +1,10 @@ #-------------------------------------LPL--------------------------------------# +# Look for ROCclr which contains elfio +find_package(ROCclr REQUIRED CONFIG + PATHS + /opt/rocm + /opt/rocm/rocclr) + add_executable(lpl lpl.cpp) set_target_properties( lpl PROPERTIES @@ -6,7 +12,12 @@ set_target_properties( CXX_STANDARD_REQUIRED ON CXX_EXTENSIONS OFF RUNTIME_OUTPUT_DIRECTORY ${PROJECT_BINARY_DIR}) -target_include_directories(lpl PUBLIC ${PROJECT_SOURCE_DIR}/src) +target_include_directories(lpl + PUBLIC + ${PROJECT_SOURCE_DIR}/src + PRIVATE + $) + target_compile_options(lpl PUBLIC -Wall) target_link_libraries(lpl PUBLIC pthread) diff --git a/projects/hip/lpl_ca/lpl.hpp b/projects/hip/lpl_ca/lpl.hpp index 84a6930753..00635d3f26 100644 --- a/projects/hip/lpl_ca/lpl.hpp +++ b/projects/hip/lpl_ca/lpl.hpp @@ -4,7 +4,7 @@ #include "clara/clara.hpp" #include "pstreams/pstream.h" -#include "hip/hcc_detail/elfio/elfio.hpp" +#include #include @@ -71,14 +71,14 @@ inline std::string make_hipcc_call(const std::vector& sources, } inline void copy_kernel_section_to_fat_binary(const std::string& tmp, const std::string& output) { - ELFIO::elfio reader; + amd::ELFIO::elfio reader; if (!reader.load(tmp)) { throw std::runtime_error{"The result of the compilation is inaccessible."}; } const auto it = std::find_if(reader.sections.begin(), reader.sections.end(), - [](const ELFIO::section* x) { return x->get_name() == kernel_section(); }); + [](const amd::ELFIO::section* x) { return x->get_name() == kernel_section(); }); std::ofstream out{output}; diff --git a/projects/hip/rocclr/CMakeLists.txt b/projects/hip/rocclr/CMakeLists.txt index e6c4225984..378307aaef 100755 --- a/projects/hip/rocclr/CMakeLists.txt +++ b/projects/hip/rocclr/CMakeLists.txt @@ -96,8 +96,6 @@ find_package(amd_comgr REQUIRED CONFIG message(STATUS "Code Object Manager found at ${amd_comgr_DIR}.") -add_definitions(-DBSD_LIBELF) - add_library(hip64 OBJECT hip_context.cpp hip_code_object.cpp @@ -129,10 +127,8 @@ target_include_directories(hip64 ${PROJECT_SOURCE_DIR}/include ${PROJECT_BINARY_DIR}/include PRIVATE - ${CMAKE_SOURCE_DIR}/elfio ${PROJECT_SOURCE_DIR} ${PROJECT_SOURCE_DIR}/amdocl - ${PROJECT_SOURCE_DIR}/include/hip/hcc_detail/elfio ${ROCR_INCLUDES} $ $) diff --git a/projects/hip/rocclr/hip_code_object.cpp b/projects/hip/rocclr/hip_code_object.cpp index 8b89647e89..26852bee6d 100755 --- a/projects/hip/rocclr/hip_code_object.cpp +++ b/projects/hip/rocclr/hip_code_object.cpp @@ -6,27 +6,12 @@ #include "hip/hip_runtime.h" #include "hip_internal.hpp" #include "platform/program.hpp" +#include namespace hip { uint64_t CodeObject::ElfSize(const void *emi) { - const Elf64_Ehdr *ehdr = (const Elf64_Ehdr*)emi; - const Elf64_Shdr *shdr = (const Elf64_Shdr*)((char*)emi + ehdr->e_shoff); - - uint64_t max_offset = ehdr->e_shoff; - uint64_t total_size = max_offset + ehdr->e_shentsize * ehdr->e_shnum; - - for (uint16_t i=0; i < ehdr->e_shnum; ++i){ - uint64_t cur_offset = static_cast(shdr[i].sh_offset); - if (max_offset < cur_offset) { - max_offset = cur_offset; - total_size = max_offset; - if(SHT_NOBITS != shdr[i].sh_type) { - total_size += static_cast(shdr[i].sh_size); - } - } - } - return total_size; + return amd::Elf::getElfSize(emi); } bool CodeObject::isCompatibleCodeObject(const std::string& codeobj_target_id, diff --git a/projects/hip/rocclr/hip_event.cpp b/projects/hip/rocclr/hip_event.cpp index 52775990be..5fa2e9ea6e 100755 --- a/projects/hip/rocclr/hip_event.cpp +++ b/projects/hip/rocclr/hip_event.cpp @@ -90,12 +90,10 @@ hipError_t Event::elapsedTime(Event& eStop, float& ms) { return hipErrorNotReady; } - // For certain HIP API's that take start and stop event - // and no hipEventRecord needs to be called - if (event_ == eStop.event_ && !recorded_ && !eStop.recorded_) { + if (event_ != eStop.event_ && recorded_ && eStop.recorded_) { ms = static_cast(static_cast(eStop.event_->profilingInfo().end_ - - event_->profilingInfo().start_))/1000000.f; - } else if (event_ == eStop.event_) { + event_->profilingInfo().end_))/1000000.f; + } else if (event_ == eStop.event_ && (recorded_ || eStop.recorded_)) { // Events are the same, which indicates the stream is empty and likely // eventRecord is called on another stream. For such cases insert and measure a // marker. @@ -107,8 +105,10 @@ hipError_t Event::elapsedTime(Event& eStop, float& ms) { command->release(); } else { + // For certain HIP API's that take both start and stop event + // or scenarios where HIP API takes one of the events and the other event is recorded with hipEventRecord ms = static_cast(static_cast(eStop.event_->profilingInfo().end_ - - event_->profilingInfo().end_))/1000000.f; + event_->profilingInfo().start_))/1000000.f; } return hipSuccess; } diff --git a/projects/hip/rocclr/hip_memory.cpp b/projects/hip/rocclr/hip_memory.cpp index a24c8f93df..da8bd5be78 100755 --- a/projects/hip/rocclr/hip_memory.cpp +++ b/projects/hip/rocclr/hip_memory.cpp @@ -150,12 +150,14 @@ hipError_t ihipMemcpy(void* dst, const void* src, size_t sizeBytes, hipMemcpyKin size_t dOffset = 0; amd::Memory *dstMemory = getMemoryObject(dst, dOffset); amd::Device* queueDevice = &queue.device(); - - if (((srcMemory == nullptr) && (dstMemory == nullptr)) || - (kind == hipMemcpyHostToHost)) { - queue.finish(); - memcpy(dst, src, sizeBytes); - return hipSuccess; + if ((srcMemory == nullptr) && (dstMemory == nullptr)) { + if ((kind == hipMemcpyHostToHost) || (kind == hipMemcpyDefault)) { + queue.finish(); + memcpy(dst, src, sizeBytes); + return hipSuccess; + } else { + return hipErrorInvalidValue; + } } else if ((srcMemory == nullptr) && (dstMemory != nullptr)) { amd::HostQueue* pQueue = &queue; if (queueDevice != dstMemory->getContext().devices()[0]) { diff --git a/projects/hip/rocclr/hip_module.cpp b/projects/hip/rocclr/hip_module.cpp index 146c6f3829..4a09cc6ed0 100755 --- a/projects/hip/rocclr/hip_module.cpp +++ b/projects/hip/rocclr/hip_module.cpp @@ -19,7 +19,7 @@ THE SOFTWARE. */ #include -#include +#include #include #include "hip_internal.hpp" @@ -45,23 +45,7 @@ const std::string& FunctionName(const hipFunction_t f) { static uint64_t ElfSize(const void *emi) { - const Elf64_Ehdr *ehdr = (const Elf64_Ehdr*)emi; - const Elf64_Shdr *shdr = (const Elf64_Shdr*)((char*)emi + ehdr->e_shoff); - - uint64_t max_offset = ehdr->e_shoff; - uint64_t total_size = max_offset + ehdr->e_shentsize * ehdr->e_shnum; - - for (uint16_t i=0; i < ehdr->e_shnum; ++i){ - uint64_t cur_offset = static_cast(shdr[i].sh_offset); - if (max_offset < cur_offset) { - max_offset = cur_offset; - total_size = max_offset; - if(SHT_NOBITS != shdr[i].sh_type) { - total_size += static_cast(shdr[i].sh_size); - } - } - } - return total_size; + return amd::Elf::getElfSize(emi); } hipError_t hipModuleUnload(hipModule_t hmod) { @@ -303,7 +287,7 @@ hipError_t ihipModuleLaunchKernel(hipFunction_t f, uint32_t globalWorkSizeX, } } - profileNDRange = (startEvent != nullptr && stopEvent != nullptr); + profileNDRange = (startEvent != nullptr || stopEvent != nullptr); // Flag set to 1 signifies that kernel can be launched in anyorder if (flags & hipExtAnyOrderLaunch) { diff --git a/projects/hip/rocclr/hip_platform.cpp b/projects/hip/rocclr/hip_platform.cpp index 2d74d8f440..dc610affdc 100755 --- a/projects/hip/rocclr/hip_platform.cpp +++ b/projects/hip/rocclr/hip_platform.cpp @@ -26,7 +26,6 @@ #include "platform/runtime.hpp" #include -#include "elfio.hpp" constexpr unsigned __hipFatMAGIC2 = 0x48495046; // "HIPF" diff --git a/projects/hip/rocclr/hip_stream.cpp b/projects/hip/rocclr/hip_stream.cpp index 379954ef5c..7f74136080 100755 --- a/projects/hip/rocclr/hip_stream.cpp +++ b/projects/hip/rocclr/hip_stream.cpp @@ -159,10 +159,14 @@ void iHipWaitActiveStreams(amd::HostQueue* blocking_queue, bool wait_null_stream (stream->Null() == wait_null_stream)) { // Get the last valid command amd::Command* command = active_queue->getLastQueuedCommand(true); - if ((command != nullptr) && - // Check the current active status - (command->status() != CL_COMPLETE)) { - eventWaitList.push_back(command); + if (command != nullptr) { + // Check the current active status + if (command->status() != CL_COMPLETE) { + command->notifyCmdQueue(); + eventWaitList.push_back(command); + } else { + command->release(); + } } } } diff --git a/projects/hip/tests/performance/module/hipPerfModuleLoad.cpp b/projects/hip/tests/performance/module/hipPerfModuleLoad.cpp new file mode 100755 index 0000000000..06f55ed89d --- /dev/null +++ b/projects/hip/tests/performance/module/hipPerfModuleLoad.cpp @@ -0,0 +1,244 @@ +/* +Copyright (c) 2015-2017 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. +*/ + +/* HIT_START + * BUILD_CMD: hipPerfModuleLoad %hc -I%S/../../src %S/%s %S/../../src/test_common.cpp -o %T/%t EXCLUDE_HIP_PLATFORM nvcc + * TEST: %t + * HIT_END + */ + +#include "test_common.h" + +#include +#include +#include +#include +#include + +#ifdef __unix__ + +#include + +//List of Download files +std::unordered_map TL_contents { + {"Kernels.so", true}, + {"TensileLibrary.yaml", true}, + {"TensileLibrary_gfx803.co", true}, + {"TensileLibrary_gfx900.co", true}, + {"TensileLibrary_gfx906.co", true}, + {"TensileLibrary_gfx908.co", true}, + {"kernel_names.txt", true} +}; + +bool GetDirectoryContents(std::unordered_map& dir_contents) { + DIR* dir = nullptr; + struct dirent* ent = nullptr; + + //Open Current Directory + if ((dir = opendir(".")) == 0) { + std::cout<<"Failed to open current working directory, check permissions"<d_name, true}); + } + + closedir(dir); + return true; +} + +bool ContentsAvailable() { + std::unordered_map dir_contents; + + //Get recent directory contents + if(!GetDirectoryContents(dir_contents)) { + std::cout<<"Failed to get directory Contents"< mload_duration = (mload_clock_stop - mload_clock_start); + std::cout<<"Time taken for hipModuleLoad : " < + (mload_duration).count()<<" nanoseconds "< kernel_vec; + while (std::getline(kernel_file, kernel_line)) { + kernel_line.erase(std::remove(kernel_line.begin(), kernel_line.end(), '\r'), + kernel_line.end()); + kernel_vec.push_back(kernel_line); + } + + //Measure the first hipModuleGetFunction + hipFunction_t hfunc = nullptr; + auto mgetf_clock_start = std::chrono::steady_clock::now(); + HIPCHECK(hipModuleGetFunction(&hfunc, Module, kernel_vec[0].c_str())); + auto mgetf_clock_stop = std::chrono::steady_clock::now(); + std::chrono::duration mgetf_duration = (mgetf_clock_stop - mgetf_clock_start); + std::cout<<"Time taken to fetch a function via hipModuleGetFunction : " + < + (mgetf_duration).count()<<" nanoseconds "< + (mgetf_duration).count()<<" nanoseconds "< 0) { + std::cout << "Time taken for Average hipModuleGetFunction : " + << (static_cast(all_duration) / static_cast(kernel_vec.size()))<<" nanoseconds "< -#include -#include -#include "hip/hip_runtime.h" -#include "test_common.h" - -#define NUM_THREADS 10 -#define NUM_ELM 1024*1024 -#define HIPTEST_TRUE 1 - -int memcpyPeersOnly = 1; -int testAllTypes = 0; -int Available_Gpus = 0; -std::atomic failureCount{0}; - -enum apiToTest {TEST_MEMCPY, TEST_MEMCPYH2D, TEST_MEMCPYD2H, TEST_MEMCPYD2D, - TEST_MEMCPYASYNC, TEST_MEMCPYH2DASYNC, TEST_MEMCPYD2HASYNC, - TEST_MEMCPYD2DASYNC, TEST_MAX}; -std::vector apiNameToTest = { "hipMemcpy", "hipMemcpyH2D", - "hipMemcpyD2H", "hipMemcpyD2D", "hipMemcpyAsync", - "hipMemcpyH2DAsync", "hipMemcpyD2HAsync", "hipMemcpyD2DAsync" }; - -// If memcpyPeersOnly is true, then checks if given gpus are peers and returns -// true if they are peers, else false -// If memcpyPeersOnly is false, then returns true always -bool gpusIsPeer(int gpu0, int gpu1) { - bool bRet = true; - if (HIPTEST_TRUE == memcpyPeersOnly) { - int CanAccessPeer1 = 0, CanAccessPeer2 = 0; - HIPCHECK(hipDeviceCanAccessPeer(&CanAccessPeer1, gpu0, gpu1)); - HIPCHECK(hipDeviceCanAccessPeer(&CanAccessPeer2, gpu1, gpu0)); - if ((CanAccessPeer1 * CanAccessPeer2) == 0) { - bRet = false; - } - } - - return bRet; -} - -template -class memcpyTests { - public: - T *A_h, *B_h; - apiToTest api; - size_t NUM_ELMTS = 0; - hipStream_t stream; - memcpyTests(apiToTest val, size_t num_elmts); - bool Memcpy_And_verify(); - ~memcpyTests(); -}; - -class Memcpy_Negative_Tests { - public: - // The following function will test negative scenarios with hipMemcpy() - bool Test_Memcpy(void); - bool Test_MemcpyAsync(void); - bool Test_MemcpyHtoD(void); - bool Test_MemcpyHtoDAsync(void); - bool Test_MemcpyDtoH(void); - bool Test_MemcpyDtoHAsync(void); - bool Test_MemcpyDtoD(void); - bool Test_MemcpyDtoDAsync(void); -}; - - -bool Memcpy_Negative_Tests::Test_Memcpy(void) { - bool IfTestPassed = true; - std::string str_out, str_err = "hipErrorInvalidValue"; - float *A_h = NULL, *B_h = NULL, *A_d = NULL, *A_d1 = NULL; - A_h = reinterpret_cast(malloc(NUM_ELM * sizeof(float))); - B_h = reinterpret_cast(malloc(NUM_ELM * sizeof(float))); - - if ((A_h == NULL) || (B_h == NULL)) { - failed("Malloc call failed!"); - } - - HIPCHECK(hipMalloc(&A_d, (NUM_ELM*sizeof(float)))); - HIPCHECK(hipMalloc(&A_d1, (NUM_ELM*sizeof(float)))); - - for ( int i = 0; i < NUM_ELM; ++i ) { - A_h[i] = 123; - B_h[i] = 0; - } - - // Copying only half the memory on device side from host - HIPCHECK(hipMemcpy(A_d, A_h, (NUM_ELM/2) * sizeof(float), hipMemcpyDefault)); - // Copying device memory to host to verify if the content is consistent - HIPCHECK(hipMemcpy(B_h, A_d, NUM_ELM * sizeof(float), hipMemcpyDefault)); - // Verifying the host content copied in the above step for consistency. - int Data_mismatch = 0; - - for (int i = 0; i < (NUM_ELM/2); ++i) { - if (B_h[i] != 123) { - Data_mismatch++; - } - } - - if (Data_mismatch != 0) { - printf("Data Mismatch for negative test\n"); - IfTestPassed = false; - } - - str_out = hipGetErrorString(hipMemcpy(NULL, A_d, (NUM_ELM/2) * sizeof(float), - hipMemcpyDefault)); - if ((str_err.compare(str_out)) != 0) { - printf("Failed for hipMemcpy with NULL for destination parameter.\n"); - printf("Error: %s\n", str_out.c_str()); - IfTestPassed = false; - } - - str_out = hipGetErrorString(hipMemcpy(A_h, NULL, (NUM_ELM/2) * sizeof(float), - hipMemcpyDefault)); - if ((str_err.compare(str_out)) != 0) { - printf("Failed for hipMemcpy with NULL for source\n"); - IfTestPassed = false; - } - - str_out = hipGetErrorString(hipMemcpy(NULL, NULL, (NUM_ELM/2) * sizeof(float), - hipMemcpyDefault)); - if ((str_err.compare(str_out)) != 0) { - printf("Failed for hipMemcpy with NULL for source and destination\n"); - IfTestPassed = false; - } - - // To check the behaviour if both the ptrs provided are same - HIPCHECK(hipMemcpy(A_d, A_d, (NUM_ELM/2) * sizeof(float), hipMemcpyDefault)); - HIPCHECK(hipMemcpy(A_h, A_h, (NUM_ELM/2) * sizeof(float), hipMemcpyDefault)); - - // To check the consistency of the data - HIPCHECK(hipMemcpy(B_h, A_d, (NUM_ELM/2) * sizeof(float), hipMemcpyDefault)); - Data_mismatch = 0; - - for (int i = 0; i < (NUM_ELM/2); ++i) { - if (B_h[i] != 123) { - Data_mismatch++; - } - } - - if (Data_mismatch != 0) { - printf("Data Mismatch after memcpy of same src and destination\n"); - IfTestPassed = false; - } - - // Memory copy on same device with two different regions - HIPCHECK(hipMemcpy(A_d1, A_d, (NUM_ELM) * sizeof(float), hipMemcpyDefault)); - - HIPCHECK(hipFree(A_d)); - HIPCHECK(hipFree(A_d1)); - free(A_h); - free(B_h); - - return IfTestPassed; -} - -bool Memcpy_Negative_Tests::Test_MemcpyAsync(void) { - bool IfTestPassed = true; - float *A_h = NULL, *B_h = NULL, *A_d = NULL, *A_d1 = NULL; - std::string str_out, str_err = "hipErrorInvalidValue"; - hipStream_t stream; - HIPCHECK(hipStreamCreate(&stream)); - A_h = reinterpret_cast(malloc(NUM_ELM * sizeof(float))); - B_h = reinterpret_cast(malloc(NUM_ELM * sizeof(float))); - - if ((A_h == nullptr) || (B_h == nullptr)) { - failed("Malloc call failed!"); - } - - HIPCHECK(hipMalloc(&A_d, (NUM_ELM * sizeof(float)))); - HIPCHECK(hipMalloc(&A_d1, (NUM_ELM * sizeof(float)))); - for (int i = 0; i < NUM_ELM; ++i) { - A_h[i] = 123; - B_h[i] = 0; - } - - // Copying host data into the device. - HIPCHECK(hipMemcpyAsync(A_d1, A_h, NUM_ELM * sizeof(float), - hipMemcpyDefault, stream)); - - // Passing null pointer: seg fault observed with the following. - str_out = hipGetErrorString(hipMemcpyAsync(NULL, A_h, NUM_ELM * sizeof(float), - hipMemcpyDefault, stream)); - if ((str_err.compare(str_out)) != 0) { - printf("Failed for hipMemcpyAsync with NULL for destination\n"); - IfTestPassed = false; - } - - str_out = hipGetErrorString(hipMemcpyAsync(A_d, NULL, NUM_ELM * sizeof(float), - hipMemcpyDefault, stream)); - if ((str_err.compare(str_out)) != 0) { - printf("Failed for hipMemcpyAsync with NULL for source\n"); - IfTestPassed = false; - } - - str_out = hipGetErrorString(hipMemcpyAsync(NULL, NULL, - NUM_ELM * sizeof(float), - hipMemcpyDefault, stream)); - if ((str_err.compare(str_out)) != 0) { - printf("Failed for hipMemcpyAsync with NULL for source and destination\n"); - IfTestPassed = false; - } - - // Passing default stream just for sanity kind of check - HIPCHECK(hipMemcpyAsync(A_d, A_h, NUM_ELM * sizeof(float), hipMemcpyDefault, - 0)); - - // Passing stream object belong to destination gpu - // which is against the suggested usage. - HIPCHECK(hipMemcpyAsync(A_d, A_d1, NUM_ELM * sizeof(float), - hipMemcpyDefault, stream)); - - // Passing incorrect memcpy kind is not allowed hence those scenarios - // are not included - - // Copying only half the memory on device side from host - HIPCHECK(hipMemcpyAsync(A_d, A_h, (NUM_ELM/2) * sizeof(float), - hipMemcpyDefault, stream)); - // Copying device memory to host to verify the content is consistent. - HIPCHECK(hipMemcpy(B_h, A_d, (NUM_ELM/2) * sizeof(float), hipMemcpyDefault)); - - // Verifying the host content copied in the above step for consistency. - int Data_mismatch = 0; - for (int i = 0; i < (NUM_ELM/2); ++i) { - if (B_h[i] != 123) { - Data_mismatch++; - } - } - - if (Data_mismatch != 0) { - printf("Data Mismatch after half the size memcpyAsync\n"); - IfTestPassed = false; - } - - // To check the behaviour if both the ptrs provided are same - HIPCHECK(hipMemcpyAsync(A_d, A_d, (NUM_ELM/2) * sizeof(float), - hipMemcpyDefault, stream)); - HIPCHECK(hipMemcpyAsync(A_h, A_h, (NUM_ELM/2) * sizeof(float), - hipMemcpyDefault, stream)); - // To check the consistency of the data - HIPCHECK(hipMemcpy(B_h, A_d, (NUM_ELM) * sizeof(float), hipMemcpyDefault)); - Data_mismatch = 0; - for (int i = 0; i < (NUM_ELM); ++i) { - if (B_h[i] != 123) { - Data_mismatch++; - } - } - - if (Data_mismatch != 0) { - printf("Data Mismatch after memcpyAsync of same src and destination\n"); - IfTestPassed = false; - } - - // Memory copy on same device with two different regions - HIPCHECK(hipMemcpyAsync(A_d1, A_d, (NUM_ELM) * sizeof(float), - hipMemcpyDefault, stream)); - - HIPCHECK(hipStreamSynchronize(stream)); - HIPCHECK(hipFree(A_d)); - HIPCHECK(hipFree(A_d1)); - free(A_h); - free(B_h); - - return IfTestPassed; -} - -bool Memcpy_Negative_Tests::Test_MemcpyHtoD(void) { - bool IfTestPassed = true; - float *A_h = NULL, *B_h = NULL, *A_d = NULL, *A_d1 = NULL; - std::string str_out, str_err = "hipErrorInvalidValue"; - A_h = reinterpret_cast(malloc(NUM_ELM * sizeof(float))); - B_h = reinterpret_cast(malloc(NUM_ELM * sizeof(float))); - if ((A_h == nullptr) || (B_h == nullptr)) { - failed("Malloc call failed!"); - } - HIPCHECK(hipMalloc(&A_d, (NUM_ELM * sizeof(float)))); - HIPCHECK(hipMalloc(&A_d1, (NUM_ELM * sizeof(float)))); - for (int i = 0; i < NUM_ELM; ++i) { - A_h[i] = 123; - B_h[i] = 0; - } - - // Passing null ptr to check the API behavior. - // Expectation: It should not crash and exit gracefully. - str_out = hipGetErrorString(hipMemcpyHtoD(NULL, A_h, - NUM_ELM * sizeof(float))); - if ((str_err.compare(str_out)) != 0) { - printf("Failed for hipMemcpyHtoD with NULL for destination\n"); - IfTestPassed = false; - } - str_out = hipGetErrorString(hipMemcpyHtoD(A_d, NULL, - NUM_ELM * sizeof(float))); - if ((str_err.compare(str_out)) != 0) { - printf("Failed for hipMemcpyHtoD with NULL for source\n"); - IfTestPassed = false; - } - str_out = hipGetErrorString(hipMemcpyHtoD(NULL, NULL, - NUM_ELM * sizeof(float))); - if ((str_err.compare(str_out)) != 0) { - printf("Failed for hipMemcpyHtoD with NULL for source and destination\n"); - IfTestPassed = false; - } - // Copy half of the allocated memory - HIPCHECK(hipMemcpyHtoD(A_d, A_h, NUM_ELM * sizeof(float) / 2)); - // copying back to host to verify - HIPCHECK(hipMemcpyDtoH(B_h, A_d, NUM_ELM * sizeof(float))); - - int Data_mismatch = 0; - for (int i = 0; i < (NUM_ELM / 2); ++i) - if (B_h[i] != 123) - Data_mismatch++; - - if (Data_mismatch != 0) { - printf("Data Mismatch after hipMemcpyHtoD with half size\n"); - IfTestPassed = false; - } - - // Setting device memory to zero - HIPCHECK(hipMemset(A_d, 0, NUM_ELM * sizeof(float))); - // Swap source and destination pointer - HIPCHECK(hipMemcpyHtoD(A_h, A_d, NUM_ELM * sizeof(float))); - - // Pass same pointers in source and destination params - HIPCHECK(hipMemcpyHtoD(A_h, A_h, NUM_ELM * sizeof(float))); - HIPCHECK(hipMemcpyHtoD(A_d, A_d, NUM_ELM * sizeof(float))); - - // Mem copy on same device with two different regions - HIPCHECK(hipMemcpyHtoD(A_d1, A_d, NUM_ELM * sizeof(float))); - - HIPCHECK(hipFree(A_d)); - HIPCHECK(hipFree(A_d1)); - free(A_h); - free(B_h); - - return IfTestPassed; -} - -bool Memcpy_Negative_Tests::Test_MemcpyHtoDAsync(void) { - bool IfTestPassed = true; - float *A_h = NULL, *B_h = NULL, *A_d = NULL, *A_d1 = NULL; - std::string str_out, str_err = "hipErrorInvalidValue"; - hipStream_t stream; - HIPCHECK(hipStreamCreate(&stream)); - A_h = reinterpret_cast(malloc(NUM_ELM * sizeof(float))); - B_h = reinterpret_cast(malloc(NUM_ELM * sizeof(float))); - if ((A_h == nullptr) || (B_h == nullptr)) { - failed("Malloc call failed!"); - } - HIPCHECK(hipMalloc(&A_d, (NUM_ELM * sizeof(float)))); - HIPCHECK(hipMalloc(&A_d1, (NUM_ELM * sizeof(float)))); - for (int i = 0; i < NUM_ELM; ++i) { - A_h[i] = 123; - B_h[i] = 0; - } - - // Passing null ptr to check the API behavior. - // Expectation: It should not crash and exit gracefully. - str_out = hipGetErrorString(hipMemcpyHtoDAsync(NULL, A_h, - NUM_ELM * sizeof(float), - stream)); - if ((str_err.compare(str_out)) != 0) { - printf("Failed for hipMemcpyHtoDAsync with NULL for destination\n"); - IfTestPassed = false; - } - - str_out = hipGetErrorString(hipMemcpyHtoDAsync(A_d, NULL, - NUM_ELM * sizeof(float), - stream)); - if ((str_err.compare(str_out)) != 0) { - printf("Failed for hipMemcpyHtoDAsync with NULL for source\n"); - IfTestPassed = false; - } - str_out = hipGetErrorString(hipMemcpyHtoDAsync(NULL, NULL, - NUM_ELM * sizeof(float), - stream)); - if ((str_err.compare(str_out)) != 0) { - printf("Failed MemcpyHtoDAsync with NULL for source and destination\n"); - IfTestPassed = false; - } - - // Copy half of the allocated memory - HIPCHECK(hipMemcpyHtoDAsync(A_d, A_h, NUM_ELM * sizeof(float)/2, stream)); - // copying back to host to verify - HIPCHECK(hipMemcpyDtoH(B_h, A_d, NUM_ELM * sizeof(float))); - int Data_mismatch = 0; - for (int i = 0; i < (NUM_ELM/2); ++i) - if (B_h[i] != 123) - Data_mismatch++; - if (Data_mismatch != 0) { - printf("Data Mismatch after hipMemcpyHtoDAsync with half size\n"); - IfTestPassed = false; - } - - // Setting device memory to zero - HIPCHECK(hipMemset(A_d, 0, NUM_ELM * sizeof(float))); - // Swap source and destination pointer - HIPCHECK(hipMemcpyHtoDAsync(B_h, A_d, NUM_ELM * sizeof(float), stream)); - HIPCHECK(hipStreamSynchronize(stream)); - if (B_h[0] != 0) { - printf("Data Mismatch after hipMemcpyHtoDAsync with memset to 0\n"); - IfTestPassed = false; - } - - // Pass same pointers in source and destination params - HIPCHECK(hipMemcpyHtoDAsync(A_h, A_h, NUM_ELM * sizeof(float), stream)); - HIPCHECK(hipMemcpyHtoDAsync(A_d, A_d, NUM_ELM * sizeof(float), stream)); - - // Mem copy on same device with two different regions - HIPCHECK(hipMemcpyHtoDAsync(A_d1, A_d, NUM_ELM * sizeof(float), stream)); - HIPCHECK(hipStreamSynchronize(stream)); - // Checking the api with null stream - HIPCHECK(hipMemcpyHtoDAsync(A_d1, A_d, NUM_ELM * sizeof(float), 0)); - HIPCHECK(hipStreamSynchronize(stream)); - - HIPCHECK(hipFree(A_d)); - HIPCHECK(hipFree(A_d1)); - free(A_h); - free(B_h); - - return IfTestPassed; -} - -bool Memcpy_Negative_Tests::Test_MemcpyDtoH(void) { - bool IfTestPassed = true; - float *A_h = NULL, *B_h = NULL, *A_d = NULL, *A_d1 = NULL; - std::string str_out, str_err = "hipErrorInvalidValue"; - A_h = reinterpret_cast(malloc(NUM_ELM * sizeof(float))); - B_h = reinterpret_cast(malloc(NUM_ELM * sizeof(float))); - if ((A_h == nullptr) || (B_h == nullptr)) { - failed("Malloc call failed!"); - } - HIPCHECK(hipMalloc(&A_d, (NUM_ELM * sizeof(float)))); - HIPCHECK(hipMalloc(&A_d1, (NUM_ELM * sizeof(float)))); - for (int i = 0; i < NUM_ELM; ++i) { - A_h[i] = 123; - B_h[i] = 0; - } - - // Copying data from host to device for further operations - HIPCHECK(hipMemcpyHtoD(A_d, A_h, NUM_ELM * sizeof(float))); - - // Passing null ptr to check the API behavior. - // Expectation: It should not crash and exit gracefully. - str_out = hipGetErrorString(hipMemcpyDtoH(NULL, A_d, - NUM_ELM * sizeof(float))); - if ((str_err.compare(str_out)) != 0) { - printf("Failed for hipMemcpyDtoH with NULL for destination\n"); - IfTestPassed = false; - } - str_out = hipGetErrorString(hipMemcpyDtoH(A_d, NULL, - NUM_ELM * sizeof(float))); - if ((str_err.compare(str_out)) != 0) { - printf("Failed for hipMemcpyDtoH with NULL for source\n"); - IfTestPassed = false; - } - str_out = hipGetErrorString(hipMemcpyDtoH(NULL, NULL, - NUM_ELM * sizeof(float))); - if ((str_err.compare(str_out)) != 0) { - printf("Failed for hipMemcpyDtoH with NULL for source and destination\n"); - IfTestPassed = false; - } - // Copy half of the allocated memory - HIPCHECK(hipMemcpyDtoH(B_h, A_d, NUM_ELM * sizeof(float)/2)); - - int Data_mismatch = 0; - for (int i = 0; i < (NUM_ELM/2); ++i) - if (B_h[i] != 123) - Data_mismatch++; - - if (Data_mismatch != 0) { - printf("Data Mismatch after hipMemcpyDtoH with half size\n"); - IfTestPassed = false; - } - - // Setting device memory to zero - HIPCHECK(hipMemset(A_d, 0, NUM_ELM * sizeof(float))); - // Swap source and destination pointer - HIPCHECK(hipMemcpyDtoH(A_d, A_h, NUM_ELM * sizeof(float))); - - // Pass same pointers in source and destination params - HIPCHECK(hipMemcpyDtoH(A_h, A_h, NUM_ELM * sizeof(float))); - HIPCHECK(hipMemcpyDtoH(A_d, A_d, NUM_ELM * sizeof(float))); - - // Mem copy on same device with two diffeent regions - HIPCHECK(hipMemcpyDtoH(A_d1, A_d, NUM_ELM * sizeof(float))); - - HIPCHECK(hipFree(A_d)); - HIPCHECK(hipFree(A_d1)); - free(A_h); - free(B_h); - - return IfTestPassed; -} - -bool Memcpy_Negative_Tests::Test_MemcpyDtoHAsync(void) { - bool IfTestPassed = true; - float *A_h = NULL, *B_h = NULL, *A_d = NULL, *A_d1 = NULL; - std::string str_out, str_err = "hipErrorInvalidValue"; - hipStream_t stream; - HIPCHECK(hipStreamCreate(&stream)); - A_h = reinterpret_cast(malloc(NUM_ELM * sizeof(float))); - B_h = reinterpret_cast(malloc(NUM_ELM * sizeof(float))); - if ((A_h == nullptr) || (B_h == nullptr)) { - failed("Malloc call failed!"); - } - HIPCHECK(hipMalloc(&A_d, (NUM_ELM * sizeof(float)))); - HIPCHECK(hipMalloc(&A_d1, (NUM_ELM * sizeof(float)))); - for (int i = 0; i < NUM_ELM; ++i) { - A_h[i] = 123; - B_h[i] = 0; - } - - // Copying data from host to device for further operations - HIPCHECK(hipMemcpyHtoDAsync(A_d, A_h, NUM_ELM * sizeof(float), stream)); - - // Passing null ptr to check the API behavior. - // Expectation: It should not crash and exit gracefully. - str_out = hipGetErrorString(hipMemcpyDtoHAsync(NULL, A_d, - NUM_ELM * sizeof(float), - stream)); - if ((str_err.compare(str_out)) != 0) { - printf("Failed for hipMemcpyDtoHAsync with NULL for destination\n"); - IfTestPassed = false; - } - str_out = hipGetErrorString(hipMemcpyDtoHAsync(A_d, NULL, - NUM_ELM * sizeof(float), - stream)); - if ((str_err.compare(str_out)) != 0) { - printf("Failed for hipMemcpyDtoHAsync with NULL for source\n"); - IfTestPassed = false; - } - str_out = hipGetErrorString(hipMemcpyDtoHAsync(NULL, NULL, - NUM_ELM * sizeof(float), - stream)); - if ((str_err.compare(str_out)) != 0) { - printf("Failed hipMemcpyDtoHAsync with NULL for source and destination\n"); - IfTestPassed = false; - } - - // Copy half of the allocated memory - HIPCHECK(hipMemcpyDtoHAsync(B_h, A_d, NUM_ELM * sizeof(float)/2, stream)); - HIPCHECK(hipStreamSynchronize(stream)); - int Data_mismatch = 0; - for (int i = 0; i < (NUM_ELM/2); ++i) - if (B_h[i] != 123) - Data_mismatch++; - - if (Data_mismatch != 0) { - printf("Data Mismatch after hipMemcpyDtoHAsync with half size\n"); - IfTestPassed = false; - } - // Checking the api with default stream - HIPCHECK(hipMemcpyDtoHAsync(B_h, A_d, NUM_ELM * sizeof(float), 0)); - // Setting device memory to zero - HIPCHECK(hipMemset(A_d, 0, NUM_ELM * sizeof(float))); - // Swap source and destination pointer - HIPCHECK(hipMemcpyDtoHAsync(A_d, A_h, NUM_ELM * sizeof(float), stream)); - - // Pass same pointers in source and destination params - HIPCHECK(hipMemcpyDtoHAsync(A_h, A_h, NUM_ELM * sizeof(float), stream)); - HIPCHECK(hipStreamSynchronize(stream)); - HIPCHECK(hipMemcpyDtoHAsync(A_d, A_d, NUM_ELM * sizeof(float), stream)); - HIPCHECK(hipStreamSynchronize(stream)); - // Mem copy on same device with two different regions - HIPCHECK(hipMemcpyDtoHAsync(A_d1, A_d, NUM_ELM * sizeof(float), stream)); - HIPCHECK(hipStreamSynchronize(stream)); - - HIPCHECK(hipFree(A_d)); - HIPCHECK(hipFree(A_d1)); - free(A_h); - free(B_h); - - return IfTestPassed; -} - -bool Memcpy_Negative_Tests::Test_MemcpyDtoD(void) { - bool IfTestPassed = true; - float *A_h = NULL, *B_h = NULL, *A_d1 = NULL, *A_d2 = NULL, *Ad1 = NULL; - std::string str_out, str_err = "hipErrorInvalidValue"; - A_h = reinterpret_cast(malloc(NUM_ELM * sizeof(float))); - B_h = reinterpret_cast(malloc(NUM_ELM * sizeof(float))); - if ((A_h == nullptr) || (B_h == nullptr)) { - failed("Malloc call failed!"); - } - HIPCHECK(hipMalloc(&A_d1, (NUM_ELM * sizeof(float)))); - HIPCHECK(hipMalloc(&Ad1, (NUM_ELM * sizeof(float)))); - HIPCHECK(hipMemset(A_d1, 0, NUM_ELM * sizeof(float))); - if (Available_Gpus > 1) { - HIPCHECK(hipSetDevice(1)); - HIPCHECK(hipMalloc(&A_d2, (NUM_ELM * sizeof(float)))); - HIPCHECK(hipMemset(A_d2, 1, NUM_ELM * sizeof(float))); - } - for (int i = 0; i < NUM_ELM; ++i) { - A_h[i] = 123; - B_h[i] = 0; - } - // Passing null pointers to check the behaviour:: - str_out = hipGetErrorString(hipMemcpyDtoD(&A_d1, NULL, - NUM_ELM * sizeof(float))); - if ((str_err.compare(str_out)) != 0) { - printf("Failed for hipMemcpyDtoD with NULL for source\n"); - IfTestPassed = false; - } - str_out = hipGetErrorString(hipMemcpyDtoD(NULL, &A_d2, - NUM_ELM * sizeof(float))); - if ((str_err.compare(str_out)) != 0) { - printf("Failed for hipMemcpyDtoD with NULL for destination\n"); - IfTestPassed = false; - } - str_out = hipGetErrorString(hipMemcpyDtoD(NULL, NULL, - NUM_ELM * sizeof(float))); - if ((str_err.compare(str_out)) != 0) { - printf("Failed for hipMemcpyDtoD with NULL for source and destination\n"); - IfTestPassed = false; - } - - // Pass real but host ptr:: The below two scenarios gives seg fault. - // Behaviour is as expected - // HIPCHECK(hipMemcpyDtoD(&A_d1, &A_h, NUM_ELM * sizeof(float))); - // HIPCHECK(hipMemcpyDtoD(&A_h, &A_d1, NUM_ELM * sizeof(float))); - int Data_mismatch = 0; - // Copying half of actually allocated memory - HIPCHECK(hipSetDevice(0)); - if (Available_Gpus > 1) { - HIPCHECK(hipMemcpyHtoD(A_d1, A_h, NUM_ELM * sizeof(float))); - if (true == gpusIsPeer(0, 1)) { - HIPCHECK(hipMemcpyDtoD(A_d2, A_d1, NUM_ELM * sizeof(float)/2)); - HIPCHECK(hipMemcpyDtoH(B_h, A_d2, NUM_ELM * sizeof(float))); - for (int i = 0; i < NUM_ELM/2; ++i) { - if (B_h[i] != 123) - Data_mismatch++; - } - if (Data_mismatch != 0) { - printf("Data mismatch hipMemcpyDtoD between devices\n"); - IfTestPassed = false; - } - } - } - - // Passing same pointers for source and destination - HIPCHECK(hipMemcpyDtoD(A_d1, A_d1, NUM_ELM * sizeof(float))); - if (Available_Gpus > 1) { - HIPCHECK(hipMemcpyDtoD(A_d2, A_d2, NUM_ELM * sizeof(float))); - } - // Memcpy on same device with two different regions - HIPCHECK(hipMemcpyDtoD(Ad1, A_d1, NUM_ELM * sizeof(float))); - - HIPCHECK(hipFree(A_d1)); - HIPCHECK(hipFree(Ad1)); - if (Available_Gpus > 1) - HIPCHECK(hipFree(A_d2)); - free(A_h); - free(B_h); - - return IfTestPassed; -} - -bool Memcpy_Negative_Tests::Test_MemcpyDtoDAsync(void) { - bool IfTestPassed = true; - float *A_h = NULL, *B_h = NULL, *A_d1 = NULL, *A_d2 = NULL, *Ad1 = NULL; - std::string str_out, str_err = "hipErrorInvalidValue"; - hipStream_t stream; - HIPCHECK(hipStreamCreate(&stream)); - A_h = reinterpret_cast(malloc(NUM_ELM * sizeof(float))); - B_h = reinterpret_cast(malloc(NUM_ELM * sizeof(float))); - if ((A_h == nullptr) || (B_h == nullptr)) { - failed("Malloc call failed!"); - } - HIPCHECK(hipMalloc(&A_d1, (NUM_ELM * sizeof(float)))); - HIPCHECK(hipMalloc(&Ad1, (NUM_ELM * sizeof(float)))); - HIPCHECK(hipMemset(A_d1, 0, NUM_ELM * sizeof(float))); - if (Available_Gpus > 1) { - HIPCHECK(hipSetDevice(1)); - HIPCHECK(hipMalloc(&A_d2, (NUM_ELM * sizeof(float)))); - HIPCHECK(hipMemset(A_d2, 1, NUM_ELM * sizeof(float))); - } - for (int i = 0; i < NUM_ELM; ++i) { - A_h[i] = 123; - B_h[i] = 0; - } - // Passing null pointers to check the behaviour:: - str_out = hipGetErrorString(hipMemcpyDtoDAsync(&A_d1, NULL, - NUM_ELM * sizeof(float), - stream)); - if ((str_err.compare(str_out)) != 0) { - printf("Failed for hipMemcpyDtoDAsync with NULL for source\n"); - IfTestPassed = false; - } - str_out = hipGetErrorString(hipMemcpyDtoDAsync(NULL, &A_d2, - NUM_ELM * sizeof(float), - stream)); - if ((str_err.compare(str_out)) != 0) { - printf("Failed for hipMemcpyDtoDAsync with NULL for destination\n"); - IfTestPassed = false; - } - str_out = hipGetErrorString(hipMemcpyDtoDAsync(NULL, NULL, - NUM_ELM * sizeof(float), - stream)); - if ((str_err.compare(str_out)) != 0) { - printf("Failed MemcpyDtoDAsync with NULL for source and destination\n"); - IfTestPassed = false; - } - - int Data_mismatch = 0; - // Copying half of actually allocated memory - HIPCHECK(hipSetDevice(0)); - if (Available_Gpus > 1) { - HIPCHECK(hipMemcpyHtoD(A_d1, A_h, NUM_ELM * sizeof(float))); - if (true == gpusIsPeer(0, 1)) { - HIPCHECK(hipMemcpyDtoDAsync(A_d2, A_d1, NUM_ELM * sizeof(float)/2, - stream)); - HIPCHECK(hipMemcpyDtoH(B_h, A_d2, NUM_ELM * sizeof(float))); - for (int i = 0; i < NUM_ELM/2; ++i) { - if (B_h[i] != 123) - Data_mismatch++; - } - if (Data_mismatch != 0) { - printf("Data mismatch hipMemcpyDtoDAsync between devices\n"); - IfTestPassed = false; - } - } - } - - // Memcpy on same device with two different regions - HIPCHECK(hipMemcpyDtoDAsync(Ad1, A_d1, NUM_ELM * sizeof(float) , stream)); - // Testing hipMemcpyDtoDAsync between two devices. - if (Available_Gpus > 1) { - if (true == gpusIsPeer(0, 1)) { - HIPCHECK(hipMemcpyDtoDAsync(A_d2, A_d1, NUM_ELM * sizeof(float), 0)); - } - } - HIPCHECK(hipStreamSynchronize(stream)); - HIPCHECK(hipFree(A_d1)); - HIPCHECK(hipFree(Ad1)); - if (Available_Gpus > 1) - HIPCHECK(hipFree(A_d2)); - free(A_h); - free(B_h); - - return IfTestPassed; -} - -template -memcpyTests::memcpyTests(apiToTest val, size_t num_elmts) { - api = val; - NUM_ELMTS = num_elmts; - printf("%zu ", NUM_ELMTS * sizeof(T)); - fflush(stdout); - A_h = reinterpret_cast(malloc(NUM_ELMTS * sizeof(T))); - B_h = reinterpret_cast(malloc(NUM_ELMTS * sizeof(T))); - if ((A_h == NULL) || (B_h == NULL)) { - exit(1); - } - - if (api >= TEST_MEMCPYD2D) { - HIPCHECK(hipStreamCreate(&stream)); - } - for (size_t i = 0; i < NUM_ELMTS; ++i) { - A_h[i] = 123; - B_h[i] = 0; - } -} - -template -bool memcpyTests::Memcpy_And_verify() { - bool bFail = false; - std::atomic Data_mismatch{0}; - T *A_d[Available_Gpus]; - for (int i = 0; i < Available_Gpus; ++i) { - HIPCHECK(hipSetDevice(i)); - HIPCHECK(hipMalloc(&A_d[i], NUM_ELMTS * sizeof(T))); - } - HIPCHECK(hipSetDevice(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(hipSetDevice(i)); - HIPCHECK(hipMemcpy(A_d[i], A_h, NUM_ELMTS * sizeof(T), - hipMemcpyHostToDevice)); - HIPCHECK(hipMemcpy(B_h, A_d[i], NUM_ELMTS * sizeof(T), - hipMemcpyDeviceToHost)); - for (int j = 0; j < NUM_ELMTS; ++j) { - if (A_h[j] != B_h[j]) { - Data_mismatch++; - } - } - - if (Data_mismatch.load() != 0) { - printf("hipMemcpy: Failed for GPU: %d\n", i); - bFail = true; - } - } - // Device to Device copying for all combinations - for (int i = 0; i < Available_Gpus; ++i) { - for (int j = 1; j < Available_Gpus; ++j) { - if (true == gpusIsPeer(i, j)) { - HIPCHECK(hipMemcpy(A_d[j], A_d[i], NUM_ELMTS * sizeof(T), - hipMemcpyDefault)); - // Copying in direction reverse of above to check if bidirectional - // access is happening without any error - HIPCHECK(hipMemcpy(A_d[i], A_d[j], NUM_ELMTS * sizeof(T), - hipMemcpyDefault)); - // Copying data to host to verify the content - HIPCHECK(hipMemcpy(B_h, A_d[j], NUM_ELMTS * sizeof(T), - hipMemcpyDefault)); - for (int i = 0; i < NUM_ELMTS; ++i) { - if (A_h[i] != B_h[i]) - Data_mismatch++; - } - - if (Data_mismatch.load() != 0) { - printf("hipMemcpy: Failed between GPU: %d and %d\n", i, j); - bFail = true; - } - } - } - } - break; - case TEST_MEMCPYH2D: // To test hipMemcpyHtoD() - for (int i = 0; i < Available_Gpus; ++i) { - HIPCHECK(hipMemcpyHtoD(A_d[i], A_h, NUM_ELMTS * sizeof(T))); - // Copying data from device to host to check data consistency - HIPCHECK(hipMemcpy(B_h, A_d[i], NUM_ELMTS * sizeof(T), - hipMemcpyDeviceToHost)); - for (size_t i = 0; i < NUM_ELMTS; ++i) { - if (A_h[i] != B_h[i]) - Data_mismatch++; - } - if (Data_mismatch.load() != 0) { - printf("hipMemcpyHtoD: failed\n"); - bFail = true; - } - } - break; - case TEST_MEMCPYD2H: // To test hipMemcpyDtoH()--done - for (int i = 0; i < Available_Gpus; ++i) { - HIPCHECK(hipMemcpy(A_d[i], A_h, NUM_ELMTS * sizeof(T), - hipMemcpyHostToDevice)); - HIPCHECK(hipMemcpyDtoH(B_h, A_d[i], NUM_ELMTS * sizeof(T))); - for (size_t i = 0; i < NUM_ELMTS; ++i) { - if (A_h[i] != B_h[i]) - Data_mismatch++; - } - if (Data_mismatch.load() != 0) { - printf("hipMemcpyDtoH: failed\n"); - bFail = true; - } - } - 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_ELMTS * sizeof(T))); - for (int i = 0; i < Available_Gpus; ++i) { - for (int j = 1; j < Available_Gpus; ++j) { - if (true == gpusIsPeer(i, j)) { - HIPCHECK(hipMemcpyHtoD(A_d[i], A_h, NUM_ELMTS * sizeof(T))); - HIPCHECK(hipMemcpyDtoD(A_d[j], A_d[i], NUM_ELMTS * sizeof(T))); - // Copying in direction reverse of above to check if bidirectional - // access is happening without any error - HIPCHECK(hipMemcpyDtoD(A_d[i], A_d[j], NUM_ELMTS * sizeof(T))); - HIPCHECK(hipMemcpy(B_h, A_d[i], NUM_ELMTS * sizeof(T), - hipMemcpyDeviceToHost)); - for (size_t i = 0; i < NUM_ELMTS; ++i) { - if (A_h[i] != B_h[i]) - Data_mismatch++; - } - if (Data_mismatch.load() != 0) { - printf("hipMemcpyDtoD: failed between GPU: %d and %d\n", i, j); - bFail = true; - } - } - } - } - } 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_ELMTS * sizeof(T), - hipMemcpyHostToHost)); - for (size_t i = 0; i < NUM_ELMTS; ++i) { - if (A_h[i] != B_h[i]) - Data_mismatch++; - } - if (Data_mismatch.load() != 0) { - printf("hipMemcpy (Host to Host): failed\n"); - bFail = true; - } - } - 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_ELMTS * sizeof(T), - hipMemcpyHostToDevice, stream)); - HIPCHECK(hipMemcpyAsync(B_h, A_d[i], NUM_ELMTS * sizeof(T), - hipMemcpyDeviceToHost, stream)); - for (size_t i = 0; i < NUM_ELMTS; ++i) { - if (A_h[i] != B_h[i]) - Data_mismatch++; - } - - if (Data_mismatch.load() != 0) { - printf("hipMemcpyAsync: failed for GPU %d\n", i); - bFail = true; - } - } - // Device to Device copying for all combinations - for (int i = 0; i < Available_Gpus; ++i) { - for (int j = 1; j < Available_Gpus; ++j) { - if (true == gpusIsPeer(i, j)) { - HIPCHECK(hipMemcpyAsync(A_d[j], A_d[i], NUM_ELMTS * sizeof(T), - hipMemcpyDefault, stream)); - // 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_ELMTS * sizeof(T), - hipMemcpyDefault, stream)); - HIPCHECK(hipMemcpy(B_h, A_d[j], NUM_ELMTS * sizeof(T), - hipMemcpyDefault)); - for (size_t i = 0; i < NUM_ELMTS; ++i) { - if (A_h[i] != B_h[i]) - Data_mismatch++; - } - - if (Data_mismatch.load() != 0) { - printf("hipMemcpyAsync: Failed between GPU: %d and %d\n", i, j); - bFail = true; - } - } - } - } - break; - case TEST_MEMCPYH2DASYNC: // To test hipMemcpyHtoDAsync() - for (int i = 0; i < Available_Gpus; ++i) { - HIPCHECK(hipMemcpyHtoDAsync(A_d[i], A_h, NUM_ELMTS * sizeof(T), - stream)); - // Copying data from device to host to check data consistency - HIPCHECK(hipMemcpy(B_h, A_d[i], NUM_ELMTS * sizeof(T), - hipMemcpyDeviceToHost)); - for (size_t i = 0; i < NUM_ELMTS; ++i) { - if (A_h[i] != B_h[i]) - Data_mismatch++; - } - if (Data_mismatch.load() != 0) { - printf("hipMemcpyHtoDAsync: failed\n"); - bFail = true; - } - } - break; - case TEST_MEMCPYD2HASYNC: // To test hipMemcpyDtoHAsync() - for (int i = 0; i < Available_Gpus; ++i) { - HIPCHECK(hipMemcpy(A_d[i], A_h, NUM_ELMTS * sizeof(T), - hipMemcpyHostToDevice)); - HIPCHECK(hipMemcpyDtoHAsync(B_h, A_d[i], NUM_ELMTS * sizeof(T), - stream)); - for (size_t i = 0; i < NUM_ELMTS; ++i) { - if (A_h[i] != B_h[i]) - Data_mismatch++; - } - if (Data_mismatch.load() != 0) { - printf("hipMemcpyDtoHAsync: failed\n"); - bFail = true; - } - } - 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 D to H - HIPCHECK(hipMemcpyHtoD(A_d[0], A_h, NUM_ELMTS * sizeof(T))); - for (int i = 0; i < Available_Gpus; ++i) { - for (int j = 1; j < Available_Gpus; ++j) { - if (true == gpusIsPeer(i, j)) { - HIPCHECK(hipSetDevice(j)); - HIPCHECK(hipMemcpyDtoDAsync(A_d[j], A_d[i], NUM_ELMTS * sizeof(T), - stream)); - // Copying in direction reverse of above to check if bidirectional - // access is happening without any error - HIPCHECK(hipMemcpyDtoDAsync(A_d[i], A_d[j], NUM_ELMTS * sizeof(T), - stream)); - HIPCHECK(hipDeviceSynchronize()); - HIPCHECK(hipMemcpy(B_h, A_d[i], NUM_ELMTS * sizeof(T), - hipMemcpyDeviceToHost)); - for (size_t i = 0; i < NUM_ELMTS; ++i) { - if (A_h[i] != B_h[i]) - Data_mismatch++; - } - if (Data_mismatch.load() != 0) { - printf("hipMemcpyDtoDAsync: failed GPU: %d and %d\n", i, j); - bFail = true; - } - } - } - } - } 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_ELMTS * sizeof(T), - hipMemcpyHostToHost)); - for (size_t i = 0; i < NUM_ELMTS; ++i) { - if (A_h[i] != B_h[i]) - Data_mismatch++; - } - if (Data_mismatch.load() != 0) { - printf("hipMemcpy (Host to Host): failed\n"); - bFail = true; - } - } - break; - default: - printf("Did not receive valid option!\n"); - break; - } - - for (int i = 0; i < Available_Gpus; ++i) { - HIPCHECK(hipFree((A_d[i]))); - } - - // Return true if test is success - if (bFail == true) { - return false; - } else { - return true; - } -} - -template -memcpyTests::~memcpyTests() { - free(A_h); - free(B_h); - if (api >= TEST_MEMCPYD2D) { - HIPCHECK(hipStreamDestroy(stream)); - } -} - -void Thread_func(int Threadid) { - for (apiToTest api = TEST_MEMCPY; api < TEST_MAX; api = apiToTest(api + 1)) { - memcpyTests obj(api, 1024*1024); - if (false == obj.Memcpy_And_verify()) { - failureCount++; - } - } -} - -int parseExtraArguments(int argc, char* argv[]) { - int i = 0; - for (i = 1; i < argc; i++) { - const char* arg = argv[i]; - if (!strcmp(arg, " ")) { - // skip NULL args. - } else if (!strcmp(arg, "--memcpyPeersOnly")) { - if (++i >= argc || !HipTest::parseInt(argv[i], &memcpyPeersOnly)) { - failed("Bad memcpyPeersOnly argument"); - } - } else if (!strcmp(arg, "--testAllTypes")) { - if (++i >= argc || !HipTest::parseInt(argv[i], &testAllTypes)) { - failed("Bad testAllTypes argument"); - } - } else { - failed("Bad argument"); - } - } - return i; -} - - -int main(int argc, char* argv[]) { - bool TestPassed = true; - int extraArgs = 0; - HIPCHECK(hipGetDeviceCount(&Available_Gpus)); - extraArgs = HipTest::parseStandardArguments(argc, argv, false); - parseExtraArguments(extraArgs, argv); - - if (p_tests == 1) { - Memcpy_Negative_Tests test; - TestPassed = test.Test_Memcpy(); - TestPassed &= test.Test_MemcpyAsync(); - TestPassed &= test.Test_MemcpyHtoD(); - TestPassed &= test.Test_MemcpyHtoDAsync(); - TestPassed &= test.Test_MemcpyDtoD(); - TestPassed &= test.Test_MemcpyDtoDAsync(); - TestPassed &= test.Test_MemcpyDtoH(); - TestPassed &= test.Test_MemcpyDtoHAsync(); - if (TestPassed) { - passed(); - } else { - failed("Test Failed!"); - } - } else if (p_tests == 2) { - failureCount = 0; - std::thread Thrd[NUM_THREADS]; - for (int i = 0; i < NUM_THREADS; i++) - Thrd[i] = std::thread(Thread_func, 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(); - - if (failureCount.load() != 0) { - failed("Failed"); - } else { - passed(); - } - } else if (p_tests == 3) { - size_t free = 0, total = 0; - HIPCHECK(hipMemGetInfo(&free, &total)); - failureCount = 0; - // Need to see if allocating all of available free memory will result in - // any issues in windows system before adding the same - std::vector NUM_ELMTS{1, 5, 10, 100, 1024, 10*1024, 100*1024, - 1024*1024, 10*1024*1024, 100*1024*1024, - 1024*1024*1024}; - - for (apiToTest api = TEST_MEMCPY; api < TEST_MAX; api = apiToTest(api+1)) { - printf("\nTesting %s for size: ", apiNameToTest[api].c_str()); - - // Check for 0 size - memcpyTests obj(api, 0); - obj.Memcpy_And_verify(); - HIPCHECK(hipDeviceSynchronize()); - - for (size_t x : NUM_ELMTS) { - if ((x * sizeof(char)) <= free) { - memcpyTests obj(api, x); - obj.Memcpy_And_verify(); - HIPCHECK(hipDeviceSynchronize()); - } - - if (HIPTEST_TRUE == testAllTypes) { - // Testing memcpy with various data types - if ((x * sizeof(int)) <= free) { - memcpyTests obj(api, x); - obj.Memcpy_And_verify(); - HIPCHECK(hipDeviceSynchronize()); - } - if ((x * sizeof(size_t)) <= free) { - memcpyTests obj(api, x); - obj.Memcpy_And_verify(); - HIPCHECK(hipDeviceSynchronize()); - } - if ((x * sizeof(long double)) <= free) { - memcpyTests obj(api, x); - obj.Memcpy_And_verify(); - HIPCHECK(hipDeviceSynchronize()); - } - } - } - } - printf("\n"); - passed(); - } else { - failed("Didnt receive any valid option\n"); - } -}