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