From 9521a15c889579d4a1ef88af47aece93d0108157 Mon Sep 17 00:00:00 2001
From: "Yaxun (Sam) Liu" <yaxun.liu@amd.com>
Date: Tue, 22 Jan 2019 09:07:18 -0800
Subject: [PATCH] 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
---
 aten/src/THC/THCAsmUtils.cuh | 2 +-
 aten/src/THC/THCAtomics.cuh  | 2 +-
 caffe2/utils/conversions.h   | 2 +-
 caffe2/utils/fixed_divisor.h | 2 +-
 caffe2/utils/math_utils.h    | 2 +-
 cmake/Dependencies.cmake     | 5 +++++
 third_party/eigen            | 2 +-
 7 files changed, 11 insertions(+), 6 deletions(-)

diff --git a/aten/src/THC/THCAsmUtils.cuh b/aten/src/THC/THCAsmUtils.cuh
index e940592..479567f 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 427133f..67849d0 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 446cd22..8c1a644 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 3c7aaf0..64e84fb 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 b7dfc6b..a97732c 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 904113d..9adbc94 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
index f59336c..b8fa8f5 160000
--- a/third_party/eigen
+++ b/third_party/eigen
@@ -1 +1 @@
-Subproject commit f59336cee358f92b959de6a0daf07c4ab2318022
+Subproject commit b8fa8f5b5a3fb4a72b1fa020843de06f490c0447
-- 
2.7.4