NAMD
Classes | Macros | Functions
CudaUtils.h File Reference
#include <stdio.h>
#include <cuda.h>
#include <cuda_runtime.h>

Go to the source code of this file.

Classes

struct  CudaStaticAssert< bool >
 
struct  CudaStaticAssert< true >
 

Macros

#define WARPSIZE   32
 
#define WARP_FULL_MASK   0xffffffff
 
#define WARP_SHUFFLE_XOR(MASK, VAR, LANE, SIZE)   __shfl_xor(VAR, LANE, SIZE)
 
#define WARP_SHUFFLE_UP(MASK, VAR, DELTA, SIZE)   __shfl_up(VAR, DELTA, SIZE)
 
#define WARP_SHUFFLE_DOWN(MASK, VAR, DELTA, SIZE)   __shfl_down(VAR, DELTA, SIZE)
 
#define WARP_SHUFFLE(MASK, VAR, LANE, SIZE)   __shfl(VAR, LANE, SIZE)
 
#define WARP_ALL(MASK, P)   __all(P)
 
#define WARP_ANY(MASK, P)   __any(P)
 
#define WARP_BALLOT(MASK, P)   __ballot(P)
 
#define WARP_SYNC(MASK)
 
#define BLOCK_SYNC   __syncthreads()
 
#define cuda_static_assert(expr)   (CudaStaticAssert<(expr) != 0>())
 
#define cudaCheck(stmt)
 

Functions

void cudaDie (const char *msg, cudaError_t err=cudaSuccess)
 
void cudaNAMD_bug (const char *msg)
 
void clear_device_array_async_T (void *data, const int ndata, cudaStream_t stream, const size_t sizeofT)
 
void clear_device_array_T (void *data, const int ndata, const size_t sizeofT)
 
template<class T >
void clear_device_array (T *data, const int ndata, cudaStream_t stream=0)
 
template<class T >
void clear_device_array_sync (T *data, const int ndata)
 
void allocate_host_T (void **pp, const int len, const size_t sizeofT)
 
template<class T >
void allocate_host (T **pp, const int len)
 
void allocate_device_T (void **pp, const int len, const size_t sizeofT)
 
template<class T >
void allocate_device (T **pp, const int len)
 
void deallocate_device_T (void **pp)
 
template<class T >
void deallocate_device (T **pp)
 
bool reallocate_device_T (void **pp, int *curlen, const int newlen, const float fac, const size_t sizeofT)
 
template<class T >
bool reallocate_device (T **pp, int *curlen, const int newlen, const float fac=1.0f)
 
bool reallocate_host_T (void **pp, int *curlen, const int newlen, const float fac, const unsigned int flag, const size_t sizeofT)
 
template<class T >
bool reallocate_host (T **pp, int *curlen, const int newlen, const float fac=1.0f, const unsigned int flag=cudaHostAllocDefault)
 
void deallocate_host_T (void **pp)
 
template<class T >
void deallocate_host (T **pp)
 
void copy_HtoD_async_T (const void *h_array, void *d_array, int array_len, cudaStream_t stream, const size_t sizeofT)
 
void copy_HtoD_T (const void *h_array, void *d_array, int array_len, const size_t sizeofT)
 
void copy_DtoH_async_T (const void *d_array, void *h_array, const int array_len, cudaStream_t stream, const size_t sizeofT)
 
void copy_DtoH_T (const void *d_array, void *h_array, const int array_len, const size_t sizeofT)
 
void copy_DtoD_async_T (const void *d_src, void *d_dst, const int array_len, cudaStream_t stream, const size_t sizeofT)
 
void copy_DtoD_T (const void *d_src, void *d_dst, const int array_len, const size_t sizeofT)
 
template<class T >
void copy_HtoD (const T *h_array, T *d_array, int array_len, cudaStream_t stream=0)
 
template<class T >
void copy_HtoD_sync (const T *h_array, T *d_array, int array_len)
 
template<class T >
void copy_DtoH (const T *d_array, T *h_array, const int array_len, cudaStream_t stream=0)
 
template<class T >
void copy_DtoH_sync (const T *d_array, T *h_array, const int array_len)
 
template<class T >
void copy_DtoD (const T *d_src, T *h_dst, const int array_len, cudaStream_t stream=0)
 
template<class T >
void copy_DtoD_sync (const T *d_src, T *h_dst, const int array_len)
 
void copy_PeerDtoD_async_T (const int src_dev, const int dst_dev, const void *d_src, void *d_dst, const int array_len, cudaStream_t stream, const size_t sizeofT)
 
template<class T >
void copy_PeerDtoD (const int src_dev, const int dst_dev, const T *d_src, T *d_dst, const int array_len, cudaStream_t stream=0)
 
void copy3D_HtoD_T (void *src_data, void *dst_data, int src_x0, int src_y0, int src_z0, size_t src_xsize, size_t src_ysize, int dst_x0, int dst_y0, int dst_z0, size_t dst_xsize, size_t dst_ysize, size_t width, size_t height, size_t depth, size_t sizeofT, cudaStream_t stream)
 
template<class T >
void copy3D_HtoD (T *src_data, T *dst_data, int src_x0, int src_y0, int src_z0, size_t src_xsize, size_t src_ysize, int dst_x0, int dst_y0, int dst_z0, size_t dst_xsize, size_t dst_ysize, size_t width, size_t height, size_t depth, cudaStream_t stream=0)
 
void copy3D_DtoH_T (void *src_data, void *dst_data, int src_x0, int src_y0, int src_z0, size_t src_xsize, size_t src_ysize, int dst_x0, int dst_y0, int dst_z0, size_t dst_xsize, size_t dst_ysize, size_t width, size_t height, size_t depth, size_t sizeofT, cudaStream_t stream)
 
template<class T >
void copy3D_DtoH (T *src_data, T *dst_data, int src_x0, int src_y0, int src_z0, size_t src_xsize, size_t src_ysize, int dst_x0, int dst_y0, int dst_z0, size_t dst_xsize, size_t dst_ysize, size_t width, size_t height, size_t depth, cudaStream_t stream=0)
 
void copy3D_DtoD_T (void *src_data, void *dst_data, int src_x0, int src_y0, int src_z0, size_t src_xsize, size_t src_ysize, int dst_x0, int dst_y0, int dst_z0, size_t dst_xsize, size_t dst_ysize, size_t width, size_t height, size_t depth, size_t sizeofT, cudaStream_t stream)
 
template<class T >
void copy3D_DtoD (T *src_data, T *dst_data, int src_x0, int src_y0, int src_z0, size_t src_xsize, size_t src_ysize, int dst_x0, int dst_y0, int dst_z0, size_t dst_xsize, size_t dst_ysize, size_t width, size_t height, size_t depth, cudaStream_t stream=0)
 
void copy3D_PeerDtoD_T (int src_dev, int dst_dev, void *src_data, void *dst_data, int src_x0, int src_y0, int src_z0, size_t src_xsize, size_t src_ysize, int dst_x0, int dst_y0, int dst_z0, size_t dst_xsize, size_t dst_ysize, size_t width, size_t height, size_t depth, size_t sizeofT, cudaStream_t stream)
 
template<class T >
void copy3D_PeerDtoD (int src_dev, int dst_dev, T *src_data, T *dst_data, int src_x0, int src_y0, int src_z0, size_t src_xsize, size_t src_ysize, int dst_x0, int dst_y0, int dst_z0, size_t dst_xsize, size_t dst_ysize, size_t width, size_t height, size_t depth, cudaStream_t stream=0)
 

Macro Definition Documentation

#define BLOCK_SYNC   __syncthreads()

Definition at line 44 of file CudaUtils.h.

#define cuda_static_assert (   expr)    (CudaStaticAssert<(expr) != 0>())

Definition at line 69 of file CudaUtils.h.

Referenced by reduceVariables().

#define cudaCheck (   stmt)
Value:
do { \
cudaError_t err = stmt; \
if (err != cudaSuccess) { \
char msg[256]; \
sprintf(msg, "%s in file %s, function %s, line %d\n", #stmt,__FILE__,__FUNCTION__,__LINE__); \
cudaDie(msg, err); \
} \
} while(0)
if(ComputeNonbondedUtil::goMethod==2)
void cudaDie(const char *msg, cudaError_t err=cudaSuccess)
Definition: CudaUtils.C:9

Definition at line 79 of file CudaUtils.h.

Referenced by allocate_device_T(), allocate_host_T(), CudaFFTCompute::backward(), batchTranspose_xyz_yzx(), batchTranspose_xyz_zxy(), CudaComputeNonbondedKernel::bindExclusions(), bindTextureObject(), ComputeBondedCUDAKernel::bondedForce(), CudaTileListKernel::buildTileLists(), clear_device_array_async_T(), clear_device_array_T(), ComputeBondedCUDAKernel::ComputeBondedCUDAKernel(), copy3D_DtoD_T(), copy3D_DtoH_T(), copy3D_HtoD_T(), copy3D_PeerDtoD_T(), copy_DtoD_async_T(), copy_DtoD_T(), copy_DtoH_async_T(), copy_DtoH_T(), copy_HtoD_async_T(), copy_HtoD_T(), copy_PeerDtoD_async_T(), CudaPmeRealSpaceCompute::copyAtoms(), CudaPmeTranspose::copyDataDeviceToDevice(), CudaPmeTranspose::copyDataDeviceToHost(), CudaPmeTranspose::copyDataHostToDevice(), createStream(), CudaComputeGBISKernel::CudaComputeGBISKernel(), CudaComputeNonbonded::CudaComputeNonbonded(), CudaComputeNonbondedKernel::CudaComputeNonbondedKernel(), CudaNonbondedTables::CudaNonbondedTables(), CudaPmeKSpaceCompute::CudaPmeKSpaceCompute(), CudaPmeRealSpaceCompute::CudaPmeRealSpaceCompute(), CudaPmeTranspose::CudaPmeTranspose(), CudaTileListKernel::CudaTileListKernel(), deallocate_device_T(), deallocate_host_T(), ComputeNonbondedCUDA::doWork(), CudaPmeKSpaceCompute::energyAndVirialSetCallback(), CudaComputeNonbonded::finishReductions(), CudaTileListKernel::finishTileList(), CudaFFTCompute::forward(), gather_force(), CudaPmeRealSpaceCompute::gatherForce(), CudaPmeRealSpaceCompute::gatherForceSetCallback(), CudaComputeGBISKernel::GBISphase1(), CudaComputeGBISKernel::GBISphase2(), CudaComputeGBISKernel::GBISphase3(), DeviceCUDA::getMaxNumBlocks(), DeviceCUDA::getMaxNumThreads(), DeviceCUDA::initialize(), CudaComputeNonbonded::initialize(), ComputePmeCUDADevice::initialize(), ComputePmeCUDAMgr::initialize_pencils(), CudaPmePencilXY::initializeDevice(), CudaPmePencilX::initializeDevice(), CudaPmePencilY::initializeDevice(), CudaPmePencilZ::initializeDevice(), CudaComputeNonbonded::launchWork(), CudaComputeNonbondedKernel::nonbondedForce(), read_CUDA_ARCH(), reallocate_device_T(), reallocate_host_T(), CudaComputeNonbondedKernel::reduceVirialEnergy(), CudaTileListKernel::reSortTileLists(), scalar_sum(), CudaPmeTranspose::setDataPtrsYZX(), CudaPmeTranspose::setDataPtrsZXY(), CudaPmeKSpaceCompute::solve(), spread_charge(), CudaPmeRealSpaceCompute::spreadCharge(), transpose_xyz_yzx(), transpose_xyz_zxy(), CudaPmeTranspose::transposeXYZtoYZX(), CudaPmeTranspose::transposeXYZtoZXY(), CudaPmeRealSpaceCompute::waitGatherForceDone(), CudaPmeTranspose::waitStreamSynchronize(), writeComplexToDisk(), writeRealToDisk(), ComputeBondedCUDAKernel::~ComputeBondedCUDAKernel(), ComputePmeCUDADevice::~ComputePmeCUDADevice(), ComputePmeCUDAMgr::~ComputePmeCUDAMgr(), CudaComputeGBISKernel::~CudaComputeGBISKernel(), CudaComputeNonbonded::~CudaComputeNonbonded(), CudaComputeNonbondedKernel::~CudaComputeNonbondedKernel(), CudaFFTCompute::~CudaFFTCompute(), CudaNonbondedTables::~CudaNonbondedTables(), CudaPmeKSpaceCompute::~CudaPmeKSpaceCompute(), CudaPmePencilX::~CudaPmePencilX(), CudaPmePencilXY::~CudaPmePencilXY(), CudaPmePencilY::~CudaPmePencilY(), CudaPmePencilZ::~CudaPmePencilZ(), CudaPmeRealSpaceCompute::~CudaPmeRealSpaceCompute(), CudaPmeTranspose::~CudaPmeTranspose(), and CudaTileListKernel::~CudaTileListKernel().

#define WARP_ALL (   MASK,
 
)    __all(P)

Definition at line 40 of file CudaUtils.h.

Referenced by if().

#define WARP_ANY (   MASK,
 
)    __any(P)

Definition at line 41 of file CudaUtils.h.

Referenced by if(), and void().

#define WARP_BALLOT (   MASK,
 
)    __ballot(P)

Definition at line 42 of file CudaUtils.h.

Referenced by gather_force(), and repackTileListsKernel().

#define WARP_FULL_MASK   0xffffffff
#define WARP_SHUFFLE (   MASK,
  VAR,
  LANE,
  SIZE 
)    __shfl(VAR, LANE, SIZE)
#define WARP_SHUFFLE_DOWN (   MASK,
  VAR,
  DELTA,
  SIZE 
)    __shfl_down(VAR, DELTA, SIZE)

Definition at line 36 of file CudaUtils.h.

#define WARP_SHUFFLE_UP (   MASK,
  VAR,
  DELTA,
  SIZE 
)    __shfl_up(VAR, DELTA, SIZE)

Definition at line 34 of file CudaUtils.h.

#define WARP_SHUFFLE_XOR (   MASK,
  VAR,
  LANE,
  SIZE 
)    __shfl_xor(VAR, LANE, SIZE)
#define WARP_SYNC (   MASK)
#define WARPSIZE   32

Definition at line 9 of file CudaUtils.h.

Function Documentation

template<class T >
void allocate_device ( T **  pp,
const int  len 
)

Definition at line 146 of file CudaUtils.h.

References allocate_device_T().

146  {
147  allocate_device_T((void **)pp, len, sizeof(T));
148 }
void allocate_device_T(void **pp, const int len, const size_t sizeofT)
Definition: CudaUtils.C:75
void allocate_device_T ( void **  pp,
const int  len,
const size_t  sizeofT 
)

Definition at line 75 of file CudaUtils.C.

References cudaCheck.

Referenced by allocate_device().

75  {
76  cudaCheck(cudaMalloc(pp, sizeofT*len));
77 }
#define cudaCheck(stmt)
Definition: CudaUtils.h:79
template<class T >
void allocate_host ( T **  pp,
const int  len 
)

Definition at line 133 of file CudaUtils.h.

References allocate_host_T().

133  {
134  allocate_host_T((void **)pp, len, sizeof(T));
135 }
void allocate_host_T(void **pp, const int len, const size_t sizeofT)
Definition: CudaUtils.C:65
void allocate_host_T ( void **  pp,
const int  len,
const size_t  sizeofT 
)

Definition at line 65 of file CudaUtils.C.

References cudaCheck.

Referenced by allocate_host().

65  {
66  cudaCheck(cudaMallocHost(pp, sizeofT*len));
67 }
#define cudaCheck(stmt)
Definition: CudaUtils.h:79
template<class T >
void clear_device_array ( T *  data,
const int  ndata,
cudaStream_t  stream = 0 
)

Definition at line 116 of file CudaUtils.h.

References clear_device_array_async_T(), and stream.

116  {
117  clear_device_array_async_T(data, ndata, stream, sizeof(T));
118 }
__thread cudaStream_t stream
void clear_device_array_async_T(void *data, const int ndata, cudaStream_t stream, const size_t sizeofT)
Definition: CudaUtils.C:51
void clear_device_array_async_T ( void data,
const int  ndata,
cudaStream_t  stream,
const size_t  sizeofT 
)

Definition at line 51 of file CudaUtils.C.

References cudaCheck.

Referenced by clear_device_array().

51  {
52  cudaCheck(cudaMemsetAsync(data, 0, sizeofT*ndata, stream));
53 }
__thread cudaStream_t stream
#define cudaCheck(stmt)
Definition: CudaUtils.h:79
template<class T >
void clear_device_array_sync ( T *  data,
const int  ndata 
)

Definition at line 121 of file CudaUtils.h.

References clear_device_array_T().

121  {
122  clear_device_array_T(data, ndata, sizeof(T));
123 }
void clear_device_array_T(void *data, const int ndata, const size_t sizeofT)
Definition: CudaUtils.C:55
void clear_device_array_T ( void data,
const int  ndata,
const size_t  sizeofT 
)

Definition at line 55 of file CudaUtils.C.

References cudaCheck.

Referenced by clear_device_array_sync().

55  {
56  cudaCheck(cudaMemset(data, 0, sizeofT*ndata));
57 }
#define cudaCheck(stmt)
Definition: CudaUtils.h:79
template<class T >
void copy3D_DtoD ( T *  src_data,
T *  dst_data,
int  src_x0,
int  src_y0,
int  src_z0,
size_t  src_xsize,
size_t  src_ysize,
int  dst_x0,
int  dst_y0,
int  dst_z0,
size_t  dst_xsize,
size_t  dst_ysize,
size_t  width,
size_t  height,
size_t  depth,
cudaStream_t  stream = 0 
)

Definition at line 360 of file CudaUtils.h.

References copy3D_DtoD_T(), and stream.

366  {
367  copy3D_DtoD_T(src_data, dst_data,
368  src_x0, src_y0, src_z0,
369  src_xsize, src_ysize,
370  dst_x0, dst_y0, dst_z0,
371  dst_xsize, dst_ysize,
372  width, height, depth,
373  sizeof(T), stream);
374 }
__thread cudaStream_t stream
void copy3D_DtoD_T(void *src_data, void *dst_data, int src_x0, int src_y0, int src_z0, size_t src_xsize, size_t src_ysize, int dst_x0, int dst_y0, int dst_z0, size_t dst_xsize, size_t dst_ysize, size_t width, size_t height, size_t depth, size_t sizeofT, cudaStream_t stream)
Definition: CudaUtils.C:275
void copy3D_DtoD_T ( void src_data,
void dst_data,
int  src_x0,
int  src_y0,
int  src_z0,
size_t  src_xsize,
size_t  src_ysize,
int  dst_x0,
int  dst_y0,
int  dst_z0,
size_t  dst_xsize,
size_t  dst_ysize,
size_t  width,
size_t  height,
size_t  depth,
size_t  sizeofT,
cudaStream_t  stream 
)

Definition at line 275 of file CudaUtils.C.

References cudaCheck.

Referenced by copy3D_DtoD().

281  {
282  cudaMemcpy3DParms parms = {0};
283 
284  parms.srcPos = make_cudaPos(sizeofT*src_x0, src_y0, src_z0);
285  parms.srcPtr = make_cudaPitchedPtr(src_data, sizeofT*src_xsize, src_xsize, src_ysize);
286 
287  parms.dstPos = make_cudaPos(sizeofT*dst_x0, dst_y0, dst_z0);
288  parms.dstPtr = make_cudaPitchedPtr(dst_data, sizeofT*dst_xsize, dst_xsize, dst_ysize);
289 
290  parms.extent = make_cudaExtent(sizeofT*width, height, depth);
291  parms.kind = cudaMemcpyDeviceToDevice;
292 
293  cudaCheck(cudaMemcpy3DAsync(&parms, stream));
294 }
__thread cudaStream_t stream
#define cudaCheck(stmt)
Definition: CudaUtils.h:79
template<class T >
void copy3D_DtoH ( T *  src_data,
T *  dst_data,
int  src_x0,
int  src_y0,
int  src_z0,
size_t  src_xsize,
size_t  src_ysize,
int  dst_x0,
int  dst_y0,
int  dst_z0,
size_t  dst_xsize,
size_t  dst_ysize,
size_t  width,
size_t  height,
size_t  depth,
cudaStream_t  stream = 0 
)

Definition at line 331 of file CudaUtils.h.

References copy3D_DtoH_T(), and stream.

337  {
338  copy3D_DtoH_T(src_data, dst_data,
339  src_x0, src_y0, src_z0,
340  src_xsize, src_ysize,
341  dst_x0, dst_y0, dst_z0,
342  dst_xsize, dst_ysize,
343  width, height, depth,
344  sizeof(T), stream);
345 }
__thread cudaStream_t stream
void copy3D_DtoH_T(void *src_data, void *dst_data, int src_x0, int src_y0, int src_z0, size_t src_xsize, size_t src_ysize, int dst_x0, int dst_y0, int dst_z0, size_t dst_xsize, size_t dst_ysize, size_t width, size_t height, size_t depth, size_t sizeofT, cudaStream_t stream)
Definition: CudaUtils.C:250
void copy3D_DtoH_T ( void src_data,
void dst_data,
int  src_x0,
int  src_y0,
int  src_z0,
size_t  src_xsize,
size_t  src_ysize,
int  dst_x0,
int  dst_y0,
int  dst_z0,
size_t  dst_xsize,
size_t  dst_ysize,
size_t  width,
size_t  height,
size_t  depth,
size_t  sizeofT,
cudaStream_t  stream 
)

Definition at line 250 of file CudaUtils.C.

References cudaCheck.

Referenced by copy3D_DtoH().

256  {
257  cudaMemcpy3DParms parms = {0};
258 
259  parms.srcPos = make_cudaPos(sizeofT*src_x0, src_y0, src_z0);
260  parms.srcPtr = make_cudaPitchedPtr(src_data, sizeofT*src_xsize, src_xsize, src_ysize);
261 
262  parms.dstPos = make_cudaPos(sizeofT*dst_x0, dst_y0, dst_z0);
263  parms.dstPtr = make_cudaPitchedPtr(dst_data, sizeofT*dst_xsize, dst_xsize, dst_ysize);
264 
265  parms.extent = make_cudaExtent(sizeofT*width, height, depth);
266  parms.kind = cudaMemcpyDeviceToHost;
267 
268  cudaCheck(cudaMemcpy3DAsync(&parms, stream));
269 }
__thread cudaStream_t stream
#define cudaCheck(stmt)
Definition: CudaUtils.h:79
template<class T >
void copy3D_HtoD ( T *  src_data,
T *  dst_data,
int  src_x0,
int  src_y0,
int  src_z0,
size_t  src_xsize,
size_t  src_ysize,
int  dst_x0,
int  dst_y0,
int  dst_z0,
size_t  dst_xsize,
size_t  dst_ysize,
size_t  width,
size_t  height,
size_t  depth,
cudaStream_t  stream = 0 
)

Definition at line 302 of file CudaUtils.h.

References copy3D_HtoD_T(), and stream.

308  {
309  copy3D_HtoD_T(src_data, dst_data,
310  src_x0, src_y0, src_z0,
311  src_xsize, src_ysize,
312  dst_x0, dst_y0, dst_z0,
313  dst_xsize, dst_ysize,
314  width, height, depth,
315  sizeof(T), stream);
316 }
void copy3D_HtoD_T(void *src_data, void *dst_data, int src_x0, int src_y0, int src_z0, size_t src_xsize, size_t src_ysize, int dst_x0, int dst_y0, int dst_z0, size_t dst_xsize, size_t dst_ysize, size_t width, size_t height, size_t depth, size_t sizeofT, cudaStream_t stream)
Definition: CudaUtils.C:225
__thread cudaStream_t stream
void copy3D_HtoD_T ( void src_data,
void dst_data,
int  src_x0,
int  src_y0,
int  src_z0,
size_t  src_xsize,
size_t  src_ysize,
int  dst_x0,
int  dst_y0,
int  dst_z0,
size_t  dst_xsize,
size_t  dst_ysize,
size_t  width,
size_t  height,
size_t  depth,
size_t  sizeofT,
cudaStream_t  stream 
)

Definition at line 225 of file CudaUtils.C.

References cudaCheck.

Referenced by copy3D_HtoD().

231  {
232  cudaMemcpy3DParms parms = {0};
233 
234  parms.srcPos = make_cudaPos(sizeofT*src_x0, src_y0, src_z0);
235  parms.srcPtr = make_cudaPitchedPtr(src_data, sizeofT*src_xsize, src_xsize, src_ysize);
236 
237  parms.dstPos = make_cudaPos(sizeofT*dst_x0, dst_y0, dst_z0);
238  parms.dstPtr = make_cudaPitchedPtr(dst_data, sizeofT*dst_xsize, dst_xsize, dst_ysize);
239 
240  parms.extent = make_cudaExtent(sizeofT*width, height, depth);
241  parms.kind = cudaMemcpyHostToDevice;
242 
243  cudaCheck(cudaMemcpy3DAsync(&parms, stream));
244 }
__thread cudaStream_t stream
#define cudaCheck(stmt)
Definition: CudaUtils.h:79
template<class T >
void copy3D_PeerDtoD ( int  src_dev,
int  dst_dev,
T *  src_data,
T *  dst_data,
int  src_x0,
int  src_y0,
int  src_z0,
size_t  src_xsize,
size_t  src_ysize,
int  dst_x0,
int  dst_y0,
int  dst_z0,
size_t  dst_xsize,
size_t  dst_ysize,
size_t  width,
size_t  height,
size_t  depth,
cudaStream_t  stream = 0 
)

Definition at line 390 of file CudaUtils.h.

References copy3D_PeerDtoD_T(), and stream.

397  {
398  copy3D_PeerDtoD_T(src_dev, dst_dev,
399  src_data, dst_data,
400  src_x0, src_y0, src_z0,
401  src_xsize, src_ysize,
402  dst_x0, dst_y0, dst_z0,
403  dst_xsize, dst_ysize,
404  width, height, depth,
405  sizeof(T), stream);
406 }
__thread cudaStream_t stream
void copy3D_PeerDtoD_T(int src_dev, int dst_dev, void *src_data, void *dst_data, int src_x0, int src_y0, int src_z0, size_t src_xsize, size_t src_ysize, int dst_x0, int dst_y0, int dst_z0, size_t dst_xsize, size_t dst_ysize, size_t width, size_t height, size_t depth, size_t sizeofT, cudaStream_t stream)
Definition: CudaUtils.C:300
void copy3D_PeerDtoD_T ( int  src_dev,
int  dst_dev,
void src_data,
void dst_data,
int  src_x0,
int  src_y0,
int  src_z0,
size_t  src_xsize,
size_t  src_ysize,
int  dst_x0,
int  dst_y0,
int  dst_z0,
size_t  dst_xsize,
size_t  dst_ysize,
size_t  width,
size_t  height,
size_t  depth,
size_t  sizeofT,
cudaStream_t  stream 
)

Definition at line 300 of file CudaUtils.C.

References cudaCheck.

Referenced by copy3D_PeerDtoD().

307  {
308  cudaMemcpy3DPeerParms parms = {0};
309 
310  parms.srcDevice = src_dev;
311  parms.dstDevice = dst_dev;
312 
313  parms.srcPos = make_cudaPos(sizeofT*src_x0, src_y0, src_z0);
314  parms.srcPtr = make_cudaPitchedPtr(src_data, sizeofT*src_xsize, src_xsize, src_ysize);
315 
316  parms.dstPos = make_cudaPos(sizeofT*dst_x0, dst_y0, dst_z0);
317  parms.dstPtr = make_cudaPitchedPtr(dst_data, sizeofT*dst_xsize, dst_xsize, dst_ysize);
318 
319  parms.extent = make_cudaExtent(sizeofT*width, height, depth);
320 
321  cudaCheck(cudaMemcpy3DPeerAsync(&parms, stream));
322 }
__thread cudaStream_t stream
#define cudaCheck(stmt)
Definition: CudaUtils.h:79
template<class T >
void copy_DtoD ( const T *  d_src,
T *  h_dst,
const int  array_len,
cudaStream_t  stream = 0 
)

Definition at line 263 of file CudaUtils.h.

References copy_DtoD_async_T(), and stream.

263  {
264  copy_DtoD_async_T(d_src, h_dst, array_len, stream, sizeof(T));
265 }
__thread cudaStream_t stream
void copy_DtoD_async_T(const void *d_src, void *d_dst, const int array_len, cudaStream_t stream, const size_t sizeofT)
Definition: CudaUtils.C:202
void copy_DtoD_async_T ( const void d_src,
void d_dst,
const int  array_len,
cudaStream_t  stream,
const size_t  sizeofT 
)

Definition at line 202 of file CudaUtils.C.

References cudaCheck.

Referenced by copy_DtoD().

203  {
204  cudaCheck(cudaMemcpyAsync(d_dst, d_src, sizeofT*array_len, cudaMemcpyDeviceToDevice, stream));
205 }
__thread cudaStream_t stream
#define cudaCheck(stmt)
Definition: CudaUtils.h:79
template<class T >
void copy_DtoD_sync ( const T *  d_src,
T *  h_dst,
const int  array_len 
)

Definition at line 271 of file CudaUtils.h.

References copy_DtoD_T().

271  {
272  copy_DtoD_T(d_src, h_dst, array_len, sizeof(T));
273 }
void copy_DtoD_T(const void *d_src, void *d_dst, const int array_len, const size_t sizeofT)
Definition: CudaUtils.C:207
void copy_DtoD_T ( const void d_src,
void d_dst,
const int  array_len,
const size_t  sizeofT 
)

Definition at line 207 of file CudaUtils.C.

References cudaCheck.

Referenced by copy_DtoD_sync().

207  {
208  cudaCheck(cudaMemcpy(d_dst, d_src, sizeofT*array_len, cudaMemcpyDeviceToDevice));
209 }
#define cudaCheck(stmt)
Definition: CudaUtils.h:79
template<class T >
void copy_DtoH ( const T *  d_array,
T *  h_array,
const int  array_len,
cudaStream_t  stream = 0 
)

Definition at line 247 of file CudaUtils.h.

References copy_DtoH_async_T(), and stream.

247  {
248  copy_DtoH_async_T(d_array, h_array, array_len, stream, sizeof(T));
249 }
__thread cudaStream_t stream
void copy_DtoH_async_T(const void *d_array, void *h_array, const int array_len, cudaStream_t stream, const size_t sizeofT)
Definition: CudaUtils.C:189
void copy_DtoH_async_T ( const void d_array,
void h_array,
const int  array_len,
cudaStream_t  stream,
const size_t  sizeofT 
)

Definition at line 189 of file CudaUtils.C.

References cudaCheck.

Referenced by copy_DtoH().

190  {
191  cudaCheck(cudaMemcpyAsync(h_array, d_array, sizeofT*array_len, cudaMemcpyDeviceToHost, stream));
192 }
__thread cudaStream_t stream
#define cudaCheck(stmt)
Definition: CudaUtils.h:79
template<class T >
void copy_DtoH_sync ( const T *  d_array,
T *  h_array,
const int  array_len 
)

Definition at line 255 of file CudaUtils.h.

References copy_DtoH_T().

255  {
256  copy_DtoH_T(d_array, h_array, array_len, sizeof(T));
257 }
void copy_DtoH_T(const void *d_array, void *h_array, const int array_len, const size_t sizeofT)
Definition: CudaUtils.C:194
void copy_DtoH_T ( const void d_array,
void h_array,
const int  array_len,
const size_t  sizeofT 
)

Definition at line 194 of file CudaUtils.C.

References cudaCheck.

Referenced by copy_DtoH_sync().

194  {
195  cudaCheck(cudaMemcpy(h_array, d_array, sizeofT*array_len, cudaMemcpyDeviceToHost));
196 }
#define cudaCheck(stmt)
Definition: CudaUtils.h:79
template<class T >
void copy_HtoD ( const T *  h_array,
T *  d_array,
int  array_len,
cudaStream_t  stream = 0 
)

Definition at line 229 of file CudaUtils.h.

References copy_HtoD_async_T(), and stream.

229  {
230  copy_HtoD_async_T(h_array, d_array, array_len, stream, sizeof(T));
231 }
__thread cudaStream_t stream
void copy_HtoD_async_T(const void *h_array, void *d_array, int array_len, cudaStream_t stream, const size_t sizeofT)
Definition: CudaUtils.C:175
void copy_HtoD_async_T ( const void h_array,
void d_array,
int  array_len,
cudaStream_t  stream,
const size_t  sizeofT 
)

Definition at line 175 of file CudaUtils.C.

References cudaCheck.

Referenced by copy_HtoD().

176  {
177  cudaCheck(cudaMemcpyAsync(d_array, h_array, sizeofT*array_len, cudaMemcpyHostToDevice, stream));
178 }
__thread cudaStream_t stream
#define cudaCheck(stmt)
Definition: CudaUtils.h:79
template<class T >
void copy_HtoD_sync ( const T *  h_array,
T *  d_array,
int  array_len 
)

Definition at line 238 of file CudaUtils.h.

References copy_HtoD_T().

238  {
239  copy_HtoD_T(h_array, d_array, array_len, sizeof(T));
240 }
void copy_HtoD_T(const void *h_array, void *d_array, int array_len, const size_t sizeofT)
Definition: CudaUtils.C:180
void copy_HtoD_T ( const void h_array,
void d_array,
int  array_len,
const size_t  sizeofT 
)

Definition at line 180 of file CudaUtils.C.

References cudaCheck.

Referenced by copy_HtoD_sync().

181  {
182  cudaCheck(cudaMemcpy(d_array, h_array, sizeofT*array_len, cudaMemcpyHostToDevice));
183 }
#define cudaCheck(stmt)
Definition: CudaUtils.h:79
template<class T >
void copy_PeerDtoD ( const int  src_dev,
const int  dst_dev,
const T *  d_src,
T *  d_dst,
const int  array_len,
cudaStream_t  stream = 0 
)

Definition at line 284 of file CudaUtils.h.

References copy_PeerDtoD_async_T(), and stream.

285  {
286  copy_PeerDtoD_async_T(src_dev, dst_dev, d_src, d_dst, array_len, stream, sizeof(T));
287 }
void copy_PeerDtoD_async_T(const int src_dev, const int dst_dev, const void *d_src, void *d_dst, const int array_len, cudaStream_t stream, const size_t sizeofT)
Definition: CudaUtils.C:215
__thread cudaStream_t stream
void copy_PeerDtoD_async_T ( const int  src_dev,
const int  dst_dev,
const void d_src,
void d_dst,
const int  array_len,
cudaStream_t  stream,
const size_t  sizeofT 
)

Definition at line 215 of file CudaUtils.C.

References cudaCheck.

Referenced by copy_PeerDtoD().

217  {
218  cudaCheck(cudaMemcpyPeerAsync(d_dst, dst_dev, d_src, src_dev, sizeofT*array_len, stream));
219 }
__thread cudaStream_t stream
#define cudaCheck(stmt)
Definition: CudaUtils.h:79
void cudaDie ( const char *  msg,
cudaError_t  err = cudaSuccess 
)

Definition at line 9 of file CudaUtils.C.

References NAMD_die().

9  {
10  char host[128];
11  gethostname(host, 128); host[127] = 0;
12  char devstr[128] = "";
13  int devnum;
14  if ( cudaGetDevice(&devnum) == cudaSuccess ) {
15  sprintf(devstr, " device %d", devnum);
16  }
17  cudaDeviceProp deviceProp;
18  if ( cudaGetDeviceProperties(&deviceProp, devnum) == cudaSuccess ) {
19  sprintf(devstr, " device %d pci %x:%x:%x", devnum,
20  deviceProp.pciDomainID, deviceProp.pciBusID, deviceProp.pciDeviceID);
21  }
22  char errmsg[1024];
23  if (err == cudaSuccess) {
24  sprintf(errmsg,"CUDA error %s on Pe %d (%s%s)", msg, CkMyPe(), host, devstr);
25  } else {
26  sprintf(errmsg,"CUDA error %s on Pe %d (%s%s): %s", msg, CkMyPe(), host, devstr, cudaGetErrorString(err));
27  }
28  NAMD_die(errmsg);
29 }
void NAMD_die(const char *err_msg)
Definition: common.C:83
void cudaNAMD_bug ( const char *  msg)

Definition at line 31 of file CudaUtils.C.

References NAMD_bug().

Referenced by CudaFFTCompute::backward(), CudaFFTCompute::forward(), gather_force(), and spread_charge().

31 {NAMD_bug(msg);}
void NAMD_bug(const char *err_msg)
Definition: common.C:123
template<class T >
void deallocate_device ( T **  pp)

Definition at line 158 of file CudaUtils.h.

References deallocate_device_T().

158  {
159  deallocate_device_T((void **)pp);
160 }
void deallocate_device_T(void **pp)
Definition: CudaUtils.C:84
void deallocate_device_T ( void **  pp)

Definition at line 84 of file CudaUtils.C.

References cudaCheck.

Referenced by deallocate_device().

84  {
85 
86  if (*pp != NULL) {
87  cudaCheck(cudaFree((void *)(*pp)));
88  *pp = NULL;
89  }
90 
91 }
#define cudaCheck(stmt)
Definition: CudaUtils.h:79
template<class T >
void deallocate_host ( T **  pp)

Definition at line 207 of file CudaUtils.h.

References deallocate_host_T().

207  {
208  deallocate_host_T((void **)pp);
209 }
void deallocate_host_T(void **pp)
Definition: CudaUtils.C:98
void deallocate_host_T ( void **  pp)

Definition at line 98 of file CudaUtils.C.

References cudaCheck.

Referenced by deallocate_host().

98  {
99 
100  if (*pp != NULL) {
101  cudaCheck(cudaFreeHost((void *)(*pp)));
102  *pp = NULL;
103  }
104 
105 }
#define cudaCheck(stmt)
Definition: CudaUtils.h:79
template<class T >
bool reallocate_device ( T **  pp,
int *  curlen,
const int  newlen,
const float  fac = 1.0f 
)

Definition at line 175 of file CudaUtils.h.

References reallocate_device_T().

175  {
176  return reallocate_device_T((void **)pp, curlen, newlen, fac, sizeof(T));
177 }
bool reallocate_device_T(void **pp, int *curlen, const int newlen, const float fac, const size_t sizeofT)
Definition: CudaUtils.C:117
bool reallocate_device_T ( void **  pp,
int *  curlen,
const int  newlen,
const float  fac,
const size_t  sizeofT 
)

Definition at line 117 of file CudaUtils.C.

References cudaCheck.

Referenced by reallocate_device().

117  {
118 
119  if (*pp != NULL && *curlen < newlen) {
120  cudaCheck(cudaFree((void *)(*pp)));
121  *pp = NULL;
122  }
123 
124  if (*pp == NULL) {
125  if (fac > 1.0f) {
126  *curlen = (int)(((double)(newlen))*(double)fac);
127  } else {
128  *curlen = newlen;
129  }
130  cudaCheck(cudaMalloc(pp, sizeofT*(*curlen)));
131  return true;
132  }
133 
134  return false;
135 }
#define cudaCheck(stmt)
Definition: CudaUtils.h:79
template<class T >
bool reallocate_host ( T **  pp,
int *  curlen,
const int  newlen,
const float  fac = 1.0f,
const unsigned int  flag = cudaHostAllocDefault 
)

Definition at line 195 of file CudaUtils.h.

References reallocate_host_T().

196  {
197  return reallocate_host_T((void **)pp, curlen, newlen, fac, flag, sizeof(T));
198 }
bool reallocate_host_T(void **pp, int *curlen, const int newlen, const float fac, const unsigned int flag, const size_t sizeofT)
Definition: CudaUtils.C:150
bool reallocate_host_T ( void **  pp,
int *  curlen,
const int  newlen,
const float  fac,
const unsigned int  flag,
const size_t  sizeofT 
)

Definition at line 150 of file CudaUtils.C.

References cudaCheck.

Referenced by reallocate_host().

151  {
152 
153  if (*pp != NULL && *curlen < newlen) {
154  cudaCheck(cudaFreeHost((void *)(*pp)));
155  *pp = NULL;
156  }
157 
158  if (*pp == NULL) {
159  if (fac > 1.0f) {
160  *curlen = (int)(((double)(newlen))*(double)fac);
161  } else {
162  *curlen = newlen;
163  }
164  cudaCheck(cudaHostAlloc(pp, sizeofT*(*curlen), flag));
165  return true;
166  }
167 
168  return false;
169 }
#define cudaCheck(stmt)
Definition: CudaUtils.h:79