Difference between raw pointer and reference

  1. Is there no significant overhead of using pointer vs reference in global and device kernels?

  2. Can references benefit from

__restrict__

as raw pointers do?

__device__ void foo( int& a, int& b){
...
}

You generally can’t use a reference as a kernel argument, unless you are using Unified Memory.

tensorflow seems to heavily use references in device and global kernels. Could you point me to a good resource of how / when to use reference? I could only find very brief mentions of reference in CUDA C programming guide.

reference is a concept that is not specific to CUDA. It is part of C++ like many other things, and CUDA mostly adheres to C++.

You’re not going to find a tutorial on references in CUDA just like you’re not going to find a tutorial on function overloading in CUDA. It’s part of C++. It should behave approximately the same way.

The general usage of references in CUDA C++, and the motivating reasons to use them, aren’t going to be any different than what they are in C++.

https://stackoverflow.com/questions/12728794/why-should-i-use-reference-variables-at-all

I’m not aware of any claims that CUDA makes about using restrict with references, or even if it is allowed or even if it makes sense.

A restrict decorator is a contract that you are making with the compiler, that essentially says that you won’t use pointer aliasing. But a reference is an alias for the variable that it is a reference to. So I’m not sure using the two together makes sense.

Since pass by value to a function makes a copy of the variable, one of the important benefits of using references is that you don’t need to make unnecessary copies. It also sorts out some vexing use cases where you would like the function to be able to modify the contents of an object or variable in the calling environment, without having to resort to using pointer-to-pointer (which is the C style approach to sorting that out.