| // Check that operator new and operator delete work. |
| |
| #include <assert.h> |
| #include <new> |
| #include <stdio.h> |
| |
| // We can not use host-side variable std::nothrow. |
| // Re-declare it as a __device__ variable. |
| // TODO: do that in clang's CUDA header wrappers. |
| # if __cplusplus >= 201103L |
| extern decltype(std::nothrow) __device__ std::nothrow; |
| #define NOTHROW_AVAILABLE 1 |
| #else |
| #define NOTHROW_AVAILABLE 0 |
| #endif |
| |
| __device__ void global_new() { |
| void* x = ::operator new(42); |
| assert(x != NULL); |
| ::operator delete(x); |
| |
| #if NOTHROW_AVAILABLE |
| x = ::operator new(42, std::nothrow); |
| assert(x != NULL); |
| ::operator delete(x, std::nothrow); |
| #endif |
| |
| x = ::operator new[](42); |
| assert(x != NULL); |
| ::operator delete[](x); |
| |
| #if NOTHROW_AVAILABLE |
| x = ::operator new[](42, std::nothrow); |
| assert(x != NULL); |
| ::operator delete[](x, std::nothrow); |
| #endif |
| } |
| |
| __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<<<1, 1>>>(); |
| cudaError_t err = cudaDeviceSynchronize(); |
| if (err != cudaSuccess) { |
| printf("CUDA error %d\n", (int)err); |
| return 1; |
| } |
| printf("Success!\n"); |
| return 0; |
| } |