From a156207c28dff1c706355a9e1eb7c51d64c58aff 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/clr commit: 451c7b37cc8d7d2528346975479ed1937330bfd1] --- .../clr/hipamd/tests/src/p2p/hipPeerToPeer_simple.cpp | 8 ++++---- 1 file changed, 4 insertions(+), 4 deletions(-) diff --git a/projects/clr/hipamd/tests/src/p2p/hipPeerToPeer_simple.cpp b/projects/clr/hipamd/tests/src/p2p/hipPeerToPeer_simple.cpp index 32cab371b8..74fefa8fdb 100644 --- a/projects/clr/hipamd/tests/src/p2p/hipPeerToPeer_simple.cpp +++ b/projects/clr/hipamd/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 215b31eba94e2b0321bb74b205130201d80c48c9 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/clr commit: df73989d693d2c19e84fba37c60bd0a50cf788a7] --- .../tests/src/runtimeApi/device/hipGetDeviceAttribute.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/projects/clr/hipamd/tests/src/runtimeApi/device/hipGetDeviceAttribute.cpp b/projects/clr/hipamd/tests/src/runtimeApi/device/hipGetDeviceAttribute.cpp index 2919939694..0b965a3ae3 100644 --- a/projects/clr/hipamd/tests/src/runtimeApi/device/hipGetDeviceAttribute.cpp +++ b/projects/clr/hipamd/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 ec4cc279867eb99649627cdc3551dcfa8c612382 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/clr commit: ca5bcb5af471212a525ec45683d3e978ec628670] --- .../include/hip/hcc_detail/hip_runtime_api.h | 49 +++++++++++++++ .../include/hip/nvcc_detail/hip_runtime_api.h | 12 ++++ projects/clr/hipamd/src/hip_memory.cpp | 59 +++++++++++++++++++ 3 files changed, 120 insertions(+) diff --git a/projects/clr/hipamd/include/hip/hcc_detail/hip_runtime_api.h b/projects/clr/hipamd/include/hip/hcc_detail/hip_runtime_api.h index 9d0757f83a..7f159572d7 100644 --- a/projects/clr/hipamd/include/hip/hcc_detail/hip_runtime_api.h +++ b/projects/clr/hipamd/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/clr/hipamd/include/hip/nvcc_detail/hip_runtime_api.h b/projects/clr/hipamd/include/hip/nvcc_detail/hip_runtime_api.h index dbd6d8b300..902e3620fa 100644 --- a/projects/clr/hipamd/include/hip/nvcc_detail/hip_runtime_api.h +++ b/projects/clr/hipamd/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/clr/hipamd/src/hip_memory.cpp b/projects/clr/hipamd/src/hip_memory.cpp index 77526cf9ac..0fb13abeb6 100644 --- a/projects/clr/hipamd/src/hip_memory.cpp +++ b/projects/clr/hipamd/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 f15543d5cf504be948edfd6329eb3ac35bb37866 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/clr commit: b8c23f979b088c535c83a001a02fb9efe1ee435b] --- .../clr/hipamd/include/hip/hcc_detail/program_state.hpp | 6 ++---- projects/clr/hipamd/src/hip_memory.cpp | 1 + projects/clr/hipamd/src/hip_module.cpp | 3 +-- projects/clr/hipamd/src/program_state.cpp | 9 ++++----- 4 files changed, 8 insertions(+), 11 deletions(-) diff --git a/projects/clr/hipamd/include/hip/hcc_detail/program_state.hpp b/projects/clr/hipamd/include/hip/hcc_detail/program_state.hpp index 02e2f1e524..f7de214f10 100644 --- a/projects/clr/hipamd/include/hip/hcc_detail/program_state.hpp +++ b/projects/clr/hipamd/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/clr/hipamd/src/hip_memory.cpp b/projects/clr/hipamd/src/hip_memory.cpp index 77526cf9ac..e1016f4af4 100644 --- a/projects/clr/hipamd/src/hip_memory.cpp +++ b/projects/clr/hipamd/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/clr/hipamd/src/hip_module.cpp b/projects/clr/hipamd/src/hip_module.cpp index 45a44b3666..d173a2f295 100644 --- a/projects/clr/hipamd/src/hip_module.cpp +++ b/projects/clr/hipamd/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/clr/hipamd/src/program_state.cpp b/projects/clr/hipamd/src/program_state.cpp index e867887da2..35785dcad5 100644 --- a/projects/clr/hipamd/src/program_state.cpp +++ b/projects/clr/hipamd/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 b347b33786604bc82f7d838a2309f888c5937e30 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/clr commit: c47ae45d0a7b9b09dcdc85704a75af69d27c931a] --- .../docs/markdown/hip_programming_guide.md | 23 +++++++++++++------ 1 file changed, 16 insertions(+), 7 deletions(-) diff --git a/projects/clr/hipamd/docs/markdown/hip_programming_guide.md b/projects/clr/hipamd/docs/markdown/hip_programming_guide.md index 5d0c1f2497..9313eb22e1 100644 --- a/projects/clr/hipamd/docs/markdown/hip_programming_guide.md +++ b/projects/clr/hipamd/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 f06408fa54db56eae72cb5b329a645b3f886b4d5 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/clr commit: 6ef0c494d1b31305da87d63995cfdc0166c5695b] --- projects/clr/hipamd/.clang-format | 20 ++++++++++++++++++++ projects/clr/hipamd/CMakeLists.txt | 26 ++++++++++++++++++++++---- 2 files changed, 42 insertions(+), 4 deletions(-) create mode 100644 projects/clr/hipamd/.clang-format diff --git a/projects/clr/hipamd/.clang-format b/projects/clr/hipamd/.clang-format new file mode 100644 index 0000000000..1793af2ba2 --- /dev/null +++ b/projects/clr/hipamd/.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/clr/hipamd/CMakeLists.txt b/projects/clr/hipamd/CMakeLists.txt index baf3b49df1..1573ddee5c 100644 --- a/projects/clr/hipamd/CMakeLists.txt +++ b/projects/clr/hipamd/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 e24c6ff5f9fe73204e722ac76ba6b7a8bd85a544 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/clr commit: cd71f55fd7c36e325cc4d5e2842d6c5895473aba] --- .../clr/hipamd/tests/src/p2p/hipPeerToPeer_simple.cpp | 8 ++++---- 1 file changed, 4 insertions(+), 4 deletions(-) diff --git a/projects/clr/hipamd/tests/src/p2p/hipPeerToPeer_simple.cpp b/projects/clr/hipamd/tests/src/p2p/hipPeerToPeer_simple.cpp index 74fefa8fdb..9181ebdb59 100644 --- a/projects/clr/hipamd/tests/src/p2p/hipPeerToPeer_simple.cpp +++ b/projects/clr/hipamd/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 9fddf100b732684c5a818802ee37287ceb640c59 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/clr commit: 0ca9591e556adab34bb65a1c1ad6783c31612587] --- projects/clr/hipamd/src/program_state.cpp | 12 ++++++++++++ 1 file changed, 12 insertions(+) diff --git a/projects/clr/hipamd/src/program_state.cpp b/projects/clr/hipamd/src/program_state.cpp index 35785dcad5..aef92d463c 100644 --- a/projects/clr/hipamd/src/program_state.cpp +++ b/projects/clr/hipamd/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.