Use cudaLaunchKernel with CUDA 7.0 and later

CUDA 7.0 introduced the cudaLaunchKernel API call similar to the
CUDA driver API and avoids the chevron notation. This has the benefit
of a slight reduction in runtime API overhead (up to 2%) partly
because two runtime API calls that precede the kernel launch are skipped
(cudaSetupArgument and cudaConfigureCall).

For future dev-testing the GMX_DISABLE_CUDALAUNCH env. var. can be used to
force the chevron-notation kernel launch.

Change-Id: Id057fb01489814b99ae290de9e4ddd9f530a04be

diff --git a/docs/user-guide/environment-variables.rst b/docs/user-guide/environment-variables.rst
index 76e25db1e3646cd06c19cdd0be8737cbdb8cad91..8cdf343617d48a6f3a80f8fbaaa0fe99d6eb4bcb 100644 (file)
--- a/docs/user-guide/environment-variables.rst
+++ b/docs/user-guide/environment-variables.rst
@@ -132,6 +132,10 @@ Performance and Run Control
         to performance loss due to a known CUDA driver bug present in API v5.0 NVIDIA drivers (pre-30x.xx).
         Cannot be set simultaneously with ``GMX_NO_CUDA_STREAMSYNC``.
 
+``GMX_DISABLE_CUDALAUNCH``
+        disable the use of the lower-latency cudaLaunchKernel API even when supported (CUDA >=v7.0).
+        Should only be used for benchmarking purposes.
+
 ``GMX_CYCLE_ALL``
         times all code during runs.  Incompatible with threads.
 

diff --git a/src/gromacs/mdlib/nbnxn_cuda/nbnxn_cuda.cu b/src/gromacs/mdlib/nbnxn_cuda/nbnxn_cuda.cu
index 1f5c865889762800589af678800c5075ed2e6ec0..13f8ef89a6f5bf090103f2fcf2ef9ce0dd75f02c 100644 (file)
--- a/src/gromacs/mdlib/nbnxn_cuda/nbnxn_cuda.cu
+++ b/src/gromacs/mdlib/nbnxn_cuda/nbnxn_cuda.cu
@@ -125,6 +125,10 @@ 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);
 
 /* XXX always/never run the energy/pruning kernels -- only for benchmarking purposes */
 static bool always_ener  = (getenv("GMX_GPU_ALWAYS_ENER") != NULL);
@@ -421,7 +425,22 @@ void nbnxn_gpu_launch_kernel(gmx_nbnxn_cuda_t       *nb,
                 shmem);
     }
 
-    nb_kernel<<< dim_grid, dim_block, shmem, stream>>> (*adat, *nbp, *plist, bCalcFshift);
+    if (bUseCudaLaunchKernel)
+    {
+        gmx_unused 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);
+    }
     CU_LAUNCH_ERR("k_calc_nb");
 
     if (bDoTime)