jlebar created this revision.
jlebar added a reviewer: tra.
jlebar added a subscriber: cfe-commits.
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.
https://reviews.llvm.org/D23239
Files:
clang/lib/Headers/__clang_cuda_runtime_wrapper.h
Index: clang/lib/Headers/__clang_cuda_runtime_wrapper.h
===================================================================
--- clang/lib/Headers/__clang_cuda_runtime_wrapper.h
+++ clang/lib/Headers/__clang_cuda_runtime_wrapper.h
@@ -312,5 +312,15 @@
#pragma pop_macro("uint3")
#pragma pop_macro("__USE_FAST_MATH__")
+// Device overrides for placement new and delete.
+__device__ inline void *operator new(__SIZE_TYPE__, void *__ptr) {
+ return __ptr;
+}
+__device__ inline void *operator new[](__SIZE_TYPE__, void *__ptr) {
+ return __ptr;
+}
+__device__ inline void operator delete(void *, void *) {}
+__device__ inline void operator delete[](void *, void *) {}
+
#endif // __CUDA__
#endif // __CLANG_CUDA_RUNTIME_WRAPPER_H__
Index: clang/lib/Headers/__clang_cuda_runtime_wrapper.h
===================================================================
--- clang/lib/Headers/__clang_cuda_runtime_wrapper.h
+++ clang/lib/Headers/__clang_cuda_runtime_wrapper.h
@@ -312,5 +312,15 @@
#pragma pop_macro("uint3")
#pragma pop_macro("__USE_FAST_MATH__")
+// Device overrides for placement new and delete.
+__device__ inline void *operator new(__SIZE_TYPE__, void *__ptr) {
+ return __ptr;
+}
+__device__ inline void *operator new[](__SIZE_TYPE__, void *__ptr) {
+ return __ptr;
+}
+__device__ inline void operator delete(void *, void *) {}
+__device__ inline void operator delete[](void *, void *) {}
+
#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
