                 grid.x = 65535;
                 }
 
-            shared_size = warp_scan<cur_tpp>::capacity*sizeof(unsigned char)*(block_size/cur_tpp);
+            if (sm < 30) shared_size = warp_scan<cur_tpp>::capacity*sizeof(unsigned char)*(block_size/cur_tpp);
 
             gpu_compute_nlist_binned_shared_kernel<2,cur_tpp><<<grid, block_size,shared_size>>>(d_nlist,
                                                                              d_n_neigh,

              bool filter_body,
              unsigned int block_size)
    {
    unsigned int shared_size;

    if (tpp == cur_tpp && cur_tpp != 0)
        {

                grid.x = 65535;
                }

            if (sm < 30) shared_size = warp_scan<cur_tpp>::capacity*sizeof(unsigned char)*(block_size/cur_tpp);

            gpu_compute_nlist_binned_shared_kernel<1,cur_tpp><<<grid, block_size,shared_size>>>(d_nlist,
                                                                             d_n_neigh,
