Lennard-Jones
[Force Calculation]

Functions

__device__ void gpuAddLJForce (FLOAT_ARRAY_TYPE *dForce, FLOAT_ARRAY_TYPE *dA, FLOAT_ARRAY_TYPE *dB, float *boxsize)
 Add Lennard-Jones force to one particle.
__device__ void gpuAddLJForceCut (FLOAT_ARRAY_TYPE *dForce, FLOAT_ARRAY_TYPE *dA, FLOAT_ARRAY_TYPE *dB, float *boxsize, float *rCutSq)
 Add Lennard-Jones force to one particle under respect of a cutoff.
__global__ void gpuLJForces (FLOAT_ARRAY_TYPE *dPos, FLOAT_ARRAY_TYPE *dForce, int N, float boxsize)
 'Brute force' Lennard-Jones force calculation kernel
__global__ void gpuLJForcesCut (FLOAT_ARRAY_TYPE *dPos, FLOAT_ARRAY_TYPE *dForce, int N, float boxsize, float rCutSq)
 'Brute force' Lennard-Jones force calculation kernel with potential cutoff
__global__ void gpuLJForcesCutRSmem (FLOAT_ARRAY_TYPE *dPos, FLOAT_ARRAY_TYPE *dForce, int N, float boxsize, float rCutSq)
 'Brute force' Lennard-Jones force calculation kernel with potential cutoff - using shared memory for reference particles
__global__ void gpuLJForcesCutSmem (FLOAT_ARRAY_TYPE *dPos, FLOAT_ARRAY_TYPE *dForce, int N, float boxsize, float rCutSq)
 'Brute force' Lennard-Jones force calculation kernel with potential cutoff - using registers for reference particles, shared memory for others
__global__ void gpuLJForcesCutSmemSmem (FLOAT_ARRAY_TYPE *dPos, FLOAT_ARRAY_TYPE *dForce, int N, float boxsize, float rCutSq)
 'Brute force' Lennard-Jones force calculation kernel with potential cutoff using shared memory, Version 2
__global__ void gpuLJForcesCutVar1 (FLOAT_ARRAY_TYPE *dPos, FLOAT_ARRAY_TYPE *dForce, int N, float boxsize, float rCutSq)
 'Brute force' Lennard-Jones force calculation kernel with potential cutoff - using registers for reference particles
__global__ void gpuLJForcesSmem (FLOAT_ARRAY_TYPE *dPos, FLOAT_ARRAY_TYPE *dForce, int N, float boxsize)
 'Brute force' Lennard-Jones force calculation kernel using shared memory
__global__ void gpuLJForcesSmemSmem (FLOAT_ARRAY_TYPE *dPos, FLOAT_ARRAY_TYPE *dForce, int N, float boxsize)
 'Brute force' Lennard-Jones force calculation kernel using shared memory for all coordinates
__global__ void gpuLJForcesVerletList (FLOAT_ARRAY_TYPE *dPos, FLOAT_ARRAY_TYPE *dForce, int N, int *dVlists, size_t vlists_pitch, int *dVcount, float boxsize, float rCutSq)
 Verlet List Lennard-Jones force calculation kernel with potential cutoff.

Detailed Description

This module contains functions to compute Lennard-Jones forces.

Function Documentation

__device__ void gpuAddLJForce ( FLOAT_ARRAY_TYPE *  dForce,
FLOAT_ARRAY_TYPE *  dA,
FLOAT_ARRAY_TYPE *  dB,
float *  boxsize 
) [inline]

The Lennard-Jones force $ \vec{F}_{LJ}(\vec{a})_{ab} $ between particle a and b is calculated and added to the f (belonging to particle a): $ \vec{F}_{LJ}(\vec{a})_{ab}=\frac{48\epsilon\vec{r}_{ab}}{r_{ab}^2}\left[\left[\frac{\sigma}{r_{ab}}\right]^{12}-\left[\frac{\sigma}{2r_{ab}}\right]^{6}\right] $ , where $ r_{ab}=\left| \vec{a} - \vec{b} \right| $ . Due to usage of Lennard-Jones units $ \epsilon = 1 $ and $ \sigma = 1 $ .

Parameters:
[out] dForce Force contribution to particle a due to Lennard-Jones interaction with particle b
[in] dA Position in device memory of particle a
[in] dB Position in device memory of particle b
[in] boxsize Edge length of the cubic simulation box.
Returns:
Returns CUDA error information.

Definition at line 27 of file gpu_forces.cu.

References FLOAT_ARRAY_TYPE, and gpuDistPBC().

Referenced by gpuLJForces(), gpuLJForcesSmem(), and gpuLJForcesSmemSmem().

__device__ void gpuAddLJForceCut ( FLOAT_ARRAY_TYPE *  dForce,
FLOAT_ARRAY_TYPE *  dA,
FLOAT_ARRAY_TYPE *  dB,
float *  boxsize,
float *  rCutSq 
) [inline]

The Lennard-Jones force $ \vec{F}_{LJ}(\vec{a})_{ab} $ between particle a and b is calculated and added to the f (belonging to particle a): $ \vec{F}_{LJ}(\vec{a})_{ab} = \left\{ \begin{array}{ll} \frac{48\epsilon\vec{r}_{ab}}{r_{ab}^2}\left[\left[\frac{\sigma}{r_{ab}}\right]^{12}-\left[\frac{\sigma}{2r_{ab}}\right]^{6}\right] & \textrm{for } r_{ab} < r_{cut} \\ 0 & \textrm{else} \end{array}\right. $ , where $ r_{ab}=\left| \vec{a} - \vec{b} \right| $ . Due to usage of Lennard-Jones units $ \epsilon = 1 $ and $ \sigma = 1 $ .

Parameters:
[out] dForce Force contribution to particle a due to Lennard-Jones interaction with particle b
[in] dA Position in device memory of particle a
[in] dB Position in device memory of particle b
[in] boxsize Edge length of the cubic simulation box
[in] rCutSq Squared cutoff radius of the Lennard-Jones potential
Returns:
Returns CUDA error information.

Definition at line 40 of file gpu_forces.cu.

References FLOAT_ARRAY_TYPE, and gpuDistPBC().

Referenced by gpuLJForcesCut(), gpuLJForcesCutRSmem(), gpuLJForcesCutSmem(), gpuLJForcesCutSmemSmem(), gpuLJForcesCutVar1(), and gpuLJForcesVerletList().

__global__ void gpuLJForces ( FLOAT_ARRAY_TYPE *  dPos,
FLOAT_ARRAY_TYPE *  dForce,
int  N,
float  boxsize 
)

This kernel calculates the Lennard-Jones force of each particle in a $N^2$-manner: The reference particle of each thread is stored in a register, all interacting particles are fetched from global memory.

Parameters:
[in] dPos Array of all particle positions
[in,out] dForce Array of all forces
[in] N Number of particle
[in] boxsize Edge length of the cubic simulation box.
Returns:
Returns CUDA error information.

Definition at line 68 of file gpu_forces.cu.

References FLOAT_ARRAY_TYPE, gpuAddLJForce(), and gpuCumulateVec().

__global__ void gpuLJForcesCut ( FLOAT_ARRAY_TYPE *  dPos,
FLOAT_ARRAY_TYPE *  dForce,
int  N,
float  boxsize,
float  rCutSq 
)

This kernel calculates the Lennard-Jones force - under respect of a potential cutoff - of each particle in a $N^2$-manner: Here, all coordinates are directly read from global memory.

Parameters:
[in] dPos Array of all particle positions
[in,out] dForce Array of all forces
[in] N Number of particle
[in] boxsize Edge length of the cubic simulation box
[in] rCutSq Squared Lennard-Jones potential cutoff radius
Returns:
Returns CUDA error information.

Definition at line 196 of file gpu_forces.cu.

References FLOAT_ARRAY_TYPE, gpuAddLJForceCut(), and gpuCumulateVec().

__global__ void gpuLJForcesCutRSmem ( FLOAT_ARRAY_TYPE *  dPos,
FLOAT_ARRAY_TYPE *  dForce,
int  N,
float  boxsize,
float  rCutSq 
)

This kernel calculates the Lennard-Jones force - under respect of a potential cutoff - of each particle in a $N^2$-manner: The reference particle of each thread is stored in shared memory, all interacting particles are fetched from global memory.

Attention:
In the kernel launch configuration you have to specify the amount of shared memory used by each thread block: numThreadsPerBlock*sizeof(FLOAT_ARRAY_TYPE)
Parameters:
[in] dPos Array of all particle positions
[in,out] dForce Array of all forces
[in] N Number of particle
[in] boxsize Edge length of the cubic simulation box
[in] rCutSq Squared Lennard-Jones potential cutoff radius
Returns:
Returns CUDA error information.

Definition at line 243 of file gpu_forces.cu.

References FLOAT_ARRAY_TYPE, gpuAddLJForceCut(), and gpuCumulateVec().

__global__ void gpuLJForcesCutSmem ( FLOAT_ARRAY_TYPE *  dPos,
FLOAT_ARRAY_TYPE *  dForce,
int  N,
float  boxsize,
float  rCutSq 
)

This kernel calculates the Lennard-Jones force - under respect of a potential cutoff - of each particle in a $N^2$-manner: The reference particle of each thread is stored in a register, all interacting particles are fetched blockwise and stored in shared memory. Then the whole thread block gets these values via broadcast from shared memory.

Attention:
In the kernel launch configuration you have to specify the amount of shared memory used by each thread block: numThreadsPerBlock*sizeof(FLOAT_ARRAY_TYPE)
Parameters:
[in] dPos Array of all particle positions
[in,out] dForce Array of all forces
[in] N Number of particle
[in] boxsize Edge length of the cubic simulation box
[in] rCutSq Squared Lennard-Jones potential cutoff radius
Returns:
Returns CUDA error information.

Definition at line 289 of file gpu_forces.cu.

References FLOAT_ARRAY_TYPE, gpuAddLJForceCut(), and gpuCumulateVec().

__global__ void gpuLJForcesCutSmemSmem ( FLOAT_ARRAY_TYPE *  dPos,
FLOAT_ARRAY_TYPE *  dForce,
int  N,
float  boxsize,
float  rCutSq 
)

This kernel calculates the Lennard-Jones force - under respect of a potential cutoff - of each particle in a $N^2$-manner: The reference particle of each thread is stored shared memory, all interacting particles are fetched from global memory blockwise and stored into shared memory. Then the whole thread block gets these values via broadcast from shared memory.

Attention:
In the kernel launch configuration you have to specify the amount of shared memory used by each thread block: 2*numThreadsPerBlock*sizeof(FLOAT_ARRAY_TYPE)
Parameters:
[in] dPos Array of all particle positions
[in,out] dForce Array of all forces
[in] N Number of particle
[in] boxsize Edge length of the cubic simulation box
[in] rCutSq Squared Lennard-Jones potential cutoff radius
Returns:
Returns CUDA error information.

Definition at line 323 of file gpu_forces.cu.

References FLOAT_ARRAY_TYPE, gpuAddLJForceCut(), and gpuCumulateVec().

__global__ void gpuLJForcesCutVar1 ( FLOAT_ARRAY_TYPE *  dPos,
FLOAT_ARRAY_TYPE *  dForce,
int  N,
float  boxsize,
float  rCutSq 
)

This kernel calculates the Lennard-Jones force - under respect of a potential cutoff - of each particle in a $N^2$-manner: The reference particle of each thread is stored in a register, all interacting particles are fetched from global memory.

Parameters:
[in] dPos Array of all particle positions
[in,out] dForce Array of all forces
[in] N Number of particle
[in] boxsize Edge length of the cubic simulation box
[in] rCutSq Squared Lennard-Jones potential cutoff radius
Returns:
Returns CUDA error information.

Definition at line 219 of file gpu_forces.cu.

References FLOAT_ARRAY_TYPE, gpuAddLJForceCut(), and gpuCumulateVec().

__global__ void gpuLJForcesSmem ( FLOAT_ARRAY_TYPE *  dPos,
FLOAT_ARRAY_TYPE *  dForce,
int  N,
float  boxsize 
)

This kernel calculates the Lennard-Jones force of each particle in a $N^2$-manner: The reference particle of each thread is stored shared memory, all interacting particles are fetched from global memory, execpt those of the same thread block, who are accessible in shared memory.

Attention:
In the kernel launch configuration you have to specify the amount of shared memory used by each thread block: numThreadsPerBlock*sizeof(FLOAT_ARRAY_TYPE)
Parameters:
[in] dPos Array of all particle positions
[in,out] dForce Array of all forces
[in] N Number of particle
[in] boxsize Edge length of the cubic simulation box.
Returns:
Returns CUDA error information.

Definition at line 91 of file gpu_forces.cu.

References FLOAT_ARRAY_TYPE, gpuAddLJForce(), and gpuCumulateVec().

__global__ void gpuLJForcesSmemSmem ( FLOAT_ARRAY_TYPE *  dPos,
FLOAT_ARRAY_TYPE *  dForce,
int  N,
float  boxsize 
)

This kernel calculates the Lennard-Jones force of each particle in a $N^2$-manner: The reference particle of each thread is stored shared memory, all interacting particles are fetched from global memory blockwise and stored into shared memory. Then the whole thread block gets these values via broadcast from shared memory.

Attention:
In the kernel launch configuration you have to specify the amount of shared memory used by each thread block: 2*numThreadsPerBlock*sizeof(FLOAT_ARRAY_TYPE)
Parameters:
[in] dPos Array of all particle positions
[in,out] dForce Array of all forces
[in] N Number of particle
[in] boxsize Edge length of the cubic simulation box.
Returns:
Returns CUDA error information.

Definition at line 133 of file gpu_forces.cu.

References FLOAT_ARRAY_TYPE, gpuAddLJForce(), and gpuCumulateVec().

__global__ void gpuLJForcesVerletList ( FLOAT_ARRAY_TYPE *  dPos,
FLOAT_ARRAY_TYPE *  dForce,
int  N,
int *  dVlists,
size_t  vlists_pitch,
int *  dVcount,
float  boxsize,
float  rCutSq 
)

The force on each particle is caculated with the help of a Verlet neighbor list generated by gpuGenerateVerletListSmem . This functions stores the reference particle position inside a register, the nieghbor particles are fetched in order of occurence in the Verlet list.

Parameters:
[in] dPos Array of all particle positions
[in,out] dForce Array of all forces
[in] N Number of particle
[in] dVlists Array holding in each column the Verlet List for one particle
[in] vlists_pitch Pitch value for accessing the vlists array via CUDA in the right way.
[in] dVcount This vector counts the number of interaction particles in the Verlet list.
[in] boxsize Edge length of the cubic simulation box
[in] rCutSq Squared Lennard-Jones potential cutoff radius
Returns:
Returns CUDA error information.

Definition at line 400 of file gpu_forces.cu.

References FLOAT_ARRAY_TYPE, gpuAddLJForceCut(), and gpuCumulateVec().

 All Files Functions Defines
Generated on Thu Jun 17 14:22:54 2010 for mdgpu - a molecular dynamic simulation program using GPUs by  doxygen 1.6.1