PME spline+spread CUDA kernel and unit tests

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

diff --git a/src/gromacs/gpu_utils/cuda_kernel_utils.cuh b/src/gromacs/gpu_utils/cuda_kernel_utils.cuh
index e96edb366f4cb06b0c7a362f4f70ab2f2b187407..2bba07dbae373cecff338a552d235dc23410bf25 100644 (file)
--- a/src/gromacs/gpu_utils/cuda_kernel_utils.cuh
+++ b/src/gromacs/gpu_utils/cuda_kernel_utils.cuh
@@ -56,5 +56,65 @@ __device__ __forceinline__ T LDG(const T* ptr)
 #endif
 }
 
+/*! \brief Fetch the value by \p index from the texture object or reference.
+ * Fetching from the object is the preferred behaviour on CC >= 3.0.
+ *
+ * \tparam[in] T        Raw data type
+ * \param[in] texObj    Table texture object
+ * \param[in] texRef    Table texture reference
+ * \param[in] index     Non-negative element index
+ * \returns             The value from the table at \p index
+ */
+template <typename T>
+static __forceinline__ __device__
+T fetchFromTexture(const cudaTextureObject_t texObj,
+                   const struct texture<T, 1, cudaReadModeElementType> texRef,
+                   int index)
+{
+    assert(index >= 0);
+    assert(!c_disableCudaTextures);
+    T result;
+#if GMX_PTX_ARCH >= 300  // Preferring texture objects on any new arch
+    GMX_UNUSED_VALUE(texRef);
+    result = tex1Dfetch<T>(texObj, index);
+#else
+    GMX_UNUSED_VALUE(texObj);
+    result = tex1Dfetch(texRef, index);
+#endif
+    return result;
+}
+
+/*! \brief Fetch the value by \p index from the parameter lookup table.
+ *
+ *  Depending on what is supported, it fetches parameters either
+ *  using direct load, texture objects, or texture references.
+ *
+ * \tparam[in] T        Raw data type
+ * \param[in] d_ptr     Device pointer to the raw table memory
+ * \param[in] texObj    Table texture object
+ * \param[in] texRef    Table texture reference
+ * \param[in] index     Non-negative element index
+ * \returns             The value from the table at \p index
+ */
+template <typename T>
+static __forceinline__ __device__
+T fetchFromParamLookupTable(const T                  *d_ptr,
+                            const cudaTextureObject_t texObj,
+                            const struct texture<T, 1, cudaReadModeElementType> texRef,
+                            int index)
+{
+    assert(index >= 0);
+    T result;
+#if DISABLE_CUDA_TEXTURES
+    GMX_UNUSED_VALUE(texObj);
+    GMX_UNUSED_VALUE(texRef);
+    result = LDG(d_ptr + index);
+#else
+    GMX_UNUSED_VALUE(d_ptr);
+    result = fetchFromTexture<T>(texObj, texRef, index);
+#endif
+    return result;
+}
+
 
 #endif /* GMX_GPU_UTILS_CUDA_KERNEL_UTILS_CUH */