From 64b531e782d309efda100c74d0144984ee386454 Mon Sep 17 00:00:00 2001 From: Alexey Chernov <4ernov@gmail.com> Date: Fri, 24 Jan 2020 13:51:49 +0300 Subject: [PATCH 1/5] Clear `HIP_PATH` before version detection (#1786) Don't allow `HIP_PATH` to be propagated to `hipconfig`, when run by CMake to detect the package version, as it leads to the wrong version is detected: when there's already HIP of some different version installed in the system and `HIP_PATH` points to its location, `hipconfig` tends to return the version of the installed HIP, rather than the value defined for the distribution. The compiled results report wrong version and spoils the rest of the stack in this case. [ROCm/clr commit: 9f5a4148ce90ea38e8fd43ab9ae6bf1871d0c85c] --- projects/clr/hipamd/CMakeLists.txt | 1 + 1 file changed, 1 insertion(+) diff --git a/projects/clr/hipamd/CMakeLists.txt b/projects/clr/hipamd/CMakeLists.txt index 9a1d940421..3eedf35b09 100644 --- a/projects/clr/hipamd/CMakeLists.txt +++ b/projects/clr/hipamd/CMakeLists.txt @@ -20,6 +20,7 @@ endmacro() # Setup version information ############################# # Determine HIP_BASE_VERSION +set(ENV{HIP_PATH} "") execute_process(COMMAND ${CMAKE_CURRENT_SOURCE_DIR}/bin/hipconfig --version OUTPUT_VARIABLE HIP_BASE_VERSION OUTPUT_STRIP_TRAILING_WHITESPACE) From adb93af8c30bfc20a9139b4c212c5ba0d3501c99 Mon Sep 17 00:00:00 2001 From: Siu Chi Chan Date: Fri, 24 Jan 2020 05:52:49 -0500 Subject: [PATCH 2/5] Fix associate code object symbols with host allocation bug (#1799) The current implementation skips this procedure for a given device object when a global symbol is found in the cache. This is incorrect: - There could be other undefined globals that have not been previously encountered further down the list - If a symbol is found in the cache, it doesn't need to be pinned again but it still need to be defined for the current executable Added special case for the printf buffer symbol (already pinned by HCC) The bug was exposed by running printf on different GPUs. [ROCm/clr commit: 8fc7cad90f948bdeffe02cb743ffb55e26d6ba39] --- projects/clr/hipamd/src/program_state.cpp | 2 +- projects/clr/hipamd/src/program_state.inl | 60 +++++++++++++------ .../tests/src/kernel/hipPrintfKernel.cpp | 9 ++- 3 files changed, 50 insertions(+), 21 deletions(-) diff --git a/projects/clr/hipamd/src/program_state.cpp b/projects/clr/hipamd/src/program_state.cpp index dbd7d3ebc4..5e9f9976be 100644 --- a/projects/clr/hipamd/src/program_state.cpp +++ b/projects/clr/hipamd/src/program_state.cpp @@ -61,7 +61,7 @@ namespace hip_impl { if (it == impl->get_globals().end()) return nullptr; else - return it->second; + return it->second.first; } hsa_executable_t program_state::load_executable(const char* data, diff --git a/projects/clr/hipamd/src/program_state.inl b/projects/clr/hipamd/src/program_state.inl index 993418de96..8861558e04 100644 --- a/projects/clr/hipamd/src/program_state.inl +++ b/projects/clr/hipamd/src/program_state.inl @@ -18,6 +18,7 @@ #include #include #include +#include "hc.hpp" #include @@ -193,7 +194,8 @@ public: std::tuple< std::once_flag, std::mutex, - std::unordered_map> globals; + // map from string to pair + std::unordered_map>> globals; using RAII_code_reader = std::unique_ptr& get_globals() { + std::unordered_map>& get_globals() { std::call_once(std::get<0>(globals), [this]() { std::get<2>(globals).reserve(get_symbol_addresses().size()); }); @@ -349,30 +351,52 @@ public: auto& g_mutex = get_globals_mutex(); for (auto&& x : undefined_symbols) { - if (g.find(x) != g.cend()) return; - const auto it1 = get_symbol_addresses().find(x); - if (it1 == get_symbol_addresses().cend()) { hip_throw(std::runtime_error{ "Global symbol: " + x + " is undefined."}); } - std::lock_guard lck{g_mutex}; + hsa_status_t status; + auto check_hsa_global_var_define_error = [&x](hsa_status_t s) { + if (s != HSA_STATUS_SUCCESS) { + const char* es; + hsa_status_string(s, &es); + hip_throw(std::runtime_error{ "Error when defining symbol " + x + " : " + es}); + } + }; - if (g.find(x) != g.cend()) return; + auto retrieve_pinned_address_from_cache = [](decltype(g) g, decltype(x) x) { + const auto& global_addr = g.find(x); + if (global_addr != g.cend()) { + return global_addr->second.second; + } + return (void*)nullptr; + }; - g.emplace(x, (void*)(it1->second.first)); - void* p = nullptr; - hsa_amd_memory_lock( - reinterpret_cast(it1->second.first), - it1->second.second, - nullptr, // All agents. - 0, - &p); - - hsa_executable_agent_global_variable_define( - executable, agent, x.c_str(), p); + void* p = retrieve_pinned_address_from_cache(g, x); + if (p == nullptr) { + std::lock_guard lck{g_mutex}; + p = retrieve_pinned_address_from_cache(g, x); + if (p == nullptr) { + if (x == "_ZN2hc13printf_bufferE") { + // This is the printf buffer, get the pinned address from HCC + p = Kalmar::getContext()->getPrintfBufferPointerVA(); + } + else { + status = hsa_amd_memory_lock(reinterpret_cast(it1->second.first), + it1->second.second, + nullptr, // All agents. + 0, &p); + check_hsa_global_var_define_error(status); + } + // cache the global address and its pinned address + g.emplace(x, std::make_pair(reinterpret_cast(it1->second.first), p)); + } + } + status = hsa_executable_agent_global_variable_define( + executable, agent, x.c_str(), p); + check_hsa_global_var_define_error(status); } } diff --git a/projects/clr/hipamd/tests/src/kernel/hipPrintfKernel.cpp b/projects/clr/hipamd/tests/src/kernel/hipPrintfKernel.cpp index 1d4fa5fe30..5675f2e6bd 100644 --- a/projects/clr/hipamd/tests/src/kernel/hipPrintfKernel.cpp +++ b/projects/clr/hipamd/tests/src/kernel/hipPrintfKernel.cpp @@ -30,7 +30,12 @@ THE SOFTWARE. __global__ void run_printf() { printf("Hello World\n"); } int main() { - hipLaunchKernelGGL(HIP_KERNEL_NAME(run_printf), dim3(1), dim3(1), 0, 0); - hipDeviceSynchronize(); + int device_count = 0; + hipGetDeviceCount(&device_count); + for (int i = 0; i < device_count; ++i) { + hipSetDevice(i); + hipLaunchKernelGGL(HIP_KERNEL_NAME(run_printf), dim3(1), dim3(1), 0, 0); + hipDeviceSynchronize(); + } passed(); } From fec5105d9cbfdf3770e504dd87567cc48918db80 Mon Sep 17 00:00:00 2001 From: paulfreddy <52053501+paulfreddy@users.noreply.github.com> Date: Fri, 24 Jan 2020 02:52:57 -0800 Subject: [PATCH 3/5] Fix install script syntax error (#1805) Fix hip-nvcc install warning on postinstall and prerm script [ROCm/clr commit: 5cd5c62f298cd1923bf86430f488cce73f7aba4d] --- projects/clr/hipamd/packaging/hip-nvcc.postinst | 2 +- projects/clr/hipamd/packaging/hip-nvcc.prerm | 2 +- 2 files changed, 2 insertions(+), 2 deletions(-) diff --git a/projects/clr/hipamd/packaging/hip-nvcc.postinst b/projects/clr/hipamd/packaging/hip-nvcc.postinst index 2f901324cb..b70cf2848f 100755 --- a/projects/clr/hipamd/packaging/hip-nvcc.postinst +++ b/projects/clr/hipamd/packaging/hip-nvcc.postinst @@ -3,6 +3,6 @@ ROCMDIR=@ROCM_PATH@ HIPDIR=$ROCMDIR/hip -if [ -d $ROCMDIR] +if [ -d $ROCMDIR] ; then ln -s -f $ROCMDIR /opt/rocm fi diff --git a/projects/clr/hipamd/packaging/hip-nvcc.prerm b/projects/clr/hipamd/packaging/hip-nvcc.prerm index 96875e4a9c..baa0e6f5c7 100755 --- a/projects/clr/hipamd/packaging/hip-nvcc.prerm +++ b/projects/clr/hipamd/packaging/hip-nvcc.prerm @@ -1,5 +1,5 @@ #!/bin/bash -if [ -L "/opt/rocm" ] +if [ -L "/opt/rocm" ] ; then unlink /opt/rocm fi From a49cbb1580c21d4028eb76003af3b5779be8e869 Mon Sep 17 00:00:00 2001 From: mshivama <47909405+mshivama@users.noreply.github.com> Date: Fri, 24 Jan 2020 16:23:28 +0530 Subject: [PATCH 4/5] SWDEV-220503: this_grid().thread_rank() gives incorrect result (#1808) * fix a minor bug while computing this.grid()::thread_rank() [ROCm/clr commit: 0605426049798cf42f124a6bde1f8699700c75d3] --- .../include/hip/hcc_detail/hip_cooperative_groups_helper.h | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/projects/clr/hipamd/include/hip/hcc_detail/hip_cooperative_groups_helper.h b/projects/clr/hipamd/include/hip/hcc_detail/hip_cooperative_groups_helper.h index b74d16d23b..9738448d94 100644 --- a/projects/clr/hipamd/include/hip/hcc_detail/hip_cooperative_groups_helper.h +++ b/projects/clr/hipamd/include/hip/hcc_detail/hip_cooperative_groups_helper.h @@ -106,7 +106,7 @@ __CG_STATIC_QUALIFIER__ uint32_t size() { } __CG_STATIC_QUALIFIER__ uint32_t thread_rank() { - // Compute global id of the workgroup to which the current threads belongs to + // Compute global id of the workgroup to which the current thread belongs to uint32_t blkIdx = (uint32_t)((hipBlockIdx_z * hipGridDim_y * hipGridDim_x) + (hipBlockIdx_y * hipGridDim_x) + @@ -115,7 +115,7 @@ __CG_STATIC_QUALIFIER__ uint32_t thread_rank() { // Compute total number of threads being passed to reach current workgroup // within grid uint32_t num_threads_till_current_workgroup = - (uint32_t)(blkIdx * (hipBlockIdx_x * hipBlockIdx_y * hipBlockIdx_z)); + (uint32_t)(blkIdx * (hipBlockDim_x * hipBlockDim_y * hipBlockDim_z)); // Compute thread local rank within current workgroup uint32_t local_thread_rank = From 6fafadae75841004059f1e911cba6ee2024e60a4 Mon Sep 17 00:00:00 2001 From: aakanksha555 <41199349+aakanksha555@users.noreply.github.com> Date: Fri, 24 Jan 2020 05:53:47 -0500 Subject: [PATCH 5/5] Fix for a syntax error in deb packages (#1814) [ROCm/clr commit: ef3dbc1f91aa3107a0561f053e38ad4316aa12f7] --- projects/clr/hipamd/hipify-clang/packaging/hipify-clang.txt | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/projects/clr/hipamd/hipify-clang/packaging/hipify-clang.txt b/projects/clr/hipamd/hipify-clang/packaging/hipify-clang.txt index 5f78e7e67e..b189eff1e6 100644 --- a/projects/clr/hipamd/hipify-clang/packaging/hipify-clang.txt +++ b/projects/clr/hipamd/hipify-clang/packaging/hipify-clang.txt @@ -48,7 +48,7 @@ endif() set(CPACK_PACKAGE_FILE_NAME ${CPACK_PACKAGE_NAME}-${CPACK_PACKAGE_VERSION_MAJOR}.${CPACK_PACKAGE_VERSION_MINOR}.${CPACK_PACKAGE_VERSION_PATCH}) set(CPACK_GENERATOR "TGZ;DEB;RPM") set(CPACK_BINARY_DEB "ON") -set(CPACK_DEBIAN_PACKAGE_DEPENDS "cuda >= 7.0") +set(CPACK_DEBIAN_PACKAGE_DEPENDS "cuda (>= 7.0)") set(CPACK_BINARY_RPM "ON") set(CPACK_RPM_PACKAGE_ARCHITECTURE "${CMAKE_SYSTEM_PROCESSOR}") set(CPACK_RPM_PACKAGE_AUTOREQPROV "NO")