          ***/
 
 
-#ifdef OLDPREFIX
-        int n_total = calc_prefix<DIM2>(prefix, tid,  nchild);
-        prefix[tid] += n_offset - nchild;
-        __syncthreads();
-#else
-        inclusive_scan_block<ADDOP<int>, int>(prefix, nchild, tid);        // inclusive scan to compute memory offset of each child
-        int n_total = prefix[blockDim.x - 1];                              // fetch total number of children, i.e. offset of the last child -1
-        __syncthreads();                                                   // thread barrier to make sure that warps completed their jobs
-        prefix[tid] += n_offset - nchild;                                  // convert inclusive into exclusive scan for referencing purpose
+        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
-#endif
 
         for (int i = n_offset; i < n_offset + n_total; i += DIM)         //nullify part of the array that will be filled with children
           nodesM[tid + i] = 0;                                          //but do not touch those parts which has already been filled

  __syncthreads();
  return x;
#else
  
  int offset = 0;
  int tid2 = tid << 1;




#define TEXTURES
#if 0
#define _ORIG_SHMEM_
#endif
