/home/julian/svn/gpucpumd/mdgpu2/gpu_tools.cuh

Go to the documentation of this file.

#ifndef GPU_TOOLS_CUH
#define GPU_TOOLS_CUH

#include <cuda.h>
#include "define.h"
#include <iostream>
#include <fstream>
#include <iomanip>
#include <string>
#include <sstream>
#include <limits>

//defining groups used for documentation
//####################################### Inline functions ####################################//




__device__ inline int gpuRnd(float b)
{
    return b<0.0f ? static_cast<int>(b-0.5f) : static_cast<int>(b+0.5f);
}

__device__ inline void gpuCumulateVec(FLOAT_ARRAY_TYPE *dVec, FLOAT_ARRAY_TYPE *dVecToAdd)
{
    (*dVec).x += (*dVecToAdd).x;
    (*dVec).y += (*dVecToAdd).y;
    (*dVec).z += (*dVecToAdd).z;
} // end of helptool

__device__ void gpuDistPBC(FLOAT_ARRAY_TYPE *dDist, FLOAT_ARRAY_TYPE *dA, FLOAT_ARRAY_TYPE *dB, float *boxSize)
{
    (*dDist).x = (*dA).x - (*dB).x;
    (*dDist).y = (*dA).y - (*dB).y;
    (*dDist).z = (*dA).z - (*dB).z;
    (*dDist).x = ( (*dDist).x - gpuRnd( (*dDist).x / (*boxSize) ) * (*boxSize) );//regard nearest image convention
    (*dDist).y = ( (*dDist).y - gpuRnd( (*dDist).y / (*boxSize) ) * (*boxSize) );
    (*dDist).z = ( (*dDist).z - gpuRnd( (*dDist).z / (*boxSize) ) * (*boxSize) );
}





__host__ __device__ inline void checkPBC(FLOAT_ARRAY_TYPE *pos, float *boxSize)
{
    (*pos).x = (*pos).x - int( (*pos).x / (*boxSize) ) * (*boxSize);
    if( (*pos).x<0.0f) (*pos).x = (*pos).x + (*boxSize);

    (*pos).y = (*pos).y - int( (*pos).y / (*boxSize) ) * (*boxSize);
    if( (*pos).y<0.0f) (*pos).y = (*pos).y + (*boxSize);

    (*pos).z = (*pos).z - int( (*pos).z / (*boxSize) ) * (*boxSize);
    if( (*pos).z<0.0f) (*pos).z = (*pos).z + (*boxSize);
} // end of pbctools




//####################################### Verlet List functions ####################################//

__global__ void gpuGenerateVerletList(FLOAT_ARRAY_TYPE *dPos, int *dVlists, size_t pitchDVlists, int *dVcount, float *dmax, float rListSq, int N, float boxSize, bool *doUpdate);


__global__ void gpuGenerateVerletListVar1(FLOAT_ARRAY_TYPE *dPos, int *dVlists, size_t pitchDVlists, int *dVcount, float *dmax, float rListSq, int N, float boxSize, bool *doUpdate);


__global__ void gpuGenerateVerletListSmem(FLOAT_ARRAY_TYPE *dPos, int *dVlists, size_t pitchDVlists, int *dVcount, float *dmax, float rListSq, int N, float boxSize, bool *doUpdate);


__global__ void gpuGenerateVerletListSmemVar1(FLOAT_ARRAY_TYPE *dPos, int *dVlists, size_t pitchDVlists, int *dVcount, float *dmax, float rListSq, int N, float boxSize, bool *doUpdate); // end of verletlistgroup
//####################################### NVT ensemble methods ####################################//
void rescaleT(int N, int dim, int numInterData, int numThreads, float T, int maxThreads, int maxBlocks, FLOAT_ARRAY_TYPE *dVel, float *sumSq);


__global__ void gpuRescaleTKernel(int N, int dim, float T, FLOAT_ARRAY_TYPE *dVel, float *vSumSq); // end of constTtools





//####################################### general MD helper functions ####################################//

void removeDrift(int N, int numInterData, int numThreads, int maxThreads, int maxBlocks, FLOAT_ARRAY_TYPE *dVel, FLOAT_ARRAY_TYPE *compSum);


void __global__ gpuRemoveDriftKernel(int N, FLOAT_ARRAY_TYPE *dVel, FLOAT_ARRAY_TYPE *compSum);
__global__ void gpuGetVelOnTime(FLOAT_ARRAY_TYPE *dVel, FLOAT_ARRAY_TYPE *dVt, FLOAT_ARRAY_TYPE *dForce, int N, float dt);
//####################################### Observable functions ####################################//



float getKineticEnergy(FLOAT_ARRAY_TYPE *dVel, int &N, float *sumSq, int &maxThreads, int &maxBlocks);


float getHarmonicPotentialEnergy(float *dEpot, FLOAT_ARRAY_TYPE *dPos, float *sum, int *cnn, size_t &pitch, int &N, int &numNN, float &k, float &boxSize, dim3 &dimGrid, dim3 &dimBlock, int &maxThreads, int &maxBlocks);


__global__ void gpuGetHarmonicPotentialEnergyKernel(float *dEpot, FLOAT_ARRAY_TYPE *dPos, int *cnn, size_t pitch, int N, int numNN, float k, float boxSize);


float getLJCutPotentialEnergy(FLOAT_ARRAY_TYPE *dPos, float *dEpot, float *sum, int &N, float &boxSize, float &rCutSq, dim3 &dimGrid, dim3 &dimBlock, int &maxThreads, int &maxBlocks);


__global__ void gpuLJCutPotentialEnergyKernel(FLOAT_ARRAY_TYPE *dPos, float *dEpot, int N, float boxSize, float rCutSq);
//####################################### CUDA helper functions ####################################//

inline static unsigned long inKB(unsigned long bytes);


inline static unsigned long inMB(unsigned long bytes);



__global__ void copyDevToDev(FLOAT_ARRAY_TYPE *dA, FLOAT_ARRAY_TYPE *dB, int N);


__global__ void gpuSimpleSqSum(FLOAT_ARRAY_TYPE *dVal, float*sqSum, int N);

__global__ void gpuSimpleComponentSum(FLOAT_ARRAY_TYPE *dVal, FLOAT_ARRAY_TYPE *sum, int N);
//####################################### Screen output functions ####################################//

bool memInfo();


void checkCudaError(const char *identifier);

void showPositions(FLOAT_ARRAY_TYPE *hPos, int N);


void showVelocities(FLOAT_ARRAY_TYPE *hVel, int &N);



void showForces(FLOAT_ARRAY_TYPE *hForce, int &N);



void showDeviceArray(FLOAT_ARRAY_TYPE *dArray, int N);


void showNeareastNeighborList(int *dNNListArray, int &N, int &numNN);


void showVerletList(int *dVlists, size_t pitchVlists, int *dVcount, int N, int NList);
//####################################### Simulation parameter functions ####################################//
bool setCudaDevice(int gpuDevNo);

bool generateStartScript(std::string exeName);


bool generateRestartScript(std::string exeName, std::string basename, int dim, int N, int numNN, float T, float k, float boxSize, float dt, int NOffset, int NSim, int NOut, int NAdjust, int gpuDevNo);


bool generateScript(std::string scriptname, std::string exeName, std::string inputfile, std::string basename, int dim, int N, int numNN, float T, float k, float boxSize, float dt, int NOffset, int NSim, int NOut, int NAdjust, int gpuDevNo);
//####################################### Parallel primitives ####################################//



__global__ void gpuReduce1(float *dInData, float *dInterData, int numInterData, int numThreads, int N);


__global__ void gpuReduce2(float *dInterData, int numInterData, int numThreads);


__global__ void gpuSqSumReduce1(FLOAT_ARRAY_TYPE *dInData, float *dInterData, int numInterData, int numThreads, int N);


__global__ void gpuComponentSumReduce1(FLOAT_ARRAY_TYPE *dInData, FLOAT_ARRAY_TYPE *dInterData, int numInterData, int numThreads, int N);


__global__ void gpuComponentSumReduce2(FLOAT_ARRAY_TYPE *dInterData, int numInterData, int numThreads);


void sumReduction(float *dInData, float *dInterData, int numInterData, int numThreads, int N);


void SqVecSumReduction(FLOAT_ARRAY_TYPE *dInData, float *dInterData, int numInterData, int numThreads, int N);


void componentSumReduction(FLOAT_ARRAY_TYPE *dInData, FLOAT_ARRAY_TYPE *dInterData, int numInterData, int numThreads, int N);
#endif //GPU_TOOLS_CUH

---