Author: jlebar Date: Tue Aug 9 20:09:14 2016 New Revision: 278194 URL: http://llvm.org/viewvc/llvm-project?rev=278194&view=rev Log: [CUDA] Add __device__ overloads for placement new and delete.
Summary: Previously these sort of worked because they didn't end up resulting in calls at the ptx layer. But I'm adding stricter checks that break placement new without these changes. Reviewers: tra Subscribers: cfe-commits Differential Revision: https://reviews.llvm.org/D23239 Modified: cfe/trunk/lib/Headers/__clang_cuda_runtime_wrapper.h Modified: cfe/trunk/lib/Headers/__clang_cuda_runtime_wrapper.h URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Headers/__clang_cuda_runtime_wrapper.h?rev=278194&r1=278193&r2=278194&view=diff ============================================================================== --- cfe/trunk/lib/Headers/__clang_cuda_runtime_wrapper.h (original) +++ cfe/trunk/lib/Headers/__clang_cuda_runtime_wrapper.h Tue Aug 9 20:09:14 2016 @@ -312,5 +312,23 @@ __device__ inline __cuda_builtin_gridDim #pragma pop_macro("uint3") #pragma pop_macro("__USE_FAST_MATH__") +// Device overrides for placement new and delete. +#pragma push_macro("CUDA_NOEXCEPT") +#if __cplusplus >= 201103L +#define CUDA_NOEXCEPT noexcept +#else +#define CUDA_NOEXCEPT +#endif + +__device__ inline void *operator new(__SIZE_TYPE__, void *__ptr) CUDA_NOEXCEPT { + return __ptr; +} +__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 {} +#pragma pop_macro("CUDA_NOEXCEPT") + #endif // __CUDA__ #endif // __CLANG_CUDA_RUNTIME_WRAPPER_H__ _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
