Bump required CUDA version to 7.0

[alexxy/gromacs.git] / src / gromacs / mdlib / nbnxn_cuda / nbnxn_cuda.cu

diff --git a/src/gromacs/mdlib/nbnxn_cuda/nbnxn_cuda.cu b/src/gromacs/mdlib/nbnxn_cuda/nbnxn_cuda.cu
index 7010ae6dc296dfab736743a1e3ba8fb7516ea80b..e18c197235b4eab756905ce4ef68032b4d58aca0 100644 (file)
--- a/src/gromacs/mdlib/nbnxn_cuda/nbnxn_cuda.cu
+++ b/src/gromacs/mdlib/nbnxn_cuda/nbnxn_cuda.cu
@@ -116,11 +116,6 @@ typedef void (*nbnxn_cu_kfunc_ptr_t)(const cu_atomdata_t,
 
 /*********************************/
 
-/* XXX switch between chevron and cudaLaunch (supported only in CUDA >=7.0)
-   -- only for benchmarking purposes */
-static const bool bUseCudaLaunchKernel =
-    (GMX_CUDA_VERSION >= 7000) && (getenv("GMX_DISABLE_CUDALAUNCH") == NULL);
-
 /*! Returns the number of blocks to be used for the nonbonded GPU kernel. */
 static inline int calc_nb_kernel_nblock(int nwork_units, const gmx_device_info_t *dinfo)
 {
@@ -432,22 +427,13 @@ void nbnxn_gpu_launch_kernel(gmx_nbnxn_cuda_t       *nb,
                 shmem);
     }
 
-    if (bUseCudaLaunchKernel)
-    {
-        gmx_unused void* kernel_args[4];
-        kernel_args[0] = adat;
-        kernel_args[1] = nbp;
-        kernel_args[2] = plist;
-        kernel_args[3] = &bCalcFshift;
+    void* kernel_args[4];
+    kernel_args[0] = adat;
+    kernel_args[1] = nbp;
+    kernel_args[2] = plist;
+    kernel_args[3] = &bCalcFshift;
 
-#if GMX_CUDA_VERSION >= 7000
-        cudaLaunchKernel((void *)nb_kernel, dim_grid, dim_block, kernel_args, shmem, stream);
-#endif
-    }
-    else
-    {
-        nb_kernel<<< dim_grid, dim_block, shmem, stream>>> (*adat, *nbp, *plist, bCalcFshift);
-    }
+    cudaLaunchKernel((void *)nb_kernel, dim_grid, dim_block, kernel_args, shmem, stream);
     CU_LAUNCH_ERR("k_calc_nb");
 
     if (bDoTime)
@@ -562,36 +548,20 @@ void nbnxn_gpu_launch_kernel_pruneonly(gmx_nbnxn_cuda_t       *nb,
                 shmem);
     }
 
-    if (bUseCudaLaunchKernel)
-    {
-        gmx_unused void* kernel_args[5];
-        kernel_args[0] = adat;
-        kernel_args[1] = nbp;
-        kernel_args[2] = plist;
-        kernel_args[3] = &numParts;
-        kernel_args[4] = &part;
+    void* kernel_args[5];
+    kernel_args[0] = adat;
+    kernel_args[1] = nbp;
+    kernel_args[2] = plist;
+    kernel_args[3] = &numParts;
+    kernel_args[4] = &part;
 
-#if GMX_CUDA_VERSION >= 7000
-        if (plist->haveFreshList)
-        {
-            cudaLaunchKernel((void *)nbnxn_kernel_prune_cuda<true>, dim_grid, dim_block, kernel_args, shmem, stream);
-        }
-        else
-        {
-            cudaLaunchKernel((void *)nbnxn_kernel_prune_cuda<false>, dim_grid, dim_block, kernel_args, shmem, stream);
-        }
-#endif
+    if (plist->haveFreshList)
+    {
+        cudaLaunchKernel((void *)nbnxn_kernel_prune_cuda<true>, dim_grid, dim_block, kernel_args, shmem, stream);
     }
     else
     {
-        if (plist->haveFreshList)
-        {
-            nbnxn_kernel_prune_cuda<true><<< dim_grid, dim_block, shmem, stream>>> (*adat, *nbp, *plist, numParts, part);
-        }
-        else
-        {
-            nbnxn_kernel_prune_cuda<false><<< dim_grid, dim_block, shmem, stream>>> (*adat, *nbp, *plist, numParts, part);
-        }
+        cudaLaunchKernel((void *)nbnxn_kernel_prune_cuda<false>, dim_grid, dim_block, kernel_args, shmem, stream);
     }
     CU_LAUNCH_ERR("k_pruneonly");