From 8cd13f320685f492875b9910704784ec4560ccc4 Mon Sep 17 00:00:00 2001 From: Phaneendr-kumar Lanka Date: Wed, 20 Dec 2017 14:10:47 +0530 Subject: [PATCH 1/8] [nvccTests] Enable hipPeerToPeer_simple on nvcc [ROCm/hip commit: 88374a8a63889ffd2a990fc2acc9a0a7bdd45785] --- projects/hip/tests/src/p2p/hipPeerToPeer_simple.cpp | 8 ++++---- 1 file changed, 4 insertions(+), 4 deletions(-) diff --git a/projects/hip/tests/src/p2p/hipPeerToPeer_simple.cpp b/projects/hip/tests/src/p2p/hipPeerToPeer_simple.cpp index 32cab371b8..74fefa8fdb 100644 --- a/projects/hip/tests/src/p2p/hipPeerToPeer_simple.cpp +++ b/projects/hip/tests/src/p2p/hipPeerToPeer_simple.cpp @@ -23,10 +23,10 @@ THE SOFTWARE. // Also serves as a template for other tests. /* HIT_START - * BUILD: %t %s ../test_common.cpp - * RUN: %t EXCLUDE_HIP_PLATFORM all - * RUN: %t --memcpyWithPeer EXCLUDE_HIP_PLATFORM all - * RUN: %t --mirrorPeers EXCLUDE_HIP_PLATFORM all + * BUILD: %t %s ../test_common.cpp + * RUN: %t EXCLUDE_HIP_PLATFORM hipcc + * RUN: %t --memcpyWithPeer EXCLUDE_HIP_PLATFORM hipcc + * RUN: %t --mirrorPeers EXCLUDE_HIP_PLATFORM hipcc * HIT_END */ From e2b7d5b9fd29d4cb460fe8aab24c22bc73a7c298 Mon Sep 17 00:00:00 2001 From: Phaneendr-kumar Lanka Date: Wed, 10 Jan 2018 10:36:25 +0530 Subject: [PATCH 2/8] [nvccTests] Enable hipGetDeviceAttribute on nvcc [ROCm/hip commit: 5642a95dfa89b67795f5cc38cb998b98d866cb47] --- .../hip/tests/src/runtimeApi/device/hipGetDeviceAttribute.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/projects/hip/tests/src/runtimeApi/device/hipGetDeviceAttribute.cpp b/projects/hip/tests/src/runtimeApi/device/hipGetDeviceAttribute.cpp index 2919939694..0b965a3ae3 100644 --- a/projects/hip/tests/src/runtimeApi/device/hipGetDeviceAttribute.cpp +++ b/projects/hip/tests/src/runtimeApi/device/hipGetDeviceAttribute.cpp @@ -23,7 +23,7 @@ THE SOFTWARE. /* HIT_START * BUILD: %t %s ../../test_common.cpp - * RUN: %t EXCLUDE_HIP_PLATFORM nvcc + * RUN: %t * HIT_END */ From 5da7dcbd3b1c4764f3926e6609c9e18753290ead Mon Sep 17 00:00:00 2001 From: Rahul Garg Date: Tue, 16 Jan 2018 11:44:19 +0530 Subject: [PATCH 3/8] Added support for - - hipMemcpyFromArray - hipMemcpyAtoH - hipMemcpyHtoA [ROCm/hip commit: 487a430b5a3679ee3d0ee02e0ce6f22acd474ec0] --- .../include/hip/hcc_detail/hip_runtime_api.h | 49 +++++++++++++++ .../include/hip/nvcc_detail/hip_runtime_api.h | 12 ++++ projects/hip/src/hip_memory.cpp | 59 +++++++++++++++++++ 3 files changed, 120 insertions(+) diff --git a/projects/hip/include/hip/hcc_detail/hip_runtime_api.h b/projects/hip/include/hip/hcc_detail/hip_runtime_api.h index 9d0757f83a..7f159572d7 100644 --- a/projects/hip/include/hip/hcc_detail/hip_runtime_api.h +++ b/projects/hip/include/hip/hcc_detail/hip_runtime_api.h @@ -1420,7 +1420,56 @@ hipError_t hipMemcpy2DToArray(hipArray* dst, size_t wOffset, size_t hOffset, con hipError_t hipMemcpyToArray(hipArray* dst, size_t wOffset, size_t hOffset, const void* src, size_t count, hipMemcpyKind kind); +/** + * @brief Copies data between host and device. + * + * @param[in] dst Destination memory address + * @param[in] srcArray Source memory address + * @param[in] woffset Source starting X offset + * @param[in] hOffset Source starting Y offset + * @param[in] count Size in bytes to copy + * @param[in] kind Type of transfer + * @return #hipSuccess, #hipErrorInvalidValue, #hipErrorInvalidPitchValue, #hipErrorInvalidDevicePointer, #hipErrorInvalidMemcpyDirection + * + * @see hipMemcpy, hipMemcpy2DToArray, hipMemcpy2D, hipMemcpyFromArray, hipMemcpyToSymbol, hipMemcpyAsync + */ +hipError_t hipMemcpyFromArray(void* dst, hipArray_const_t srcArray, size_t wOffset, size_t hOffset, + size_t count, hipMemcpyKind kind); +/** + * @brief Copies data between host and device. + * + * @param[in] dst Destination memory address + * @param[in] srcArray Source array + * @param[in] srcoffset Offset in bytes of source array + * @param[in] count Size of memory copy in bytes + * @return #hipSuccess, #hipErrorInvalidValue, #hipErrorInvalidPitchValue, #hipErrorInvalidDevicePointer, #hipErrorInvalidMemcpyDirection + * + * @see hipMemcpy, hipMemcpy2DToArray, hipMemcpy2D, hipMemcpyFromArray, hipMemcpyToSymbol, hipMemcpyAsync + */ +hipError_t hipMemcpyAtoH(void* dst, hipArray* srcArray, size_t srcOffset, size_t count); + +/** + * @brief Copies data between host and device. + * + * @param[in] dstArray Destination memory address + * @param[in] dstOffset Offset in bytes of destination array + * @param[in] srcHost Source host pointer + * @param[in] count Size of memory copy in bytes + * @return #hipSuccess, #hipErrorInvalidValue, #hipErrorInvalidPitchValue, #hipErrorInvalidDevicePointer, #hipErrorInvalidMemcpyDirection + * + * @see hipMemcpy, hipMemcpy2DToArray, hipMemcpy2D, hipMemcpyFromArray, hipMemcpyToSymbol, hipMemcpyAsync + */ +hipError_t hipMemcpyHtoA(hipArray* dstArray, size_t dstOffset, const void* srcHost, size_t count); + +/** + * @brief Copies data between host and device. + * + * @param[in] p 3D memory copy parameters + * @return #hipSuccess, #hipErrorInvalidValue, #hipErrorInvalidPitchValue, #hipErrorInvalidDevicePointer, #hipErrorInvalidMemcpyDirection + * + * @see hipMemcpy, hipMemcpy2DToArray, hipMemcpy2D, hipMemcpyFromArray, hipMemcpyToSymbol, hipMemcpyAsync + */ hipError_t hipMemcpy3D(const struct hipMemcpy3DParms *p); // doxygen end Memory diff --git a/projects/hip/include/hip/nvcc_detail/hip_runtime_api.h b/projects/hip/include/hip/nvcc_detail/hip_runtime_api.h index dbd6d8b300..902e3620fa 100644 --- a/projects/hip/include/hip/nvcc_detail/hip_runtime_api.h +++ b/projects/hip/include/hip/nvcc_detail/hip_runtime_api.h @@ -483,6 +483,18 @@ inline static hipError_t hipMemcpyToArray(hipArray* dst, size_t wOffset, size_t return hipCUDAErrorTohipError(cudaMemcpyToArray(dst, wOffset, hOffset, src, count, hipMemcpyKindToCudaMemcpyKind(kind))); } +inline static hipError_t hipMemcpyFromArray(void* dst, hipArray_const_t srcArray, size_t wOffset, size_t hOffset, size_t count, hipMemcpyKind kind) { + return hipCUDAErrorTohipError(cudaMemcpyFromArray(dst, srcArray, wOffset, hOffset, count, hipMemcpyKindToCudaMemcpyKind(kind))); +} + +inline static hipError_t hipMemcpyAtoH(void* dst, hipArray* srcArray, size_t srcOffset, size_t count) { + return hipCUResultTohipError(cuMemcpyAtoH(dst, (CUarray)srcArray, srcOffset, count)); +} + +inline static hipError_t hipMemcpyHtoA(hipArray* dstArray, size_t dstOffset, const void* srcHost, size_t count) { + return hipCUResultTohipError(cuMemcpyHtoA((CUarray)dstArray, dstOffset, srcHost, count)); +} + inline static hipError_t hipDeviceSynchronize() { return hipCUDAErrorTohipError(cudaDeviceSynchronize()); } diff --git a/projects/hip/src/hip_memory.cpp b/projects/hip/src/hip_memory.cpp index 77526cf9ac..0fb13abeb6 100644 --- a/projects/hip/src/hip_memory.cpp +++ b/projects/hip/src/hip_memory.cpp @@ -1411,6 +1411,65 @@ hipError_t hipMemcpyToArray(hipArray* dst, size_t wOffset, size_t hOffset, return ihipLogStatus(e); } +hipError_t hipMemcpyFromArray(void* dst, hipArray_const_t srcArray, size_t wOffset, size_t hOffset, + size_t count, hipMemcpyKind kind) { + + HIP_INIT_SPECIAL_API((TRACE_MCMD), dst, srcArray, wOffset, hOffset, count, kind); + + hipStream_t stream = ihipSyncAndResolveStream(hipStreamNull); + + hc::completion_future marker; + + hipError_t e = hipSuccess; + + try { + stream->locked_copySync((char *)dst, (char*)srcArray->data + wOffset, count, kind); + } + catch (ihipException &ex) { + e = ex._code; + } + + return ihipLogStatus(e); +} + +hipError_t hipMemcpyHtoA(hipArray* dstArray, size_t dstOffset, const void* srcHost, size_t count) +{ + HIP_INIT_SPECIAL_API((TRACE_MCMD), dstArray, dstOffset, srcHost, count); + + hipStream_t stream = ihipSyncAndResolveStream(hipStreamNull); + + hc::completion_future marker; + + hipError_t e = hipSuccess; + try { + stream->locked_copySync((char *)dstArray->data + dstOffset, srcHost, count, hipMemcpyHostToDevice); + } catch (ihipException &ex) { + e = ex._code; + } + + return ihipLogStatus(e); +} + +hipError_t hipMemcpyAtoH(void* dst, hipArray* srcArray, size_t srcOffset, size_t count) +{ + HIP_INIT_SPECIAL_API((TRACE_MCMD), dst, srcArray, srcOffset, count); + + hipStream_t stream = ihipSyncAndResolveStream(hipStreamNull); + + hc::completion_future marker; + + hipError_t e = hipSuccess; + + try { + stream->locked_copySync((char *)dst, (char*)srcArray->data + srcOffset, count, hipMemcpyDeviceToHost); + } + catch (ihipException &ex) { + e = ex._code; + } + + return ihipLogStatus(e); +} + hipError_t hipMemcpy3D(const struct hipMemcpy3DParms *p) { HIP_INIT_SPECIAL_API((TRACE_MCMD), p); From e2ade308cf7a084fdd682cbda8826bd5d502facc Mon Sep 17 00:00:00 2001 From: Rahul Garg Date: Tue, 30 Jan 2018 18:06:31 +0530 Subject: [PATCH 4/8] Fixed host allocated globals address lookup for host usage Fixed texture driver APIs failure [ROCm/hip commit: 24ab820a11210ad9e4c0d29dee87e357b30ed4e2] --- projects/hip/include/hip/hcc_detail/program_state.hpp | 6 ++---- projects/hip/src/hip_memory.cpp | 1 + projects/hip/src/hip_module.cpp | 3 +-- projects/hip/src/program_state.cpp | 9 ++++----- 4 files changed, 8 insertions(+), 11 deletions(-) diff --git a/projects/hip/include/hip/hcc_detail/program_state.hpp b/projects/hip/include/hip/hcc_detail/program_state.hpp index 02e2f1e524..f7de214f10 100644 --- a/projects/hip/include/hip/hcc_detail/program_state.hpp +++ b/projects/hip/include/hip/hcc_detail/program_state.hpp @@ -69,18 +69,16 @@ namespace hip_impl } }; - using RAII_global = std::unique_ptr; - const std::unordered_map< hsa_agent_t, std::vector>& executables(); const std::unordered_map< std::uintptr_t, std::vector>>& functions(); const std::unordered_map& function_names(); - std::unordered_map& globals(); + std::unordered_map& globals(); hsa_executable_t load_executable( const std::string& file, hsa_executable_t executable, hsa_agent_t agent); -} // Namespace hip_impl. \ No newline at end of file +} // Namespace hip_impl. diff --git a/projects/hip/src/hip_memory.cpp b/projects/hip/src/hip_memory.cpp index 77526cf9ac..e1016f4af4 100644 --- a/projects/hip/src/hip_memory.cpp +++ b/projects/hip/src/hip_memory.cpp @@ -453,6 +453,7 @@ hipError_t hipArrayCreate ( hipArray** array, const HIP_ARRAY_DESCRIPTOR* pAlloc array[0]->width = pAllocateArray->width; array[0]->height = pAllocateArray->height; array[0]->isDrv = true; + array[0]->textureType = hipTextureType2D; void ** ptr = &array[0]->data; if (ctx) { const unsigned am_flags = 0; diff --git a/projects/hip/src/hip_module.cpp b/projects/hip/src/hip_module.cpp index 45a44b3666..d173a2f295 100644 --- a/projects/hip/src/hip_module.cpp +++ b/projects/hip/src/hip_module.cpp @@ -568,7 +568,6 @@ hipError_t hipModuleGetTexRef( const auto it = globals().find(name); if (it == globals().end()) return ihipLogStatus(hipErrorInvalidValue); - *texRef = static_cast(it->second.get()); - + *texRef = reinterpret_cast(it->second); return ihipLogStatus(hipSuccess); } diff --git a/projects/hip/src/program_state.cpp b/projects/hip/src/program_state.cpp index e867887da2..35785dcad5 100644 --- a/projects/hip/src/program_state.cpp +++ b/projects/hip/src/program_state.cpp @@ -169,7 +169,7 @@ namespace lock_guard lck{mtx}; if (globals().find(x) != globals().cend()) return; - + globals().emplace(x, (void*)(it1->second.first)); void* p = nullptr; hsa_amd_memory_lock( reinterpret_cast(it1->second.first), @@ -181,7 +181,6 @@ namespace hsa_executable_agent_global_variable_define( executable, agent, x.c_str(), p); - globals().emplace(x, RAII_global{p, hsa_amd_memory_unlock}); } } @@ -462,9 +461,9 @@ namespace hip_impl return r; } - unordered_map& globals() + unordered_map& globals() { - static unordered_map r; + static unordered_map r; static once_flag f; call_once(f, []() { r.reserve(symbol_addresses().size()); }); @@ -491,4 +490,4 @@ namespace hip_impl return executable; } -} // Namespace hip_impl. \ No newline at end of file +} // Namespace hip_impl. From 88e41efc44b8f507b6f688a3f8d2845414474411 Mon Sep 17 00:00:00 2001 From: Rahul Garg Date: Tue, 6 Feb 2018 14:29:04 +0530 Subject: [PATCH 5/8] Update the programming guide with environemnt variables names and default threshold values used. [ROCm/hip commit: b1ce616a024eb4ff8689cfc27e13bab5062484d2] --- .../docs/markdown/hip_programming_guide.md | 23 +++++++++++++------ 1 file changed, 16 insertions(+), 7 deletions(-) diff --git a/projects/hip/docs/markdown/hip_programming_guide.md b/projects/hip/docs/markdown/hip_programming_guide.md index 5d0c1f2497..9313eb22e1 100644 --- a/projects/hip/docs/markdown/hip_programming_guide.md +++ b/projects/hip/docs/markdown/hip_programming_guide.md @@ -54,7 +54,7 @@ A stronger system-level fence can be specified when the event is created with hi - HIP/ROCm also supports the ability to cache host memory in the GPU using the "Non-Coherent" host memory allocations. This can provide performance benefit, but care must be taken to use the correct synchronization. -## Unpinned Memory Transfer Optimizations +## Unpinned Memory Transfer Optimization Please note that this document lists possible ways for experimenting with HIP stack to gain performance. Performance may vary from platform to platform. ### On Small BAR Setup @@ -79,11 +79,20 @@ stage the copy through an optimized pinned staging buffer, to implement H2D and PinInPlace is another algorithm which pins the host memory "in-place", and copies it with the DMA engine. -By default staging buffers are used for unpinned memory transfers. Environment variables allow control over the unpinned copy algorithm and parameters: +Unpinned memory transfer mode can be controlled using environment variable HCC_UNPINNED_COPY_MODE. -- HIP_PININPLACE - This environment variable forces the use of PinInPlace logic for all unpinned memory copies +By default HCC_UNPINNED_COPY_MODE is set to 0, which uses default threshold values to decide which transfer way to use based on data size. -- HIP_OPTIMAL_MEM_TRANSFER- This environment variable enables a hybrid memory copy logic based on thresholds. These thresholds can be managed with following environment variables: - - HIP_H2D_MEM_TRANSFER_THRESHOLD_STAGING_OR_PININPLACE - Threshold in bytes for H2D copy. For sizes smaller than threshold staging buffers logic would be used else PinInPlace logic. - - HIP_H2D_MEM_TRANSFER_THRESHOLD_DIRECT_OR_STAGING - Threshold in bytes for H2D copy. For sizes smaller than threshold direct copy logic would be used else staging buffers logic. - - HIP_D2H_MEM_TRANSFER_THRESHOLD - Threshold in bytes for D2H copy. For sizes smaller than threshold staging buffer logic would be used else PinInPlace logic. +Setting HCC_UNPINNED_COPY_MODE = 1, forces all unpinned transfer to use PinInPlace logic. + +Setting HCC_UNPINNED_COPY_MODE = 2, forces all unpinned transfer to use Staging buffers. + +Setting HCC_UNPINNED_COPY_MODE = 3, forces all unpinned transfer to use direct memcpy on large BAR systems. + +Following environment variables can be used to control the transfer thresholds: + +- HCC_H2D_STAGING_THRESHOLD - Threshold in KB for H2D copy. For sizes smaller than threshold direct copy logic would be used else staging buffers logic. By default it is set to 64. + +- HCC_H2D_PININPLACE_THRESHOLD - Threshold in KB for H2D copy. For sizes smaller than threshold staging buffers logic would be used else PinInPlace logic. By default it is set to 4096. + +- HCC_D2H_PININPLACE_THRESHOLD - Threshold in KB for D2H copy. For sizes smaller than threshold staging buffer logic would be used else PinInPlace logic. By default it is set to 1024. From 00aeb8514f5385f5415297b618cddc4120c94161 Mon Sep 17 00:00:00 2001 From: Maneesh Gupta Date: Mon, 12 Feb 2018 10:21:27 +0530 Subject: [PATCH 6/8] [build] Add clangformat target to cmake - Added a new target "clangformat" which formats the code using the clang-format tool found in HCC package. - Renamed target "static_check" to "cppcheck". - Use find_program() macro to make the build system more robust. Change-Id: Ifefbf36f23ff7ef27a870120d5b9170d0cc8aa52 [ROCm/hip commit: 1e559d19c39b33bfaab5a0c4f28535d25b238253] --- projects/hip/.clang-format | 20 ++++++++++++++++++++ projects/hip/CMakeLists.txt | 26 ++++++++++++++++++++++---- 2 files changed, 42 insertions(+), 4 deletions(-) create mode 100644 projects/hip/.clang-format diff --git a/projects/hip/.clang-format b/projects/hip/.clang-format new file mode 100644 index 0000000000..1793af2ba2 --- /dev/null +++ b/projects/hip/.clang-format @@ -0,0 +1,20 @@ +--- +Language: Cpp +BasedOnStyle: Google +AlignEscapedNewlinesLeft: false +ColumnLimit: 100 +DerivePointerAlignment: false +IndentWrappedFunctionNames: false +MaxEmptyLinesToKeep: 2 +SortIncludes: false +IndentWidth: 4 +--- +Language: ObjC +BasedOnStyle: Google +AlignEscapedNewlinesLeft: false +ColumnLimit: 100 +DerivePointerAlignment: false +IndentWrappedFunctionNames: false +MaxEmptyLinesToKeep: 2 +SortIncludes: false +IndentWidth: 4 diff --git a/projects/hip/CMakeLists.txt b/projects/hip/CMakeLists.txt index baf3b49df1..1573ddee5c 100644 --- a/projects/hip/CMakeLists.txt +++ b/projects/hip/CMakeLists.txt @@ -220,8 +220,11 @@ endif() file(WRITE "${PROJECT_BINARY_DIR}/.hipVersion" ${_versionInfo}) # Build doxygen documentation -add_custom_target(doc COMMAND HIP_PATH=${CMAKE_CURRENT_SOURCE_DIR} doxygen ${CMAKE_CURRENT_SOURCE_DIR}/docs/doxygen-input/doxy.cfg - WORKING_DIRECTORY ${CMAKE_CURRENT_SOURCE_DIR}/docs) +find_program(DOXYGEN_EXE doxygen) +if(DOXYGEN_EXE) + add_custom_target(doc COMMAND HIP_PATH=${CMAKE_CURRENT_SOURCE_DIR} ${DOXYGEN_EXE} ${CMAKE_CURRENT_SOURCE_DIR}/docs/doxygen-input/doxy.cfg + WORKING_DIRECTORY ${CMAKE_CURRENT_SOURCE_DIR}/docs) +endif() ############################# # Install steps @@ -364,8 +367,23 @@ endif() ############################# # Code analysis ############################# -# Target: static_check -add_custom_target(static_check COMMAND cppcheck --force --quiet --enable=warning,performance,portability,information,missingInclude src include -I /opt/rocm/include/hcc -I /opt/rocm/include --suppress=*:/opt/rocm/include/hcc/hc.hpp WORKING_DIRECTORY ${CMAKE_CURRENT_SOURCE_DIR}) +# Target: cppcheck +find_program(CPPCHECK_EXE cppcheck) +if(CPPCHECK_EXE) + add_custom_target(cppcheck COMMAND ${CPPCHECK_EXE} --force --quiet --enable=warning,performance,portability,information,missingInclude src include -I /opt/rocm/include/hcc -I /opt/rocm/include --suppress=*:/opt/rocm/include/hcc/hc.hpp + WORKING_DIRECTORY ${CMAKE_CURRENT_SOURCE_DIR}) +endif() + +############################# +# Code formatting +############################# +# Target: clangformat +find_program(CLANGFORMAT_EXE clang-format PATHS ${HCC_HOME}/bin) +if(CLANGFORMAT_EXE) + file(GLOB_RECURSE FORMAT_SOURCE_FILE_LIST *.cpp *.hpp *.h) + add_custom_target(clangformat COMMAND ${CLANGFORMAT_EXE} -style=file -i ${FORMAT_SOURCE_FILE_LIST} + WORKING_DIRECTORY ${CMAKE_CURRENT_SOURCE_DIR}) +endif() ############################# # Testing steps From cb9b6d3613c66935cbecf9dbfdc4dc6426f70c80 Mon Sep 17 00:00:00 2001 From: Maneesh Gupta Date: Wed, 14 Feb 2018 09:13:50 +0530 Subject: [PATCH 7/8] [dtests] Fix HIT block in hipPeerToPeer_simple.cpp [ROCm/hip commit: 2ff55124fa64c19e42c1cce3755dd762e787fdd2] --- projects/hip/tests/src/p2p/hipPeerToPeer_simple.cpp | 8 ++++---- 1 file changed, 4 insertions(+), 4 deletions(-) diff --git a/projects/hip/tests/src/p2p/hipPeerToPeer_simple.cpp b/projects/hip/tests/src/p2p/hipPeerToPeer_simple.cpp index 74fefa8fdb..9181ebdb59 100644 --- a/projects/hip/tests/src/p2p/hipPeerToPeer_simple.cpp +++ b/projects/hip/tests/src/p2p/hipPeerToPeer_simple.cpp @@ -23,10 +23,10 @@ THE SOFTWARE. // Also serves as a template for other tests. /* HIT_START - * BUILD: %t %s ../test_common.cpp - * RUN: %t EXCLUDE_HIP_PLATFORM hipcc - * RUN: %t --memcpyWithPeer EXCLUDE_HIP_PLATFORM hipcc - * RUN: %t --mirrorPeers EXCLUDE_HIP_PLATFORM hipcc + * BUILD: %t %s ../test_common.cpp + * RUN: %t EXCLUDE_HIP_PLATFORM hcc + * RUN: %t --memcpyWithPeer EXCLUDE_HIP_PLATFORM hcc + * RUN: %t --mirrorPeers EXCLUDE_HIP_PLATFORM hcc * HIT_END */ From afc71e3bada78343a6e12f16f9ad587a269d573f Mon Sep 17 00:00:00 2001 From: Siu Chi Chan Date: Sun, 18 Feb 2018 14:19:21 -0500 Subject: [PATCH 8/8] make HIP to load the GPU objects and to setup the function symbol map on startup [ROCm/hip commit: ebc92e597553360fe6389be64de99370713bef10] --- projects/hip/src/program_state.cpp | 12 ++++++++++++ 1 file changed, 12 insertions(+) diff --git a/projects/hip/src/program_state.cpp b/projects/hip/src/program_state.cpp index 35785dcad5..aef92d463c 100644 --- a/projects/hip/src/program_state.cpp +++ b/projects/hip/src/program_state.cpp @@ -490,4 +490,16 @@ namespace hip_impl return executable; } + + // To force HIP to load the kernels and to setup the function + // symbol map on program startup + class startup_kernel_loader { + private: + startup_kernel_loader() { functions(); } + startup_kernel_loader(const startup_kernel_loader&) = delete; + startup_kernel_loader& operator= (const startup_kernel_loader&) = delete; + static startup_kernel_loader skl; + }; + startup_kernel_loader startup_kernel_loader::skl; + } // Namespace hip_impl.