Pointer of local variable can not be send to nested kernel?

Accelerated Computing CUDA CUDA Programming and Performance
1

global void B(float3* pointer)
{
pointer; // address is 0xxxxxxff7800, content is empty rather than f3
}
global void A()
{
…
float3 f3[3] = {…}; // f3 address is: 00ff7800
B<<<1, 1>>>(f3); // wht can not do this?
}

but the guide says parent kernel exits after all sub-kernel exiting. why the local memory f3 seems released?

2

9.2.2. Memory Model
Parent and child grids share the same global and constant memory storage, but have distinct local and shared memory.

You cannot share pointers to local variables between kernels.

1 Like
3

Alternatives are to use global memory (possibly with a call-specific index)
or to call __device__ functions instead of other __global__ sub-kernels from your kernel.