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

Broke nvptx compilation on files including <complex>

This reverts commit 12da97e.

Author: JonChesterfield

clang/lib/Driver/ToolChains/Clang.cpp

@@ -1256,8 +1256,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.

clang/lib/Headers/__clang_hip_cmath.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__")

clang/lib/Headers/__clang_hip_math.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

clang/lib/Headers/openmp_wrappers/__clang_openmp_device_functions.h

@@ -75,58 +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__
-
-// Define overloads otherwise which are absent
-#define __DEVICE__ static constexpr __attribute__((always_inline, nothrow))
-
-__DEVICE__ float acos(float __x) { return ::acosf(__x); }
-__DEVICE__ float acosh(float __x) { return ::acoshf(__x); }
-__DEVICE__ float asin(float __x) { return ::asinf(__x); }
-__DEVICE__ float asinh(float __x) { return ::asinhf(__x); }
-__DEVICE__ float atan(float __x) { return ::atanf(__x); }
-__DEVICE__ float atan2(float __x, float __y) { return ::atan2f(__x, __y); }
-__DEVICE__ float atanh(float __x) { return ::atanhf(__x); }
-__DEVICE__ float cbrt(float __x) { return ::cbrtf(__x); }
-__DEVICE__ float cosh(float __x) { return ::coshf(__x); }
-__DEVICE__ float erf(float __x) { return ::erff(__x); }
-__DEVICE__ float erfc(float __x) { return ::erfcf(__x); }
-__DEVICE__ float exp2(float __x) { return ::exp2f(__x); }
-__DEVICE__ float expm1(float __x) { return ::expm1f(__x); }
-__DEVICE__ float fdim(float __x, float __y) { return ::fdimf(__x, __y); }
-__DEVICE__ float hypot(float __x, float __y) { return ::hypotf(__x, __y); }
-__DEVICE__ int ilogb(float __x) { return ::ilogbf(__x); }
-__DEVICE__ float ldexp(float __arg, int __exp) {
-  return ::ldexpf(__arg, __exp);
-}
-__DEVICE__ float lgamma(float __x) { return ::lgammaf(__x); }
-__DEVICE__ float log1p(float __x) { return ::log1pf(__x); }
-__DEVICE__ float logb(float __x) { return ::logbf(__x); }
-__DEVICE__ float nextafter(float __x, float __y) {
-  return ::nextafterf(__x, __y);
-}
-__DEVICE__ float remainder(float __x, float __y) {
-  return ::remainderf(__x, __y);
-}
-__DEVICE__ float scalbn(float __x, int __y) { return ::scalbnf(__x, __y); }
-__DEVICE__ float sinh(float __x) { return ::sinhf(__x); }
-__DEVICE__ float tan(float __x) { return ::tanf(__x); }
-__DEVICE__ float tanh(float __x) { return ::tanhf(__x); }
-__DEVICE__ float tgamma(float __x) { return ::tgammaf(__x); }
-
-#undef __DEVICE__
-
-#pragma omp end declare variant
-#endif // __AMDGCN__
-
 #endif

clang/lib/Headers/openmp_wrappers/cmath

@@ -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

clang/lib/Headers/openmp_wrappers/math.h

@@ -27,4 +27,3 @@ float abs(float __x) { return fabs(__x); }
 double abs(double __x) { return fabs(__x); }
 
 }
-

clang/test/Headers/Inputs/include/algorithm

@@ -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;
 }