diff --git a/common/Membrane.cpp b/common/Membrane.cpp index e21b64c3..e2d51081 100644 --- a/common/Membrane.cpp +++ b/common/Membrane.cpp @@ -768,7 +768,7 @@ void Membrane::RecvD3Q7AA(double *dist){ // Unpack the distributions on the device //................................................................................... //...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)................................ 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)................................ ScaLBL_D3Q7_Membrane_Unpack(5,dvcRecvDist_Z, recvbuf_Z,recvCount_Z,dist,Np,coefficient_Z); //.................................................................................. - + MPI_COMM_SCALBL.barrier(); //................................................................................... Lock=false; // unlock the communicator after communications complete diff --git a/cuda/Ion.cu b/cuda/Ion.cu index 8db86a70..559ef7ec 100644 --- a/cuda/Ion.cu +++ b/cuda/Ion.cu @@ -91,7 +91,7 @@ __global__ void dvc_ScaLBL_D3Q7_Membrane_AssignLinkCoef_halo( // swap rule means that the distributions in recvbuf are OPPOSITE of q // 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 psiLocal, psiNonlocal, membranePotential; double ap,aq; // coefficient diff --git a/hip/Ion.hip b/hip/Ion.hip index 8f72a500..4da89201 100644 --- a/hip/Ion.hip +++ b/hip/Ion.hip @@ -3,7 +3,7 @@ #include "hip/hip_runtime.h" #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, 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 // 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 psiLocal, psiNonlocal, membranePotential; double ap,aq; // coefficient /* 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>>( + dvc_ScaLBL_D3Q7_Membrane_AssignLinkCoef_halo<<>>( Cqx, Cqy, Cqz, Map, Distance, Psi, Threshold, MassFractionIn, MassFractionOut, ThresholdMassFractionIn, ThresholdMassFractionOut, 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, double *dist, int N, double *coef){ - int GRID = count / NTHREADS + 1; - - dvc_ScaLBL_D3Q7_Membrane_Unpack<<>>(q, d3q7_recvlist, recvbuf,count, + dvc_ScaLBL_D3Q7_Membrane_Unpack<<>>(q, d3q7_recvlist, recvbuf,count, dist, N, coef); hipError_t err = hipGetLastError();