From: Justin Lebar <jlebar@google.com>
Date: Wed, 10 Aug 2016 01:09:14 +0000 (+0000)
Subject: [CUDA] Add __device__ overloads for placement new and delete.
X-Git-Url: https://granicus.if.org/sourcecode?a=commitdiff_plain;h=984f42a083bf8d99237355fc8b83830cdd12bea8;p=clang

[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

git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@278194 91177308-0d34-0410-b5e6-96231b3b80d8
---

diff --git a/lib/Headers/__clang_cuda_runtime_wrapper.h b/lib/Headers/__clang_cuda_runtime_wrapper.h
index 6445f9b76b..05a85fa3d5 100644
--- a/lib/Headers/__clang_cuda_runtime_wrapper.h
+++ b/lib/Headers/__clang_cuda_runtime_wrapper.h
@@ -312,5 +312,23 @@ __device__ inline __cuda_builtin_gridDim_t::operator dim3() const {
 #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__