diff options
author | Artem Belevich <tra@google.com> | 2020-11-19 10:06:57 -0800 |
---|---|---|
committer | Tom Stellard <tstellar@redhat.com> | 2020-12-09 12:42:33 -0500 |
commit | aa29049404efdc0134066839bc14d135d69ec225 (patch) | |
tree | f3b23153957816bcec05da3d68313135888106d6 | |
parent | [X86][AVX] Only share broadcasts of different widths from the same SDValue of... (diff) | |
download | llvm-project-aa29049404efdc0134066839bc14d135d69ec225.tar.gz llvm-project-aa29049404efdc0134066839bc14d135d69ec225.tar.bz2 llvm-project-aa29049404efdc0134066839bc14d135d69ec225.zip |
[CUDA] Unbreak CUDA compilation with -std=c++20
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.
Differential Revision: https://reviews.llvm.org/D91807
(cherry picked from commit 9a465057a64dba8a8614424d26136f5c0452bcc3)
-rw-r--r-- | clang/lib/Headers/cuda_wrappers/new | 38 |
1 files changed, 24 insertions, 14 deletions
diff --git a/clang/lib/Headers/cuda_wrappers/new b/clang/lib/Headers/cuda_wrappers/new index f49811c5a57c..47690f1152fe 100644 --- a/clang/lib/Headers/cuda_wrappers/new +++ b/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 |