[CUDA] Unbreak CUDA compilation with -std=c++20

[tra]

Summary

Standard libc++ headers in stdc++ mode include <new> which picks up
cuda_wrappers/new before any of the CUDA macros have been defined.

We can not include CUDA headers that early, so the work-around is to
postpone the CUDA-specific parts until after <new> is included from user
sources and all CUDA-related macros and declarations are available.

PR48228 https://bugs.llvm.org/show_bug.cgi?id=48228

[jlebar]

How fun. :)

[tra]

Reopening the tracker for a better fix.

[tra]

Better fix for the early inclusion of <new>

[tra]

@jlebar PTAL. The previous fix did not quite fix the problem. __device__ operator new/delete was still not defined, but the issue was not reported due to deferred diags if the functions were not used.

[tra]

Removed forgotten pop_macro(DEVICE)

clang/lib/Headers/cuda_wrappers/new

  *===-----------------------------------------------------------------------===
  */
 
 #ifndef __CLANG_CUDA_WRAPPERS_NEW
 #define __CLANG_CUDA_WRAPPERS_NEW
 
 #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
 #else
 #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("__device__")
 #pragma pop_macro("CUDA_NOEXCEPT")
 
+#endif // __device__
 #endif // include guard