When I compile another nonsense program:
__shared__ int shared[8];
__device__ void share1(void)
{
int i = threadIdx.x;
__shared__ int j[8];
j[i] = shared[i+1];
shared[i] = j[i+1];
}
__device__ void share2(void)
{
int i = threadIdx.x;
__shared__ int j[8];
j[i] = shared[i+1];
shared[i] += j[i+1];
}
typedef struct {void (*fn)();} Test;
__constant__ Test t[2] = {share1, share2};
__global__ void kernel(int i)
{
t[i].fn();
}
I get the error message from nvcc:
Now that I understand that there is no stack on the device, and hence no functions, and nvcc inlines everything, this sounds reasonable. I really would have saved some head scratching if you told us up front in as many words. Also a more friendly error message would help!
Thanks, Eric