Overloading new[] doesn't work in device code

Accelerated Computing CUDA CUDA Programming and Performance
1

Hi there,

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

//-*-c++-*-

#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()

{

  arraynew<<<1,1>>>();

  return(0);

}

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.

Best,

Eray

2

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.

3

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.

4

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

5

Yes!