cuda_common.h
No OneTemporary
Actions

Subscribers

None

File Metadata

Created: Tue, Nov 12, 11:41

cuda_common.h
View Options

	/* ----------------------------------------------------------------------
	LAMMPS - Large-scale Atomic/Molecular Massively Parallel Simulator

	Original Version:
	http://lammps.sandia.gov, Sandia National Laboratories
	Steve Plimpton, sjplimp@sandia.gov

	See the README file in the top-level LAMMPS directory.

	-----------------------------------------------------------------------

	USER-CUDA Package and associated modifications:
	https://sourceforge.net/projects/lammpscuda/

	Christian Trott, christian.trott@tu-ilmenau.de
	Lars Winterfeld, lars.winterfeld@tu-ilmenau.de
	Theoretical Physics II, University of Technology Ilmenau, Germany

	See the README file in the USER-CUDA directory.

	This software is distributed under the GNU General Public License.
	------------------------------------------------------------------------- */

	#ifndef _CUDA_COMMON_H_
	#define _CUDA_COMMON_H_

	//#include "cutil.h"
	#include "cuda_precision.h"
	#include "cuda_wrapper_cu.h"

	#define CUDA_MAX_TYPES_PLUS_ONE 12 //for pair styles which use constant space for parameters, this needs to be one larger than the number of atom types
	//this can not be arbitrarly large, since constant space is limited.
	//in principle one could alter potentials to use global memory for parameters, some du that already since the first examples I encountered had a high number (20+) of atom types
	//Christian
	#define CUDA_MAX_TYPES2 (CUDA_MAX_TYPES_PLUS_ONE * CUDA_MAX_TYPES_PLUS_ONE)
	#define CUDA_MAX_NSPECIAL 25

	// define some easy-to-use debug and emulation macros
	#ifdef _DEBUG
	#define MYDBG(a) a
	#else
	#define MYDBG(a)
	#endif

	#if __DEVICE_EMULATION__
	#define MYEMU(a) a
	#else
	#define MYEMU(a)
	#endif

	#define MYEMUDBG(a) MYEMU(MYDBG(a))

	// Add Prefix (needed as workaround, same constant's names in different files causes conflict)
	#define MY_ADD_PREFIX(prefix, var) prefix##_##var
	#define MY_ADD_PREFIX2(prefix, var) MY_ADD_PREFIX(prefix, var)
	#define MY_AP(var) MY_ADD_PREFIX2(MY_PREFIX, var)

	#define MY_VAR_TO_STR(var) #var
	#define MY_VAR_TO_STR2(var) MY_VAR_TO_STR(var)
	//#define &MY_AP(var) (MY_VAR_TO_STR2(MY_PREFIX) "_" MY_VAR_TO_STR2(var))
	//#define &MY_AP(var) &(MY_AP(var))
	#define CUDA_USE_TEXTURE
	#define CUDA_USE_FLOAT4

	//constants used by many classes

	//domain
	#define _boxhi MY_AP(boxhi)
	#define _boxlo MY_AP(boxlo)
	#define _subhi MY_AP(subhi)
	#define _sublo MY_AP(sublo)
	#define _box_size MY_AP(box_size)
	#define _prd MY_AP(prd)
	#define _periodicity MY_AP(periodicity)
	#define _triclinic MY_AP(triclinic)
	#define _boxhi_lamda MY_AP(boxhi_lamda)
	#define _boxlo_lamda MY_AP(boxlo_lamda)
	#define _prd_lamda MY_AP(prd_lamda)
	#define _h MY_AP(h)
	#define _h_inv MY_AP(h_inv)
	#define _h_rate MY_AP(h_rate)
	__device__ __constant__ X_FLOAT _boxhi[3];
	__device__ __constant__ X_FLOAT _boxlo[3];
	__device__ __constant__ X_FLOAT _subhi[3];
	__device__ __constant__ X_FLOAT _sublo[3];
	__device__ __constant__ X_FLOAT _box_size[3];
	__device__ __constant__ X_FLOAT _prd[3];
	__device__ __constant__ int _periodicity[3];
	__device__ __constant__ int _triclinic;
	__device__ __constant__ X_FLOAT _boxhi_lamda[3];
	__device__ __constant__ X_FLOAT _boxlo_lamda[3];
	__device__ __constant__ X_FLOAT _prd_lamda[3];
	__device__ __constant__ X_FLOAT _h[6];
	__device__ __constant__ X_FLOAT _h_inv[6];
	__device__ __constant__ V_FLOAT _h_rate[6];


	//atom properties
	#define _x MY_AP(x)
	#define _v MY_AP(v)
	#define _f MY_AP(f)
	#define _tag MY_AP(tag)
	#define _type MY_AP(type)
	#define _mask MY_AP(mask)
	#define _image MY_AP(image)
	#define _q MY_AP(q)
	#define _mass MY_AP(mass)
	#define _rmass MY_AP(rmass)
	#define _rmass_flag MY_AP(rmass_flag)
	#define _eatom MY_AP(eatom)
	#define _vatom MY_AP(vatom)
	#define _x_type MY_AP(x_type)
	#define _radius MY_AP(radius)
	#define _density MY_AP(density)
	#define _omega MY_AP(omega)
	#define _torque MY_AP(torque)
	#define _special MY_AP(special)
	#define _maxspecial MY_AP(maxspecial)
	#define _nspecial MY_AP(nspecial)
	#define _special_flag MY_AP(special_flag)
	#define _molecule MY_AP(molecule)
	#define _v_radius MY_AP(v_radius)
	#define _omega_rmass MY_AP(omega_rmass)
	#define _freeze_group_bit MY_AP(freeze_group_bit)
	#define _map_array MY_AP(map_array)
	__device__ __constant__ X_FLOAT* _x; //holds pointer to positions
	__device__ __constant__ V_FLOAT* _v;
	__device__ __constant__ F_FLOAT* _f;
	__device__ __constant__ int* _tag;
	__device__ __constant__ int* _type;
	__device__ __constant__ int* _mask;
	__device__ __constant__ int* _image;
	__device__ __constant__ V_FLOAT* _mass;
	__device__ __constant__ F_FLOAT* _q;
	__device__ __constant__ V_FLOAT* _rmass;
	__device__ __constant__ int _rmass_flag;
	__device__ __constant__ ENERGY_FLOAT* _eatom;
	__device__ __constant__ ENERGY_FLOAT* _vatom;
	__device__ __constant__ X_FLOAT4* _x_type; //holds pointer to positions
	__device__ __constant__ X_FLOAT* _radius;
	__device__ __constant__ F_FLOAT* _density;
	__device__ __constant__ V_FLOAT* _omega;
	__device__ __constant__ F_FLOAT* _torque;
	__device__ __constant__ int* _special;
	__device__ __constant__ int _maxspecial;
	__device__ __constant__ int* _nspecial;
	__device__ __constant__ int _special_flag[4];
	__device__ __constant__ int* _molecule;
	__device__ __constant__ V_FLOAT4* _v_radius; //holds pointer to positions
	__device__ __constant__ V_FLOAT4* _omega_rmass; //holds pointer to positions
	__device__ __constant__ int _freeze_group_bit;
	__device__ __constant__ int* _map_array;

	#ifdef CUDA_USE_TEXTURE

	#define _x_tex MY_AP(x_tex)
	#if X_PRECISION == 1
	texture<float> _x_tex;
	#else
	texture<int2, 1> _x_tex;
	#endif

	#define _type_tex MY_AP(type_tex)
	texture<int> _type_tex;

	#define _x_type_tex MY_AP(x_type_tex)
	#if X_PRECISION == 1
	texture<float4, 1> _x_type_tex;
	#else
	texture<int4, 1> _x_type_tex;
	#endif

	#define _v_radius_tex MY_AP(v_radius_tex)
	#if V_PRECISION == 1
	texture<float4, 1> _v_radius_tex;
	#else
	texture<int4, 1> _v_radius_tex;
	#endif

	#define _omega_rmass_tex MY_AP(omega_rmass_tex)
	#if V_PRECISION == 1
	texture<float4, 1> _omega_rmass_tex;
	#else
	texture<int4, 1> _omega_rmass_tex;
	#endif

	#define _q_tex MY_AP(q_tex)
	#if F_PRECISION == 1
	texture<float> _q_tex;
	#else
	texture<int2, 1> _q_tex;
	#endif

	#endif

	//neighbor
	#ifdef IncludeCommonNeigh
	#define _inum MY_AP(inum)
	#define _inum_border MY_AP(inum_border)
	#define _ilist MY_AP(ilist)
	#define _ilist_border MY_AP(ilist_border)
	#define _numneigh MY_AP(numneigh)
	#define _numneigh_border MY_AP(numneigh_border)
	#define _numneigh_inner MY_AP(numneigh_inner)
	#define _firstneigh MY_AP(firstneigh)
	#define _neighbors MY_AP(neighbors)
	#define _neighbors_border MY_AP(neighbors_border)
	#define _neighbors_inner MY_AP(neighbors_inner)
	#define _reneigh_flag MY_AP(reneigh_flag)
	#define _triggerneighsq MY_AP(triggerneighsq)
	#define _xhold MY_AP(xhold)
	#define _maxhold MY_AP(maxhold)
	#define _dist_check MY_AP(dist_check)
	#define _neighbor_maxlocal MY_AP(neighbor_maxlocal)
	#define _maxneighbors MY_AP(maxneighbors)
	#define _overlap_comm MY_AP(overlap_comm)
	__device__ __constant__ int _inum;
	__device__ __constant__ int* _inum_border;
	__device__ __constant__ int* _ilist;
	__device__ __constant__ int* _ilist_border;
	__device__ __constant__ int* _numneigh;
	__device__ __constant__ int* _numneigh_border;
	__device__ __constant__ int* _numneigh_inner;
	__device__ __constant__ int** _firstneigh;
	__device__ __constant__ int* _neighbors;
	__device__ __constant__ int* _neighbors_border;
	__device__ __constant__ int* _neighbors_inner;
	__device__ __constant__ int* _reneigh_flag;
	__device__ __constant__ X_FLOAT _triggerneighsq;
	__device__ __constant__ X_FLOAT* _xhold; //holds pointer to positions
	__device__ __constant__ int _maxhold;
	__device__ __constant__ int _dist_check;
	__device__ __constant__ int _neighbor_maxlocal;
	__device__ __constant__ int _maxneighbors;
	__device__ __constant__ int _overlap_comm;
	#endif

	//system properties
	#define _nall MY_AP(nall)
	#define _nghost MY_AP(nghost)
	#define _nlocal MY_AP(nlocal)
	#define _nmax MY_AP(nmax)
	#define _cuda_ntypes MY_AP(cuda_ntypes)
	#define _dtf MY_AP(dtf)
	#define _dtv MY_AP(dtv)
	#define _factor MY_AP(factor)
	#define _virial MY_AP(virial)
	#define _eng_vdwl MY_AP(eng_vdwl)
	#define _eng_coul MY_AP(eng_coul)
	#define _molecular MY_AP(molecular)
	__device__ __constant__ unsigned _nall;
	__device__ __constant__ unsigned _nghost;
	__device__ __constant__ unsigned _nlocal;
	__device__ __constant__ unsigned _nmax;
	__device__ __constant__ unsigned _cuda_ntypes;
	__device__ __constant__ V_FLOAT _dtf;
	__device__ __constant__ X_FLOAT _dtv;
	__device__ __constant__ V_FLOAT _factor;
	__device__ __constant__ ENERGY_FLOAT* _virial;
	__device__ __constant__ ENERGY_FLOAT* _eng_vdwl;
	__device__ __constant__ ENERGY_FLOAT* _eng_coul;
	__device__ __constant__ int _molecular;

	//other general constants
	#define _buffer MY_AP(buffer)
	#define _flag MY_AP(flag)
	#define _debugdata MY_AP(debugdata)
	__device__ __constant__ void* _buffer;
	__device__ __constant__ int* _flag;
	__device__ __constant__ int* _debugdata;

	// pointers to data fields on GPU are hold in constant space
	// -> reduces register usage and number of parameters for kernelcalls
	// will be variables of file scope in cuda files




	// maybe used to output cudaError_t
	#define MY_OUTPUT_RESULT(result) \
	switch(result) \
	{ \
	case cudaSuccess: printf(" => cudaSuccess\n"); break; \
	case cudaErrorInvalidValue: printf(" => cudaErrorInvalidValue\n"); break; \
	case cudaErrorInvalidSymbol: printf(" => cudaErrorInvalidSymbol\n"); break; \
	case cudaErrorInvalidDevicePointer: printf(" => cudaErrorInvalidDevicePointer\n"); break; \
	case cudaErrorInvalidMemcpyDirection: printf(" => cudaErrorInvalidMemcpyDirection\n"); break; \
	default: printf(" => unknown\n"); break; \
	}

	#ifdef _DEBUG
	# define CUT_CHECK_ERROR(errorMessage) { \
	cudaError_t err = cudaGetLastError(); \
	if( cudaSuccess != err) { \
	fprintf(stderr, "Cuda error: %s in file '%s' in line %i : %s.\n", \
	errorMessage, __FILE__, __LINE__, cudaGetErrorString( err) );\
	exit(EXIT_FAILURE); \
	} \
	err = cudaThreadSynchronize(); \
	if( cudaSuccess != err) { \
	fprintf(stderr, "Cuda error: %s in file '%s' in line %i : %s.\n", \
	errorMessage, __FILE__, __LINE__, cudaGetErrorString( err) );\
	exit(EXIT_FAILURE); \
	} \
	}
	#else
	# define CUT_CHECK_ERROR(errorMessage) { \
	cudaError_t err = cudaGetLastError(); \
	if( cudaSuccess != err) { \
	fprintf(stderr, "Cuda error: %s in file '%s' in line %i : %s.\n", \
	errorMessage, __FILE__, __LINE__, cudaGetErrorString( err) );\
	exit(EXIT_FAILURE); \
	} \
	}
	#endif

	# define CUDA_SAFE_CALL_NO_SYNC( call) { \
	cudaError err = call; \
	if( cudaSuccess != err) { \
	fprintf(stderr, "Cuda error in file '%s' in line %i : %s.\n", \
	__FILE__, __LINE__, cudaGetErrorString( err) ); \
	exit(EXIT_FAILURE); \
	} }

	# define CUDA_SAFE_CALL( call) CUDA_SAFE_CALL_NO_SYNC(call);

	#define X_MASK 1
	#define V_MASK 2
	#define F_MASK 4
	#define TAG_MASK 8
	#define TYPE_MASK 16
	#define MASK_MASK 32
	#define IMAGE_MASK 64
	#define Q_MASK 128
	#define MOLECULE_MASK 256
	#define RMASS_MASK 512
	#define RADIUS_MASK 1024
	#define DENSITY_MASK 2048
	#define OMEGA_MASK 4096
	#define TORQUE_MASK 8192



	#endif // #ifdef _CUDA_COMMON_H_

cuda_common.hNo OneTemporaryActions

File Metadata

cuda_common.hView Options

Event Timeline

cuda_common.h
No OneTemporary
Actions

cuda_common.h
View Options