Registers occupation in __device__ functions

Accelerated Computing CUDA CUDA Programming and Performance
1

If my kernel invokes a device function to which it passes certain parameters, does this function need additional memory registers to store these parameters, or do they reside in the same registers as in the kernel.
In other words, is code 1 equivalent to code 2 in terms of registers occupation?

Code 1:
global void myKernel(int a, intb){
int c;

c = a + b;

}

Code 2:
device int sum(int s1, int s2){
int result;

result = s1+s2;
return(result);
}

global void myKernel(int a, intb){
int c;

c = sum(a, B);

}

2

So far every device function has been inlined by the compiler for me.
Check your .ptx file, int sum() shouldn’t even show up there.

3

The compiler performs extremely aggressive register usage reductions. You shouldn’t see any increase in register usage by calling device functions. I have one unrolled loop that calls a device function 27 times, and the total register usage of the kernel is only 22.