[cfe-users] Operator new in CUDA kernels

[Justin Lebar via cfe-users]

Thank you for the bug report.  This looks like a real bug, we should
fix it.  I will add it to our list.

Sorry for the long delay in getting back to you -- I don't read the
mailing lists habitually, and just happened to do a search today that
showed me that I'd missed many emails I should have responded to.

-Justin

On Sun, Jun 18, 2017 at 4:11 PM, Ralph Kube via cfe-users
<cfe-users at lists.llvm.org> wrote:
> 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);
> }
> _______________________________________________
> cfe-users mailing list
> cfe-users at lists.llvm.org
> http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-users