NAMD
CudaPmeSolverUtilKernel.h
Go to the documentation of this file.
1 #ifndef CUDAPMESOLVERUTILKERNEL_H
2 #define CUDAPMESOLVERUTILKERNEL_H
3 #include "HipDefines.h"
4 #if defined(NAMD_CUDA) || defined(NAMD_HIP)
5 
6 void spread_charge(const float4 *atoms, const int numAtoms,
7  const int nfftx, const int nffty, const int nfftz,
8  const int xsize, const int ysize, const int zsize,
9  const int xdim, const int y00, const int z00,
10  const bool periodicY, const bool periodicZ,
11  float* data, const int order, cudaStream_t stream);
12 
13 void scalar_sum(const bool orderXYZ, const int nfft1, const int nfft2, const int nfft3,
14  const int size1, const int size2, const int size3, const double kappa,
15  const float recip1x, const float recip1y, const float recip1z,
16  const float recip2x, const float recip2y, const float recip2z,
17  const float recip3x, const float recip3y, const float recip3z,
18  const double volume,
19  const float* prefac1, const float* prefac2, const float* prefac3,
20  const int k2_00, const int k3_00,
21  const bool doEnergyVirial, double* energy, double* virial, float2* data,
22  // const int cuda_arch,
23  cudaStream_t stream);
24 
25 void gather_force(const float4 *atoms, const int numAtoms,
26  // const float recip11, const float recip22, const float recip33,
27  const int nfftx, const int nffty, const int nfftz,
28  const int xsize, const int ysize, const int zsize,
29  const int xdim, const int y00, const int z00,
30  const bool periodicY, const bool periodicZ,
31  const float* data, const int order, float3* force,
32 #ifdef NAMD_CUDA
33  const cudaTextureObject_t gridTexObj,
34 #endif
35  cudaStream_t stream);
36 
37 // void calc_sum_charge_squared(const float4 *atoms, const int numAtoms, double* sum_charge_squared,
38 // cudaStream_t stream);
39 
40 template <typename T>
42  T* data_in;
44  int nx;
45  int zsize_out;
46  int xsize_out;
47  int ysize_out;
48 };
49 
51  const int nx, const int ny, const int nz,
52  const int xsize_in, const int ysize_in,
53  const int ysize_out, const int zsize_out,
54  const float2* data_in, float2* data_out, cudaStream_t stream);
55 
57  const int numBatches, TransposeBatch<float2>* batches,
58  const int max_nx, const int ny, const int nz,
59  const int xsize_in, const int ysize_in, cudaStream_t stream);
60 
62  const int nx, const int ny, const int nz,
63  const int xsize_in, const int ysize_in,
64  const int zsize_out, const int xsize_out,
65  const float2* data_in, float2* data_out, cudaStream_t stream);
66 
68  const int numBatches, TransposeBatch<float2>* batches,
69  const int max_nx, const int ny, const int nz,
70  const int xsize_in, const int ysize_in,
71  cudaStream_t stream);
72 
73 #endif // NAMD_CUDA
74 
75 #endif // CUDAPMESOLVERUTILKERNEL_H
void batchTranspose_xyz_yzx(const int numBatches, TransposeBatch< float2 > *batches, const int max_nx, const int ny, const int nz, const int xsize_in, const int ysize_in, cudaStream_t stream)
void scalar_sum(const bool orderXYZ, const int nfft1, const int nfft2, const int nfft3, const int size1, const int size2, const int size3, const double kappa, const float recip1x, const float recip1y, const float recip1z, const float recip2x, const float recip2y, const float recip2z, const float recip3x, const float recip3y, const float recip3z, const double volume, const float *prefac1, const float *prefac2, const float *prefac3, const int k2_00, const int k3_00, const bool doEnergyVirial, double *energy, double *virial, float2 *data, cudaStream_t stream)
static __thread atom * atoms
void transpose_xyz_zxy(const int nx, const int ny, const int nz, const int xsize_in, const int ysize_in, const int zsize_out, const int xsize_out, const float2 *data_in, float2 *data_out, cudaStream_t stream)
__thread cudaStream_t stream
#define order
Definition: PmeRealSpace.C:235
void batchTranspose_xyz_zxy(const int numBatches, TransposeBatch< float2 > *batches, const int max_nx, const int ny, const int nz, const int xsize_in, const int ysize_in, cudaStream_t stream)
__global__ void gather_force(const float4 *xyzq, const int ncoord, const int nfftx, const int nffty, const int nfftz, const int xsize, const int ysize, const int zsize, const int xdim, const int y00, const int z00, const float *data, const cudaTextureObject_t gridTexObj, const int stride, FT *force)
void spread_charge(const float4 *atoms, const int numAtoms, const int nfftx, const int nffty, const int nfftz, const int xsize, const int ysize, const int zsize, const int xdim, const int y00, const int z00, const bool periodicY, const bool periodicZ, float *data, const int order, cudaStream_t stream)
void transpose_xyz_yzx(const int nx, const int ny, const int nz, const int xsize_in, const int ysize_in, const int ysize_out, const int zsize_out, const float2 *data_in, float2 *data_out, cudaStream_t stream)