OSchip
/
llvm-project
mirror of https://github.com/THU-DSP-LAB/llvm-project.git
13
0
Fork 6
Code Issues Packages Projects Releases Wiki Activity

[CUDA] Another attempt to fix early inclusion of <new> from libstdc++ 

Previous patch (9a465057a6) 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
This commit is contained in:
Artem Belevich 2020-12-04 11:27:39 -08:00
parent 4b5dc150b9
commit 4326792942
1 changed files with 22 additions and 24 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

46
clang/lib/Headers/cuda_wrappers/new
View File

@ -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