 
 
 template<int DIM2, int SHIFT>
-__device__ float4 approximate_gravity(int DIM2x, int DIM2y,
+__device__ __forceinline__ float4 approximate_gravity(int DIM2x, int DIM2y,
     int tx, int ty,
     int body_i, float4 pos_i,
     real4 group_pos,

#pragma unroll 16
            for (int j = 0; j < DIMx; j++)
              acc_i = add_acc(acc_i, pos_i, sh_mass[offs + j], sh_pos[offs + j], eps2);
#if 0
            direCount += DIMx;
#endif
            __syncthreads();
#endif
          }

#pragma unroll
    for (int j = 0; j < DIMx; j++) 
      acc_i = add_acc(acc_i, pos_i, sh_mass[offs + j], sh_pos[offs + j], eps2);
    direCount += DIMx;
    __syncthreads();
  }

