2021-03-25 06:49:28 +01:00
|
|
|
/*
|
|
|
|
* =======================================================================================
|
|
|
|
*
|
|
|
|
* Author: Jan Eitzinger (je), jan.eitzinger@fau.de
|
|
|
|
* Copyright (c) 2021 RRZE, University Erlangen-Nuremberg
|
|
|
|
*
|
|
|
|
* This file is part of MD-Bench.
|
|
|
|
*
|
|
|
|
* MD-Bench is free software: you can redistribute it and/or modify it
|
|
|
|
* under the terms of the GNU Lesser General Public License as published
|
|
|
|
* by the Free Software Foundation, either version 3 of the License, or
|
|
|
|
* (at your option) any later version.
|
|
|
|
*
|
|
|
|
* MD-Bench is distributed in the hope that it will be useful, but WITHOUT ANY
|
|
|
|
* WARRANTY; without even the implied warranty of MERCHANTABILITY or FITNESS FOR A
|
|
|
|
* PARTICULAR PURPOSE. See the GNU Lesser General Public License for more
|
|
|
|
* details.
|
|
|
|
*
|
|
|
|
* You should have received a copy of the GNU Lesser General Public License along
|
|
|
|
* with MD-Bench. If not, see <https://www.gnu.org/licenses/>.
|
|
|
|
* =======================================================================================
|
|
|
|
*/
|
2021-11-11 15:23:46 +01:00
|
|
|
#include <math.h>
|
2021-11-09 08:37:37 +01:00
|
|
|
#include <stdio.h>
|
|
|
|
#include <stdlib.h>
|
2021-11-10 16:02:05 +01:00
|
|
|
#include <stddef.h>
|
2021-11-08 20:32:12 +01:00
|
|
|
#include <cuda_runtime.h>
|
2021-11-10 16:02:05 +01:00
|
|
|
#include <device_launch_parameters.h>
|
|
|
|
|
2021-11-11 14:49:29 +01:00
|
|
|
extern "C" {
|
|
|
|
#include <likwid-marker.h>
|
2021-03-25 06:49:28 +01:00
|
|
|
|
2021-11-11 14:49:29 +01:00
|
|
|
#include <timing.h>
|
|
|
|
#include <neighbor.h>
|
|
|
|
#include <parameter.h>
|
|
|
|
#include <atom.h>
|
|
|
|
}
|
2021-03-25 06:49:28 +01:00
|
|
|
|
2021-11-11 20:14:30 +01:00
|
|
|
void checkError(const char *msg, cudaError_t err)
|
|
|
|
{
|
|
|
|
if (err != cudaSuccess)
|
|
|
|
{
|
|
|
|
//print a human readable error message
|
|
|
|
printf("[CUDA ERROR %s]: %s\r\n", msg, cudaGetErrorString(err));
|
|
|
|
exit(-1);
|
|
|
|
}
|
|
|
|
}
|
|
|
|
|
2021-11-10 16:02:05 +01:00
|
|
|
// cuda kernel
|
|
|
|
__global__ void calc_force(
|
2021-11-11 15:23:46 +01:00
|
|
|
Atom a,
|
2021-11-10 16:02:05 +01:00
|
|
|
MD_FLOAT xtmp, MD_FLOAT ytmp, MD_FLOAT ztmp,
|
|
|
|
MD_FLOAT *fix, MD_FLOAT *fiy, MD_FLOAT *fiz,
|
2021-11-11 20:14:30 +01:00
|
|
|
MD_FLOAT cutforcesq, MD_FLOAT sigma6, MD_FLOAT epsilon,
|
2021-11-10 16:02:05 +01:00
|
|
|
int i, int numneighs, int *neighs) {
|
|
|
|
|
|
|
|
// Calculate idx k from thread information
|
|
|
|
const long long k = blockIdx.x * blockDim.x + threadIdx.x;
|
|
|
|
if( k >= numneighs ) {
|
|
|
|
return;
|
|
|
|
}
|
|
|
|
|
2021-11-11 15:23:46 +01:00
|
|
|
Atom *atom = &a;
|
|
|
|
|
2021-11-11 20:14:30 +01:00
|
|
|
const int j = neighs[k];
|
2021-11-10 16:02:05 +01:00
|
|
|
MD_FLOAT delx = xtmp - atom_x(j);
|
|
|
|
MD_FLOAT dely = ytmp - atom_y(j);
|
|
|
|
MD_FLOAT delz = ztmp - atom_z(j);
|
|
|
|
MD_FLOAT rsq = delx * delx + dely * dely + delz * delz;
|
|
|
|
|
2021-11-11 20:14:30 +01:00
|
|
|
#ifdef EXPLICIT_TYPES
|
2021-11-10 16:02:05 +01:00
|
|
|
const int type_i = atom->type[i];
|
|
|
|
const int type_j = atom->type[j];
|
|
|
|
const int type_ij = type_i * atom->ntypes + type_j;
|
|
|
|
const MD_FLOAT cutforcesq = atom->cutforcesq[type_ij];
|
|
|
|
const MD_FLOAT sigma6 = atom->sigma6[type_ij];
|
|
|
|
const MD_FLOAT epsilon = atom->epsilon[type_ij];
|
2021-11-11 20:14:30 +01:00
|
|
|
#endif
|
2021-11-10 16:02:05 +01:00
|
|
|
|
|
|
|
if(rsq < cutforcesq) {
|
|
|
|
MD_FLOAT sr2 = 1.0 / rsq;
|
|
|
|
MD_FLOAT sr6 = sr2 * sr2 * sr2 * sigma6;
|
|
|
|
MD_FLOAT force = 48.0 * sr6 * (sr6 - 0.5) * sr2 * epsilon;
|
2021-11-11 20:14:30 +01:00
|
|
|
fix[k] = delx * force;
|
|
|
|
fiy[k] = dely * force;
|
|
|
|
fiz[k] = delz * force;
|
2021-11-10 16:02:05 +01:00
|
|
|
}
|
|
|
|
}
|
|
|
|
|
2021-11-11 14:49:29 +01:00
|
|
|
extern "C" {
|
|
|
|
|
2021-10-26 09:11:17 +02:00
|
|
|
double computeForce(
|
|
|
|
Parameter *param,
|
|
|
|
Atom *atom,
|
|
|
|
Neighbor *neighbor
|
|
|
|
)
|
|
|
|
{
|
2021-03-25 06:49:28 +01:00
|
|
|
int Nlocal = atom->Nlocal;
|
|
|
|
int* neighs;
|
2021-10-26 09:11:17 +02:00
|
|
|
MD_FLOAT* fx = atom->fx;
|
|
|
|
MD_FLOAT* fy = atom->fy;
|
|
|
|
MD_FLOAT* fz = atom->fz;
|
|
|
|
#ifndef EXPLICIT_TYPES
|
2021-03-25 06:49:28 +01:00
|
|
|
MD_FLOAT cutforcesq = param->cutforce * param->cutforce;
|
2021-11-11 20:14:30 +01:00
|
|
|
MD_FLOAT sigma6 = param->sigma6;
|
|
|
|
MD_FLOAT epsilon = param->epsilon;
|
2021-10-26 09:11:17 +02:00
|
|
|
#endif
|
2021-03-25 06:49:28 +01:00
|
|
|
|
|
|
|
for(int i = 0; i < Nlocal; i++) {
|
|
|
|
fx[i] = 0.0;
|
|
|
|
fy[i] = 0.0;
|
|
|
|
fz[i] = 0.0;
|
|
|
|
}
|
|
|
|
|
2021-11-11 20:14:30 +01:00
|
|
|
Atom c_atom;
|
|
|
|
c_atom.Natoms = atom->Natoms;
|
|
|
|
c_atom.Nlocal = atom->Nlocal;
|
|
|
|
c_atom.Nghost = atom->Nghost;
|
|
|
|
c_atom.Nmax = atom->Nmax;
|
|
|
|
c_atom.ntypes = atom->ntypes;
|
|
|
|
|
|
|
|
size_t available, total;
|
|
|
|
cudaMemGetInfo(&available, &total);
|
|
|
|
printf("Available memory: %ldGB\r\n", available / 1024 / 1024 / 1024);
|
|
|
|
cudaDeviceSetLimit(cudaLimitMallocHeapSize, available);
|
|
|
|
|
|
|
|
// HINT: Run with cuda-memcheck ./MDBench-NVCC in case of error
|
|
|
|
// HINT: Only works for data layout = AOS!!!
|
|
|
|
|
|
|
|
checkError( "Malloc1", cudaMalloc((void**)&(c_atom.x), sizeof(MD_FLOAT) * atom->Nmax * 3) );
|
|
|
|
checkError( "Memcpy1", cudaMemcpy((void*)(c_atom.x), atom->x, sizeof(MD_FLOAT) * atom->Nmax * 3, cudaMemcpyHostToDevice) );
|
|
|
|
|
|
|
|
checkError( "Malloc4", cudaMalloc((void**)&(c_atom.type), sizeof(int) * atom->Nmax) );
|
|
|
|
checkError( "Memcpy4", cudaMemcpy(c_atom.type, atom->type, sizeof(int) * atom->Nmax, cudaMemcpyHostToDevice) );
|
|
|
|
|
|
|
|
checkError( "Malloc5", cudaMalloc((void**)&(c_atom.epsilon), sizeof(MD_FLOAT) * atom->ntypes * atom->ntypes) );
|
|
|
|
checkError( "Memcpy5", cudaMemcpy(c_atom.epsilon, atom->epsilon, sizeof(MD_FLOAT) * atom->ntypes * atom->ntypes, cudaMemcpyHostToDevice) );
|
|
|
|
|
|
|
|
checkError( "Malloc6", cudaMalloc((void**)&(c_atom.sigma6), sizeof(MD_FLOAT) * atom->ntypes * atom->ntypes) );
|
|
|
|
checkError( "Memcpy6", cudaMemcpy(c_atom.sigma6, atom->sigma6, sizeof(MD_FLOAT) * atom->ntypes * atom->ntypes, cudaMemcpyHostToDevice) );
|
|
|
|
|
|
|
|
checkError( "Malloc7", cudaMalloc((void**)&(c_atom.cutforcesq), sizeof(MD_FLOAT) * atom->ntypes * atom->ntypes) );
|
|
|
|
checkError( "Memcpy7", cudaMemcpy(c_atom.cutforcesq, atom->cutforcesq, sizeof(MD_FLOAT) * atom->ntypes * atom->ntypes, cudaMemcpyHostToDevice) );
|
|
|
|
|
2021-11-11 14:49:29 +01:00
|
|
|
double S = getTimeStamp();
|
2021-06-30 13:44:02 +02:00
|
|
|
LIKWID_MARKER_START("force");
|
2021-10-12 22:39:54 +02:00
|
|
|
|
2021-11-11 20:14:30 +01:00
|
|
|
// #pragma omp parallel for
|
2021-10-26 09:11:17 +02:00
|
|
|
for(int i = 0; i < Nlocal; i++) {
|
|
|
|
neighs = &neighbor->neighbors[i * neighbor->maxneighs];
|
|
|
|
int numneighs = neighbor->numneigh[i];
|
|
|
|
MD_FLOAT xtmp = atom_x(i);
|
|
|
|
MD_FLOAT ytmp = atom_y(i);
|
|
|
|
MD_FLOAT ztmp = atom_z(i);
|
|
|
|
|
|
|
|
#ifdef EXPLICIT_TYPES
|
|
|
|
const int type_i = atom->type[i];
|
|
|
|
#endif
|
2021-10-12 22:39:54 +02:00
|
|
|
|
2021-11-09 16:40:25 +01:00
|
|
|
int *c_neighs;
|
2021-11-11 20:29:14 +01:00
|
|
|
checkError( "c_neighs malloc", cudaMalloc((void**)&c_neighs, sizeof(int) * numneighs) );
|
|
|
|
checkError( "c_neighs memcpy", cudaMemcpy(c_neighs, neighs, sizeof(int) * numneighs, cudaMemcpyHostToDevice) );
|
2021-11-09 08:37:37 +01:00
|
|
|
|
2021-11-09 16:40:25 +01:00
|
|
|
MD_FLOAT *c_fix, *c_fiy, *c_fiz;
|
2021-11-11 20:29:14 +01:00
|
|
|
checkError( "c_fix malloc", cudaMalloc((void**)&c_fix, sizeof(MD_FLOAT) * numneighs) );
|
|
|
|
checkError( "c_fiy malloc", cudaMalloc((void**)&c_fiy, sizeof(MD_FLOAT) * numneighs) );
|
|
|
|
checkError( "c_fiz malloc", cudaMalloc((void**)&c_fiz, sizeof(MD_FLOAT) * numneighs) );
|
2021-11-10 16:02:05 +01:00
|
|
|
|
|
|
|
const int num_blocks = 64;
|
2021-11-11 15:23:46 +01:00
|
|
|
const int num_threads_per_block = ceil((float)numneighs / (float)num_blocks);
|
|
|
|
// printf("numneighs: %d => num-blocks: %d, num_threads_per_block => %d\r\n", numneighs, num_blocks, num_threads_per_block);
|
2021-11-10 16:02:05 +01:00
|
|
|
|
|
|
|
// launch cuda kernel
|
2021-11-11 20:14:30 +01:00
|
|
|
calc_force <<< num_blocks, num_threads_per_block >>> (c_atom, xtmp, ytmp, ztmp, c_fix, c_fiy, c_fiz, cutforcesq, sigma6, epsilon, i, numneighs, c_neighs);
|
|
|
|
checkError( "PeekAtLastError", cudaPeekAtLastError() );
|
|
|
|
checkError( "DeviceSync", cudaDeviceSynchronize() );
|
|
|
|
|
2021-11-10 16:02:05 +01:00
|
|
|
// sum result
|
2021-11-11 20:14:30 +01:00
|
|
|
MD_FLOAT *d_fix = (MD_FLOAT*)malloc(sizeof(MD_FLOAT) * numneighs);
|
|
|
|
MD_FLOAT *d_fiy = (MD_FLOAT*)malloc(sizeof(MD_FLOAT) * numneighs);
|
|
|
|
MD_FLOAT *d_fiz = (MD_FLOAT*)malloc(sizeof(MD_FLOAT) * numneighs);
|
2021-11-11 20:29:14 +01:00
|
|
|
checkError( "d_fix copy to host", cudaMemcpy(d_fix, c_fix, sizeof(MD_FLOAT) * numneighs, cudaMemcpyDeviceToHost) );
|
|
|
|
checkError( "d_fiy copy to host", cudaMemcpy(d_fiy, c_fiy, sizeof(MD_FLOAT) * numneighs, cudaMemcpyDeviceToHost) );
|
|
|
|
checkError( "d_fiz copy to host", cudaMemcpy(d_fiz, c_fiz, sizeof(MD_FLOAT) * numneighs, cudaMemcpyDeviceToHost) );
|
2021-11-11 20:14:30 +01:00
|
|
|
|
2021-10-26 09:11:17 +02:00
|
|
|
for(int k = 0; k < numneighs; k++) {
|
2021-11-10 16:02:05 +01:00
|
|
|
fx[i] += d_fix[k];
|
|
|
|
fy[i] += d_fiy[k];
|
|
|
|
fz[i] += d_fiz[k];
|
2021-04-07 00:46:51 +02:00
|
|
|
}
|
2021-10-26 09:11:17 +02:00
|
|
|
|
2021-11-11 20:29:14 +01:00
|
|
|
checkError( "cudaFree c_fix", cudaFree(c_fix) );
|
|
|
|
checkError( "cudaFree c_fiy", cudaFree(c_fiy) );
|
|
|
|
checkError( "cudaFree c_fiz", cudaFree(c_fiz) );
|
|
|
|
checkError( "cudaFree c_neighs", cudaFree(c_neighs) );
|
|
|
|
free(d_fix);
|
|
|
|
free(d_fiy);
|
|
|
|
free(d_fiz);
|
2021-03-25 06:49:28 +01:00
|
|
|
}
|
2021-10-12 22:39:54 +02:00
|
|
|
|
2021-11-11 20:29:14 +01:00
|
|
|
cudaFree(c_atom.x);
|
|
|
|
cudaFree(c_atom.type);
|
|
|
|
cudaFree(c_atom.epsilon);
|
|
|
|
cudaFree(c_atom.sigma6);
|
|
|
|
cudaFree(c_atom.cutforcesq);
|
|
|
|
|
2021-06-30 13:44:02 +02:00
|
|
|
LIKWID_MARKER_STOP("force");
|
2021-11-11 14:49:29 +01:00
|
|
|
double E = getTimeStamp();
|
2021-10-11 16:57:02 +02:00
|
|
|
|
2021-11-11 14:49:29 +01:00
|
|
|
return E-S;
|
2021-03-25 06:49:28 +01:00
|
|
|
}
|
2021-11-11 14:49:29 +01:00
|
|
|
}
|