CUDA 5.5 introduced steam priorities with 2 levels. We make use of this
feature by launching the non-local non-bonded kernel in a high priority
stream. As a consequence, the non-local kernel will preempt the local
one and finish first. This will improve performance in multi-node runs
by reducing the possibility of late arrival of non-local forces.
Change-Id: I4efc65546e4135f12006c0422e1fca42a788129f
init_plist(nb->plist[eintLocal]);
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, 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]);
/* 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 >= 5500
+ {
+ 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");
stat = cudaStreamCreate(&nb->stream[eintNonlocal]);
CU_RET_ERR(stat, "cudaStreamCreate on stream[eintNonlocal] failed");
}
/* init events for sychronization (timing disabled for performance reasons!) */
}
/* init events for sychronization (timing disabled for performance reasons!) */
stat = cudaEventCreateWithFlags(&nb->misc_ops_done, cudaEventDisableTiming);
CU_RET_ERR(stat, "cudaEventCreate on misc_ops_one failed");
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
/* 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