Move CUDA texture setup code from NB CUDA module to cudautils.cu

[alexxy/gromacs.git] / src / gromacs / gpu_utils / cudautils.cu

diff --git a/src/gromacs/gpu_utils/cudautils.cu b/src/gromacs/gpu_utils/cudautils.cu
index 6021331ffb4740684113555b319579d683fda89a..75d87fa8ceae5f6b9f46a0d56d893ab196b6f4b8 100644 (file)
--- a/src/gromacs/gpu_utils/cudautils.cu
+++ b/src/gromacs/gpu_utils/cudautils.cu
@@ -37,8 +37,10 @@
 
 #include "cudautils.cuh"
 
-#include <stdlib.h>
+#include <cassert>
+#include <cstdlib>
 
+#include "gromacs/gpu_utils/cuda_arch_utils.cuh"
 #include "gromacs/utility/smalloc.h"
 
 /*** Generic CUDA data operation wrappers ***/
@@ -246,3 +248,97 @@ void cu_realloc_buffered(void **d_dest, void *h_src,
         }
     }
 }
+
+bool use_texobj(const gmx_device_info_t *dev_info)
+{
+    assert(!c_disableCudaTextures);
+    /* Only device CC >= 3.0 (Kepler and later) support texture objects */
+    return (dev_info->prop.major >= 3);
+}
+
+/*! \brief Set up texture object for an array of type T.
+ *
+ * Set up texture object for an array of type T and bind it to the device memory
+ * \p d_ptr points to.
+ *
+ * \tparam[in] T        Raw data type
+ * \param[out] texObj   texture object to initialize
+ * \param[in]  d_ptr    pointer to device global memory to bind \p texObj to
+ * \param[in]  sizeInBytes  size of memory area to bind \p texObj to
+ */
+template <typename T>
+static void setup1DTexture(cudaTextureObject_t &texObj,
+                           void                *d_ptr,
+                           size_t               sizeInBytes)
+{
+    assert(!c_disableCudaTextures);
+
+    cudaError_t      stat;
+    cudaResourceDesc rd;
+    cudaTextureDesc  td;
+
+    memset(&rd, 0, sizeof(rd));
+    rd.resType                = cudaResourceTypeLinear;
+    rd.res.linear.devPtr      = d_ptr;
+    rd.res.linear.desc        = cudaCreateChannelDesc<T>();
+    rd.res.linear.sizeInBytes = sizeInBytes;
+
+    memset(&td, 0, sizeof(td));
+    td.readMode                 = cudaReadModeElementType;
+    stat = cudaCreateTextureObject(&texObj, &rd, &td, NULL);
+    CU_RET_ERR(stat, "cudaCreateTextureObject failed");
+}
+
+/*! \brief Set up texture reference for an array of type T.
+ *
+ * Set up texture object for an array of type T and bind it to the device memory
+ * \p d_ptr points to.
+ *
+ * \tparam[in] T        Raw data type
+ * \param[out] texObj   texture reference to initialize
+ * \param[in]  d_ptr    pointer to device global memory to bind \p texObj to
+ * \param[in]  sizeInBytes  size of memory area to bind \p texObj to
+ */
+template <typename T>
+static void setup1DTexture(const struct texture<T, 1, cudaReadModeElementType> *texRef,
+                           const void                                          *d_ptr,
+                           size_t                                              sizeInBytes)
+{
+    assert(!c_disableCudaTextures);
+
+    cudaError_t           stat;
+    cudaChannelFormatDesc cd;
+
+    cd   = cudaCreateChannelDesc<T>();
+    stat = cudaBindTexture(nullptr, texRef, d_ptr, &cd, sizeInBytes);
+    CU_RET_ERR(stat, "cudaBindTexture failed");
+}
+
+template <typename T>
+void initParamLookupTable(T                        * &d_ptr,
+                          cudaTextureObject_t       &texObj,
+                          const struct texture<T, 1, cudaReadModeElementType> *texRef,
+                          const T                   *h_ptr,
+                          int                        numElem,
+                          const gmx_device_info_t   *devInfo)
+{
+    const size_t sizeInBytes = numElem * sizeof(*d_ptr);
+    cudaError_t  stat        = cudaMalloc((void **)&d_ptr, sizeInBytes);
+    CU_RET_ERR(stat, "cudaMalloc failed in initParamLookupTable");
+    cu_copy_H2D(d_ptr, (void *)h_ptr, sizeInBytes);
+
+    if (!c_disableCudaTextures)
+    {
+        if (use_texobj(devInfo))
+        {
+            setup1DTexture<T>(texObj, d_ptr, sizeInBytes);
+        }
+        else
+        {
+            setup1DTexture<T>(texRef, d_ptr, sizeInBytes);
+        }
+    }
+}
+
+//! Add explicit instantiations of initParamLookupTable() here as needed
+template void initParamLookupTable<float>(float * &, cudaTextureObject_t &, const texture<float, 1, cudaReadModeElementType> *, const float *, int, const gmx_device_info_t *);