Fix one error and compiler warnings with Cuda & clang-3.6

Clang-3.6 on OS X can now be used by nvcc. clang found one
error related to || being used instead of | to set flag bits,
and a handful of warnings variables in headers not being used.
The latter is caused by declaring constants in headers, and
making then static to avoid clashing symbols. However, this emits
them in every single compile unit that includes the header. Fixed
by either moving names to a cpp file, or changing to defines.

Change-Id: Ib4d59c40aa8caffc667cc202a3efe45891b2abe3

diff --git a/src/gromacs/gmxlib/cuda_tools/pmalloc_cuda.cu b/src/gromacs/gmxlib/cuda_tools/pmalloc_cuda.cu
index eeef88fd41e02197b191968da5a0c7c7d422bb1b..96d21e1a139ec8a698d81e530df8ad9f13f09074 100644 (file)
--- a/src/gromacs/gmxlib/cuda_tools/pmalloc_cuda.cu
+++ b/src/gromacs/gmxlib/cuda_tools/pmalloc_cuda.cu
@@ -78,7 +78,7 @@ void pmalloc_wc(void **h_ptr, size_t nbytes)
 {
     cudaError_t stat;
     char        strbuf[STRLEN];
-    int         flag = cudaHostAllocDefault || cudaHostAllocWriteCombined;
+    int         flag = cudaHostAllocDefault | cudaHostAllocWriteCombined;
 
     if (nbytes == 0)
     {

diff --git a/src/gromacs/gmxlib/gmx_detect_hardware.cpp b/src/gromacs/gmxlib/gmx_detect_hardware.cpp
index 96ea3c3db17b10c0d3a6d24fe914a94e730d061b..ff2e416d785e2725b83b0b378b6ebbab187b10e0 100644 (file)
--- a/src/gromacs/gmxlib/gmx_detect_hardware.cpp
+++ b/src/gromacs/gmxlib/gmx_detect_hardware.cpp
@@ -82,6 +82,12 @@ const gmx_bool bGPUBinary = TRUE;
 const gmx_bool bGPUBinary = FALSE;
 #endif
 
+/* Names of the GPU detection/check results (see e_gpu_detect_res_t in hw_info.h). */
+const char * const gpu_detect_res_str[egpuNR] =
+{
+    "compatible", "inexistent", "incompatible", "insane"
+};
+
 static const char * invalid_gpuid_hint =
     "A delimiter-free sequence of valid numeric IDs of available GPUs is expected.";
 

diff --git a/src/gromacs/legacyheaders/types/hw_info.h b/src/gromacs/legacyheaders/types/hw_info.h
index 51c3b9812f00ff650ef008d0075bd47804a131c3..e94d550b7d2e021ef27c0e6f1278c9b8ad74197f 100644 (file)
--- a/src/gromacs/legacyheaders/types/hw_info.h
+++ b/src/gromacs/legacyheaders/types/hw_info.h
@@ -55,14 +55,11 @@ struct gmx_device_info_t;
  * incompatible driver/runtime. */
 typedef enum
 {
-    egpuCompatible = 0,  egpuNonexistent,  egpuIncompatible, egpuInsane
+    egpuCompatible = 0,  egpuNonexistent,  egpuIncompatible, egpuInsane, egpuNR
 } e_gpu_detect_res_t;
 
-/* Textual names of the GPU detection/check results (see e_gpu_detect_res_t). */
-static const char * const gpu_detect_res_str[] =
-{
-    "compatible", "inexistent", "incompatible", "insane"
-};
+/* Names of the GPU detection/check results */
+extern const char * const gpu_detect_res_str[egpuNR];
 
 /* GPU device information -- for now with only CUDA devices
  * The gmx_hardware_detect module initializes it. */

diff --git a/src/gromacs/mdlib/nbnxn_consts.h b/src/gromacs/mdlib/nbnxn_consts.h
index 3c5da829321578e51da66c7e089a7640fb313a5a..3ada608f60379ae1a37a46848ca02087a2c15a62 100644 (file)
--- a/src/gromacs/mdlib/nbnxn_consts.h
+++ b/src/gromacs/mdlib/nbnxn_consts.h
@@ -98,15 +98,15 @@ extern "C" {
  * Bit i*CJ_SIZE + j tells if atom i and j interact.
  */
 /* All interaction mask is the same for all kernels */
-static const unsigned int NBNXN_INTERACTION_MASK_ALL       = 0xffffffffU;
+#define NBNXN_INTERACTION_MASK_ALL        0xffffffffU
 /* 4x4 kernel diagonal mask */
-static const unsigned int NBNXN_INTERACTION_MASK_DIAG      = 0x08ceU;
+#define NBNXN_INTERACTION_MASK_DIAG       0x08ceU
 /* 4x2 kernel diagonal masks */
-static const unsigned int NBNXN_INTERACTION_MASK_DIAG_J2_0 = 0x0002U;
-static const unsigned int NBNXN_INTERACTION_MASK_DIAG_J2_1 = 0x002fU;
+#define NBNXN_INTERACTION_MASK_DIAG_J2_0  0x0002U
+#define NBNXN_INTERACTION_MASK_DIAG_J2_1  0x002fU
 /* 4x8 kernel diagonal masks */
-static const unsigned int NBNXN_INTERACTION_MASK_DIAG_J8_0 = 0xf0f8fcfeU;
-static const unsigned int NBNXN_INTERACTION_MASK_DIAG_J8_1 = 0x0080c0e0U;
+#define NBNXN_INTERACTION_MASK_DIAG_J8_0  0xf0f8fcfeU
+#define NBNXN_INTERACTION_MASK_DIAG_J8_1  0x0080c0e0U
 
 
 #ifdef __cplusplus

diff --git a/src/gromacs/mdlib/nbnxn_cuda/nbnxn_cuda_kernel.cuh b/src/gromacs/mdlib/nbnxn_cuda/nbnxn_cuda_kernel.cuh
index d6793ca658750a1581d94259a8e233c63bcd8c58..01d848a93b2fd41df06b21b0c93643bccf1a6e87 100644 (file)
--- a/src/gromacs/mdlib/nbnxn_cuda/nbnxn_cuda_kernel.cuh
+++ b/src/gromacs/mdlib/nbnxn_cuda/nbnxn_cuda_kernel.cuh
@@ -335,7 +335,10 @@ __global__ void NB_KERNEL_FUNC_NAME(nbnxn_kernel, _F_cuda)
 #endif
             for (jm = 0; jm < NBNXN_GPU_JGROUP_SIZE; jm++)
             {
-                if (imask & (supercl_interaction_mask << (jm * NCL_PER_SUPERCL)))
+                /* ((1U << NCL_PER_SUPERCL) - 1U) is the i-cluster interaction
+                 * mask for a super-cluster with all NCL_PER_SUPERCL bits set.
+                 */
+                if (imask & (((1U << NCL_PER_SUPERCL) - 1U) << (jm * NCL_PER_SUPERCL)))
                 {
                     mask_ji = (1U << (jm * NCL_PER_SUPERCL));
 

diff --git a/src/gromacs/mdlib/nbnxn_cuda/nbnxn_cuda_kernel_utils.cuh b/src/gromacs/mdlib/nbnxn_cuda/nbnxn_cuda_kernel_utils.cuh
index 2b7b4ddcad599acfb16b4396919e3ac9fbf5f648..4f59dc262867a7272b4e4c78610174b9806bc785 100644 (file)
--- a/src/gromacs/mdlib/nbnxn_cuda/nbnxn_cuda_kernel_utils.cuh
+++ b/src/gromacs/mdlib/nbnxn_cuda/nbnxn_cuda_kernel_utils.cuh
@@ -59,9 +59,6 @@
 #define ONE_TWELVETH_F  0.08333333f
 
 
-/*! i-cluster interaction mask for a super-cluster with all NCL_PER_SUPERCL bits set */
-const unsigned supercl_interaction_mask = ((1U << NCL_PER_SUPERCL) - 1U);
-
 /*! Apply force switch,  force + energy version. */
 static inline __device__
 void calculate_force_switch_F(const  cu_nbparam_t nbparam,