diff options
author | 2020-12-04 11:27:39 -0800 | |
---|---|---|
committer | 2020-12-09 12:42:33 -0500 | |
commit | 59012b685fd69d7350eb55166a8817688e413db8 (patch) | |
tree | acc9f126f79b09292254e2b5b3e190e5d53fe048 | |
parent | [CUDA] Unbreak CUDA compilation with -std=c++20 (diff) | |
download | llvm-project-59012b685fd69d7350eb55166a8817688e413db8.tar.gz llvm-project-59012b685fd69d7350eb55166a8817688e413db8.tar.bz2 llvm-project-59012b685fd69d7350eb55166a8817688e413db8.zip |
[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)
-rw-r--r-- | clang/lib/Headers/cuda_wrappers/new | 46 |
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 |