     return x;
   }
 }
-#endif
+
+
+template<int DIM2>
+__device__ int calc_prefix(int N, int* prefix_in, int tid) 
+{
+  const int DIM = 1 << DIM2;
+  const int N2  = N >> DIM2;
+
+  int y = calc_prefix<DIM2>(prefix_in, tid, prefix_in[tid]);
+  if (N2 == 0) return y;
+
+  for (int p = DIM; p < N; p += DIM) 
+  {
+    int *prefix = &prefix_in[p];
+    const int y1 = calc_prefix<DIM2>(prefix, tid, prefix[tid]);
+    prefix[tid] += y;
+    y += y1;
+  }
+  __syncthreads();
+
+  return y;
+} 
+
 
 
   template<int SHIFT>

  x = prefix[DIM - 1];
  __syncthreads();
  return x;
#else

  int offset = 0;
  int tid2 = tid << 1;

#pragma unroll
  for (int d = DIM >> 1; d > 0; d >>= 1) {
    __syncthreads();

    int iflag = BTEST(tid < d);
    int ai = (((tid2 + 1) << offset) - 1) & iflag;
    int bi = (((tid2 + 2) << offset) - 1) & iflag;

    prefix[bi] += prefix[ai] & iflag;
    offset++;
  }

  // clear the last element
  if (tid == 0) prefix[DIM - 1] = 0;

  // traverse down the tree building the scan in place
#pragma unroll
  for (int d = 1; d < DIM; d <<= 1) {
    offset--;
    __syncthreads();

    int iflag = BTEST(tid < d);
    int ai = (((tid2 + 1) << offset) - 1) & iflag;
    int bi = (((tid2 + 2) << offset) - 1) & iflag;

    int t       = prefix[ai];
    if (tid < d) {
      prefix[ai]  = (prefix[bi] & iflag) + (t & BTEST(tid >= d));
      prefix[bi] += t & iflag;
    }
  }
  __syncthreads();

  prefix[tid] += value;
  __syncthreads();

  x = prefix[DIM - 1];
  __syncthreads();
  return x;
#endif
}

#if 1
__device__ __forceinline__ uint shfl_scan_add_step(uint partial, uint up_offset)
{
  uint result;

        if (flag && nodesM[offset + 6] == 0) nodesM[offset + 6] = child + 6;
        if (flag && nodesM[offset + 7] == 0) nodesM[offset + 7] = child + 7;
        __syncthreads();

        n_offset += n_total;    //Increase the offset in the array by the number of newly added nodes

