 /********* 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;
 }

        bool leaf       = node_pos.w <= 0;  //Small AND equal incase of a 1 particle cell       //Check if it is a leaf
        //         split = true;

        uint mask    = BTEST((split && !leaf) && use_node);               // mask = #FFFFFFFF if use_node+split+not_a_leaf==true, otherwise zero
        int child    =    node_data & 0x0FFFFFFF;                         //Index to the first child of the node
        int nchild   = (((node_data & 0xF0000000) >> 28)) & mask;         //The number of children this node has




template<const int SHIFT, const int BLOCKDIM2, const int NI>
__device__ 
#if 0
__noinline__
#else
__forceinline__ 
#endif
void approximate_gravity(
    float4 pos_i[NI],
    real4 group_pos,
    float eps2,
