Question about shared array in kernel


I have a question about declaring an array as shared. I though that the following code should output the same (i.e. all zeros), whether the array hd is defined as shared or as a simple float, but the code with the array declared as shared outputs 0, 1, 2, 3, … Why is this happening? I thought that even when declared as shared, the arrays were still private to each thread. Is this true?


__global__ void test(int n)
  //__shared__ float hd[12];
  float hd[12];
  int t = threadIdx.x;
  int k;
  for ( k=0; k<12; k++ ) *(hd+k) = 0;
  if (t==0) {
      for (k=0; k<12; k++ ) printf("%.10f\n",*(hd+k));
  for (k=0;k<12;k++) *(hd+k) = k;

shared arrays are not private to each thread. That is the whole point of calling them shared. They are shared among all threads in a threadblock.