diff --git a/projects/hip/INSTALL.md b/projects/hip/INSTALL.md
index 7b12bd42ef..df7cdd1de5 100644
--- a/projects/hip/INSTALL.md
+++ b/projects/hip/INSTALL.md
@@ -36,37 +36,32 @@ 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.
-
-* 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
+ * 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.
* 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 +94,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
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.
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{};
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);