         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
+#if 0
+        inclusive_scan_block<ADDOP<int>, int, DIM2>(prefix, nchild, tid);        // inclusive scan to compute memory offset of each child
+#else
+        int v0 = inclusive_scan_block<ADDOP<int>, int, DIM2>(nchild, tid);        // inclusive scan to compute memory offset of each child
+        prefix[tid] = v0;
+        __syncthreads();
+#endif
         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

texture<float4, 1, cudaReadModeElementType> texMultipole;
texture<float4, 1, cudaReadModeElementType> texBody;

template<class T>
 struct ADDOP {
  __device__ static inline T identity()           {return (T)(0);}

          __syncthreads();

          //step 2:
#if 1 //def OLDPREFIX
          int nl = calc_prefix<DIM2>(nb, &body_list[n_direct], tid);
#else
          int nl = inclusive_scan_array<ADDOP<int>, int>              // inclusive scan to compute number of leaves to process
