OilfieldToolsNet/LBPM
4
0
Fork 0
Code Issues Pull Requests Packages Projects Releases Wiki Activity

multiple nvidia gpu working with membrane

Browse Source
This commit is contained in:
James E McClure 2022-05-12 20:50:05 -04:00
parent cf3bc417ce
commit b6227dd823
3 changed files with 10 additions and 14 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

4
common/Membrane.cpp
View File

@ -768,7 +768,7 @@ void Membrane::RecvD3Q7AA(double *dist){
// Unpack the distributions on the device // Unpack the distributions on the device
//................................................................................... //...................................................................................
//...Unpacking for x face(q=2)................................ //...Unpacking for x face(q=2)................................
ScaLBL_D3Q7_Membrane_Unpack(2,dvcRecvDist_x, recvbuf_x,recvCount_x,dist,Np,coefficient_x); ScaLBL_D3Q7_Membrane_Unpack(2,dvcRecvDist_x, recvbuf_x,recvCount_x,dist,Np,coefficient_x);
//................................................................................... //...................................................................................
//...Packing for X face(q=1)................................ //...Packing for X face(q=1)................................
ScaLBL_D3Q7_Membrane_Unpack(1,dvcRecvDist_X, recvbuf_X,recvCount_X,dist,Np,coefficient_X); ScaLBL_D3Q7_Membrane_Unpack(1,dvcRecvDist_X, recvbuf_X,recvCount_X,dist,Np,coefficient_X);
@ -784,7 +784,7 @@ void Membrane::RecvD3Q7AA(double *dist){
//...Packing for Z face(q=5)................................ //...Packing for Z face(q=5)................................
ScaLBL_D3Q7_Membrane_Unpack(5,dvcRecvDist_Z, recvbuf_Z,recvCount_Z,dist,Np,coefficient_Z); ScaLBL_D3Q7_Membrane_Unpack(5,dvcRecvDist_Z, recvbuf_Z,recvCount_Z,dist,Np,coefficient_Z);
//.................................................................................. //..................................................................................
MPI_COMM_SCALBL.barrier(); MPI_COMM_SCALBL.barrier();
//................................................................................... //...................................................................................
Lock=false; // unlock the communicator after communications complete Lock=false; // unlock the communicator after communications complete

2
cuda/Ion.cu
View File

@ -91,7 +91,7 @@ __global__ void dvc_ScaLBL_D3Q7_Membrane_AssignLinkCoef_halo(
// swap rule means that the distributions in recvbuf are OPPOSITE of q // swap rule means that the distributions in recvbuf are OPPOSITE of q
// dist may be even or odd distributions stored by stream layout // dist may be even or odd distributions stored by stream layout
//.................................................................................... //....................................................................................
int n, idx, nqm, npm, i, j, k; int n, idx, nqm, npm, label, i, j, k;
double distanceLocal, distanceNonlocal; double distanceLocal, distanceNonlocal;
double psiLocal, psiNonlocal, membranePotential; double psiLocal, psiNonlocal, membranePotential;
double ap,aq; // coefficient double ap,aq; // coefficient

18
hip/Ion.hip
View File

@ -3,7 +3,7 @@
#include "hip/hip_runtime.h" #include "hip/hip_runtime.h"
#define NBLOCKS 1024 #define NBLOCKS 1024
#define NTHREADS 256 #define NTHREADS 512
extern "C" void Membrane_D3Q19_Unpack(int q, int *list, int *links, int start, int linkCount, extern "C" void Membrane_D3Q19_Unpack(int q, int *list, int *links, int start, int linkCount,
double *recvbuf, double *dist, int N) { double *recvbuf, double *dist, int N) {
@ -91,13 +91,13 @@ __global__ void dvc_ScaLBL_D3Q7_Membrane_AssignLinkCoef_halo(
// swap rule means that the distributions in recvbuf are OPPOSITE of q // swap rule means that the distributions in recvbuf are OPPOSITE of q
// dist may be even or odd distributions stored by stream layout // dist may be even or odd distributions stored by stream layout
//.................................................................................... //....................................................................................
int n, idx, nqm, npm, i, j, k; int n, idx, nqm, npm, label, i, j, k;
double distanceLocal, distanceNonlocal; double distanceLocal, distanceNonlocal;
double psiLocal, psiNonlocal, membranePotential; double psiLocal, psiNonlocal, membranePotential;
double ap,aq; // coefficient double ap,aq; // coefficient
/* second enforce custom rule for membrane links */ /* second enforce custom rule for membrane links */
int S = (count-nlinks)/NBLOCKS/NTHREADS + 1; int S = count/NBLOCKS/NTHREADS + 1;
for (int s=0; s<S; s++){ for (int s=0; s<S; s++){
//........Get 1-D index for this thread.................... //........Get 1-D index for this thread....................
idx = S*blockIdx.x*blockDim.x + s*blockDim.x + threadIdx.x; idx = S*blockIdx.x*blockDim.x + s*blockDim.x + threadIdx.x;
@ -174,8 +174,8 @@ __global__ void dvc_ScaLBL_D3Q7_Membrane_Unpack(int q,
fp = recvbuf[idx]; fp = recvbuf[idx];
fqq = (1-aq)*fq+ap*fp; fqq = (1-aq)*fq+ap*fp;
dist[q * N + n] = fqq; dist[q * N + n] = fqq;
} }
} }
} }
} }
__global__ void dvc_ScaLBL_D3Q7_Membrane_IonTransport(int *membrane, double *coef, __global__ void dvc_ScaLBL_D3Q7_Membrane_IonTransport(int *membrane, double *coef,
@ -932,9 +932,7 @@ extern "C" void ScaLBL_D3Q7_Membrane_AssignLinkCoef_halo(
int *d3q7_recvlist, int *d3q7_linkList, double *coef, int start, int nlinks, int count, int *d3q7_recvlist, int *d3q7_linkList, double *coef, int start, int nlinks, int count,
const int N, const int Nx, const int Ny, const int Nz) { const int N, const int Nx, const int Ny, const int Nz) {
int GRID = count / 1024 + 1; dvc_ScaLBL_D3Q7_Membrane_AssignLinkCoef_halo<<<NBLOCKS,NTHREADS>>>(
dvc_ScaLBL_D3Q7_Membrane_AssignLinkCoef_halo<<<GRID,1024 >>>(
Cqx, Cqy, Cqz, Map, Distance, Psi, Threshold, Cqx, Cqy, Cqz, Map, Distance, Psi, Threshold,
MassFractionIn, MassFractionOut, ThresholdMassFractionIn, ThresholdMassFractionOut, MassFractionIn, MassFractionOut, ThresholdMassFractionIn, ThresholdMassFractionOut,
d3q7_recvlist, d3q7_linkList, coef, start, nlinks, count, N, Nx, Ny, Nz); d3q7_recvlist, d3q7_linkList, coef, start, nlinks, count, N, Nx, Ny, Nz);
@ -950,9 +948,7 @@ extern "C" void ScaLBL_D3Q7_Membrane_Unpack(int q,
int *d3q7_recvlist, double *recvbuf, int count, int *d3q7_recvlist, double *recvbuf, int count,
double *dist, int N, double *coef){ double *dist, int N, double *coef){
int GRID = count / NTHREADS + 1; dvc_ScaLBL_D3Q7_Membrane_Unpack<<<NBLOCKS,NTHREADS >>>(q, d3q7_recvlist, recvbuf,count,
dvc_ScaLBL_D3Q7_Membrane_Unpack<<<GRID,NTHREADS >>>(q, d3q7_recvlist, recvbuf,count,
dist, N, coef); dist, N, coef);
hipError_t err = hipGetLastError(); hipError_t err = hipGetLastError();