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

update hip

Browse Source
This commit is contained in:
James McClure
2022-05-12 06:58:10 -04:00
parent 0e769186a5
commit ffe4bdd917
1 changed files with 56 additions and 76 deletions
Download Patch File Download Diff File Expand all files Collapse all files

132
hip/Ion.hip
View File

@@ -104,52 +104,52 @@ __global__ void dvc_ScaLBL_D3Q7_Membrane_AssignLinkCoef_halo(
if (idx < count) { if (idx < count) {
n = d3q7_recvlist[idx]; n = d3q7_recvlist[idx];
label = d3q7_linkList[idx]; label = d3q7_linkList[idx];
ap = 1.0; // regular streaming rule ap = 1.0; // regular streaming rule
aq = 1.0; aq = 1.0;
if (label > 0 && !(n < 0)){ if (label > 0 && !(n < 0)){
nqm = Map[n]; nqm = Map[n];
distanceLocal = Distance[nqm]; distanceLocal = Distance[nqm];
psiLocal = Psi[nqm]; psiLocal = Psi[nqm];
// Get the 3-D indices from the send process // Get the 3-D indices from the send process
k = nqm/(Nx*Ny); j = (nqm-Nx*Ny*k)/Nx; i = nqm-Nx*Ny*k-Nx*j; k = nqm/(Nx*Ny); j = (nqm-Nx*Ny*k)/Nx; i = nqm-Nx*Ny*k-Nx*j;
// Streaming link the non-local distribution // Streaming link the non-local distribution
i -= Cqx; j -= Cqy; k -= Cqz; i -= Cqx; j -= Cqy; k -= Cqz;
npm = k*Nx*Ny + j*Nx + i; npm = k*Nx*Ny + j*Nx + i;
distanceNonlocal = Distance[npm]; distanceNonlocal = Distance[npm];
psiNonlocal = Psi[npm]; psiNonlocal = Psi[npm];
membranePotential = psiLocal - psiNonlocal; membranePotential = psiLocal - psiNonlocal;
aq = MassFractionIn; aq = MassFractionIn;
ap = MassFractionOut; ap = MassFractionOut;
/* link is inside membrane */ /* link is inside membrane */
if (distanceLocal > 0.0){ if (distanceLocal > 0.0){
if (membranePotential < Threshold*(-1.0)){ if (membranePotential < Threshold*(-1.0)){
ap = MassFractionIn; ap = MassFractionIn;
aq = MassFractionOut; aq = MassFractionOut;
} }
else { else {
ap = ThresholdMassFractionIn; ap = ThresholdMassFractionIn;
aq = ThresholdMassFractionOut; aq = ThresholdMassFractionOut;
} }
} }
else if (membranePotential > Threshold){ else if (membranePotential > Threshold){
aq = ThresholdMassFractionIn; aq = ThresholdMassFractionIn;
ap = ThresholdMassFractionOut; ap = ThresholdMassFractionOut;
} }
} }
coef[2*idx]=aq; coef[2*idx]=aq;
coef[2*idx+1]=ap; coef[2*idx+1]=ap;
} }
} }
} }
__global__ void dvc_ScaLBL_D3Q7_Membrane_Unpack(int q, __global__ void dvc_ScaLBL_D3Q7_Membrane_Unpack(int q,
int *d3q7_recvlist, int *d3q7_linkList, int start, int nlinks, int count, int *d3q7_recvlist, double *recvbuf, int count,
double *recvbuf, double *dist, int N, double *coef) { double *dist, int N, double *coef) {
//.................................................................................... //....................................................................................
// Unack distribution from the recv buffer // Unack distribution from the recv buffer
// Distribution q matche Cqx, Cqy, Cqz // Distribution q matche Cqx, Cqy, Cqz
@@ -163,41 +163,21 @@ __global__ void dvc_ScaLBL_D3Q7_Membrane_Unpack(int q,
int S = count/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....................
link = S*blockIdx.x*blockDim.x + s*blockDim.x + threadIdx.x; idx = S*blockIdx.x*blockDim.x + s*blockDim.x + threadIdx.x;
if (idx < count){
/* First unpack the regular links */ n = d3q7_recvlist[idx];
if (link < nlinks) { // update link based on mass transfer coefficients
// get the index for the recv list (deal with reordering of links) if (!(n < 0)){
idx = d3q7_linkList[start+link]; aq = coef[2*idx];
// get the distribution index ap = coef[2*idx+1];
n = d3q7_recvlist[start+idx]; fq = dist[q * N + n];
if (n > N){ fp = recvbuf[idx];
printf(" MEMBRANE ERROR: q=%i, idx=%i, n=%i, Np=%i \n",q,idx,n,N); fqq = (1-aq)*fq+ap*fp;
} dist[q * N + n] = fqq;
else if (!(n < 0)){ }
fp = recvbuf[start + idx]; }
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;
// }
// }
} }
} }
__global__ void dvc_ScaLBL_D3Q7_Membrane_IonTransport(int *membrane, double *coef, __global__ void dvc_ScaLBL_D3Q7_Membrane_IonTransport(int *membrane, double *coef,
double *dist, double *Den, int memLinks, int Np){ double *dist, double *Den, int memLinks, int Np){
int link,iq,ip,nq,np; int link,iq,ip,nq,np;
@@ -967,13 +947,13 @@ extern "C" void ScaLBL_D3Q7_Membrane_AssignLinkCoef_halo(
extern "C" void ScaLBL_D3Q7_Membrane_Unpack(int q, extern "C" void ScaLBL_D3Q7_Membrane_Unpack(int q,
int *d3q7_recvlist, int *d3q7_linkList, int start, int nlinks, int count, int *d3q7_recvlist, double *recvbuf, int count,
double *recvbuf, double *dist, int N, double *coef) { double *dist, int N, double *coef){
int GRID = count / 1024 + 1; int GRID = count / NTHREADS + 1;
dvc_ScaLBL_D3Q7_Membrane_Unpack<<<GRID,1024 >>>(q, d3q7_recvlist, d3q7_linkList, start, nlinks, count, dvc_ScaLBL_D3Q7_Membrane_Unpack<<<GRID,NTHREADS >>>(q, d3q7_recvlist, recvbuf,count,
recvbuf, dist, N, coef) ; dist, N, coef);
hipError_t err = hipGetLastError(); hipError_t err = hipGetLastError();
if (hipSuccess != err){ if (hipSuccess != err){