Help from coders

[ET_]

I have this snippet of code:

__global__ void stuff(float a)
{
    a = 9;
}

called from this function:

void caller(float &a)
{
    cout >> a;  //actually a has value 3

    stuff<<<blocks, threads>>>(a);

    cout >> a;  // a has still value 9
}

that doesn't work.

As I am not familiar with C++ parameter passing, I ask: how would the code look like to permit to the __global__ stuff to modify the "a" value in the caller scope?
I tried all thepossible exchanges with "*" and "&" operators, but I can't find a way.

Would you help me?

Thanks

Luigi

[Qubit]

Proper disclosure: I am absolutely unfamiliar with CUDA, but anyway this seems like a C++ question.

It seems as if it should print 3 instead of 9, but the desired output is unclear from your wording.

Anyway, if instead you write

__global__ void stuff(float& a)

it should print 9.

The '&' is the reference operator in C++. Writing "float &a"/"float& a"/"float & a" indicates that you're passing the parameter a by reference.
Without the reference operator, you'd be passing the parameter a by value, which will essentially create a copy of a which would be destroyed upon returning from the function.

[chalsall]

Originally Posted by ET_

Would you help me?

I'll try... Give:

...
stuff<<<blocks, threads>>>(&a);
...

and

__global__ void stuff(float *a)
{
    *a = 9;
}

...a try. While your function "stuff" may be defined as "global", your variable "a" was passed in as a scaler to the function.

At least under "C", "&" says "give me the address of this object", while "*" says "give me the object this address points to".

Of course, I may very well be completely wrong.

[TheJudger]

Hello Luigi,

unless my mind is screwed up:
CPU vs. GPU context!
Executing a GPU kernel from CPU context is asynchronous (CPU continues before GPU has finished the kernel), thus never can return values to the caller.
I guess you'll need explicit

• synchronisation (let the CPU wait for the GPU to finish)
• data transfers CPU<->GPU

when using the CUDA API.

Oliver

[jasonp]

Communication from CPU to GPU and back has to go through memory; just passing the address of something to a GPU will not work, because CPU and GPU have separate address spaces. If you want a GPU to write something and then have the CPU see it, then you have to do a memory transfer into the CPU.

Except when using the latest CUDA toolkit, where you can force a range of memory to be identical between CPU and GPU, so that either one automatically can see the latest version of something in that memory range, no matter who wrote it. I believe this only works on linux though.

[ET_]

@Chalsall: That was my reasoning also, but the __global__ functions on device don't allow reference arguments. That's the reason of my question

@Oliver: It was not a matter of synchronization of threads, your hint didn't work either.

@Jason: Correct. I cudaMemcopied the value to the device, and after some struggling with pointers it worked!

Thanks to you all!

Luigi

P.S. And yes, I made typo on my forst message . Before che change, the result in main SHOULD be 9, but it was still 3 instead. Sorry for the confusion.