Unified Memory in CUDA 6

Technical Blogs & Events Technical Blog
81

How come deriving from Managed does String becomes pass-by-reference allowed? I know I'm missing something important.
On the other hand, DataElem::name of type String is not a pointer (i.e. not String*), therefore when we do DataElem* elem = new DataElem we are not allocating memory for DataElem::name::data.
So how/when does that allocation happens in the GPU?

82

I think the examples in the post are a bit simplified. You can see the real code here and run it yourself: https://github.com/parallel...

83

I tried it and get an error 77 (illegal address) after executing launch_by_pointer(e);
My guess is, it is caused to String::data not being allocated using cudaMallocManaged becuse in the DataElem we have 'String name' instead of 'String* name'.

84

I changed the code to reflect what I suggest in my previous post. It works fine for launch_by_pointer and launch_by_ref but not for launch_by_value. I have a page fault error when String::~String() is called. I put your code modified into a personal Github in case you happen to have the change to take a look at it.
https://github.com/BRabbit2...

Thanks for the help !

85

Just in case someone runs into the same issue, I just tried Mark's code in a P100 and works just fine, no errors. I need to re-read the Unified Memory documentation because the same code in a GTX 760 I still get the Illegal Address errors (cudaError = 77) but if I use the heap to allocate memory for the name (String* name) I get a memory page error when calling the destructor from the launch_by_value. I'll try to find out why and post back.

86

I am new to CUDA but experienced with C++. I find CUDA very promising,
I suggest NVIDIA should make CUDA interface more STL-like instead of strange syntax like “func<<…>>(…)". You could wrap all the functionalities into an STL-like library and do not introduce a foreign syntax into C++. If you do that you do not also need a CUDA compiler.
Of course I might be missing something, but I would like very much to understand, if I do.

87

Hi, please check, I think there's a typo here:
// pass data to kernel by reference
kernel_by_ref<<<1,1>>>(*data); -> (&data); ?
Thanks!

88

data is a pointer. So to pass by reference you have to dereference it by applying the *. Taking the address with & and passing that will result in a compiler error.