From 409aa49218d7ba2b7fecd89aef3c5568882ee35a 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. --- INSTALL.md | 20 ++++++++++---------- 1 file changed, 10 insertions(+), 10 deletions(-) diff --git a/INSTALL.md b/INSTALL.md index 7b12bd42ef..f8f67ac71f 100644 --- a/INSTALL.md +++ b/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 ed2fbb27e4fed42d43d0ebc23fd99c8980bae3c1 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 --- INSTALL.md | 2 -- 1 file changed, 2 deletions(-) diff --git a/INSTALL.md b/INSTALL.md index f8f67ac71f..df7cdd1de5 100644 --- a/INSTALL.md +++ b/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 98b9e92908ebbc1a76231b83c208c84b0c4a46d3 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. --- src/hip_clang.cpp | 8 +++++--- src/hip_hcc_internal.h | 2 ++ src/hip_module.cpp | 26 ++++++++++++++++++++++---- 3 files changed, 29 insertions(+), 7 deletions(-) diff --git a/src/hip_clang.cpp b/src/hip_clang.cpp index 9af4f5a7a2..82e181ca69 100644 --- a/src/hip_clang.cpp +++ b/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/src/hip_hcc_internal.h b/src/hip_hcc_internal.h index fb1588cc54..0e0db0c628 100644 --- a/src/hip_hcc_internal.h +++ b/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/src/hip_module.cpp b/src/hip_module.cpp index e81204d86b..c00004656a 100644 --- a/src/hip_module.cpp +++ b/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 7c221450fa991d61a84cdf8bf03eab90196ab58c 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 --- hipify-clang/CMakeLists.txt | 4 +++- hipify-clang/README.md | 3 ++- 2 files changed, 5 insertions(+), 2 deletions(-) diff --git a/hipify-clang/CMakeLists.txt b/hipify-clang/CMakeLists.txt index 81b593a25a..e1fba870e4 100644 --- a/hipify-clang/CMakeLists.txt +++ b/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/hipify-clang/README.md b/hipify-clang/README.md index ee25b23f2d..5fc14a97c1 100644 --- a/hipify-clang/README.md +++ b/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 8bca9f1d7a70a4389ef2821b373e7d080baf8ad1 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 --- tests/hipify-clang/unit_tests/libraries/cuSPARSE/cuSPARSE_10.cu | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/tests/hipify-clang/unit_tests/libraries/cuSPARSE/cuSPARSE_10.cu b/tests/hipify-clang/unit_tests/libraries/cuSPARSE/cuSPARSE_10.cu index 6690a7c296..2ab2d605de 100644 --- a/tests/hipify-clang/unit_tests/libraries/cuSPARSE/cuSPARSE_10.cu +++ b/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);