 #ifdef OLDPREFIX
         calc_prefix<DIM2>(prefix0, tid, flag);
 #else
-        inclusive_scan_block<ADDOP<int>, int>(prefix0, (int)flag, tid);       // inclusive scan on flags to construct array
+#if 0
+        inclusive_scan_block<ADDOP<int>, int, DIM2>(prefix0, (int)flag, tid);       // inclusive scan on flags to construct array
+#else
+        v0 = inclusive_scan_block<ADDOP<int>, int, DIM2>((int)flag, tid);       // inclusive scan on flags to construct array
+        prefix0[tid] = v0;
+        __syncthreads();
+#endif
 #endif
 
         if (flag) prefix1[prefix0[tid] - 1] = tid;                             //with tidś whose leaves have to be opened

        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

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

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