 /********* SEGMENTED SCAN ***********/
 /************************************/
 
+__device__ __forceinline__ int ShflSegScanStepB(
+            int partial,
+            uint distance,
+            uint up_offset)
+{
+  asm(
+      "{.reg .u32 r0;"
+      ".reg .pred p;"
+      "shfl.up.b32 r0, %1, %2, 0;"
+      "setp.le.u32 p, %2, %3;"
+      "@p add.u32 %1, r0, %1;"
+      "mov.u32 %0, %1;}"
+      : "=r"(partial) : "r"(partial), "r"(up_offset), "r"(distance));
+  return partial;
+}
   template<const int SIZE2>
 __device__ __forceinline__ int inclusive_segscan_warp(int value, const int distance)
 {
 
+#if 0
   const int SIZE = 1 << SIZE2; 
 
   for (int i = 0; i < SIZE2; i++)
     value += __shfl_up(value, 1 << i, SIZE) & BTEST(laneId >= (1<<i)) & BTEST((1<<i) <= distance);
+#else
+  for (int i = 0; i < SIZE2; i++)
+    value = ShflSegScanStepB(value, distance, 1<<i);
+#endif
 
   return value;
 }

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

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



        int n_total = calc_prefix(prefix,  nchild);              // inclusive scan to compute memory offset of each child (return total # of children)
        int offset  = prefix[laneId];
        for (int i = n_offset; i < n_offset + n_total; i += DIM)         //nullify part of the array that will be filled with children
          nodesM[laneId + i] = 0;                                          //but do not touch those parts which has already been filled
#if 1
        offset     += n_offset - nchild;                                  // convert inclusive into exclusive scan for referencing purpose

        if (flag) nodesM[offset] = child;                            //Thread with the node that is about to be split
        //writes the first child in the array of nodes
        /*** in the following 8 lines, we calculate indexes of all the children that have to be walked from the index of the first child***/
