From c0682dda969de0549e85ad11324984117dc707ba Mon Sep 17 00:00:00 2001 From: Yaxun Sam Liu Date: Sat, 30 Mar 2019 08:24:49 -0400 Subject: [PATCH 1/5] hip-clang: update installation guide. [ROCm/hip commit: 409aa49218d7ba2b7fecd89aef3c5568882ee35a] --- projects/hip/INSTALL.md | 20 ++++++++++---------- 1 file changed, 10 insertions(+), 10 deletions(-) diff --git a/projects/hip/INSTALL.md b/projects/hip/INSTALL.md index 7b12bd42ef..f8f67ac71f 100644 --- a/projects/hip/INSTALL.md +++ b/projects/hip/INSTALL.md @@ -36,37 +36,34 @@ apt-get install hip_hcc * By default HIP is installed into /opt/rocm/hip (can be overridden by setting HIP_PATH environment variable). * Optionally, consider adding /opt/rocm/bin to your PATH to make it easier to use the tools. -## AMD-clang +## HIP-clang * Using clang to compile HIP program for AMD GPU is under development. Users need to build LLVM, clang, lld, ROCm device library, and HIP from source. * Install the [rocm](http://gpuopen.com/getting-started-with-boltzmann-components-platforms-installation/) packages. ROCm will install some of the necessary components, including the kernel driver, HSA runtime, etc. -* Build LLVM/clang/lld by using the following repository and branch and following the general LLVM/clang build procedure: +* Build LLVM/clang/lld by using the following repository and branch and following the general LLVM/clang build procedure. It is recommended to use -DCMAKE_INSTALL_PREFIX=/opt/rocm/llvm with cmake so that LLVM/clang/lld are installed to the default path expected by hipcc. * LLVM: https://github.com/RadeonOpenCompute/llvm.git amd-common branch - * clang: https://github.com/RadeonOpenCompute/clang checkout amd-hip-upstream branch, then merge with amd-common branch to match LLVM/lld + * clang: https://github.com/RadeonOpenCompute/clang amd-common branch * lld: https://github.com/RadeonOpenCompute/lld amd-common branch * Build Rocm device library - * Checkout https://github.com/RadeonOpenCompute/ROCm-Device-Libs.git amd-hip branch and build it with clang built from the last step. + * Checkout https://github.com/RadeonOpenCompute/ROCm-Device-Libs.git master branch and build it with clang built from the last step. * Build HIP - * Checkout https://github.com/ROCm-Developer-Tools/HIP.git hip-clang branch and build it with HCC installed with ROCm packages. + * Checkout https://github.com/ROCm-Developer-Tools/HIP.git master branch and build it with HCC installed with ROCm packages. Please use -DHIP_COMPILER=clang with cmake to enable hip-clang. * Environment variables to let hipcc to use clang to compile HIP program - By default hipcc uses hcc to compile HIP program for AMD GPU. To let hipcc to use clang to compile HIP program, the following environment variables must be set: - - * HIP_CLANG_PATH - Path to clang - * DEVICE_LIB_PATH - Path to the device library - * Default paths and environment variables: * By default HIP looks for HSA in /opt/rocm/hsa (can be overridden by setting HSA_PATH environment variable) * By default HIP is installed into /opt/rocm/hip (can be overridden by setting HIP_PATH environment variable). + * By default HIP looks for clang in /opt/rocm/llvm/bin (can be overridden by setting HIP_CLANG_PATH environment variable) + * By default HIP looks for device library in /opt/rocm/lib (can be overriden by setting DEVICE_LIB_PATH environment variable). * Optionally, consider adding /opt/rocm/bin to your PATH to make it easier to use the tools. * Optionally, set HIPCC_VERBOSE=7 to output the command line for compilation to make sure clang is used instead of hcc. @@ -99,6 +96,9 @@ HIP source code is available and the project can be built from source on the HCC 2. Download HIP source code (from the [GitHub repot](https://github.com/ROCm-Developer-Tools/HIP).) 3. Install HIP build-time dependencies using ```sudo apt-get install libelf-dev```. 4. Build and install HIP (This is the simple version assuming default paths ; see below for additional options.) + +By default, HIP uses HCC to compile programs. To use HIP-Clang, add -DHIP_COMPILER=clang to cmake command line. + ``` cd HIP mkdir build From 595d95cfe8d8a7de6d0641af3bc52df004668af9 Mon Sep 17 00:00:00 2001 From: "Yaxun (Sam) Liu" Date: Sat, 30 Mar 2019 08:29:08 -0400 Subject: [PATCH 2/5] Update INSTALL.md [ROCm/hip commit: ed2fbb27e4fed42d43d0ebc23fd99c8980bae3c1] --- projects/hip/INSTALL.md | 2 -- 1 file changed, 2 deletions(-) diff --git a/projects/hip/INSTALL.md b/projects/hip/INSTALL.md index f8f67ac71f..df7cdd1de5 100644 --- a/projects/hip/INSTALL.md +++ b/projects/hip/INSTALL.md @@ -56,8 +56,6 @@ apt-get install hip_hcc * Checkout https://github.com/ROCm-Developer-Tools/HIP.git master branch and build it with HCC installed with ROCm packages. Please use -DHIP_COMPILER=clang with cmake to enable hip-clang. -* Environment variables to let hipcc to use clang to compile HIP program - * Default paths and environment variables: * By default HIP looks for HSA in /opt/rocm/hsa (can be overridden by setting HSA_PATH environment variable) From 5072c98f327d9eb5108cc04761a254aca445cf8e Mon Sep 17 00:00:00 2001 From: Yaxun Sam Liu Date: Sat, 30 Mar 2019 07:50:42 -0400 Subject: [PATCH 3/5] hip-clang: fix kernel not found on multi-gpu __hipRegisterFunction is called during by .init functions during program initialization. It calls hipModuleGetFunction to locate kernel symbol in code objects. hipModuleGetFunction assumes current device when locating kernel symbols. This works for HCC but not for hip-clang, since hip-clang needs to locate kernel symbols for different devices without switching between devices. This patch introduces a new hsa agent parameter to ihipModuleGetFunction, which allows __hipRegisterFunction to choose the correct hsa agent when locating kernel symbols. By default it uses this_agent(), therefore this patch has no impact on HCC. [ROCm/hip commit: 98b9e92908ebbc1a76231b83c208c84b0c4a46d3] --- projects/hip/src/hip_clang.cpp | 8 +++++--- projects/hip/src/hip_hcc_internal.h | 2 ++ projects/hip/src/hip_module.cpp | 26 ++++++++++++++++++++++---- 3 files changed, 29 insertions(+), 7 deletions(-) diff --git a/projects/hip/src/hip_clang.cpp b/projects/hip/src/hip_clang.cpp index 9af4f5a7a2..82e181ca69 100644 --- a/projects/hip/src/hip_clang.cpp +++ b/projects/hip/src/hip_clang.cpp @@ -132,11 +132,13 @@ extern "C" void __hipRegisterFunction( assert(modules && modules->size() >= g_deviceCnt); for (int deviceId = 0; deviceId < g_deviceCnt; ++deviceId) { hipFunction_t function; - if ((hipSuccess == hipModuleGetFunction(&function, modules->at(deviceId), deviceName) || + hsa_agent_t agent = g_allAgents[deviceId + 1]; + if ((hipSuccess == hipModuleGetFunctionEx(&function, modules->at(deviceId), deviceName, &agent) || // With code-object-v3, we need to match the kernel descriptor symbol name - (hipSuccess == hipModuleGetFunction( + (hipSuccess == hipModuleGetFunctionEx( &function, modules->at(deviceId), - (std::string(deviceName) + std::string(".kd")).c_str() + (std::string(deviceName) + std::string(".kd")).c_str(), + &agent ))) && function != nullptr) { functions[deviceId] = function; } diff --git a/projects/hip/src/hip_hcc_internal.h b/projects/hip/src/hip_hcc_internal.h index fb1588cc54..0e0db0c628 100644 --- a/projects/hip/src/hip_hcc_internal.h +++ b/projects/hip/src/hip_hcc_internal.h @@ -943,6 +943,8 @@ extern hipError_t ihipDeviceSetState(); extern ihipDevice_t* ihipGetDevice(int); ihipCtx_t* ihipGetPrimaryCtx(unsigned deviceIndex); +hipError_t hipModuleGetFunctionEx(hipFunction_t* hfunc, hipModule_t hmod, + const char* name, hsa_agent_t *agent); hipStream_t ihipSyncAndResolveStream(hipStream_t); diff --git a/projects/hip/src/hip_module.cpp b/projects/hip/src/hip_module.cpp index e81204d86b..c00004656a 100644 --- a/projects/hip/src/hip_module.cpp +++ b/projects/hip/src/hip_module.cpp @@ -360,13 +360,14 @@ inline hsa_status_t copy_agent_global_variables(hsa_executable_t, hsa_agent_t ag return HSA_STATUS_SUCCESS; } -hsa_executable_symbol_t find_kernel_by_name(hsa_executable_t executable, const char* kname) { +hsa_executable_symbol_t find_kernel_by_name(hsa_executable_t executable, const char* kname, + hsa_agent_t agent) { using namespace hip_impl; pair r{kname, {}}; hsa_executable_iterate_agent_symbols( - executable, this_agent(), + executable, agent, [](hsa_executable_t, hsa_agent_t, hsa_executable_symbol_t x, void* s) { auto p = static_cast*>(s); @@ -384,6 +385,11 @@ hsa_executable_symbol_t find_kernel_by_name(hsa_executable_t executable, const c return r.second; } +hsa_executable_symbol_t find_kernel_by_name(hsa_executable_t executable, const char* kname) { + return find_kernel_by_name(executable, kname, hip_impl::this_agent()); +} + + string read_elf_file_as_string(const void* file) { // Precondition: file points to an ELF image that was BITWISE loaded // into process accessible memory, and not one loaded by @@ -437,7 +443,8 @@ namespace hip_impl { } } // Namespace hip_impl. -hipError_t ihipModuleGetFunction(hipFunction_t* func, hipModule_t hmod, const char* name) { +hipError_t ihipModuleGetFunction(hipFunction_t* func, hipModule_t hmod, const char* name, + hsa_agent_t *agent = nullptr) { using namespace hip_impl; if (!func || !name) return hipErrorInvalidValue; @@ -450,7 +457,10 @@ hipError_t ihipModuleGetFunction(hipFunction_t* func, hipModule_t hmod, const ch if (!*func) return hipErrorInvalidValue; - auto kernel = find_kernel_by_name(hmod->executable, name); + if (!agent) + *agent = this_agent(); + + auto kernel = find_kernel_by_name(hmod->executable, name, *agent); if (kernel.handle == 0u) return hipErrorNotFound; @@ -462,11 +472,19 @@ hipError_t ihipModuleGetFunction(hipFunction_t* func, hipModule_t hmod, const ch return hipSuccess; } +// Get kernel for the current hsa agent. hipError_t hipModuleGetFunction(hipFunction_t* hfunc, hipModule_t hmod, const char* name) { HIP_INIT_API(hipModuleGetFunction, hfunc, hmod, name); return ihipLogStatus(ihipModuleGetFunction(hfunc, hmod, name)); } +// Get kernel for the given hsa agent. Internal use only. +hipError_t hipModuleGetFunctionEx(hipFunction_t* hfunc, hipModule_t hmod, + const char* name, hsa_agent_t *agent) { + HIP_INIT_API(hipModuleGetFunctionEx, hfunc, hmod, name); + return ihipLogStatus(ihipModuleGetFunction(hfunc, hmod, name, agent)); +} + namespace { hipFuncAttributes make_function_attributes(const amd_kernel_code_t& header) { hipFuncAttributes r{}; From 9e3258475390101c6cc154892a0a92a8b9cf8ada Mon Sep 17 00:00:00 2001 From: Evgeny Mankov Date: Mon, 1 Apr 2019 19:44:52 +0300 Subject: [PATCH 4/5] [HIPIFY][cmake] Update CMakeLists and Readme cause CUDA 10.1 and clang 8.0.0 are released [ROCm/hip commit: 7c221450fa991d61a84cdf8bf03eab90196ab58c] --- projects/hip/hipify-clang/CMakeLists.txt | 4 +++- projects/hip/hipify-clang/README.md | 3 ++- 2 files changed, 5 insertions(+), 2 deletions(-) diff --git a/projects/hip/hipify-clang/CMakeLists.txt b/projects/hip/hipify-clang/CMakeLists.txt index 81b593a25a..e1fba870e4 100644 --- a/projects/hip/hipify-clang/CMakeLists.txt +++ b/projects/hip/hipify-clang/CMakeLists.txt @@ -108,7 +108,7 @@ if (HIPIFY_CLANG_TESTS) (CUDA_VERSION VERSION_GREATER "7.5" AND LLVM_PACKAGE_VERSION VERSION_LESS "4.0") OR (CUDA_VERSION VERSION_GREATER "8.0" AND LLVM_PACKAGE_VERSION VERSION_LESS "6.0") OR (CUDA_VERSION VERSION_GREATER "9.0" AND LLVM_PACKAGE_VERSION VERSION_LESS "7.0") OR - CUDA_VERSION VERSION_EQUAL "10.0") + (CUDA_VERSION VERSION_GREATER "10.1" AND LLVM_PACKAGE_VERSION VERSION_LESS "8.0")) message(SEND_ERROR "CUDA ${CUDA_VERSION} is not supported by clang ${LLVM_PACKAGE_VERSION}.") if (CUDA_VERSION VERSION_LESS "7.0") message(STATUS "Please install CUDA 7.0 or higher.") @@ -120,6 +120,8 @@ if (HIPIFY_CLANG_TESTS) message(STATUS "Please install clang 6.0 or higher.") elseif ((CUDA_VERSION VERSION_EQUAL "9.1") OR (CUDA_VERSION VERSION_EQUAL "9.2")) message(STATUS "Please install clang 7.0 or higher.") + elseif ((CUDA_VERSION VERSION_EQUAL "10.0") OR (CUDA_VERSION VERSION_EQUAL "10.1")) + message(STATUS "Please install clang 8.0 or higher.") endif() endif() diff --git a/projects/hip/hipify-clang/README.md b/projects/hip/hipify-clang/README.md index ee25b23f2d..5fc14a97c1 100644 --- a/projects/hip/hipify-clang/README.md +++ b/projects/hip/hipify-clang/README.md @@ -52,7 +52,8 @@ | 7.0.0 | 9.2 | -
not working due to
the clang's bug [38811](https://bugs.llvm.org/show_bug.cgi?id=38811) | -
not working due to
the clang's bug [36384](https://bugs.llvm.org/show_bug.cgi?id=36384) | | 7.0.1 | 9.2 | -
not working due to
the clang's bug [38811](https://bugs.llvm.org/show_bug.cgi?id=38811) | -
not working due to
the clang's bug [36384](https://bugs.llvm.org/show_bug.cgi?id=36384) | | 7.1.0 | 9.2 (?) | -
LLVM 7.1.0
is not yet released | -
LLVM 7.1.0
is not yet released | -| 8.0.0 | 10.0 (?) | -
LLVM 8.0.0
is not yet released | -
LLVM 8.0.0
is not yet released | +| 8.0.0 | 10.1 | -
not working due to
the clang's bug [38811](https://bugs.llvm.org/show_bug.cgi?id=38811) | -
not working due to
the clang's bug [36384](https://bugs.llvm.org/show_bug.cgi?id=36384) | +| 8.0.1 | 10.1 | -
LLVM 8.0.1
is not yet released | -
LLVM 8.0.1
is not yet released | In most cases, you can get a suitable version of LLVM+CLANG with your package manager. From 7bbf825c74c9d92a86ee45c4af74ac3331783eb3 Mon Sep 17 00:00:00 2001 From: Evgeny Mankov Date: Mon, 1 Apr 2019 19:52:08 +0300 Subject: [PATCH 5/5] [HIPIFY][tests] Fix typo in test for CUDA 10.x [ROCm/hip commit: 8bca9f1d7a70a4389ef2821b373e7d080baf8ad1] --- .../hipify-clang/unit_tests/libraries/cuSPARSE/cuSPARSE_10.cu | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/projects/hip/tests/hipify-clang/unit_tests/libraries/cuSPARSE/cuSPARSE_10.cu b/projects/hip/tests/hipify-clang/unit_tests/libraries/cuSPARSE/cuSPARSE_10.cu index 6690a7c296..2ab2d605de 100644 --- a/projects/hip/tests/hipify-clang/unit_tests/libraries/cuSPARSE/cuSPARSE_10.cu +++ b/projects/hip/tests/hipify-clang/unit_tests/libraries/cuSPARSE/cuSPARSE_10.cu @@ -131,7 +131,7 @@ int main(int argc, char*argv[]) // CHECK: assert(HIPSPARSE_STATUS_SUCCESS == status); assert(CUSPARSE_STATUS_SUCCESS == status); // CHECK: cublasStat = hipblasCreate(&cublasH); - cublasStat = cublasCreate(cublasH); + cublasStat = cublasCreate(&cublasH); // CHECK: assert(HIPBLAS_STATUS_SUCCESS == cublasStat); assert(CUBLAS_STATUS_SUCCESS == cublasStat); // CHECK: cublasStat = hipblasSetStream(cublasH, stream);