   
   const int SIZE = 1 << SIZE2; 
 
-#if 0
+#if 0  /* this one forces lmem usage to 236 bytes of lmem */
   for (int i = 1; i <= SIZE; i <<= 1) 
     value += __shfl_up(value, i, SIZE) & BTEST(laneId >= i && i <= distance);
-#else
+#else  /* this one uses only 40 bytes of lmem, any idea why ?!? */
   for (int i = 0; i < SIZE2; i++)
     value += __shfl_up(value, 1 << i, SIZE) & BTEST(laneId >= (1<<i)) & BTEST((1<<i) <= distance);
 #endif



    const int blockDim2 = NTHREAD2;
    __shared__ int shmem[11*(1 << blockDim2)];
    //    __shared__ int shmem[24*(1 << blockDim2)]; is possible on FERMI
    //    int             lmem[LMEM_STACK_SIZE];



  int *directS = shmem;                              //  0*DIM,  1*DIM,  1*DIM
  int *nodesS  = directS + DIM;                      //  1*DIM, 10*DIM,  9*DIM
  int *prefix  = nodesS  + DIM*8;                    //  9*DIM, 10*DIM,  1*DIM

  const int NJMAX = DIM*3;
  int    *body_list = (int*   )&nodesS   [  DIM]; //  2*DIM,   5*DIM,  2*DIM
