Merge branch release-2016

[alexxy/gromacs.git] / src / gromacs / gpu_utils / cuda_arch_utils.cuh

diff --git a/src/gromacs/gpu_utils/cuda_arch_utils.cuh b/src/gromacs/gpu_utils/cuda_arch_utils.cuh
index b6a23047267fd4b9428b8f006ecaebc829c7e9b8..a60bf9ebb80537871ba7f5b52100458ae62c6509 100644 (file)
--- a/src/gromacs/gpu_utils/cuda_arch_utils.cuh
+++ b/src/gromacs/gpu_utils/cuda_arch_utils.cuh
@@ -134,4 +134,41 @@ T gmx_shfl_down_sync(const unsigned int activeMask,
 #endif
 }
 
+/*! \brief Allow disabling CUDA textures using the GMX_DISABLE_CUDA_TEXTURES macro.
+ *
+ *  This option will not influence functionality. All features using textures ought
+ *  to have fallback for texture-less reads (direct/LDG loads), all new code needs
+ *  to provide fallback code.
+ */
+#if defined GMX_DISABLE_CUDA_TEXTURES
+#define DISABLE_CUDA_TEXTURES 1
+#else
+#define DISABLE_CUDA_TEXTURES 0
+#endif
+
+/* CUDA architecture technical characteristics. Needs macros because it is used
+ * in the __launch_bounds__ function qualifiers and might need it in preprocessor
+ * conditionals.
+ *
+ */
+#if GMX_PTX_ARCH > 0
+    #if   GMX_PTX_ARCH <= 210  // CC 2.x
+        #define GMX_CUDA_MAX_BLOCKS_PER_MP   8
+        #define GMX_CUDA_MAX_THREADS_PER_MP  1536
+    #elif GMX_PTX_ARCH <= 370  // CC 3.x
+        #define GMX_CUDA_MAX_BLOCKS_PER_MP   16
+        #define GMX_CUDA_MAX_THREADS_PER_MP  2048
+    #else // CC 5.x, 6.x
+          /* Note that this final branch covers all future architectures (current gen
+           * is 6.x as of writing), hence assuming that these *currently defined* upper
+           * limits will not be lowered.
+           */
+        #define GMX_CUDA_MAX_BLOCKS_PER_MP   32
+        #define GMX_CUDA_MAX_THREADS_PER_MP  2048
+    #endif
+#else
+        #define GMX_CUDA_MAX_BLOCKS_PER_MP   0
+        #define GMX_CUDA_MAX_THREADS_PER_MP  0
+#endif
+
 #endif /* CUDA_ARCH_UTILS_CUH_ */