Merge 'master' into 'amd-master'
Change-Id: I4369829444ebd150e59f07f5db44e4584d9cc396
[ROCm/hip commit: 87ce7c278a]
Tento commit je obsažen v:
Jenkins
@@ -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
|
||||
|
||||
@@ -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()
|
||||
|
||||
|
||||
@@ -52,7 +52,8 @@
|
||||
| 7.0.0 | 9.2 | - <br/> not working due to <br/> the clang's bug [38811](https://bugs.llvm.org/show_bug.cgi?id=38811) | - <br/> not working due to <br/> the clang's bug [36384](https://bugs.llvm.org/show_bug.cgi?id=36384) |
|
||||
| 7.0.1 | 9.2 | - <br/> not working due to <br/> the clang's bug [38811](https://bugs.llvm.org/show_bug.cgi?id=38811) | - <br/> not working due to <br/> the clang's bug [36384](https://bugs.llvm.org/show_bug.cgi?id=36384) |
|
||||
| 7.1.0 | 9.2 (?) | - <br/> LLVM 7.1.0 <br/> is not yet released | - <br/> LLVM 7.1.0 <br/> is not yet released |
|
||||
| 8.0.0 | 10.0 (?) | - <br/> LLVM 8.0.0 <br/> is not yet released | - <br/> LLVM 8.0.0 <br/> is not yet released |
|
||||
| 8.0.0 | 10.1 | - <br/> not working due to <br/> the clang's bug [38811](https://bugs.llvm.org/show_bug.cgi?id=38811) | - <br/> not working due to <br/> the clang's bug [36384](https://bugs.llvm.org/show_bug.cgi?id=36384) |
|
||||
| 8.0.1 | 10.1 | - <br/> LLVM 8.0.1 <br/> is not yet released | - <br/> LLVM 8.0.1 <br/> is not yet released |
|
||||
|
||||
In most cases, you can get a suitable version of LLVM+CLANG with your package manager.
|
||||
|
||||
|
||||
@@ -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;
|
||||
}
|
||||
|
||||
@@ -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);
|
||||
|
||||
@@ -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<const char*, hsa_executable_symbol_t> 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<pair<const char*, hsa_executable_symbol_t>*>(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{};
|
||||
|
||||
@@ -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);
|
||||
|
||||
Odkázat v novém úkolu
Zablokovat Uživatele