[CUDA] Add __device__ overloads for placement new and delete.

[jlebar]

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.

[tra]

I think we need to add noexcept for these in c++11.

[jlebar]

Add noexcept.

[jlebar]

In D23239#508863, @tra wrote:

I think we need to add noexcept for these in c++11.

Argh, you are completely right.

Thank you for the review.

[tra]

LGTM.

cfe/trunk/lib/Headers/__clang_cuda_runtime_wrapper.h

 #pragma push_macro("uint3")
 #define dim3 __cuda_builtin_blockDim_t
 #define uint3 __cuda_builtin_threadIdx_t
 #include "curand_mtgp32_kernel.h"
 #pragma pop_macro("dim3")
 #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__