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

diff --git a/src/gromacs/mdlib/nbnxn_cuda/nbnxn_cuda_kernel.cuh b/src/gromacs/mdlib/nbnxn_cuda/nbnxn_cuda_kernel.cuh
index b24915c1f922cefed0f1f0b9b48969e74c167600..4afcd4e7cfa3b3192065c753fd42da610606a0cd 100644 (file)
--- 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