[CUDA] Another attempt to fix early inclusion of <new> from libstdc++

Previous patch (9a465057a64dba) did not fix the problem.
https://bugs.llvm.org/show_bug.cgi?id=48228

If the <new> is included too early, before CUDA-specific defines are available,
just include-next the standard <new> and undo the include guard.  CUDA-specific
variants of operator new/delete will be declared if/when <new> is used from the
CUDA source itself, when all CUDA-related macros are available.

Differential Revision: https://reviews.llvm.org/D91807

(cherry picked from commit 43267929423bf768bbbcc65e47a07e37af7f4e22)

1 files changed, 22 insertions, 24 deletions

diff --git a/clang/lib/Headers/cuda_wrappers/new b/clang/lib/Headers/cuda_wrappers/new
index 47690f1152fe..7f255314056a 100644
--- a/clang/lib/Headers/cuda_wrappers/new
+++ b/clang/lib/Headers/cuda_wrappers/new
@@ -26,6 +26,13 @@
 
 #include_next <new>
 
+#if !defined(__device__)
+// The header has been included too early from the standard C++ library
+// and CUDA-specific macros are not available yet.
+// Undo the include guard and try again later.
+#undef __CLANG_CUDA_WRAPPERS_NEW
+#else
+
 #pragma push_macro("CUDA_NOEXCEPT")
 #if __cplusplus >= 201103L
 #define CUDA_NOEXCEPT noexcept
@@ -33,76 +40,67 @@
 #define CUDA_NOEXCEPT
 #endif
 
-#pragma push_macro("__DEVICE__")
-#if defined __device__
-#define __DEVICE__ __device__
-#else
-// <new> has been included too early from the standard libc++ headers and the
-// standard CUDA macros are not available yet. We have to define our own.
-#define __DEVICE__ __attribute__((device))
-#endif
-
 // Device overrides for non-placement new and delete.
-__DEVICE__ inline void *operator new(__SIZE_TYPE__ size) {
+__device__ inline void *operator new(__SIZE_TYPE__ size) {
   if (size == 0) {
     size = 1;
   }
   return ::malloc(size);
 }
-__DEVICE__ inline void *operator new(__SIZE_TYPE__ size,
+__device__ inline void *operator new(__SIZE_TYPE__ size,
                                      const std::nothrow_t &) CUDA_NOEXCEPT {
   return ::operator new(size);
 }
 
-__DEVICE__ inline void *operator new[](__SIZE_TYPE__ size) {
+__device__ inline void *operator new[](__SIZE_TYPE__ size) {
   return ::operator new(size);
 }
-__DEVICE__ inline void *operator new[](__SIZE_TYPE__ size,
+__device__ inline void *operator new[](__SIZE_TYPE__ size,
                                        const std::nothrow_t &) {
   return ::operator new(size);
 }
 
-__DEVICE__ inline void operator delete(void* ptr) CUDA_NOEXCEPT {
+__device__ inline void operator delete(void* ptr) CUDA_NOEXCEPT {
   if (ptr) {
     ::free(ptr);
   }
 }
-__DEVICE__ inline void operator delete(void *ptr,
+__device__ inline void operator delete(void *ptr,
                                        const std::nothrow_t &) CUDA_NOEXCEPT {
   ::operator delete(ptr);
 }
 
-__DEVICE__ inline void operator delete[](void* ptr) CUDA_NOEXCEPT {
+__device__ inline void operator delete[](void* ptr) CUDA_NOEXCEPT {
   ::operator delete(ptr);
 }
-__DEVICE__ inline void operator delete[](void *ptr,
+__device__ inline void operator delete[](void *ptr,
                                          const std::nothrow_t &) CUDA_NOEXCEPT {
   ::operator delete(ptr);
 }
 
 // Sized delete, C++14 only.
 #if __cplusplus >= 201402L
-__DEVICE__ inline void operator delete(void *ptr,
+__device__ inline void operator delete(void *ptr,
                                        __SIZE_TYPE__ size) CUDA_NOEXCEPT {
   ::operator delete(ptr);
 }
-__DEVICE__ inline void operator delete[](void *ptr,
+__device__ inline void operator delete[](void *ptr,
                                          __SIZE_TYPE__ size) CUDA_NOEXCEPT {
   ::operator delete(ptr);
 }
 #endif
 
 // Device overrides for placement new and delete.
-__DEVICE__ inline void *operator new(__SIZE_TYPE__, void *__ptr) CUDA_NOEXCEPT {
+__device__ inline void *operator new(__SIZE_TYPE__, void *__ptr) CUDA_NOEXCEPT {
   return __ptr;
 }
-__DEVICE__ inline void *operator new[](__SIZE_TYPE__, void *__ptr) CUDA_NOEXCEPT {
+__device__ inline void *operator new[](__SIZE_TYPE__, void *__ptr) CUDA_NOEXCEPT {
   return __ptr;
 }
-__DEVICE__ inline void operator delete(void *, void *) CUDA_NOEXCEPT {}
-__DEVICE__ inline void operator delete[](void *, void *) CUDA_NOEXCEPT {}
+__device__ inline void operator delete(void *, void *) CUDA_NOEXCEPT {}
+__device__ inline void operator delete[](void *, void *) CUDA_NOEXCEPT {}
 
-#pragma pop_macro("__DEVICE__")
 #pragma pop_macro("CUDA_NOEXCEPT")
 
+#endif // __device__
 #endif // include guard