Homec4science
Diffusion lammps (lammps-icms) Browse

lammps/lib/gpu/device.cu
9ee36a3c2298lammps-icms

Last Change
Raw File

device.cu
View Options

// **************************************************************************
// device.cu
// -------------------
// W. Michael Brown (ORNL)
//
// Device code for device information
//
// __________________________________________________________________________
// This file is part of the LAMMPS Accelerator Library (LAMMPS_AL)
// __________________________________________________________________________
//
// begin :
// email : brownw@ornl.gov
// ***************************************************************************/
/*************************************************************************
Preprocessor Definitions
Note: It is assumed that constants with the same names are defined with
the same values in all files.
ARCH
Definition: Architecture number for accelerator
MEM_THREADS
Definition: Number of threads with sequential ids accessing memory
simultaneously on multiprocessor
WARP_SIZE:
Definition: Number of threads guaranteed to be on the same instruction
THREADS_PER_ATOM
Definition: Default number of threads assigned per atom for pair styles
Restructions: Must be power of 2; THREADS_PER_ATOM<=WARP_SIZE
THREADS_PER_CHARGE
Definition: Default number of threads assigned per atom for pair styles
with charge
Restructions: Must be power of 2; THREADS_PER_ATOM<=WARP_SIZE
PPPM_MAX_SPLINE
Definition: Maximum order for splines in PPPM
PPPM_BLOCK_1D
Definition: Thread block size for PPPM kernels
Restrictions: PPPM_BLOCK_1D>=PPPM_MAX_SPLINE*PPPM_MAX_SPLINE
PPPM_BLOCK_1D%32==0
BLOCK_PAIR
Definition: Default thread block size for pair styles
Restrictions:
MAX_SHARED_TYPES 8
Definition: Max number of atom type params can be stored in shared memory
Restrictions: MAX_SHARED_TYPES*MAX_SHARED_TYPES<=BLOCK_PAIR
BLOCK_CELL_2D
Definition: Default block size in each dimension for cell list builds
and matrix transpose
BLOCK_CELL_ID
Definition: Default block size for binning atoms in cell list builds
BLOCK_NBOR_BUILD
Definition: Default block size for neighbor list builds
BLOCK_BIO_PAIR
Definition: Default thread block size for "bio" pair styles
MAX_BIO_SHARED_TYPES
Definition: Max number of atom type params can be stored in shared memory
Restrictions: MAX_BIO_SHARED_TYPES<=BLOCK_BIO_PAIR*2 &&
MAX_BIO_SHARED_TYPES>=BLOCK_BIO_PAIR
*************************************************************************/
#ifdef NV_KERNEL
#include "nv_kernel_def.h"
#else
#define GLOBAL_ID_X get_global_id(0)
#define ARCH 0
#define DRIVER 0
#define MEM_THREADS 16
#define WARP_SIZE 1
#define THREADS_PER_ATOM 1
#define THREADS_PER_CHARGE 1
#define BLOCK_PAIR 64
#define MAX_SHARED_TYPES 8
#define BLOCK_NBOR_BUILD 64
#define BLOCK_BIO_PAIR 64
#endif
#define PPPM_MAX_SPLINE 8
#define PPPM_BLOCK_1D 64
#define BLOCK_CELL_2D 8
#define BLOCK_CELL_ID 128
#define MAX_BIO_SHARED_TYPES 128
__kernel void kernel_zero(__global int *mem, int numel) {
int ii=GLOBAL_ID_X;
if (ii<numel)
mem[ii]=0;
}
__kernel void kernel_info(__global int *info) {
info[0]=ARCH;
info[1]=MEM_THREADS;
info[2]=WARP_SIZE;
info[3]=THREADS_PER_ATOM;
info[4]=PPPM_MAX_SPLINE;
info[5]=PPPM_BLOCK_1D;
info[6]=BLOCK_PAIR;
info[7]=MAX_SHARED_TYPES;
info[8]=BLOCK_CELL_2D;
info[9]=BLOCK_CELL_ID;
info[10]=BLOCK_NBOR_BUILD;
info[11]=BLOCK_BIO_PAIR;
info[12]=MAX_BIO_SHARED_TYPES;
info[13]=THREADS_PER_CHARGE;
}
c4science · Help