/home/julian/svn/gpucpumd/mdgpu2/gpu_integrator.cu

Go to the documentation of this file.

#include "gpu_integrator.cuh"

__global__ void gpuIntegratorLeapFrog(FLOAT_ARRAY_TYPE *dPos, FLOAT_ARRAY_TYPE *dVel, FLOAT_ARRAY_TYPE *dForce, int N, float dt, float boxsize, float *dmax, float maxDisplacement, bool *needListUpdate)
{
    int tid = threadIdx.x;
    int idx = blockDim.x * blockIdx.x + tid;
    if (idx < N)
    {
        float dx, dy, dz;
        //update velocities at half time step
        dVel[idx].x = dVel[idx].x + dt * dForce[idx].x;
        dVel[idx].y = dVel[idx].y + dt * dForce[idx].y;
        dVel[idx].z = dVel[idx].z + dt * dForce[idx].z;
        //update positions
        dx = dt * dVel[idx].x;
        dy = dt * dVel[idx].y;
        dz = dt * dVel[idx].z;
        dPos[idx].x = dPos[idx].x + dx;
        dPos[idx].y = dPos[idx].y + dy;
        dPos[idx].z = dPos[idx].z + dz;
        //calculate displacement
        dmax[idx] += sqrt(dx * dx + dy * dy + dz * dz);
        if ( dmax[idx] > maxDisplacement )
        {
            (*needListUpdate) = true; //here no barrier is needed -> if condition is once fulfilled the flag is set
        }
        //force periodic boundary conditions
        checkPBC( &dPos[idx], &boxsize);
    }
}

__global__ void gpuInitialBackstep(FLOAT_ARRAY_TYPE *dVel, FLOAT_ARRAY_TYPE *dForce, int N, float dt)
{
    int tid = threadIdx.x;
    int idx = blockDim.x * blockIdx.x + tid;
    if (idx < N)
    {
        //update velocities at half time step
        dVel[idx].x = dVel[idx].x - 0.5f * dt * dForce[idx].x;
        dVel[idx].y = dVel[idx].y - 0.5f * dt * dForce[idx].y;
        dVel[idx].z = dVel[idx].z - 0.5f * dt * dForce[idx].z;
    }
}

---