         static unsigned int max_block_size_1 = UINT_MAX;
         static unsigned int max_block_size_2 = UINT_MAX;
 
-		cudaFuncAttributes attr1;
-		cudaFuncGetAttributes(&attr1, gpu_compute_eam_tex_inter_forces_kernel<0>);
-
-		cudaFuncAttributes attr2;
-		cudaFuncGetAttributes(&attr2, gpu_compute_eam_tex_inter_forces_kernel_2<0>);
-
-		max_block_size_1 = attr1.maxThreadsPerBlock;
-		max_block_size_2 = attr2.maxThreadsPerBlock;
+        cudaFuncAttributes attr1;
+        cudaFuncGetAttributes(&attr1, gpu_kernel_1<0>);
+
+        cudaFuncAttributes attr2;
+        cudaFuncGetAttributes(&attr2, gpu_kernel_2<0>);
+
+        max_block_size_1 = attr1.maxThreadsPerBlock;
+        max_block_size_2 = attr2.maxThreadsPerBlock;
 
         unsigned int run_block_size_1 = min(eam_data.block_size, max_block_size_1);
         unsigned int run_block_size_2 = min(eam_data.block_size, max_block_size_2);

        dim3 grid_2((int) ceil((double) N / (double) run_block_size_2), 1, 1);
        dim3 threads_2(run_block_size_2, 1, 1);

        gpu_compute_eam_tex_inter_forces_kernel<0> <<<grid_1, threads_1>>>(d_force,
                d_virial, virial_pitch, N, d_pos, box, d_n_neigh, d_nlist,
                d_head_list, d_F, d_rho, d_rphi, d_dF, d_drho, d_drphi, d_dFdP);
        gpu_compute_eam_tex_inter_forces_kernel_2<0> <<<grid_2, threads_2>>>(d_force,
                d_virial, virial_pitch, N, d_pos, box, d_n_neigh, d_nlist,
                d_head_list, d_F, d_rho, d_rphi, d_dF, d_drho, d_drphi, d_dFdP);
    }

    return cudaSuccess;
}
            return error;
    }

    if (compute_capability < 350) {
        tex_F.normalized = false;
        tex_F.filterMode = cudaFilterModePoint;
        error = cudaBindTexture(0, tex_F, d_F, sizeof(Scalar4) * eam_data.nrho * eam_data.ntypes);
