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

[Artem Belevich via Phabricator via cfe-commits]

tra created this revision.
tra added a reviewer: jlebar.
Herald added subscribers: bixia, yaxunl.
Herald added a project: clang.
tra requested review of this revision.

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 define
__device__ in the wrapper header itself.

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


Repository:
  rG LLVM Github Monorepo

https://reviews.llvm.org/D91807

Files:
  clang/lib/Headers/cuda_wrappers/new


Index: clang/lib/Headers/cuda_wrappers/new
===================================================================
--- clang/lib/Headers/cuda_wrappers/new
+++ clang/lib/Headers/cuda_wrappers/new
@@ -33,66 +33,76 @@
 #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 // include guard


-------------- next part --------------
A non-text attachment was scrubbed...
Name: D91807.306454.patch
Type: text/x-patch
Size: 3488 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/cfe-commits/attachments/20201119/e46025d1/attachment.bin>