 #define WARP_SIZE2 5
 #define WARP_SIZE  32
 
-#if 1
+#if 0
+#define tid    (threadIdx.x)
+#else
+#define tid    (threadIdx.y*blockDim.x + threadIdx.x)
+#endif
 #define laneId (threadIdx.x & (WARP_SIZE - 1))
 #define warpId (threadIdx.x >> WARP_SIZE2)
-#endif
 
 __forceinline__ __device__ float Wkernel(const float q)
 {

         ***/


        int n_total = calc_prefix<DIM2>(prefix, tid,  nchild);              // inclusive scan to compute memory offset of each child (return total # of children)
        int offset  = prefix[tid];
        offset     += n_offset - nchild;                                  // convert inclusive into exclusive scan for referencing purpose
        __syncthreads();                                                   // thread barrier


        lmem = &MEM_BUF[gridDim.x*(LMEM_STACK_SIZE*blockDim.x + LMEM_EXTRA_SIZE)];    //Use the extra large buffer
        apprCount = direCount = 0;
        acc_i = approximate_gravity<blockDim2, 8>( DIM2x, DIM2y, tx, ty,
            body_i, pos_i, group_pos,
            eps2, node_begend,
            multipole_data, body_pos,
