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

[alexxy/gromacs.git] / src / gromacs / mdlib / nbnxn_cuda / nbnxn_cuda_data_mgmt.cu

diff --git a/src/gromacs/mdlib/nbnxn_cuda/nbnxn_cuda_data_mgmt.cu b/src/gromacs/mdlib/nbnxn_cuda/nbnxn_cuda_data_mgmt.cu
index 79ef8e842e53f188aab4b5ecab761094e24907a6..24f316d5498c1d2162f8e47f57eb1d89f95a94bf 100644 (file)
--- a/src/gromacs/mdlib/nbnxn_cuda/nbnxn_cuda_data_mgmt.cu
+++ b/src/gromacs/mdlib/nbnxn_cuda/nbnxn_cuda_data_mgmt.cu
@@ -86,19 +86,6 @@ static void nbnxn_cuda_clear_e_fshift(gmx_nbnxn_cuda_t *nb);
 static void nbnxn_cuda_free_nbparam_table(cu_nbparam_t            *nbparam,
                                           const gmx_device_info_t *dev_info);
 
-
-/*! \brief Return whether texture objects are used on this device.
- *
- * \param[in]   pointer to the GPU device info structure to inspect for texture objects support
- * \return      true if texture objects are used on this device
- */
-static 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 Return whether combination rules are used.
  *
  * \param[in]   pointer to nonbonded paramter struct
@@ -110,104 +97,6 @@ static inline bool useLjCombRule(const cu_nbparam_t  *nbparam)
             nbparam->vdwtype == evdwCuCUTCOMBLB);
 }
 
-/*! \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");
-}
-
-/*! \brief Initialize parameter lookup table.
- *
- * Initializes device memory, copies data from host and binds
- * a texture to allocated device memory to be used for LJ/Ewald/... parameter
- * lookup.
- *
- * \tparam[in] T         Raw data type
- * \param[out] d_ptr     device pointer to the memory to be allocated
- * \param[out] texObj    texture object to be initialized
- * \param[out] texRef    texture reference to be initialized
- * \param[in]  h_ptr     pointer to the host memory to be uploaded to the device
- * \param[in]  numElem   number of elements in the h_ptr
- * \param[in]  devInfo   pointer to the info struct of the device in use
- */
-template <typename T>
-static 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);
-        }
-    }
-}
-
 /*! \brief Initialized the Ewald Coulomb correction GPU table.
 
     Tabulates the Ewald Coulomb force and initializes the size/scale