Fix malformed CUDA version macro check

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

diff --git a/src/mdlib/nbnxn_cuda/nbnxn_cuda_data_mgmt.cu b/src/mdlib/nbnxn_cuda/nbnxn_cuda_data_mgmt.cu
index cf9b0f772579e4c3e97077afb9a9e0d873a422cb..a35bfedb818b393dc12c769ccf722fc62eef66be 100644 (file)
--- a/src/mdlib/nbnxn_cuda/nbnxn_cuda_data_mgmt.cu
+++ b/src/mdlib/nbnxn_cuda/nbnxn_cuda_data_mgmt.cu
@@ -458,7 +458,7 @@ static int pick_nbnxn_kernel_version(FILE            *fplog,
     }
 
     bCUDA32 = bCUDA40 = false;
-#if CUDA_VERSION == 3200
+#if CUDA_VERSION == 3020
     bCUDA32 = true;
     sprintf(sbuf, "3.2");
 #elif CUDA_VERSION == 4000
@@ -505,7 +505,9 @@ static int pick_nbnxn_kernel_version(FILE            *fplog,
 
 void nbnxn_cuda_init(FILE *fplog,
                      nbnxn_cuda_ptr_t *p_cu_nb,
-                     const gmx_gpu_info_t *gpu_info, int my_gpu_index,
+                     const gmx_gpu_info_t *gpu_info,
+                     const gmx_gpu_opt_t *gpu_opt,
+                     int my_gpu_index,
                      gmx_bool bLocalAndNonlocal)
 {
     cudaError_t stat;
@@ -539,14 +541,36 @@ void nbnxn_cuda_init(FILE *fplog,
 
     init_plist(nb->plist[eintLocal]);
 
+    /* set device info, just point it to the right GPU among the detected ones */
+    nb->dev_info = &gpu_info->cuda_dev[get_gpu_device_id(gpu_info, gpu_opt, my_gpu_index)];
+
     /* local/non-local GPU streams */
     stat = cudaStreamCreate(&nb->stream[eintLocal]);
     CU_RET_ERR(stat, "cudaStreamCreate on stream[eintLocal] failed");
     if (nb->bUseTwoStreams)
     {
         init_plist(nb->plist[eintNonlocal]);
+
+        /* CUDA stream priority available in the CUDA RT 5.5 API.
+         * Note that the device we're running on does not have to support
+         * priorities, because we are querying the priority range which in this
+         * case will be a single value.
+         */
+#if CUDA_VERSION >= 5050
+        {
+            int highest_priority;
+            stat = cudaDeviceGetStreamPriorityRange(NULL, &highest_priority);
+            CU_RET_ERR(stat, "cudaDeviceGetStreamPriorityRange failed");
+
+            stat = cudaStreamCreateWithPriority(&nb->stream[eintNonlocal],
+                                                cudaStreamDefault,
+                                                highest_priority);
+            CU_RET_ERR(stat, "cudaStreamCreateWithPriority on stream[eintNonlocal] failed");
+        }
+#else
         stat = cudaStreamCreate(&nb->stream[eintNonlocal]);
         CU_RET_ERR(stat, "cudaStreamCreate on stream[eintNonlocal] failed");
+#endif
     }
 
     /* init events for sychronization (timing disabled for performance reasons!) */
@@ -555,9 +579,6 @@ void nbnxn_cuda_init(FILE *fplog,
     stat = cudaEventCreateWithFlags(&nb->misc_ops_done, cudaEventDisableTiming);
     CU_RET_ERR(stat, "cudaEventCreate on misc_ops_one failed");
 
-    /* set device info, just point it to the right GPU among the detected ones */
-    nb->dev_info = &gpu_info->cuda_dev[get_gpu_device_id(gpu_info, my_gpu_index)];
-
     /* On GPUs with ECC enabled, cudaStreamSynchronize shows a large overhead
      * (which increases with shorter time/step) caused by a known CUDA driver bug.
      * To work around the issue we'll use an (admittedly fragile) memory polling
@@ -582,7 +603,7 @@ void nbnxn_cuda_init(FILE *fplog,
     bTMPIAtomics = false;
 #endif
 
-#if defined(i386) || defined(__x86_64__)
+#ifdef GMX_TARGET_X86
     bX86 = true;
 #else
     bX86 = false;
@@ -607,7 +628,7 @@ void nbnxn_cuda_init(FILE *fplog,
          *   - GPUs are not being shared.
          */
         bool bShouldUsePollSync = (bX86 && bTMPIAtomics &&
-                                   (gmx_count_gpu_dev_shared(gpu_info) < 1));
+                                   (gmx_count_gpu_dev_shared(gpu_opt) < 1));
 
         if (bStreamSync)
         {
@@ -875,7 +896,7 @@ void nbnxn_cuda_init_atomdata(nbnxn_cuda_ptr_t cu_nb,
     }
 }
 
-void nbnxn_cuda_free(FILE *fplog, nbnxn_cuda_ptr_t cu_nb)
+void nbnxn_cuda_free(nbnxn_cuda_ptr_t cu_nb)
 {
     cudaError_t     stat;
     cu_atomdata_t   *atdat;