From: Aleksei Iupinov Date: Fri, 22 Sep 2017 13:35:28 +0000 (+0200) Subject: Template/move CUDA texture cleanup code from NB CUDA module to cudautils.cu X-Git-Url: http://biod.pnpi.spb.ru/gitweb/?a=commitdiff_plain;h=510fa2e8ef8d915c3496fbd7b2ceaa9e9a3c940d;p=alexxy%2Fgromacs.git 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 75d87fa8ce..36ee46f4e8 100644 --- 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 +void destroyParamLookupTable(T *d_ptr, + cudaTextureObject_t texObj, + const struct texture *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() during texture setup + * looks reasonable, when instantiating the templates for new types - just in case. + */ template void initParamLookupTable(float * &, cudaTextureObject_t &, const texture *, const float *, int, const gmx_device_info_t *); +template void destroyParamLookupTable(float *, cudaTextureObject_t, const texture *, const gmx_device_info_t *); diff --git a/src/gromacs/gpu_utils/cudautils.cuh b/src/gromacs/gpu_utils/cudautils.cuh index cfe5c2a381..43b40a93fc 100644 --- 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 +void destroyParamLookupTable(T *d_ptr, + cudaTextureObject_t texObj, + const struct texture *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 24f316d549..991e558631 100644 --- 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);