diff options
| author | Yaxun (Sam) Liu <yaxun.liu@amd.com> | 2019-01-22 09:07:18 -0800 |
|---|---|---|
| committer | Facebook Github Bot <facebook-github-bot@users.noreply.github.com> | 2019-01-22 09:09:48 -0800 |
| commit | 9521a15c889579d4a1ef88af47aece93d0108157 (patch) | |
| tree | 3c1cd953e5feff554c08ce79b3281dc6e0e18549 | |
| parent | 4cf76574b958b41ae1df7fba6e87af101414549e (diff) | |
| download | pytorch-9521a15c889579d4a1ef88af47aece93d0108157.tar.gz pytorch-9521a15c889579d4a1ef88af47aece93d0108157.tar.bz2 pytorch-9521a15c889579d4a1ef88af47aece93d0108157.zip | |
hip-clang enablement (#16085)
Summary:
Initial enabling of the upcoming hip-clang compiler for the PyTorch source base.
Changes:
* update the Eigen submodule to a version including our upstreamed hip-clang enabling there
* modify a few ifdef guards with the `__HIP__` macro used by hip-clang
* use `__lane_id` instead of `hc::__lane_id`
* add Debug flags for ROCm to the cmake infrastructure
Pull Request resolved: https://github.com/pytorch/pytorch/pull/16085
Differential Revision: D13709459
Pulled By: ezyang
fbshipit-source-id: 1b7b33fe810a0434766180580d4443ea177eb7c7
| -rw-r--r-- | aten/src/THC/THCAsmUtils.cuh | 2 | ||||
| -rw-r--r-- | aten/src/THC/THCAtomics.cuh | 2 | ||||
| -rw-r--r-- | caffe2/utils/conversions.h | 2 | ||||
| -rw-r--r-- | caffe2/utils/fixed_divisor.h | 2 | ||||
| -rw-r--r-- | caffe2/utils/math_utils.h | 2 | ||||
| -rw-r--r-- | cmake/Dependencies.cmake | 5 | ||||
| m--------- | third_party/eigen | 0 |
7 files changed, 10 insertions, 5 deletions
diff --git a/aten/src/THC/THCAsmUtils.cuh b/aten/src/THC/THCAsmUtils.cuh index e940592d1c..479567fba5 100644 --- a/aten/src/THC/THCAsmUtils.cuh +++ b/aten/src/THC/THCAsmUtils.cuh @@ -84,7 +84,7 @@ struct Bitfield<uint64_t> { __device__ __forceinline__ int getLaneId() { #if defined(__HIP_PLATFORM_HCC__) - return hc::__lane_id(); + return __lane_id(); #else int laneId; asm("mov.s32 %0, %laneid;" : "=r"(laneId) ); diff --git a/aten/src/THC/THCAtomics.cuh b/aten/src/THC/THCAtomics.cuh index 427133f003..67849d05a6 100644 --- a/aten/src/THC/THCAtomics.cuh +++ b/aten/src/THC/THCAtomics.cuh @@ -133,7 +133,7 @@ static inline __device__ void atomicAdd(double *address, double val) { } while (assumed != old); } #elif !defined(__CUDA_ARCH__) && (CUDA_VERSION < 8000) || defined(__HIP_PLATFORM_HCC__) -#if defined(__HIP_PLATFORM_HCC__) && __hcc_workweek__ < 18312 +#if defined(__HIP_PLATFORM_HCC__) && __hcc_workweek__ < 18312 && !__HIP__ // This needs to be defined for the host side pass static inline __device__ void atomicAdd(double *address, double val) { } #endif diff --git a/caffe2/utils/conversions.h b/caffe2/utils/conversions.h index 446cd22293..8c1a644913 100644 --- a/caffe2/utils/conversions.h +++ b/caffe2/utils/conversions.h @@ -11,7 +11,7 @@ #include <caffe2/core/hip/common_gpu.h> #endif -#if defined(__CUDA_ARCH__) || defined(__HIP_DEVICE_COMPILE__) +#if defined(__CUDA_ARCH__) || defined(__HIP_DEVICE_COMPILE__) || defined(__HIP__) #define CONVERSIONS_DECL __host__ __device__ inline #else #define CONVERSIONS_DECL inline diff --git a/caffe2/utils/fixed_divisor.h b/caffe2/utils/fixed_divisor.h index 3c7aaf016a..64e84fb7da 100644 --- a/caffe2/utils/fixed_divisor.h +++ b/caffe2/utils/fixed_divisor.h @@ -5,7 +5,7 @@ #include <cstdio> #include <cstdlib> -#if defined(__CUDA_ARCH__) || defined(__HIP_ARCH__) +#if defined(__CUDA_ARCH__) || defined(__HIP_ARCH__) || defined(__HIP__) #define FIXED_DIVISOR_DECL inline __host__ __device__ #else #define FIXED_DIVISOR_DECL inline diff --git a/caffe2/utils/math_utils.h b/caffe2/utils/math_utils.h index b7dfc6bf45..a97732cb01 100644 --- a/caffe2/utils/math_utils.h +++ b/caffe2/utils/math_utils.h @@ -3,7 +3,7 @@ #include "caffe2/core/common.h" -#if defined(__CUDA_ARCH__) || defined(__HIP_DEVICE_COMPILE__) +#if defined(__CUDA_ARCH__) || defined(__HIP_DEVICE_COMPILE__) || defined(__HIP__) #define MATH_UTILS_DECL inline __host__ __device__ #else #define MATH_UTILS_DECL inline diff --git a/cmake/Dependencies.cmake b/cmake/Dependencies.cmake index 904113d338..9adbc94e73 100644 --- a/cmake/Dependencies.cmake +++ b/cmake/Dependencies.cmake @@ -747,6 +747,11 @@ if(USE_ROCM) list(APPEND HIP_CXX_FLAGS -Wno-duplicate-decl-specifier) list(APPEND HIP_CXX_FLAGS -DCAFFE2_USE_MIOPEN) + if(CMAKE_BUILD_TYPE MATCHES Debug) + list(APPEND HIP_CXX_FLAGS -g) + list(APPEND HIP_CXX_FLAGS -O0) + endif(CMAKE_BUILD_TYPE MATCHES Debug) + set(HIP_HCC_FLAGS ${HIP_CXX_FLAGS}) # Ask hcc to generate device code during compilation so we can use # host linker to link. diff --git a/third_party/eigen b/third_party/eigen -Subproject f59336cee358f92b959de6a0daf07c4ab231802 +Subproject b8fa8f5b5a3fb4a72b1fa020843de06f490c044 |