Spill stores, spill loads

Accelerated Computing CUDA CUDA Programming and Performance
1

Hi guys,
I have some questions.
If I am right spill stores and spill loads means the compiler is spilling register values to local memory. When compiling the source file I get this ptxas info:

3>  ptxas : info : Function properties for _ZN6Helper7getFluxEff6float3S0_S0_
3>      168 bytes stack frame, 24 bytes spill stores, 24 bytes spill loads

__device__ float3 Helper::getFlux(const float ap, const float am, const float3 fp, const float3 fm, const float3 diff)
{
	if (ap - am > 0.0f)
	{
		return (ap * fm - am * fp + ap * am * diff) / (ap - am);
	}
	else
	{
		return make_float3(.0f, .0f, .0f);
	}
}

There are 24 bytes spill stores, 24 bytes spill loads, how should I interpet this?
Why is the compiler spilling registers?
Can it be avoided somehow or that’s ok?
Why has the stack frame 168 bytes?

And one more question.
How is pass-by-reference working on the GPU. It is using shared memory in this case too?
Thank you for the help.

3 Likes