Template/move CUDA texture cleanup code from NB CUDA module to cudautils.cu

Noted TODO: easy transformation into a GPU table class.

Change-Id: I20d684221fa8304d01ab7fd4a19f2c4469110142

diff --git a/src/gromacs/gpu_utils/cudautils.cu b/src/gromacs/gpu_utils/cudautils.cu
index 75d87fa8ceae5f6b9f46a0d56d893ab196b6f4b8..36ee46f4e8098d1c38cc9431a9da9be1a373ebaa 100644 (file)
--- a/src/gromacs/gpu_utils/cudautils.cu
+++ b/src/gromacs/gpu_utils/cudautils.cu
@@ -249,7 +249,12 @@ void cu_realloc_buffered(void **d_dest, void *h_src,
     }
 }
 
-bool use_texobj(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 inline bool use_texobj(const gmx_device_info_t *dev_info)
 {
     assert(!c_disableCudaTextures);
     /* Only device CC >= 3.0 (Kepler and later) support texture objects */
@@ -340,5 +345,29 @@ void initParamLookupTable(T                        * &d_ptr,
     }
 }
 
-//! Add explicit instantiations of initParamLookupTable() here as needed
+template <typename T>
+void destroyParamLookupTable(T                       *d_ptr,
+                             cudaTextureObject_t      texObj,
+                             const struct texture<T, 1, cudaReadModeElementType> *texRef,
+                             const gmx_device_info_t *devInfo)
+{
+    if (!c_disableCudaTextures)
+    {
+        if (use_texobj(devInfo))
+        {
+            CU_RET_ERR(cudaDestroyTextureObject(texObj), "cudaDestroyTextureObject on texObj failed");
+        }
+        else
+        {
+            CU_RET_ERR(cudaUnbindTexture(texRef), "cudaUnbindTexture on texRef failed");
+        }
+    }
+    CU_RET_ERR(cudaFree(d_ptr), "cudaFree failed");
+}
+
+/*! \brief Add explicit instantiations of init/destroyParamLookupTable() here as needed.
+ * One should also verify that the result of cudaCreateChannelDesc<T>() during texture setup
+ * looks reasonable, when instantiating the templates for new types - just in case.
+ */
 template void initParamLookupTable<float>(float * &, cudaTextureObject_t &, const texture<float, 1, cudaReadModeElementType> *, const float *, int, const gmx_device_info_t *);
+template void destroyParamLookupTable<float>(float *, cudaTextureObject_t, const texture<float, 1, cudaReadModeElementType> *, const gmx_device_info_t *);

diff --git a/src/gromacs/gpu_utils/cudautils.cuh b/src/gromacs/gpu_utils/cudautils.cuh
index cfe5c2a381506ae3febaec07f84956bafcd1cd42..43b40a93fcd09af7c10bdcd3000be644fe48a93b 100644 (file)
--- a/src/gromacs/gpu_utils/cudautils.cuh
+++ b/src/gromacs/gpu_utils/cudautils.cuh
@@ -165,14 +165,8 @@ float cu_event_elapsed(cudaEvent_t /*start*/, cudaEvent_t /*end*/);
 /*! Waits for event end to complete and calculates the time between start and end. */
 int cu_wait_event_time(cudaEvent_t /*end*/, cudaEvent_t /*begin*/, float * /*time*/);
 
-/*! \brief Return whether texture objects are used on this device.
- *
- * \todo This should be static in cudautils.cu, as soon as texture destruction code is moved there as well
- *
- * \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
- */
-bool use_texobj(const gmx_device_info_t *dev_info);
+// TODO: the 2 functions below are pretty much a constructor/destructor of a simple
+// GPU table object. We just need to add a templated __device__ table data fetching to complete it.
 
 /*! \brief Initialize parameter lookup table.
  *
@@ -195,4 +189,20 @@ void initParamLookupTable(T                        * &d_ptr,
                           int                        numElem,
                           const gmx_device_info_t   *devInfo);
 
+/*! \brief Destroy parameter lookup table.
+ *
+ * Unbinds texture reference/object, deallocates device memory.
+ *
+ * \tparam[in] T         Raw data type
+ * \param[in]  d_ptr     Device pointer to the memory to be deallocated
+ * \param[in]  texObj    Texture object to be deinitialized
+ * \param[in]  texRef    Texture reference to be deinitialized
+ * \param[in]  devInfo   Pointer to the info struct of the device in use
+ */
+template <typename T>
+void destroyParamLookupTable(T                         *d_ptr,
+                             cudaTextureObject_t        texObj,
+                             const struct texture<T, 1, cudaReadModeElementType> *texRef,
+                             const gmx_device_info_t   *devInfo);
+
 #endif

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 24f316d5498c1d2162f8e47f57eb1d89f95a94bf..991e55863188490bea2c22e092ef8bf8c4dec151 100644 (file)
--- a/src/gromacs/mdlib/nbnxn_cuda/nbnxn_cuda_data_mgmt.cu
+++ b/src/gromacs/mdlib/nbnxn_cuda/nbnxn_cuda_data_mgmt.cu
@@ -772,26 +772,10 @@ void nbnxn_gpu_init_atomdata(gmx_nbnxn_cuda_t              *nb,
 static void nbnxn_cuda_free_nbparam_table(cu_nbparam_t            *nbparam,
                                           const gmx_device_info_t *dev_info)
 {
-    cudaError_t stat;
-
     if (nbparam->eeltype == eelCuEWALD_TAB || nbparam->eeltype == eelCuEWALD_TAB_TWIN)
     {
-        if (!c_disableCudaTextures)
-        {
-            /* Only device CC >= 3.0 (Kepler and later) support texture objects */
-            if (use_texobj(dev_info))
-            {
-                stat = cudaDestroyTextureObject(nbparam->coulomb_tab_texobj);
-                CU_RET_ERR(stat, "cudaDestroyTextureObject on coulomb_tab_texobj failed");
-            }
-            else
-            {
-                GMX_UNUSED_VALUE(dev_info);
-                stat = cudaUnbindTexture(nbnxn_cuda_get_coulomb_tab_texref());
-                CU_RET_ERR(stat, "cudaUnbindTexture on coulomb_tab_texref failed");
-            }
-        }
-        cu_free_buffered(nbparam->coulomb_tab);
+        destroyParamLookupTable(nbparam->coulomb_tab, nbparam->coulomb_tab_texobj,
+                                &nbnxn_cuda_get_coulomb_tab_texref(), dev_info);
     }
 }
 
@@ -868,40 +852,15 @@ void nbnxn_gpu_free(gmx_nbnxn_cuda_t *nb)
 
     if (!useLjCombRule(nb->nbparam))
     {
-        if (!c_disableCudaTextures)
-        {
-            /* Only device CC >= 3.0 (Kepler and later) support texture objects */
-            if (use_texobj(nb->dev_info))
-            {
-                stat = cudaDestroyTextureObject(nbparam->nbfp_texobj);
-                CU_RET_ERR(stat, "cudaDestroyTextureObject on nbfp_texobj failed");
-            }
-            else
-            {
-                stat = cudaUnbindTexture(nbnxn_cuda_get_nbfp_texref());
-                CU_RET_ERR(stat, "cudaUnbindTexture on nbfp_texref failed");
-            }
-        }
-        cu_free_buffered(nbparam->nbfp);
+        destroyParamLookupTable(nbparam->nbfp, nbparam->nbfp_texobj,
+                                &nbnxn_cuda_get_nbfp_texref(), nb->dev_info);
+
     }
 
     if (nbparam->vdwtype == evdwCuEWALDGEOM || nbparam->vdwtype == evdwCuEWALDLB)
     {
-        if (!c_disableCudaTextures)
-        {
-            /* Only device CC >= 3.0 (Kepler and later) support texture objects */
-            if (use_texobj(nb->dev_info))
-            {
-                stat = cudaDestroyTextureObject(nbparam->nbfp_comb_texobj);
-                CU_RET_ERR(stat, "cudaDestroyTextureObject on nbfp_comb_texobj failed");
-            }
-            else
-            {
-                stat = cudaUnbindTexture(nbnxn_cuda_get_nbfp_comb_texref());
-                CU_RET_ERR(stat, "cudaUnbindTexture on nbfp_comb_texref failed");
-            }
-        }
-        cu_free_buffered(nbparam->nbfp_comb);
+        destroyParamLookupTable(nbparam->nbfp_comb, nbparam->nbfp_comb_texobj,
+                                &nbnxn_cuda_get_nbfp_comb_texref(), nb->dev_info);
     }
 
     stat = cudaFree(atdat->shift_vec);