#define DISABLE_CUDA_TEXTURES 0
#endif
+/*! \brief True if the use of texture fetch in the CUDA kernels is disabled. */
+static const bool c_disableCudaTextures = DISABLE_CUDA_TEXTURES;
+
+
/* CUDA architecture technical characteristics. Needs macros because it is used
* in the __launch_bounds__ function qualifiers and might need it in preprocessor
* conditionals.
#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 ***/
}
}
}
+
+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 *);
/*! 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);
+
+/*! \brief Initialize parameter lookup table.
+ *
+ * Initializes device memory, copies data from host and binds
+ * a texture to allocated device memory to be used for 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>
+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);
+
#endif
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
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
/*! \brief cluster size = number of atoms per cluster. */
static const int c_clSize = c_nbnxnGpuClusterSize;
-/*! \brief True if the use of texture fetch in the CUDA kernels is disabled. */
-static const bool c_disableCudaTextures = DISABLE_CUDA_TEXTURES;
-
-
#ifdef __cplusplus
extern "C" {
#endif
biod.pnpi.spb.ru / alexxy / gromacs.git / commitdiff