2 * This file is part of the GROMACS molecular simulation package.
4 * Copyright (c) 2012,2014,2015,2016,2017,2018, by the GROMACS development team, led by
5 * Mark Abraham, David van der Spoel, Berk Hess, and Erik Lindahl,
6 * and including many others, as listed in the AUTHORS file in the
7 * top-level source directory and at http://www.gromacs.org.
9 * GROMACS is free software; you can redistribute it and/or
10 * modify it under the terms of the GNU Lesser General Public License
11 * as published by the Free Software Foundation; either version 2.1
12 * of the License, or (at your option) any later version.
14 * GROMACS is distributed in the hope that it will be useful,
15 * but WITHOUT ANY WARRANTY; without even the implied warranty of
16 * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the GNU
17 * Lesser General Public License for more details.
19 * You should have received a copy of the GNU Lesser General Public
20 * License along with GROMACS; if not, see
21 * http://www.gnu.org/licenses, or write to the Free Software Foundation,
22 * Inc., 51 Franklin Street, Fifth Floor, Boston, MA 02110-1301 USA.
24 * If you want to redistribute modifications to GROMACS, please
25 * consider that scientific software is very special. Version
26 * control is crucial - bugs must be traceable. We will be happy to
27 * consider code for inclusion in the official distribution, but
28 * derived work must not be called official GROMACS. Details are found
29 * in the README & COPYING files - if they are missing, get the
30 * official version at http://www.gromacs.org.
32 * To help us fund GROMACS development, we humbly ask that you cite
33 * the research papers on the package. Check out http://www.gromacs.org.
38 #include "cudautils.cuh"
43 #include "gromacs/gpu_utils/cuda_arch_utils.cuh"
44 #include "gromacs/gpu_utils/gpu_utils.h"
45 #include "gromacs/utility/gmxassert.h"
47 /*** Generic CUDA data operation wrappers ***/
49 // TODO: template on transferKind to avoid runtime conditionals
50 int cu_copy_D2H(void *h_dest, void *d_src, size_t bytes,
51 GpuApiCallBehavior transferKind, cudaStream_t s = nullptr)
55 if (h_dest == nullptr || d_src == nullptr || bytes == 0)
62 case GpuApiCallBehavior::Async:
63 GMX_ASSERT(isHostMemoryPinned(h_dest), "Destination buffer was not pinned for CUDA");
64 stat = cudaMemcpyAsync(h_dest, d_src, bytes, cudaMemcpyDeviceToHost, s);
65 CU_RET_ERR(stat, "DtoH cudaMemcpyAsync failed");
68 case GpuApiCallBehavior::Sync:
69 stat = cudaMemcpy(h_dest, d_src, bytes, cudaMemcpyDeviceToHost);
70 CU_RET_ERR(stat, "DtoH cudaMemcpy failed");
80 int cu_copy_D2H_sync(void * h_dest, void * d_src, size_t bytes)
82 return cu_copy_D2H(h_dest, d_src, bytes, GpuApiCallBehavior::Sync);
86 * The copy is launched in stream s or if not specified, in stream 0.
88 int cu_copy_D2H_async(void * h_dest, void * d_src, size_t bytes, cudaStream_t s = nullptr)
90 return cu_copy_D2H(h_dest, d_src, bytes, GpuApiCallBehavior::Async, s);
93 // TODO: template on transferKind to avoid runtime conditionals
94 int cu_copy_H2D(void *d_dest, void *h_src, size_t bytes,
95 GpuApiCallBehavior transferKind, cudaStream_t s = nullptr)
99 if (d_dest == nullptr || h_src == nullptr || bytes == 0)
104 switch (transferKind)
106 case GpuApiCallBehavior::Async:
107 GMX_ASSERT(isHostMemoryPinned(h_src), "Source buffer was not pinned for CUDA");
108 stat = cudaMemcpyAsync(d_dest, h_src, bytes, cudaMemcpyHostToDevice, s);
109 CU_RET_ERR(stat, "HtoD cudaMemcpyAsync failed");
112 case GpuApiCallBehavior::Sync:
113 stat = cudaMemcpy(d_dest, h_src, bytes, cudaMemcpyHostToDevice);
114 CU_RET_ERR(stat, "HtoD cudaMemcpy failed");
124 int cu_copy_H2D_sync(void * d_dest, void * h_src, size_t bytes)
126 return cu_copy_H2D(d_dest, h_src, bytes, GpuApiCallBehavior::Sync);
130 * The copy is launched in stream s or if not specified, in stream 0.
132 int cu_copy_H2D_async(void * d_dest, void * h_src, size_t bytes, cudaStream_t s = nullptr)
134 return cu_copy_H2D(d_dest, h_src, bytes, GpuApiCallBehavior::Async, s);
137 /*! \brief Return whether texture objects are used on this device.
139 * \param[in] pointer to the GPU device info structure to inspect for texture objects support
140 * \return true if texture objects are used on this device
142 static inline bool use_texobj(const gmx_device_info_t *dev_info)
144 assert(!c_disableCudaTextures);
145 /* Only device CC >= 3.0 (Kepler and later) support texture objects */
146 return (dev_info->prop.major >= 3);
149 /*! \brief Set up texture object for an array of type T.
151 * Set up texture object for an array of type T and bind it to the device memory
152 * \p d_ptr points to.
154 * \tparam[in] T Raw data type
155 * \param[out] texObj texture object to initialize
156 * \param[in] d_ptr pointer to device global memory to bind \p texObj to
157 * \param[in] sizeInBytes size of memory area to bind \p texObj to
159 template <typename T>
160 static void setup1DTexture(cudaTextureObject_t &texObj,
164 assert(!c_disableCudaTextures);
170 memset(&rd, 0, sizeof(rd));
171 rd.resType = cudaResourceTypeLinear;
172 rd.res.linear.devPtr = d_ptr;
173 rd.res.linear.desc = cudaCreateChannelDesc<T>();
174 rd.res.linear.sizeInBytes = sizeInBytes;
176 memset(&td, 0, sizeof(td));
177 td.readMode = cudaReadModeElementType;
178 stat = cudaCreateTextureObject(&texObj, &rd, &td, nullptr);
179 CU_RET_ERR(stat, "cudaCreateTextureObject failed");
182 template <typename T>
183 void initParamLookupTable(T * &d_ptr,
184 cudaTextureObject_t &texObj,
187 const gmx_device_info_t *devInfo)
189 const size_t sizeInBytes = numElem * sizeof(*d_ptr);
190 cudaError_t stat = cudaMalloc((void **)&d_ptr, sizeInBytes);
191 CU_RET_ERR(stat, "cudaMalloc failed in initParamLookupTable");
192 cu_copy_H2D_sync(d_ptr, (void *)h_ptr, sizeInBytes);
194 if (!c_disableCudaTextures)
196 if (use_texobj(devInfo))
198 setup1DTexture<T>(texObj, d_ptr, sizeInBytes);
203 template <typename T>
204 void destroyParamLookupTable(T *d_ptr,
205 cudaTextureObject_t texObj,
206 const gmx_device_info_t *devInfo)
208 if (!c_disableCudaTextures)
210 if (use_texobj(devInfo))
212 CU_RET_ERR(cudaDestroyTextureObject(texObj), "cudaDestroyTextureObject on texObj failed");
215 CU_RET_ERR(cudaFree(d_ptr), "cudaFree failed");
218 /*! \brief Add explicit instantiations of init/destroyParamLookupTable() here as needed.
219 * One should also verify that the result of cudaCreateChannelDesc<T>() during texture setup
220 * looks reasonable, when instantiating the templates for new types - just in case.
222 template void initParamLookupTable<float>(float * &, cudaTextureObject_t &, const float *, int, const gmx_device_info_t *);
223 template void destroyParamLookupTable<float>(float *, cudaTextureObject_t, const gmx_device_info_t *);
224 template void initParamLookupTable<int>(int * &, cudaTextureObject_t &, const int *, int, const gmx_device_info_t *);
225 template void destroyParamLookupTable<int>(int *, cudaTextureObject_t, const gmx_device_info_t *);