           {
             //We overwrote our current stack
             apprCount = -1; 
-            return acc_i;	 
+            return;
           }
         }
 

      real4   *body_vel,
      int     *MEM_BUF) 
{
    const int blockDim2 = NTHREAD2;
    const int shMemSize = 10 * (1 << blockDim2);
    __shared__ int shmem_pool[shMemSize];

    const int nWarps2 = blockDim2 - WARP_SIZE2;
    const int sh_offs = (shMemSize >> nWarps2) * warpId;
    int *shmem = shmem_pool + sh_offs;

    /*********** check if this block is linked to a leaf **********/

    int *lmem = &MEM_BUF[blockIdx.x*(LMEM_STACK_SIZE*blockDim.x + LMEM_EXTRA_SIZE)];
    int  bid  = gridDim.x * blockIdx.y + blockIdx.x;

    while(true)
    {

#if 0
      if(threadIdx.x == 0)
      {
        bid         = atomicAdd(&active_inout[n_bodies], 1);
        shmem[0]    = bid;
      }
#else
      if(laneId == 0)
      {
        bid         = atomicAdd(&active_inout[n_bodies], 1);
        shmem[0]    = bid;
      }
#endif

      bid   = shmem[0];

      if (bid >= n_active_groups) return;

      int grpOffset = 0;

      /*********** set necessary thread constants **********/
#ifdef DO_BLOCK_TIMESTEP
      real4 curGroupSize    = groupSizeInfo[active_groups[bid + grpOffset]];
#else
      real4 curGroupSize    = groupSizeInfo[bid + grpOffset];
#endif
      int   groupData       = __float_as_int(curGroupSize.w);
      uint body_i           =   groupData & CRITMASK;
      uint nb_i             = ((groupData & INVCMASK) >> CRITBIT) + 1;

#ifdef DO_BLOCK_TIMESTEP
      real4 group_pos       = groupCenterInfo[active_groups[bid + grpOffset]];
#else
      real4 group_pos       = groupCenterInfo[bid + grpOffset];
#endif

      body_i += laneId%nb_i;

      float4 pos_i = body_pos[body_i];

      int ngb_i;

      float4 acc_i = {0.0f, 0.0f, 0.0f, 0.0f};

      const float group_eps  = 0;

      int apprCount = 0;
      int direCount = 0;


      acc_i = approximate_gravity<0, blockDim2>(
          body_i, pos_i, group_pos,
          eps2, node_begend,
          multipole_data, body_pos,
          shmem, lmem, ngb_i, apprCount, direCount, boxSizeInfo, curGroupSize, boxCenterInfo,
          group_eps, body_vel);

#if 0
      if(apprCount < 0)
      {

        //Try to get access to the big stack, only one block per time is allowed
        if(laneId == 0)
        {
          int res = atomicExch(&active_inout[n_bodies+1], 1); //If the old value (res) is 0 we can go otherwise sleep
          int waitCounter  = 0;
          while(res != 0)
          {
            //Sleep
            for(int i=0; i < (1024); i++)
            {
              waitCounter += 1;
            }
            //Test again
            shmem[0] = waitCounter;
            res = atomicExch(&active_inout[n_bodies+1], 1); 
          }
        }

        lmem = &MEM_BUF[gridDim.x*(LMEM_STACK_SIZE*blockDim.x + LMEM_EXTRA_SIZE)];    //Use the extra large buffer
        apprCount = direCount = 0;
        acc_i = approximate_gravity<8, blockDim2>(
            body_i, pos_i, group_pos,
            eps2, node_begend,
            multipole_data, body_pos,
            shmem, lmem, ngb_i, apprCount, direCount, boxSizeInfo, curGroupSize, boxCenterInfo,
            group_eps, body_vel);
        lmem = &MEM_BUF[blockIdx.x*(LMEM_STACK_SIZE*blockDim.x + LMEM_EXTRA_SIZE)];

        if(threadIdx.x == 0)
        {
          atomicExch(&active_inout[n_bodies+1], 0); //Release the lock
        }
      }//end if apprCount < 0
#endif

      if (laneId < nb_i) 
      {
        acc_out     [body_i] = acc_i;
        ngb_out     [body_i] = ngb_i;
        active_inout[body_i] = 1;
        interactions[body_i].x = apprCount;
        interactions[body_i].y = direCount ;
      }


    }     //end while

  }


#pragma unroll 16
    for (int j = 0; j < WARP_SIZE; j++) 
      for (int k = 0; k < NI; k++)
        acc_i[k] = add_acc(acc_i[k], pos_i[k], sh_mass[offs + j], sh_pos[offs + j], eps2);
    direCount += WARP_SIZE*NI;
  }
}


extern "C" __global__ void
#if 0 /* casues 164 bytes spill to lmem with NTHREAD = 128 */
__launch_bounds__(NTHREAD)
#endif
  dev_approximate_gravity(
      const int n_active_groups,
