[CUDA] Tests for device-side overloads of non-placement new/delete.

test-suite/trunk/External/CUDA/CMakeLists.txt

 # Create targets for CUDA tests that are part of the test suite.
 macro(create_local_cuda_tests VariantSuffix)
   create_one_local_test(assert assert.cu)
   create_one_local_test(axpy axpy.cu)
   create_one_local_test(algorithm algorithm.cu)
   create_one_local_test(cmath cmath.cu)
   create_one_local_test(complex complex.cu)
   create_one_local_test(math_h math_h.cu)
+  create_one_local_test(new new.cu)
   create_one_local_test(empty empty.cu)
   create_one_local_test(printf printf.cu)
   create_one_local_test(future future.cu)
 endmacro()
 
 macro(thrust_make_test_name TestName TestSourcePath)
   string(REPLACE "${THRUST_PATH}/testing/" "" _tmp ${TestSourcePath})
   string(REPLACE "/" "-" _tmp ${_tmp})

test-suite/trunk/External/CUDA/new.cu

+// Check that operator new and operator delete work.
+
+#include <assert.h>
+#include <new>
+#include <stdio.h>
+
+__device__ void global_new() {
+  void* x = ::operator new(42);
+  assert(x != NULL);
+  ::operator delete(x);
+
+  x = ::operator new(42, std::nothrow);
+  assert(x != NULL);
+  ::operator delete(x, std::nothrow);
+
+  x = ::operator new[](42);
+  assert(x != NULL);
+  ::operator delete[](x);
+
+  x = ::operator new[](42, std::nothrow);
+  assert(x != NULL);
+  ::operator delete[](x, std::nothrow);
+}
+
+__device__ void sized_delete() {
+#if __cplusplus>= 201402L
+  void* x = ::operator new(42);
+  assert(x != NULL);
+  ::operator delete(x, 42);
+
+  x = ::operator new[](42);
+  assert(x != NULL);
+  ::operator delete[](x, 42);
+#endif
+}
+
+__device__ void int_new() {
+  int* x = new int();
+  assert(*x == 0);
+  delete x;
+}
+
+struct Foo {
+  __device__ Foo() : x(42) {}
+  int x;
+};
+__device__ void class_new() {
+  Foo* foo = new Foo();
+  assert(foo->x == 42);
+  delete foo;
+}
+
+__global__ void kernel() {
+  global_new();
+  sized_delete();
+  int_new();
+  class_new();
+}
+
+int main() {
+  kernel<<<32, 32>>>();
+  cudaError_t err = cudaDeviceSynchronize();
+  if (err != cudaSuccess) {
+    printf("CUDA error %d\n", (int)err);
+    return 1;
+  }
+  printf("Success!\n");
+  return 0;
+}

test-suite/trunk/External/CUDA/new.reference_output

+Success!
+exit 0