       int direCount = 0;
 
 
-      acc_i = approximate_gravity<blockDim2, 0>( DIM2x, DIM2y, tid, tx, ty,
+      acc_i = approximate_gravity<blockDim2, 0>( DIM2x, DIM2y, tx, ty,
           body_i, pos_i, group_pos,
           eps2, node_begend,
           multipole_data, body_pos,

#define BTEST(x) (-(int)(x))

  template<int DIM2>
__device__ int calc_prefix1(int* prefix, int tid, int value)
{
  int  x;


  return mysum;
}

#else /* uses code from CUDA 4.2 Manual Appendix B13 */

  template <const int BLOCKDIM2>
__device__ __forceinline__ int inclusive_scan_warp(int value)
{
  const int BLOCKDIM = 1 << BLOCKDIM2;
  for (int i=1; i<=BLOCKDIM-1; i <<= 1)
  {
#if 0  /* uses *if* version, generate too manu spills, 240 bytes */
    int n = __shfl_up(value, i, BLOCKDIM);
    if (laneId >= i)
      value += n;
#else  /* uses masking via BTEST, only 32 bytes spills to lmem */
    value += __shfl_up(value, i, BLOCKDIM) & BTEST(laneId >= i);
#endif
  }

  return value;
}
#endif /* inlined assebly */

template<int DIM2>
__device__ __forceinline__ int calc_prefix(int* prefix, int value) 
{
  if (DIM2 != 6)  /* should never be called */
    return calc_prefix1<DIM2>(prefix, value);
  else
  {
    prefix[tid] = inclusive_scan_warp<WARP_SIZE2>(value);
