From 3f7fb7f22a9293461f252881342a3264c0b7dfad Mon Sep 17 00:00:00 2001 From: Maximilian Gaul Date: Tue, 9 Nov 2021 16:40:25 +0100 Subject: [PATCH] cudaMemcpy of Atom and other properties, first draft implementation of CUDA kernel --- src/force.c | 89 +++++++++++++++++++++++++++++++++-------------------- 1 file changed, 56 insertions(+), 33 deletions(-) diff --git a/src/force.c b/src/force.c index 5534100..320ddbb 100644 --- a/src/force.c +++ b/src/force.c @@ -78,51 +78,44 @@ double computeForce( atom->sigma6 = allocate(ALIGNMENT, atom->ntypes * atom->ntypes * sizeof(MD_FLOAT)); atom->cutforcesq = allocate(ALIGNMENT, atom->ntypes * atom->ntypes * sizeof(MD_FLOAT)); atom->cutneighsq = allocate(ALIGNMENT, atom->ntypes * atom->ntypes * sizeof(MD_FLOAT)); + atom->type = (int *) reallocate(atom->type, ALIGNMENT, atom->Nmax * sizeof(int), nold * sizeof(int)); */ - MD_FLOAT *c_xtmp; - cudaMalloc((void**)&c_xtmp, sizeof(MD_FLOAT)); - cudaMemcpy(c_xtmp, &xtmp, sizeof(MD_FLOAT), cudaMemcpyHostToDevice); + Atom *c_atom; + cudaMalloc((void**)&c_atom, sizeof(Atom)); + cudaMemcpy(c_atom, atom, sizeof(Atom), cudaMemcpyHostToDevice); - MD_FLOAT *c_ytmp; - cudaMalloc((void**)&c_ytmp, sizeof(MD_FLOAT)); - cudaMemcpy(c_ytmp, &ytmp, sizeof(MD_FLOAT), cudaMemcpyHostToDevice); + cudaMalloc((void**)&c_atom->x, sizeof(MD_FLOAT) * atom->Nmax * 3); + cudaMemcpy(c_atom->x, atom->x, sizeof(MD_FLOAT) * atom->Nmax * 3, cudaMemcpyHostToDevice); - MD_FLOAT *c_ztmp; - cudaMalloc((void**)&c_ztmp, sizeof(MD_FLOAT)); - cudaMemcpy(c_ztmp, &ztmp, sizeof(MD_FLOAT), cudaMemcpyHostToDevice); + cudaMalloc((void**)&c_atom->y, sizeof(MD_FLOAT) * atom->Nmax * 3); + cudaMemcpy(c_atom->y, atom->y, sizeof(MD_FLOAT) * atom->Nmax * 3, cudaMemcpyHostToDevice); - int *c_atom_ntypes; - cudaMalloc((void**)&c_atom_ntypes, sizeof(int)); - cudaMemcpy(c_atom_ntypes, &(atom->ntypes), sizeof(int), cudaMemcpyHostToDevice); + cudaMalloc((void**)&c_atom->z, sizeof(MD_FLOAT) * atom->Nmax * 3); + cudaMemcpy(c_atom->z, atom->z, sizeof(MD_FLOAT) * atom->Nmax * 3, cudaMemcpyHostToDevice); - int *c_neighbors; - cudaMalloc((void**)&c_neighbors, sizeof(int) * numneighs); - cudaMemcpy(c_neighbors, neighs, sizeof(int) * numneighs, cudaMemcpyHostToDevice); + cudaMalloc((void**)&c_atom->type, sizeof(int) * atom->Nmax); + cudaMemcpy(c_atom->type, atom->type, sizeof(int) * atom->Nmax, cudaMemcpyHostToDevice); - MD_FLOAT *c_atom_x; - cudaMalloc((void**)&c_atom_x, sizeof(MD_FLOAT) * atom->Nmax * 3); - cudaMemcpy(c_atom_x, atom->x, sizeof(MD_FLOAT) * atom->Nmax * 3, cudaMemcpyHostToDevice); + cudaMalloc((void**)&c_atom->epsilon, sizeof(MD_FLOAT) * atom->ntypes * atom->ntypes); + cudaMemcpy(c_atom->epsilon, atom->epsilon, sizeof(MD_FLOAT) * atom->ntypes * atom->ntypes, cudaMemcpyHostToDevice); - MD_FLOAT *c_atom_y; - cudaMalloc((void**)&c_atom_y, sizeof(MD_FLOAT) * atom->Nmax * 3); - cudaMemcpy(c_atom_y, atom->y, sizeof(MD_FLOAT) * atom->Nmax * 3, cudaMemcpyHostToDevice); + cudaMalloc((void**)&c_atom->sigma6, sizeof(MD_FLOAT) * atom->ntypes * atom->ntypes); + cudaMemcpy(c_atom->sigma6, atom->sigma6, sizeof(MD_FLOAT) * atom->ntypes * atom->ntypes, cudaMemcpyHostToDevice); - MD_FLOAT *c_atom_z; - cudaMalloc((void**)&c_atom_z, sizeof(MD_FLOAT) * atom->Nmax * 3); - cudaMemcpy(c_atom_z, atom->z, sizeof(MD_FLOAT) * atom->Nmax * 3, cudaMemcpyHostToDevice); + cudaMalloc((void**)&c_atom->cutforcesq, sizeof(MD_FLOAT) * atom->ntypes * atom->ntypes); + cudaMemcpy(c_atom->cutforcesq, atom->cutforcesq, sizeof(MD_FLOAT) * atom->ntypes * atom->ntypes, cudaMemcpyHostToDevice); - MD_FLOAT *c_atom_epsilon; - cudaMalloc((void**)&c_atom_epsilon, sizeof(MD_FLOAT) * atom->ntypes * atom->ntypes); - cudaMemcpy(c_atom_epsilon, atom->epsilon, sizeof(MD_FLOAT) * atom->ntypes * atom->ntypes, cudaMemcpyHostToDevice); + int *c_neighs; + cudaMalloc((void**)&c_neighs, sizeof(int) * numneighs); + cudaMemcpy(c_neighs, neighs, sizeof(int) * numneighs, cudaMemcpyHostToDevice); - MD_FLOAT *c_sigma6; - cudaMalloc((void**)&c_sigma6, sizeof(MD_FLOAT) * atom->ntypes * atom->ntypes); - cudaMemcpy(c_sigma6, atom->epsilon, sizeof(MD_FLOAT) * atom->ntypes * atom->ntypes, cudaMemcpyHostToDevice); + const int num_elems = numneighs; - MD_FLOAT *c_cutforcesq; - cudaMalloc((void**)&c_cutforcesq, sizeof(MD_FLOAT) * atom->ntypes * atom->ntypes); - cudaMemcpy(c_cutforcesq, atom->epsilon, sizeof(MD_FLOAT) * atom->ntypes * atom->ntypes, cudaMemcpyHostToDevice); + MD_FLOAT *c_fix, *c_fiy, *c_fiz; + cudaMalloc((void**)&c_fix, sizeof(MD_FLOAT) * num_elems); + cudaMalloc((void**)&c_fiy, sizeof(MD_FLOAT) * num_elems); + cudaMalloc((void**)&c_fiz, sizeof(MD_FLOAT) * num_elems); for(int k = 0; k < numneighs; k++) { int j = neighs[k]; @@ -159,3 +152,33 @@ double computeForce( return E-S; } + +// cuda kernel +__global__ void calc_force( + Atom *atom, + MD_FLOAT xtmp, MD_FLOAT ytmp, MD_FLOAT ztmp, + MD_FLOAT *fix, MD_FLOAT *fiy, MD_FLOAT *fiz, + int i, int k, int *neighs) { + + int j = neighs[k]; + 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; + + 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]; + + 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; + fix[j] += delx * force; + fiy[j] += dely * force; + fiz[j] += delz * force; + } +}