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 8 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)

Definition at line 24 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)

Definition at line 13 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)

Definition at line 35 of file ComputePmeCUDAKernel.h.

Function Documentation

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

Definition at line 47 of file ComputePmeCUDAKernel.cu.

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

Referenced by ComputePmeMgr::initialize_computes().

                                                                {
  float *coeffs = new float[order*order];
  float *dcoeffs = new float[order*order];
  double divisor;
  static const float *scoeffs;
  switch ( order ) {
  case 4:
    scoeffs = &order_4_coeffs[0][0];
    divisor = 6;
    break;
  case 6:
    scoeffs = &order_6_coeffs[0][0];
    divisor = 120;
    break;
  case 8:
    scoeffs = &order_8_coeffs[0][0];
    divisor = 5040;
    break;
  default:
    NAMD_die("unsupported PMEInterpOrder");
  }
  double sum = 0;
  for ( int i=0, p=order-1; i<order; ++i,--p ) {
    for ( int j=0; j<order; ++j ) {
      double c = scoeffs[i*order+(order-1-j)];  // reverse order
      sum += c;
      c /= divisor;
      coeffs[i*order+j] = c;
      dcoeffs[i*order+j] = (double)p * c;
      // printf("%d %d %f %f\n", i, j, c, (double)p*c);
    }
  }
  // printf("checksum: %f %f\n", sum, divisor);
  if ( sum != divisor )
    NAMD_die("cuda_init_bspline_coeffs static data checksum error");
  cudaMalloc((void**) c, order*order*sizeof(float));
  cudaMalloc((void**) dc, order*order*sizeof(float));
  cudaMemcpy(*c, coeffs, order*order*sizeof(float), cudaMemcpyHostToDevice);
  cudaMemcpy(*dc, dcoeffs, order*order*sizeof(float), cudaMemcpyHostToDevice);
  delete [] coeffs;
  delete [] dcoeffs;
}

Variable Documentation

CUDA_PME_CHARGES_BATCHED_PROTOTYPE

Definition at line 32 of file ComputePmeCUDAKernel.h.

CUDA_PME_CHARGES_PROTOTYPE

Definition at line 21 of file ComputePmeCUDAKernel.h.

CUDA_PME_FORCES_PROTOTYPE

Definition at line 44 of file ComputePmeCUDAKernel.h.