Author: Artem Belevich Date: 2020-12-09T12:42:33-05:00 New Revision: 59012b685fd69d7350eb55166a8817688e413db8
URL: https://github.com/llvm/llvm-project/commit/59012b685fd69d7350eb55166a8817688e413db8 DIFF: https://github.com/llvm/llvm-project/commit/59012b685fd69d7350eb55166a8817688e413db8.diff LOG: [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) Added: Modified: clang/lib/Headers/cuda_wrappers/new Removed: ################################################################################ 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 _______________________________________________ llvm-branch-commits mailing list llvm-branch-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/llvm-branch-commits