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@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@lists.llvm.org > http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-users _______________________________________________ cfe-users mailing list cfe-users@lists.llvm.org http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-users
