Revert "[OpenMP][AMDGCN] Initial math headers support"

This reverts commit 968899a.
llvm · Jul 21, 2021 · d71062f · d71062f
1 parent a733bbb
commit d71062f
Show file tree

Hide file tree

Showing 11 changed files with 92 additions and 231 deletions.
diff --git a/clang/lib/Driver/ToolChains/Clang.cpp b/clang/lib/Driver/ToolChains/Clang.cpp
@@ -1255,8 +1255,7 @@ void Clang::AddPreprocessingOptions(Compilation &C, const JobAction &JA,
   // If we are offloading to a target via OpenMP we need to include the
   // openmp_wrappers folder which contains alternative system headers.
   if (JA.isDeviceOffloading(Action::OFK_OpenMP) &&
-      (getToolChain().getTriple().isNVPTX() ||
-       getToolChain().getTriple().isAMDGCN())) {
+      getToolChain().getTriple().isNVPTX()){
     if (!Args.hasArg(options::OPT_nobuiltininc)) {
       // Add openmp_wrappers/* to our system include path.  This lets us wrap
       // standard library headers.

diff --git a/clang/lib/Headers/__clang_hip_cmath.h b/clang/lib/Headers/__clang_hip_cmath.h
diff --git a/clang/lib/Headers/__clang_hip_math.h b/clang/lib/Headers/__clang_hip_math.h
@@ -9,7 +9,7 @@
 #ifndef __CLANG_HIP_MATH_H__
 #define __CLANG_HIP_MATH_H__
 
-#if !defined(__HIP__) && !defined(__OPENMP_AMDGCN__)
+#if !defined(__HIP__)
 #error "This file is for HIP and OpenMP AMDGCN device compilation only."
 #endif
 
@@ -19,27 +19,18 @@
 #endif
 #include <limits.h>
 #include <stdint.h>
-#endif // !defined(__HIPCC_RTC__)
+#endif // __HIPCC_RTC__
 
 #pragma push_macro("__DEVICE__")
-
-#ifdef __OPENMP_AMDGCN__
-#define __DEVICE__ static inline __attribute__((always_inline, nothrow))
-#else
 #define __DEVICE__ static __device__ inline __attribute__((always_inline))
-#endif
 
 // A few functions return bool type starting only in C++11.
 #pragma push_macro("__RETURN_TYPE")
-#ifdef __OPENMP_AMDGCN__
-#define __RETURN_TYPE int
-#else
 #if defined(__cplusplus)
 #define __RETURN_TYPE bool
 #else
 #define __RETURN_TYPE int
 #endif
-#endif // __OPENMP_AMDGCN__
 
 #if defined (__cplusplus) && __cplusplus < 201103L
 // emulate static_assert on type sizes
@@ -1271,15 +1262,15 @@ float min(float __x, float __y) { return fminf(__x, __y); }
 __DEVICE__
 double min(double __x, double __y) { return fmin(__x, __y); }
 
-#if !defined(__HIPCC_RTC__) && !defined(__OPENMP_AMDGCN__)
+#if !defined(__HIPCC_RTC__)
 __host__ inline static int min(int __arg1, int __arg2) {
   return std::min(__arg1, __arg2);
 }
 
 __host__ inline static int max(int __arg1, int __arg2) {
   return std::max(__arg1, __arg2);
 }
-#endif // !defined(__HIPCC_RTC__) && !defined(__OPENMP_AMDGCN__)
+#endif // __HIPCC_RTC__
 #endif
 
 #pragma pop_macro("__DEVICE__")

diff --git a/clang/lib/Headers/openmp_wrappers/__clang_openmp_device_functions.h b/clang/lib/Headers/openmp_wrappers/__clang_openmp_device_functions.h
@@ -14,13 +14,13 @@
 #error "This file is for OpenMP compilation only."
 #endif
 
+#pragma omp begin declare variant match(                                       \
+    device = {arch(nvptx, nvptx64)}, implementation = {extension(match_any)})
+
 #ifdef __cplusplus
 extern "C" {
 #endif
 
-#pragma omp begin declare variant match(                                       \
-    device = {arch(nvptx, nvptx64)}, implementation = {extension(match_any)})
-
 #define __CUDA__
 #define __OPENMP_NVPTX__
 
@@ -33,32 +33,12 @@ extern "C" {
 #undef __OPENMP_NVPTX__
 #undef __CUDA__
 
-#pragma omp end declare variant
-
-#pragma omp begin declare variant match(device = {arch(amdgcn)})
-
-// Import types which will be used by __clang_hip_libdevice_declares.h
-#ifndef __cplusplus
-#include <stdbool.h>
-#include <stdint.h>
-#endif
-
-#define __OPENMP_AMDGCN__
-#pragma push_macro("__device__")
-#define __device__
-
-/// Include declarations for libdevice functions.
-#include <__clang_hip_libdevice_declares.h>
-
-#pragma pop_macro("__device__")
-#undef __OPENMP_AMDGCN__
-
-#pragma omp end declare variant
-
 #ifdef __cplusplus
 } // extern "C"
 #endif
 
+#pragma omp end declare variant
+
 // Ensure we make `_ZdlPv`, aka. `operator delete(void*)` available without the
 // need to `include <new>` in C++ mode.
 #ifdef __cplusplus

diff --git a/clang/lib/Headers/openmp_wrappers/cmath b/clang/lib/Headers/openmp_wrappers/cmath
@@ -75,19 +75,4 @@ __DEVICE__ float tgamma(float __x) { return ::tgammaf(__x); }
 
 #pragma omp end declare variant
 
-#ifdef __AMDGCN__
-#pragma omp begin declare variant match(device = {arch(amdgcn)})
-
-#pragma push_macro("__constant__")
-#define __constant__ __attribute__((constant))
-#define __OPENMP_AMDGCN__
-
-#include <__clang_hip_cmath.h>
-
-#pragma pop_macro("__constant__")
-#undef __OPENMP_AMDGCN__
-
-#pragma omp end declare variant
-#endif // __AMDGCN__
-
 #endif
diff --git a/clang/lib/Headers/openmp_wrappers/math.h b/clang/lib/Headers/openmp_wrappers/math.h
@@ -48,12 +48,4 @@
 
 #pragma omp end declare variant
 
-#pragma omp begin declare variant match(device = {arch(amdgcn)})
-
-#define __OPENMP_AMDGCN__
-#include <__clang_hip_math.h>
-#undef __OPENMP_AMDGCN__
-
-#pragma omp end declare variant
-
 #endif
diff --git a/clang/test/Headers/Inputs/include/algorithm b/clang/test/Headers/Inputs/include/algorithm
diff --git a/clang/test/Headers/Inputs/include/cstdlib b/clang/test/Headers/Inputs/include/cstdlib
@@ -21,13 +21,9 @@ abs(long __i) { return __builtin_labs(__i); }
 inline long long
 abs(long long __x) { return __builtin_llabs (__x); }
 
-// amdgcn already provides definition of fabs
-#ifndef __AMDGCN__
 float fabs(float __x) { return __builtin_fabs(__x); }
-#endif
 
 float abs(float __x) { return fabs(__x); }
 double abs(double __x) { return fabs(__x); }
 
 }
-
diff --git a/clang/test/Headers/Inputs/include/utility b/clang/test/Headers/Inputs/include/utility
diff --git a/clang/test/Headers/amdgcn_openmp_device_math.c b/clang/test/Headers/amdgcn_openmp_device_math.c
diff --git a/clang/test/Headers/openmp_device_math_isnan.cpp b/clang/test/Headers/openmp_device_math_isnan.cpp
@@ -21,14 +21,14 @@
 double math(float f, double d) {
   double r = 0;
   // INT_RETURN: call i32 @__nv_isnanf(float
-  // AMD_INT_RETURN: call i32 @__ocml_isnan_f32(float
+  // AMD_INT_RETURN: call i32 @_{{.*}}isnanf(float
   // BOOL_RETURN: call i32 @__nv_isnanf(float
-  // AMD_BOOL_RETURN: call i32 @__ocml_isnan_f32(float
+  // AMD_BOOL_RETURN: call zeroext i1 @_{{.*}}isnanf(float
   r += std::isnan(f);
   // INT_RETURN: call i32 @__nv_isnand(double
-  // AMD_INT_RETURN: call i32 @__ocml_isnan_f64(double
+  // AMD_INT_RETURN: call i32 @_{{.*}}isnand(double
   // BOOL_RETURN: call i32 @__nv_isnand(double
-  // AMD_BOOL_RETURN: call i32 @__ocml_isnan_f64(double
+  // AMD_BOOL_RETURN: call zeroext i1 @_{{.*}}isnand(double
   r += std::isnan(d);
   return r;
 }