               unsigned int tpp,
               bool filter_diameter,
               bool filter_body,
-              unsigned int n_blocks,
               unsigned int block_size)
     { }
 

                                     bool filter_diameter,
                                     const Scalar3& ghost_width)
    {
    int n_blocks = N/(block_size/threads_per_particle);
    if (N % (block_size/threads_per_particle)) ++n_blocks;

    #if __CUDA_ARCH__ <= 300
    // bind the position texture
    cell_xyzf_1d_tex.normalized = false;

                                                                             r_maxsq,
                                                                             sqrtf(r_maxsq),
                                                                             ghost_width);
            }
        else if (filter_diameter && filter_body)
            {
            static int max_block_size = -1;
            if (max_block_size == -1)
                max_block_size = get_max_block_size(gpu_compute_nlist_binned_shared_kernel<3,cur_tpp>);
            block_size = block_size < max_block_size ? block_size : max_block_size;
            n_blocks = N / (block_size/tpp) + 1;
            shared_size = warp_scan<cur_tpp>::capacity*sizeof(unsigned char)*(block_size/cur_tpp);


            gpu_compute_nlist_binned_shared_kernel<3,cur_tpp><<<n_blocks, block_size,shared_size>>>(d_nlist,
                                                                             d_n_neigh,
                                                                             d_last_updated_pos,
                                                                             d_conditions,
