From 4889a40a5c735b116556556b7286075fbb1f2171 Mon Sep 17 00:00:00 2001 From: Szilard Pall Date: Wed, 24 Jun 2015 23:16:49 +0200 Subject: [PATCH] Add missing macro undef in CUDA NB kernel Harmless as different architecture code-path get generated such that they don't end up in the same compilation unit, so the two macros remaining defined did not affect code where the __CUDA_ARCH__ >= 300 is not true. Change-Id: Ic6911e5c13781ac8a2835c3aef1457df6da60412 --- src/gromacs/mdlib/nbnxn_cuda/nbnxn_cuda_kernel.cuh | 5 +++++ 1 file changed, 5 insertions(+) diff --git a/src/gromacs/mdlib/nbnxn_cuda/nbnxn_cuda_kernel.cuh b/src/gromacs/mdlib/nbnxn_cuda/nbnxn_cuda_kernel.cuh index b24915c1f9..4afcd4e7cf 100644 --- a/src/gromacs/mdlib/nbnxn_cuda/nbnxn_cuda_kernel.cuh +++ b/src/gromacs/mdlib/nbnxn_cuda/nbnxn_cuda_kernel.cuh @@ -53,6 +53,7 @@ */ #if __CUDA_ARCH__ >= 300 +/* Note: convenience macros, need to be undef-ed at the end of the file. */ #define REDUCE_SHUFFLE /* On Kepler pre-loading i-atom types to shmem gives a few %, but on Fermi it does not */ @@ -79,6 +80,7 @@ #define LJ_EWALD #endif + /* Kernel launch parameters: - #blocks = #pair lists, blockId = pair list Id @@ -575,6 +577,9 @@ __global__ void NB_KERNEL_FUNC_NAME(nbnxn_kernel, _F_cuda) #endif } +#undef REDUCE_SHUFFLE +#undef IATYPE_SHMEM + #undef EL_EWALD_ANY #undef EXCLUSION_FORCES #undef LJ_EWALD -- 2.22.0