             case 2:
                 {
                 static int max_block_size = -1;
+                static int sm = -1;
                 if (max_block_size == -1)
                     max_block_size = get_max_block_size(gpu_compute_pair_forces_shared_kernel<evaluator, 2, 0>);
+                if (sm == -1)
+                    sm = get_compute_capability(gpu_compute_pair_forces_shared_kernel<evaluator, 2, 0>);
+
                 block_size = block_size < max_block_size ? block_size : max_block_size;
-                n_blocks = pair_args.N / (block_size/tpp) + 1;
+                dim3 grid(pair_args.N / (block_size/tpp) + 1, 1, 1);
+                if (sm < 30 && grid.x > 65535)
+                    {
+                    grid.y = grid.x/65535 + 1;
+                    grid.x = 65535;
+                    }
+
                 shared_bytes += sizeof(Scalar)*block_size;
 
                 gpu_compute_pair_forces_shared_kernel<evaluator, 2, 0>
-                  <<<n_blocks, block_size, shared_bytes>>>(pair_args.d_force, pair_args.d_virial,
+                  <<<grid, block_size, shared_bytes>>>(pair_args.d_force, pair_args.d_virial,
                   pair_args.virial_pitch, pair_args.N, pair_args.d_pos, pair_args.d_diameter,
                   pair_args.d_charge, pair_args.box, pair_args.d_n_neigh, pair_args.d_nlist,
                   pair_args.nli, d_params, pair_args.d_rcutsq, pair_args.d_ronsq, pair_args.ntypes,

            case 1:
                {
                static int max_block_size = -1;
                if (max_block_size == -1)
                    max_block_size = get_max_block_size(gpu_compute_pair_forces_shared_kernel<evaluator, 1, 0>);
                block_size = block_size < max_block_size ? block_size : max_block_size;
                n_blocks = pair_args.N / (block_size/tpp) + 1;
                shared_bytes += sizeof(Scalar)*block_size;

                gpu_compute_pair_forces_shared_kernel<evaluator, 1, 0>
                  <<<n_blocks, block_size, shared_bytes>>>(pair_args.d_force, pair_args.d_virial,
                  pair_args.virial_pitch, pair_args.N, pair_args.d_pos, pair_args.d_diameter,
                  pair_args.d_charge, pair_args.box, pair_args.d_n_neigh, pair_args.d_nlist,
                  pair_args.nli, d_params, pair_args.d_rcutsq, pair_args.d_ronsq, pair_args.ntypes,

            case 0:
                {
                static int max_block_size = -1;
                static int sm = -1;
                if (max_block_size == -1)
                    max_block_size = get_max_block_size(gpu_compute_pair_forces_shared_kernel<evaluator, 0, 0>);
                if (sm == -1)
                    sm = get_compute_capability(gpu_compute_pair_forces_shared_kernel<evaluator, 0, 0>);

                block_size = block_size < max_block_size ? block_size : max_block_size;
                dim3 grid(pair_args.N / (block_size/tpp) + 1, 1, 1);
                if (sm < 30 && grid.x > 65535)
                    {
                    grid.y = grid.x/65535 + 1;
                    grid.x = 65535;
                    }

                shared_bytes += sizeof(Scalar)*block_size;

                gpu_compute_pair_forces_shared_kernel<evaluator, 0, 0>
                  <<<grid, block_size, shared_bytes>>>(pair_args.d_force, pair_args.d_virial,
                  pair_args.virial_pitch, pair_args.N, pair_args.d_pos, pair_args.d_diameter,
                  pair_args.d_charge, pair_args.box, pair_args.d_n_neigh, pair_args.d_nlist,
                  pair_args.nli, d_params, pair_args.d_rcutsq, pair_args.d_ronsq, pair_args.ntypes,
