                      tpp,
                      filter_diameter,
                      filter_body,
-                     n_blocks,
                      block_size
                      );
         }

                                                                             r_maxsq,
                                                                             sqrtf(r_maxsq),
                                                                             ghost_width);
        else if (filter_diameter && filter_body)
             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,

                                                                             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<1,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<1,cur_tpp><<<n_blocks, block_size,shared_size>>>(d_nlist,
                                                                             d_n_neigh,
                                                                             d_last_updated_pos,
