NAMD
Macros | Functions | Variables
ComputePmeCUDAKernel.h File Reference

Go to the source code of this file.

Macros

#define __align__(X)
 
#define CUDA_PME_CHARGES_PROTOTYPE
 
#define CUDA_PME_CHARGES_BATCHED_PROTOTYPE
 
#define CUDA_PME_FORCES_PROTOTYPE
 

Functions

void cuda_init_bspline_coeffs (float **c, float **dc, int order)
 

Variables

 CUDA_PME_CHARGES_PROTOTYPE
 
 CUDA_PME_CHARGES_BATCHED_PROTOTYPE
 
 CUDA_PME_FORCES_PROTOTYPE
 

Macro Definition Documentation

#define __align__ (   X)

Definition at line 5 of file ComputePmeCUDAKernel.h.

#define CUDA_PME_CHARGES_BATCHED_PROTOTYPE
Value:
void cuda_pme_charges_batched( \
const float *coeffs, \
float * const *q_arr, int *f_arr, int *fz_arr, \
float **a_data_ptr, int* n_atoms_ptr, \
int* K1_ptr, int* K2_ptr, int* K3_ptr, \
int order, int numPatches, int n_max_atoms, cudaStream_t stream)
__global__ void const int const TileList *__restrict__ TileExcl *__restrict__ const int *__restrict__ const int const float2 *__restrict__ const int *__restrict__ const float3 const float3 const float3 const float4 *__restrict__ const float cudaTextureObject_t cudaTextureObject_t cudaTextureObject_t const int const float const PatchPairRecord *__restrict__ const int *__restrict__ const int2 *__restrict__ const unsigned int *__restrict__ unsigned int *__restrict__ int *__restrict__ int *__restrict__ TileListStat *__restrict__ const BoundingBox *__restrict__ float4 *__restrict__ float4 *__restrict__ float *__restrict__ float *__restrict__ float *__restrict__ float *__restrict__ float *__restrict__ float *__restrict__ float *__restrict__ float *__restrict__ const int numPatches
__thread cudaStream_t stream
#define order
Definition: PmeRealSpace.C:235

Definition at line 21 of file ComputePmeCUDAKernel.h.

#define CUDA_PME_CHARGES_PROTOTYPE
Value:
void cuda_pme_charges( \
const float *coeffs, \
float * const *q_arr, int *f_arr, int *fz_arr, \
float *a_data, int n_atoms, \
int K1, int K2, int K3, \
int order, cudaStream_t stream)
__thread cudaStream_t stream
#define order
Definition: PmeRealSpace.C:235

Definition at line 10 of file ComputePmeCUDAKernel.h.

#define CUDA_PME_FORCES_PROTOTYPE
Value:
void cuda_pme_forces( \
const float *coeffs, \
float * const *q_arr, \
float * const *afn, int dimy, int maxn, \
/* float *a_data, float *f_data, int n_atoms, */ \
int K1, int K2, int K3, \
int order, cudaStream_t stream)
__thread cudaStream_t stream
#define order
Definition: PmeRealSpace.C:235

Definition at line 32 of file ComputePmeCUDAKernel.h.

Function Documentation

void cuda_init_bspline_coeffs ( float **  c,
float **  dc,
int  order 
)

Definition at line 43 of file ComputePmeCUDAKernel.cu.

References NAMD_die(), order, order_4_coeffs, order_6_coeffs, and order_8_coeffs.

Referenced by ComputePmeMgr::initialize_computes().

43  {
44  float *coeffs = new float[order*order];
45  float *dcoeffs = new float[order*order];
46  double divisor;
47  static const float *scoeffs;
48  switch ( order ) {
49  case 4:
50  scoeffs = &order_4_coeffs[0][0];
51  divisor = 6;
52  break;
53  case 6:
54  scoeffs = &order_6_coeffs[0][0];
55  divisor = 120;
56  break;
57  case 8:
58  scoeffs = &order_8_coeffs[0][0];
59  divisor = 5040;
60  break;
61  default:
62  NAMD_die("unsupported PMEInterpOrder");
63  }
64  double sum = 0;
65  for ( int i=0, p=order-1; i<order; ++i,--p ) {
66  for ( int j=0; j<order; ++j ) {
67  double c = scoeffs[i*order+(order-1-j)]; // reverse order
68  sum += c;
69  c /= divisor;
70  coeffs[i*order+j] = c;
71  dcoeffs[i*order+j] = (double)p * c;
72  // printf("%d %d %f %f\n", i, j, c, (double)p*c);
73  }
74  }
75  // printf("checksum: %f %f\n", sum, divisor);
76  if ( sum != divisor )
77  NAMD_die("cuda_init_bspline_coeffs static data checksum error");
78  cudaMalloc((void**) c, order*order*sizeof(float));
79  cudaMalloc((void**) dc, order*order*sizeof(float));
80  cudaMemcpy(*c, coeffs, order*order*sizeof(float), cudaMemcpyHostToDevice);
81  cudaMemcpy(*dc, dcoeffs, order*order*sizeof(float), cudaMemcpyHostToDevice);
82  delete [] coeffs;
83  delete [] dcoeffs;
84 }
static const float order_8_coeffs[8][8]
#define order
Definition: PmeRealSpace.C:235
void NAMD_die(const char *err_msg)
Definition: common.C:83
static const float order_4_coeffs[4][4]
static const float order_6_coeffs[6][6]

Variable Documentation

CUDA_PME_CHARGES_BATCHED_PROTOTYPE

Definition at line 29 of file ComputePmeCUDAKernel.h.

CUDA_PME_CHARGES_PROTOTYPE

Definition at line 18 of file ComputePmeCUDAKernel.h.

CUDA_PME_FORCES_PROTOTYPE

Definition at line 41 of file ComputePmeCUDAKernel.h.