Overloading new[] doesn't work in device code

Hi there,

I’m writing some container classes for the device side, and I’ve stumbled upon this error. When compiling this code:


#include "stdlib.h"

#include "stdio.h"

__device__ inline

void *operator new(size_t size) {return malloc(size);}

__device__ __host__ inline

void operator delete(void *p) {free(p);}

__device__ inline

void *operator new[](size_t size) {return malloc(size);}

__device__ inline

void operator delete[](void *p) {free(p);}

__global__ void arraynew()


  char* ptr = new char[100];

  printf("ptr is %p\n", ptr);

  delete[] ptr;

struct Bucket {

//default constructor

    __device__ inline Bucket() : key(0), data(0){}

    __device__ inline Bucket(int k, int d) : key(k), data(d) {}

int key;

    int data;


Bucket* ptr2 = new Bucket[100];

  printf("ptr2 is %p\n", ptr2);

  delete[] ptr2;


int main()





It first looks as if I am able to overload the new operator but then I get:

Centauri-2:test malfunct$ nvcc -arch sm_20 arraynew.cu -keep

./arraynew.cu(32): Error: External calls are not supported (found non-inlined call to __cxa_vec_new2)

Centauri-2:test malfunct$ grep __cxa_vec_new2 *

arraynew.cpp2.i:__cuda_local_var_28154_11_non_const_ptr2 = ((struct _ZZ8arraynewvE6Bucket *)(__cxa_vec_new2(100UL, 8UL, 0UL, _ZZ8arraynewvEN6BucketC1Ev, ((void (*)(void *))0), _Znam, _ZdaPv)));

arraynew.cpp3.i:__cuda_local_var_28154_11_non_const_ptr2 = ((struct _ZZ8arraynewvE6Bucket *)(__cxa_vec_new2(100UL, 8UL, 0UL, _ZZ8arraynewvEN6BucketC1Ev, ((void (*)(void *))0), _Znam, _ZdaPv)));

arraynew.cudafe1.c:extern void *__cxa_vec_new2();

arraynew.cudafe1.gpu:__cuda_local_var_28154_11_non_const_ptr2 = ((struct _ZZ8arraynewvE6Bucket *)(__cxa_vec_new2(100UL, 8UL, 0UL, _ZZ8arraynewvEN6BucketC1Ev, ((void (*)(void *))0), _Znam, _ZdaPv)));

arraynew.cudafe2.c:extern int __cxa_vec_new2();

arraynew.cudafe2.gpu:__cuda_local_var_28154_11_non_const_ptr2 = ((struct _ZZ8arraynewvE6Bucket *)(__cxa_vec_new2(100UL, 8UL, 0UL, _ZZ8arraynewvEN6BucketC1Ev, ((void (*)(void *))0), _Znam, _ZdaPv)));

Is some workaround possible? I suppose that library routine is not yet included in the device compiler, but I don’t think there is any reason it cannot be implemented.



Neither could I get placement new to work, so I wasn’t able to implement generic container classes. Right now, I’m forcing a default init() member function for used classes but that is not a general or desirable solution.

I have found an acceptable workaround for my code, but I haven’t been able to use overload new, although it is possible to overload the new operator, which seems like an inconsistent behavior.

Do you get the compilation error if you turn off generating GPU debug information?