[cfe-users] Operator new in CUDA kernels

[Ralph Kube via cfe-users]

Hi,
I’m trying to port my CUDA project to clang. Thanks for the fantastic work,
clang is a charm to work with.

One problem however, maybe you guys can help me out.

I need to allocate memory for a class in a CUDA kernel on the heap. 
So I created a myclass** in the host program, pass it to the kernel and attempt to
allocate with new. Works fine in nvcc, in clang I get the error

[1] % clang++ -std=c++14 -o test_new_device test_new_device.cu -L/Developer/NVIDIA/CUDA-8.0/lib -lcudart
ptxas fatal   : Unresolved extern function '_Znwm'
clang-4.0: error: ptxas command failed with exit code 255 (use -v to see invocation)

I’m on osx 10.12.5, using
% clang++ --version
clang version 4.0.0 (tags/RELEASE_400/final 297808)
Target: x86_64-apple-darwin16.6.0
Thread model: posix

My analysis of the situation is:
* :: operator new() gets resolved wrong and tries to call the host-side malloc.
* ptxas cannot resolve the invoked host-side malloc _Znwm as a device function.

So, is there a way to tell clang that ::operator new() in device function should call
the device version of malloc. As described here:
http://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#memory-allocation-and-lifetime



Below is a minimal program that reproduces the error.


#include <iostream>
#include <cuda_runtime_api.h>

class myclass
{
public:
    __host__ __device__ myclass(const double _data) : data(_data) {}
    __host__ __device__ ~myclass() 
    {   
        printf("Deleting myclass\n");
    }   
    __host__ __device__ double get_data() const {return(data);}

private:
    const double data;
};



__global__
void init_myclass(myclass** mycs_ptr)
{
    (*mycs_ptr) = new myclass(14.0);
}


__global__
void access_myclass(myclass** mycs_ptr)
{
    printf("I am using data with value = %f\n", (*mycs_ptr) -> get_data());
}

__global__
void delete_myclass(myclass** mycs_ptr)
{
    delete (*mycs_ptr);
}


int main(void)
{
    myclass** myclass_ptr{nullptr};

        
    init_myclass<<<1, 1>>>(myclass_ptr);

    access_myclass<<<1, 1>>>(myclass_ptr);

    delete_myclass<<<1, 1>>>(myclass_ptr);

    return(0);
}