diff --git a/hip/Ion.hip b/hip/Ion.hip index d7e863bb..18be1aee 100644 --- a/hip/Ion.hip +++ b/hip/Ion.hip @@ -169,7 +169,7 @@ __global__ void dvc_ScaLBL_D3Q7_Membrane_Unpack(int q, /* First unpack the regular links */ if (link < nlinks) { // get the index for the recv list (deal with reordering of links) - idx = d3q7_linkList[link]; + idx = d3q7_linkList[start+link]; // get the distribution index n = d3q7_recvlist[start+idx]; if (!(n < 0)){ @@ -177,22 +177,22 @@ __global__ void dvc_ScaLBL_D3Q7_Membrane_Unpack(int q, dist[q * N + n] = fp; } } - else if (link < count){ - /* second enforce custom rule for membrane links */ - // get the index for the recv list (deal with reordering of links) - idx = d3q7_linkList[link]; - // get the distribution index - n = d3q7_recvlist[start+idx]; - // update link based on mass transfer coefficients - if (!(n < 0)){ - aq = coef[2*(link-nlinks)]; - ap = coef[2*(link-nlinks)+1]; - fq = dist[q * N + n]; - fp = recvbuf[start + idx]; - fqq = (1-aq)*fq+ap*fp; - dist[q * N + n] = fqq; - } - } +// else if (link < count){ +// /* second enforce custom rule for membrane links */ +// // get the index for the recv list (deal with reordering of links) +// idx = d3q7_linkList[link]; +// // get the distribution index +// n = d3q7_recvlist[start+idx]; +// // update link based on mass transfer coefficients +// if (!(n < 0)){ +// aq = coef[2*(link-nlinks)]; +// ap = coef[2*(link-nlinks)+1]; +// fq = dist[q * N + n]; +// fp = recvbuf[start + idx]; +// fqq = (1-aq)*fq+ap*fp; +// dist[q * N + n] = fqq; +// } +// } } } @@ -707,7 +707,9 @@ extern "C" void ScaLBL_D3Q7_Membrane_Unpack(int q, int *d3q7_recvlist, int *d3q7_linkList, int start, int nlinks, int count, double *recvbuf, double *dist, int N, double *coef) { - dvc_ScaLBL_D3Q7_Membrane_Unpack<<>>(q, d3q7_recvlist, d3q7_linkList, start, nlinks, count, + int GRID = count / 1024 + 1; + + dvc_ScaLBL_D3Q7_Membrane_Unpack<<>>(q, d3q7_recvlist, d3q7_linkList, start, nlinks, count, recvbuf, dist, N, coef) ; hipError_t err = hipGetLastError(); diff --git a/hip/Poisson.hip b/hip/Poisson.hip index e07c4d07..2808ffa1 100644 --- a/hip/Poisson.hip +++ b/hip/Poisson.hip @@ -124,7 +124,8 @@ __global__ void dvc_ScaLBL_D3Q7_AAodd_Poisson(int *neighborList, int *Map, doub //Load data //When Helmholtz-Smoluchowski slipping velocity BC is used, the bulk fluid is considered as electroneutral //and thus the net space charge density is zero. - rho_e = (UseSlippingVelBC==1) ? 0.0 : Den_charge[n] / epsilon_LB; + //rho_e = (UseSlippingVelBC==1) ? 0.0 : Den_charge[n] / epsilon_LB; + rho_e = Den_charge[n] / epsilon_LB; idx=Map[n]; psi = Psi[idx]; @@ -204,8 +205,10 @@ __global__ void dvc_ScaLBL_D3Q7_AAeven_Poisson(int *Map, double *dist, double * //Load data //When Helmholtz-Smoluchowski slipping velocity BC is used, the bulk fluid is considered as electroneutral //and thus the net space charge density is zero. - rho_e = (UseSlippingVelBC==1) ? 0.0 : Den_charge[n] / epsilon_LB; + rho_e = Den_charge[n] / epsilon_LB; + // rho_e = (UseSlippingVelBC==1) ? 0.0 : Den_charge[n] / epsilon_LB; idx=Map[n]; + psi = Psi[idx]; f0 = dist[n];