   {
     if (tid < n_direct) 
     {
-      const float4 posj = body_pos[direct[tid]];
+      const float4 posj = body_pos[directM[tid]];
 #if 0
       const float4 posj  = tex1Dfetch(texBody, direct[tid]);
 #endif
       sh_mass[tid] = posj.w;
       sh_pos [tid] = make_float3(posj.x, posj.y, posj.z);
-      sh_jid [tid] = direct[tid];
     } else {
       sh_mass[tid] = 0.0f;
       sh_pos [tid] = make_float3(1.0e10f, 1.0e10f, 1.0e10f);
-      sh_jid [tid] = -1;
     }
 
     __syncthreads();
 #pragma unroll
     for (int j = 0; j < DIMx; j++) 
-      if ((sh_jid[offs + j] >= 0)) 
-        acc_i = add_acc(acc_i, pos_i, sh_mass[offs + j], sh_pos[offs + j], eps2);
+      acc_i = add_acc(acc_i, pos_i, sh_mass[offs + j], sh_pos[offs + j], eps2);
 #if 0
     direCount += DIMx;
 #endif

        /******       DIRECT          ******/
        /***********************************/

        int *sh_body = &approx[DIM];

        flag         = split && leaf && use_node;                                //flag = split + leaf + use_node
        int  jbody   = node_data & BODYMASK;                                     //the first body in the leaf
        int  nbody   = (((node_data & INVBMASK) >> LEAFBIT)+1) & BTEST(flag);    //number of bodies in the leaf masked with the flag

        body_list[tid] = direct[tid];                                            //copy list of bodies from previous pass to body_list
        sh_body  [tid] = jbody;                                                  //store the leafs first body id into shared memory



  {
    if (tid < n_approx) 
    {
      const int address = (approxM[tid] << 1) + approxM[tid];
#ifndef TEXTURES
      float4 monopole  = multipole_data[address    ];
      float4 octopole0 = multipole_data[address + 1];
