Revision e07333e63a376a103d320cc43c22df0b78b7f2a8 authored by Tim Haines on 30 October 2018, 03:15:40 UTC, committed by Tim Haines on 11 November 2018, 10:27:44 UTC
This requires https://charm.cs.illinois.edu/gerrit/#/c/cosmo/changa/+/3393/ https://charm.cs.illinois.edu/gerrit/#/c/cosmo/changa/+/3602/ https://charm.cs.illinois.edu/gerrit/#/c/charm/+/4753/ https://charm.cs.illinois.edu/gerrit/#/c/charm/+/4755/ Change-Id: I4d0adea6284941a2d1a62ace98faaeea39a73b3c
1 parent 2fbebec
EwaldCUDA.h
#ifndef _EWALD_CUDA_H_
#define _EWALD_CUDA_H_
#include "HostCUDA.h"
/* defines for Hybrid API buffer indices */
#define EWALD_READ_ONLY_DATA 0
#define EWALD_TABLE 1
/* Defines for place in the hapiBufferInfo array */
#define PARTICLE_TABLE_IDX 0
#define EWALD_READ_ONLY_DATA_IDX 1
#define EWALD_TABLE_IDX 2
#define NEWH 80
#define BLOCK_SIZE 128
/* See "defines for Hybrid API buffer indices" in HostCUDA.h for this
* number. */
#define NUM_GRAVITY_BUFS 5
/** @brief Data for the Ewald h loop in the CUDA kernel
*/
typedef struct {
cudatype hx, hy, hz;
cudatype hCfac, hSfac;
} EwtData;
/** @brief CUDA version of complete MultipoleMoments for Ewald
*/
typedef struct {
#ifndef HEXADECAPOLE
cudatype xx, xy, xz, yy, yz, zz;
#endif
cudatype totalMass;
cudatype cmx, cmy, cmz;
} MultipoleMomentsData;
/** @brief CUDA version of MOMC for Ewald
*/
typedef struct {
cudatype m;
cudatype xx,yy,xy,xz,yz;
cudatype xxx,xyy,xxy,yyy,xxz,yyz,xyz;
cudatype xxxx,xyyy,xxxy,yyyy,xxxz,yyyz,xxyy,xxyz,xyyz;
cudatype zz;
cudatype xzz,yzz,zzz;
cudatype xxzz,xyzz,xzzz,yyzz,yzzz,zzzz;
} MomcData;
/** @brief Parameters and data for Ewald in the CUDA kernel
*/
typedef struct {
MultipoleMomentsData mm;
MomcData momcRoot;
int n, nReps, nEwReps, nEwhLoop;
cudatype L, fEwCut, alpha, alpha2, k1, ka, fEwCut2, fInner2;
} EwaldReadOnlyData;
/// @brief structure to hold information specific to GPU Ewald
typedef struct {
int EwaldRange[2]; /**< First and last particle on the
* GPU; only used for small phase */
int *EwaldMarkers; /**< indices of active particles */
EwtData *ewt; /**< h-loop table */
EwaldReadOnlyData *cachedData; /**< Root moment and other Ewald parameters */
} EwaldData;
void EwaldHostMemorySetup(EwaldData *h_idata, int size, int nEwhLoop, int largephase);
void EwaldHostMemoryFree(EwaldData *h_idata, int largephase);
#ifdef HAPI_INSTRUMENT_WRS
void EwaldHost(EwaldData *h_idata, void *cb, int myIndex, char phase, int largephase);
#else
void EwaldHost(EwaldData *h_idata, void *cb, int myIndex, int largephase);
#endif
__global__ void EwaldKernel(CompactPartData *particleCores, VariablePartData *particleVars, int *markers, int largephase, int First, int Last);
#endif
Computing file changes ...