I think I hit an nvcc bug… some code which works in toolkit 2.3 won’t even compile in 3.0. The difference seems to be some compiler confusion when two different functions both declare a pointer to the dynamic shared memory using the [font=“Arial Black”]extern shared int mydata;[/font] style of declaration. In toolkit 2.3, these declarations are local, but in 3.0 they seem to have become global and therefore can collide and kill the compiler.
I use this in templates a lot when I know the first (n) bytes of my dynamic shared memory is scratch space, so I can call functions which use that dynamic buffer for their own temporary storage.
This is the troubling definition.
template<typename T>
__device__ T sumReduce(T val)
{
extern __shared__ T array[];
....
}
Attached is a short demo program that sums the numbers from 0 to 255 both as ints and floats. It compiles and runs fine in toolkit 2.3 but in toolkit 3.0 you get a compilation error:
*** start of expression ***
[lvalue] variable: name = array, node type: array [0] of float
*** end of expression ***
externtemplate.cu(23): internal error: check_type_of_variable_node: enk_variable has wrong type
1 catastrophic error detected in the compilation of "/tmp/tmpxft_00001b37_00000000-4_externtemplate.cpp1.ii".
I was about to simply report this to NV as a bug, but before I did, I realized that the description of the extern shared declaration in the programming guide doesn’t really talk about what namespace the declaration lives in. Perhaps my scratch space definitions themselves are illegal but they just happen to work in 2.3?
externtemplate.cu (2.16 KB)