         }
     }
 
+//! determine maximum possible block size
+template<typename T>
+int get_max_block_size(T func)
+    {
+    cudaFuncAttributes attr;
+    cudaFuncGetAttributes(&attr, func);
+    int max_threads = attr.maxThreadsPerBlock;
+    // number of threads has to be multiple of warp size
+    max_threads -= max_threads % max_threads_per_particle; 
+    return max_threads;
+    }
+
+
 //! recursive template to launch neighborlist with given template parameters
 /* \tparam cur_tpp Number of threads per particle (assumed to be power of two) */
 template<int cur_tpp>

                                                                             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
        {
