From 3ee2d276012042dcacb92a83cb0dbcba2888369a Mon Sep 17 00:00:00 2001 From: Szilard Pall Date: Fri, 4 Oct 2013 02:33:04 +0200 Subject: [PATCH] make use of CUDA stream priorities 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 --- src/mdlib/nbnxn_cuda/nbnxn_cuda_data_mgmt.cu | 25 +++++++++++++++++--- 1 file changed, 22 insertions(+), 3 deletions(-) diff --git a/src/mdlib/nbnxn_cuda/nbnxn_cuda_data_mgmt.cu b/src/mdlib/nbnxn_cuda/nbnxn_cuda_data_mgmt.cu index cf9b0f7725..f5e3e02f74 100644 --- a/src/mdlib/nbnxn_cuda/nbnxn_cuda_data_mgmt.cu +++ b/src/mdlib/nbnxn_cuda/nbnxn_cuda_data_mgmt.cu @@ -539,14 +539,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, 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 >= 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"); +#endif } /* init events for sychronization (timing disabled for performance reasons!) */ @@ -555,9 +577,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 -- 2.22.0