From 5822b640b44e910b8accc6727e661907b5e47ffe Mon Sep 17 00:00:00 2001 From: Erik Lindahl Date: Fri, 12 Jun 2015 22:17:47 +0200 Subject: [PATCH] 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 --- src/gromacs/gmxlib/cuda_tools/pmalloc_cuda.cu | 2 +- src/gromacs/gmxlib/gmx_detect_hardware.cpp | 6 ++++++ src/gromacs/legacyheaders/types/hw_info.h | 9 +++------ src/gromacs/mdlib/nbnxn_consts.h | 12 ++++++------ src/gromacs/mdlib/nbnxn_cuda/nbnxn_cuda_kernel.cuh | 5 ++++- .../mdlib/nbnxn_cuda/nbnxn_cuda_kernel_utils.cuh | 3 --- 6 files changed, 20 insertions(+), 17 deletions(-) diff --git a/src/gromacs/gmxlib/cuda_tools/pmalloc_cuda.cu b/src/gromacs/gmxlib/cuda_tools/pmalloc_cuda.cu index eeef88fd41..96d21e1a13 100644 --- 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 96ea3c3db1..ff2e416d78 100644 --- 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 51c3b9812f..e94d550b7d 100644 --- 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 3c5da82932..3ada608f60 100644 --- 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 d6793ca658..01d848a93b 100644 --- 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 2b7b4ddcad..4f59dc2628 100644 --- 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, -- 2.22.0