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

Merge branch 'master' of github.com:JamesEMcClure/LBPM-WIA into merge

Browse Source
This commit is contained in:
James McClure 2023-01-21 08:48:25 -05:00
commit bf076ca633
18 changed files with 819 additions and 49 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

15
common/Domain.cpp
View File

@ -217,18 +217,25 @@ void Domain::read_swc(const std::string &Filename) {
double z = k*voxel_length + start_z; double z = k*voxel_length + start_z;
double distance; double distance;
double s = ((x-xp)*alpha+(y-yp)*beta+(z-zp)*gamma) / (alpha*alpha + beta*beta + gamma*gamma); double s = ((x-xp)*alpha+(y-yp)*beta+(z-zp)*gamma) / (alpha*alpha + beta*beta + gamma*gamma);
if (s > length){
distance = ri - sqrt((x-xi)*(x-xi) + (y-yi)*(y-yi) + (z-zi)*(z-zi)); double di = ri - sqrt((x-xi)*(x-xi) + (y-yi)*(y-yi) + (z-zi)*(z-zi));
double dp = rp - sqrt((x-xp)*(x-xp) + (y-yp)*(y-yp) + (z-zp)*(z-zp));
if (s > length ){
distance = di;
} }
else if (s < 0.0){ else if (s < 0.0){
distance = rp - sqrt((x-xp)*(x-xp) + (y-yp)*(y-yp) + (z-zp)*(z-zp)); distance = dp;
} }
else { else {
// linear variation for radius // linear variation for radius
double radius = rp + (ri - rp)*s/length; double radius = rp + (ri - rp)*s/length;
distance = radius - sqrt((x-xp-alpha*s)*(x-xp-alpha*s) + (y-yp-beta*s)*(y-yp-beta*s) + (z-zp-gamma*s)*(z-zp-gamma*s)); distance = radius - sqrt((x-xp-alpha*s)*(x-xp-alpha*s) + (y-yp-beta*s)*(y-yp-beta*s) + (z-zp-gamma*s)*(z-zp-gamma*s));
} }
if (distance < di) distance = di;
if (distance < dp) distance = dp;
if ( distance > 0.0 ){ if ( distance > 0.0 ){
/* label the voxel */ /* label the voxel */
//id[k*Nx*Ny + j*Nx + i] = label; //id[k*Nx*Ny + j*Nx + i] = label;

68
common/Membrane.cpp
View File

@ -667,6 +667,74 @@ int Membrane::Create(DoubleArray &Distance, IntArray &Map){
return mlink; return mlink;
} }
void Membrane::Write(string filename){
int mlink = membraneLinkCount;
std::ofstream ofs (filename, std::ofstream::out);
/* Create local copies of membrane data structures */
double *tmpMembraneCoef; // mass transport coefficient for the membrane
tmpMembraneCoef = new double [2*mlink*sizeof(double)];
ScaLBL_CopyToHost(tmpMembraneCoef, MembraneCoef, 2*mlink*sizeof(double));
int i,j,k;
for (int m=0; m<mlink; m++){
double a1 = tmpMembraneCoef[2*m];
double a2 = tmpMembraneCoef[2*m+1];
int m1 = membraneLinks[2*m]%Np;
int m2 = membraneLinks[2*m+1]%Np;
// map index to global i,j,k
k = m1/(Nx*Ny); j = (m1-Nx*Ny*k)/Nx; i = m1-Nx*Ny*k-Nx*j;
ofs << i << " " << j << " " << k << " "<< a1 ;
k = m2/(Nx*Ny); j = (m2-Nx*Ny*k)/Nx; i = m2-Nx*Ny*k-Nx*j;
ofs << i << " " << j << " " << k << " "<< a2 << endl;
}
ofs.close();
/*FILE *VELX_FILE;
sprintf(LocalRankFilename, "Velocity_X.%05i.raw", rank);
VELX_FILE = fopen(LocalRankFilename, "wb");
fwrite(PhaseField.data(), 8, N, VELX_FILE);
fclose(VELX_FILE);
*/
delete [] tmpMembraneCoef;
}
void Membrane::Read(string filename){
int mlink = membraneLinkCount;
/* Create local copies of membrane data structures */
double *tmpMembraneCoef; // mass transport coefficient for the membrane
tmpMembraneCoef = new double [2*mlink*sizeof(double)];
FILE *fid = fopen(filename.c_str(), "r");
INSIST(fid != NULL, "Error opening membrane file \n");
//........read the spheres..................
// We will read until a blank like or end-of-file is reached
int count = 0;
int i,j,k;
int ii,jj,kk;
double a1, a2;
while (fscanf(fid, "%i,%i,%i,%lf,%i,%i,%i,%lf,\n", &i, &j, &k, &a1, &ii, &jj, &kk, &a2) == 8){
printf("%i, %i, %i, %lf \n", i,j,k, a2);
count++;
}
if (count != mlink){
printf("WARNING (Membrane::Read): number of file lines does not match number of links \n");
}
fclose(fid);
ScaLBL_CopyToDevice(MembraneCoef, tmpMembraneCoef, 2*mlink*sizeof(double));
/*FILE *VELX_FILE;
sprintf(LocalRankFilename, "Velocity_X.%05i.raw", rank);
VELX_FILE = fopen(LocalRankFilename, "wb");
fwrite(PhaseField.data(), 8, N, VELX_FILE);
fclose(VELX_FILE);
*/
delete [] tmpMembraneCoef;
}
int Membrane::D3Q7_MapRecv(int Cqx, int Cqy, int Cqz, int *d3q19_recvlist, int Membrane::D3Q7_MapRecv(int Cqx, int Cqy, int Cqz, int *d3q19_recvlist,
int count, int *membraneRecvLabels, DoubleArray &Distance, int *dvcMap){ int count, int *membraneRecvLabels, DoubleArray &Distance, int *dvcMap){

12
common/Membrane.h
View File

@ -92,6 +92,18 @@ public:
* @param Map - mapping between regular layout and compact layout * @param Map - mapping between regular layout and compact layout
*/ */
int Create(DoubleArray &Distance, IntArray &Map); int Create(DoubleArray &Distance, IntArray &Map);
/**
* \brief Write membrane data to output file
* @param filename - name of file to save
*/
void Write(string filename);
/**
* \brief Read membrane data from input file
* @param filename - name of file to save
*/
void Read(string filename);
void SendD3Q7AA(double *dist); void SendD3Q7AA(double *dist);
void RecvD3Q7AA(double *dist); void RecvD3Q7AA(double *dist);

23
common/ScaLBL.cpp
View File

@ -2472,6 +2472,29 @@ void ScaLBL_Communicator::D3Q7_Poisson_Potential_BC_Z(int *neighborList, double
} }
} }
void ScaLBL_Communicator::D3Q19_Poisson_Potential_BC_z(int *neighborList, double *fq, double Vin, int time){
if (kproc == 0) {
if (time%2==0){
ScaLBL_D3Q19_AAeven_Poisson_Potential_BC_z(dvcSendList_z, fq, Vin, sendCount_z, N);
}
else{
ScaLBL_D3Q19_AAodd_Poisson_Potential_BC_z(neighborList, dvcSendList_z, fq, Vin, sendCount_z, N);
}
}
}
void ScaLBL_Communicator::D3Q19_Poisson_Potential_BC_Z(int *neighborList, double *fq, double Vout, int time){
if (kproc == nprocz-1){
if (time%2==0){
ScaLBL_D3Q19_AAeven_Poisson_Potential_BC_Z(dvcSendList_Z, fq, Vout, sendCount_Z, N);
}
else{
ScaLBL_D3Q19_AAodd_Poisson_Potential_BC_Z(neighborList, dvcSendList_Z, fq, Vout, sendCount_Z, N);
}
}
}
void ScaLBL_Communicator::Poisson_D3Q7_BC_z(int *Map, double *Psi, double Vin){ void ScaLBL_Communicator::Poisson_D3Q7_BC_z(int *Map, double *Psi, double Vin){
if (kproc == 0) { if (kproc == 0) {
ScaLBL_Poisson_D3Q7_BC_z(dvcSendList_z, Map, Psi, Vin, sendCount_z); ScaLBL_Poisson_D3Q7_BC_z(dvcSendList_z, Map, Psi, Vin, sendCount_z);

8
common/ScaLBL.h
View File

@ -396,6 +396,12 @@ extern "C" void ScaLBL_D3Q19_AAeven_Poisson(int *Map, double *dist,
extern "C" void ScaLBL_D3Q19_Poisson_getElectricField(double *dist, double *ElectricField, double tau, int Np); extern "C" void ScaLBL_D3Q19_Poisson_getElectricField(double *dist, double *ElectricField, double tau, int Np);
extern "C" void ScaLBL_D3Q19_AAodd_Poisson_Potential_BC_Z(int *d_neighborList, int *list, double *dist, double Vin, int count, int Np);
extern "C" void ScaLBL_D3Q19_AAodd_Poisson_Potential_BC_z(int *d_neighborList, int *list, double *dist, double Vin, int count, int Np);
extern "C" void ScaLBL_D3Q19_AAeven_Poisson_Potential_BC_Z(int *list, double *dist, double Vout, int count, int Np);
extern "C" void ScaLBL_D3Q19_AAeven_Poisson_Potential_BC_z(int *list, double *dist, double Vout, int count, int Np);
// LBM Stokes Model (adapted from MRT model) // LBM Stokes Model (adapted from MRT model)
extern "C" void ScaLBL_D3Q19_AAeven_StokesMRT(double *dist, double *Velocity, double *ChargeDensity, double *ElectricField, double rlx_setA, double rlx_setB, extern "C" void ScaLBL_D3Q19_AAeven_StokesMRT(double *dist, double *Velocity, double *ChargeDensity, double *ElectricField, double rlx_setA, double rlx_setB,
double Gx, double Gy, double Gz,double rho0, double den_scale, double h, double time_conv, bool UseSlippingVelBC, int start, int finish, int Np); double Gx, double Gy, double Gz,double rho0, double den_scale, double h, double time_conv, bool UseSlippingVelBC, int start, int finish, int Np);
@ -806,6 +812,8 @@ public:
double D3Q19_Flux_BC_z(int *neighborList, double *fq, double flux, int time); double D3Q19_Flux_BC_z(int *neighborList, double *fq, double flux, int time);
void D3Q7_Poisson_Potential_BC_z(int *neighborList, double *fq, double Vin, int time); void D3Q7_Poisson_Potential_BC_z(int *neighborList, double *fq, double Vin, int time);
void D3Q7_Poisson_Potential_BC_Z(int *neighborList, double *fq, double Vout, int time); void D3Q7_Poisson_Potential_BC_Z(int *neighborList, double *fq, double Vout, int time);
void D3Q19_Poisson_Potential_BC_z(int *neighborList, double *fq, double Vin, int time);
void D3Q19_Poisson_Potential_BC_Z(int *neighborList, double *fq, double Vout, int time);
void Poisson_D3Q7_BC_z(int *Map, double *Psi, double Vin); void Poisson_D3Q7_BC_z(int *Map, double *Psi, double Vin);
void Poisson_D3Q7_BC_Z(int *Map, double *Psi, double Vout); void Poisson_D3Q7_BC_Z(int *Map, double *Psi, double Vout);
void D3Q7_Ion_Concentration_BC_z(int *neighborList, double *fq, double Cin, int time); void D3Q7_Ion_Concentration_BC_z(int *neighborList, double *fq, double Cin, int time);

3
cpu/D3Q7BC.cpp
View File

@ -181,6 +181,7 @@ extern "C" void ScaLBL_D3Q7_AAodd_Poisson_Potential_BC_z(int *d_neighborList,
} }
} }
extern "C" void ScaLBL_D3Q7_AAodd_Poisson_Potential_BC_Z(int *d_neighborList, extern "C" void ScaLBL_D3Q7_AAodd_Poisson_Potential_BC_Z(int *d_neighborList,
int *list, int *list,
double *dist, double *dist,
@ -213,6 +214,8 @@ extern "C" void ScaLBL_D3Q7_AAodd_Poisson_Potential_BC_Z(int *d_neighborList,
} }
} }
extern "C" void ScaLBL_Poisson_D3Q7_BC_z(int *list, int *Map, double *Psi, extern "C" void ScaLBL_Poisson_D3Q7_BC_z(int *list, int *Map, double *Psi,
double Vin, int count) { double Vin, int count) {
int idx, n, nm; int idx, n, nm;

88
cpu/Poisson.cpp
View File

@ -631,6 +631,8 @@ extern "C" void ScaLBL_D3Q19_AAeven_Poisson_ElectricPotential(
} }
} }
extern "C" void ScaLBL_D3Q19_AAodd_Poisson(int *neighborList, int *Map, extern "C" void ScaLBL_D3Q19_AAodd_Poisson(int *neighborList, int *Map,
double *dist, double *Den_charge, double *dist, double *Den_charge,
double *Psi, double *ElectricField, double *Psi, double *ElectricField,
@ -915,6 +917,92 @@ extern "C" void ScaLBL_D3Q19_AAeven_Poisson(int *Map, double *dist,
} }
} }
extern "C" void ScaLBL_D3Q19_AAeven_Poisson_Potential_BC_z(int *list, double *dist, double Vin, int count, int Np) {
double W1 = 1.0/24.0;
double W2 = 1.0/48.0;
int nread, nr5;
for (int idx = 0; idx < count; idx++) {
int n = list[idx];
dist[6 * Np + n] = W1*Vin;
dist[12 * Np + n] = W2*Vin;
dist[13 * Np + n] = W2*Vin;
dist[16 * Np + n] = W2*Vin;
dist[17 * Np + n] = W2*Vin;
}
}
extern "C" void ScaLBL_D3Q19_AAeven_Poisson_Potential_BC_Z(int *list,
double *dist,
double Vout,
int count, int Np) {
double W1 = 1.0/24.0;
double W2 = 1.0/48.0;
for (int idx = 0; idx < count; idx++) {
int n = list[idx];
dist[5 * Np + n] = W1*Vout;
dist[11 * Np + n] = W2*Vout;
dist[14 * Np + n] = W2*Vout;
dist[15 * Np + n] = W2*Vout;
dist[18 * Np + n] = W2*Vout;
}
}
extern "C" void ScaLBL_D3Q19_AAodd_Poisson_Potential_BC_z(int *d_neighborList,
int *list,
double *dist,
double Vin, int count,
int Np) {
double W1 = 1.0/24.0;
double W2 = 1.0/48.0;
int nr5, nr11, nr14, nr15, nr18;
for (int idx = 0; idx < count; idx++) {
int n = list[idx];
// Unknown distributions
nr5 = d_neighborList[n + 4 * Np];
nr11 = d_neighborList[n + 10 * Np];
nr15 = d_neighborList[n + 14 * Np];
nr14 = d_neighborList[n + 13 * Np];
nr18 = d_neighborList[n + 17 * Np];
dist[nr5] = W1*Vin;
dist[nr11] = W2*Vin;
dist[nr15] = W2*Vin;
dist[nr14] = W2*Vin;
dist[nr18] = W2*Vin;
}
}
extern "C" void ScaLBL_D3Q19_AAodd_Poisson_Potential_BC_Z(int *d_neighborList, int *list, double *dist, double Vout, int count, int Np) {
double W1 = 1.0/24.0;
double W2 = 1.0/48.0;
int nr6, nr12, nr13, nr16, nr17;
for (int idx = 0; idx < count; idx++) {
int n = list[idx];
// unknown distributions
nr6 = d_neighborList[n + 5 * Np];
nr12 = d_neighborList[n + 11 * Np];
nr16 = d_neighborList[n + 15 * Np];
nr17 = d_neighborList[n + 16 * Np];
nr13 = d_neighborList[n + 12 * Np];
dist[nr6] = W1*Vout;
dist[nr12] = W2*Vout;
dist[nr16] = W2*Vout;
dist[nr17] = W2*Vout;
dist[nr13] = W2*Vout;
}
}
extern "C" void ScaLBL_D3Q19_Poisson_Init(int *Map, double *dist, double *Psi, extern "C" void ScaLBL_D3Q19_Poisson_Init(int *Map, double *dist, double *Psi,
int start, int finish, int Np) { int start, int finish, int Np) {
int n; int n;

148
cuda/Poisson.cu
View File

@ -3,7 +3,7 @@
//#include <cuda_profiler_api.h> //#include <cuda_profiler_api.h>
#define NBLOCKS 1024 #define NBLOCKS 1024
#define NTHREADS 256 #define NTHREADS 512
__global__ void dvc_ScaLBL_D3Q7_AAodd_Poisson_ElectricPotential(int *neighborList,int *Map, double *dist, double *Psi, int start, int finish, int Np){ __global__ void dvc_ScaLBL_D3Q7_AAodd_Poisson_ElectricPotential(int *neighborList,int *Map, double *dist, double *Psi, int start, int finish, int Np){
int n; int n;
@ -328,7 +328,7 @@ __global__ void dvc_ScaLBL_D3Q19_AAodd_Poisson(int *neighborList, int *Map,
f16, f17, f18; f16, f17, f18;
int nr1, nr2, nr3, nr4, nr5, nr6, nr7, nr8, nr9, nr10, nr11, nr12, nr13, int nr1, nr2, nr3, nr4, nr5, nr6, nr7, nr8, nr9, nr10, nr11, nr12, nr13,
nr14, nr15, nr16, nr17, nr18; nr14, nr15, nr16, nr17, nr18;
double error,sum_q; double sum_q;
double rlx = 1.0 / tau; double rlx = 1.0 / tau;
int idx; int idx;
@ -421,7 +421,7 @@ __global__ void dvc_ScaLBL_D3Q19_AAodd_Poisson(int *neighborList, int *Map,
f18 = dist[nr18]; f18 = dist[nr18];
sum_q = f1+f2+f3+f4+f5+f6+f7+f8+f9+f10+f11+f12+f13+f14+f15+f16+f17+f18; sum_q = f1+f2+f3+f4+f5+f6+f7+f8+f9+f10+f11+f12+f13+f14+f15+f16+f17+f18;
error = 8.0*(sum_q - f0) + rho_e; //error = 8.0*(sum_q - f0) + rho_e;
psi = 2.0*(f0*(1.0 - rlx) + rlx*(sum_q + 0.125*rho_e)); psi = 2.0*(f0*(1.0 - rlx) + rlx*(sum_q + 0.125*rho_e));
@ -545,6 +545,12 @@ __global__ void dvc_ScaLBL_D3Q19_AAeven_Poisson(int *Map, double *dist,
f17 = dist[18 * Np + n]; f17 = dist[18 * Np + n];
f18 = dist[17 * Np + n]; f18 = dist[17 * Np + n];
/* Ex = (f1 - f2) * rlx *
4.0; //NOTE the unit of electric field here is V/lu
Ey = (f3 - f4) * rlx *
4.0; //factor 4.0 is D3Q7 lattice squared speed of sound
Ez = (f5 - f6) * rlx * 4.0;
*/
Ex = (f1 - f2 + 0.5*(f7 - f8 + f9 - f10 + f11 - f12 + f13 - f14))*4.0; //NOTE the unit of electric field here is V/lu Ex = (f1 - f2 + 0.5*(f7 - f8 + f9 - f10 + f11 - f12 + f13 - f14))*4.0; //NOTE the unit of electric field here is V/lu
Ey = (f3 - f4 + 0.5*(f7 - f8 - f9 + f10 + f15 - f16 + f17 - f18))*4.0; Ey = (f3 - f4 + 0.5*(f7 - f8 - f9 + f10 + f15 - f16 + f17 - f18))*4.0;
Ez = (f5 - f6 + 0.5*(f11 - f12 - f13 + f14 + f15 - f16 - f17 + f18))*4.0; Ez = (f5 - f6 + 0.5*(f11 - f12 - f13 + f14 + f15 - f16 - f17 + f18))*4.0;
@ -554,12 +560,14 @@ __global__ void dvc_ScaLBL_D3Q19_AAeven_Poisson(int *Map, double *dist,
sum_q = f1+f2+f3+f4+f5+f6+f7+f8+f9+f10+f11+f12+f13+f14+f15+f16+f17+f18; sum_q = f1+f2+f3+f4+f5+f6+f7+f8+f9+f10+f11+f12+f13+f14+f15+f16+f17+f18;
error = 8.0*(sum_q - f0) + rho_e; error = 8.0*(sum_q - f0) + rho_e;
Error[n] = error;
psi = 2.0*(f0*(1.0 - rlx) + rlx*(sum_q + 0.125*rho_e)); psi = 2.0*(f0*(1.0 - rlx) + rlx*(sum_q + 0.125*rho_e));
idx = Map[n]; idx = Map[n];
Psi[idx] = psi; Psi[idx] = psi;
// q = 0 // q = 0
dist[n] = W0*psi;// dist[n] = W0*psi;//
@ -593,7 +601,6 @@ __global__ void dvc_ScaLBL_D3Q19_AAeven_Poisson(int *Map, double *dist,
dist[16 * Np + n] = W2*psi;//f16 * (1.0 - rlx) +W2* (rlx * psi) - (1.0-0.5*rlx)*0.02777777777777778*rho_e; dist[16 * Np + n] = W2*psi;//f16 * (1.0 - rlx) +W2* (rlx * psi) - (1.0-0.5*rlx)*0.02777777777777778*rho_e;
dist[17 * Np + n] = W2*psi;//f17 * (1.0 - rlx) +W2* (rlx * psi) - (1.0-0.5*rlx)*0.02777777777777778*rho_e; dist[17 * Np + n] = W2*psi;//f17 * (1.0 - rlx) +W2* (rlx * psi) - (1.0-0.5*rlx)*0.02777777777777778*rho_e;
dist[18 * Np + n] = W2*psi;//f18 * (1.0 - rlx) +W2* (rlx * psi) - (1.0-0.5*rlx)*0.02777777777777778*rho_e; dist[18 * Np + n] = W2*psi;//f18 * (1.0 - rlx) +W2* (rlx * psi) - (1.0-0.5*rlx)*0.02777777777777778*rho_e;
//........................................................................ //........................................................................
} }
} }
@ -635,6 +642,131 @@ __global__ void dvc_ScaLBL_D3Q19_Poisson_Init(int *Map, double *dist, double *P
} }
} }
__global__ void dvc_ScaLBL_D3Q19_AAeven_Poisson_Potential_BC_z(int *list, double *dist, double Vin, int count, int Np) {
double W1 = 1.0/24.0;
double W2 = 1.0/48.0;
int idx = blockIdx.x*blockDim.x + threadIdx.x;
if (idx < count){
int n = list[idx];
dist[6 * Np + n] = W1*Vin;
dist[12 * Np + n] = W2*Vin;
dist[13 * Np + n] = W2*Vin;
dist[16 * Np + n] = W2*Vin;
dist[17 * Np + n] = W2*Vin;
}
}
__global__ void dvc_ScaLBL_D3Q19_AAeven_Poisson_Potential_BC_Z(int *list, double *dist, double Vout, int count, int Np) {
double W1 = 1.0/24.0;
double W2 = 1.0/48.0;
int idx = blockIdx.x*blockDim.x + threadIdx.x;
if (idx < count){
int n = list[idx];
dist[5 * Np + n] = W1*Vout;
dist[11 * Np + n] = W2*Vout;
dist[14 * Np + n] = W2*Vout;
dist[15 * Np + n] = W2*Vout;
dist[18 * Np + n] = W2*Vout;
}
}
__global__ void dvc_ScaLBL_D3Q19_AAodd_Poisson_Potential_BC_z(int *d_neighborList, int *list, double *dist, double Vin, int count, int Np) {
double W1 = 1.0/24.0;
double W2 = 1.0/48.0;
int nr5, nr11, nr14, nr15, nr18;
int idx = blockIdx.x*blockDim.x + threadIdx.x;
if (idx < count){
int n = list[idx];
// Unknown distributions
nr5 = d_neighborList[n + 4 * Np];
nr11 = d_neighborList[n + 10 * Np];
nr15 = d_neighborList[n + 14 * Np];
nr14 = d_neighborList[n + 13 * Np];
nr18 = d_neighborList[n + 17 * Np];
dist[nr5] = W1*Vin;
dist[nr11] = W2*Vin;
dist[nr15] = W2*Vin;
dist[nr14] = W2*Vin;
dist[nr18] = W2*Vin;
}
}
__global__ void dvc_ScaLBL_D3Q19_AAodd_Poisson_Potential_BC_Z(int *d_neighborList, int *list, double *dist, double Vout, int count, int Np) {
double W1 = 1.0/24.0;
double W2 = 1.0/48.0;
int nr6, nr12, nr13, nr16, nr17;
int idx = blockIdx.x*blockDim.x + threadIdx.x;
if (idx < count){
int n = list[idx];
// unknown distributions
nr6 = d_neighborList[n + 5 * Np];
nr12 = d_neighborList[n + 11 * Np];
nr16 = d_neighborList[n + 15 * Np];
nr17 = d_neighborList[n + 16 * Np];
nr13 = d_neighborList[n + 12 * Np];
dist[nr6] = W1*Vout;
dist[nr12] = W2*Vout;
dist[nr16] = W2*Vout;
dist[nr17] = W2*Vout;
dist[nr13] = W2*Vout;
}
}
/* wrapper functions to launch kernels */
extern "C" void ScaLBL_D3Q19_AAeven_Poisson_Potential_BC_z(int *list, double *dist, double Vin, int count, int Np){
int GRID = count / 512 + 1;
dvc_ScaLBL_D3Q19_AAeven_Poisson_Potential_BC_z<<<GRID,512>>>(list, dist, Vin, count, Np);
cudaError_t err = cudaGetLastError();
if (cudaSuccess != err){
printf("CUDA error in ScaLBL_D3Q19_AAeven_Poisson_Potential_BC_z (kernel): %s \n",cudaGetErrorString(err));
}
}
//
extern "C" void ScaLBL_D3Q19_AAeven_Poisson_Potential_BC_Z(int *list, double *dist, double Vout, int count, int Np){
int GRID = count / 512 + 1;
dvc_ScaLBL_D3Q19_AAeven_Poisson_Potential_BC_Z<<<GRID,512>>>(list, dist, Vout, count, Np);
cudaError_t err = cudaGetLastError();
if (cudaSuccess != err){
printf("CUDA error in ScaLBL_D3Q19_AAeven_Poisson_Potential_BC_Z (kernel): %s \n",cudaGetErrorString(err));
}
}
extern "C" void ScaLBL_D3Q19_AAodd_Poisson_Potential_BC_z(int *d_neighborList, int *list, double *dist, double Vin, int count,int Np) {
int GRID = count / 512 + 1;
dvc_ScaLBL_D3Q19_AAodd_Poisson_Potential_BC_z<<<GRID,512>>>(d_neighborList, list, dist, Vin, count, Np);
cudaError_t err = cudaGetLastError();
if (cudaSuccess != err){
printf("CUDA error in ScaLBL_D3Q19_AAodd_Poisson_Potential_BC_z (kernel): %s \n",cudaGetErrorString(err));
}
}
//
extern "C" void ScaLBL_D3Q19_AAodd_Poisson_Potential_BC_Z(int *d_neighborList, int *list, double *dist, double Vout, int count, int Np) {
int GRID = count / 512 + 1;
dvc_ScaLBL_D3Q19_AAodd_Poisson_Potential_BC_Z<<<GRID,512>>>(d_neighborList, list, dist, Vout, count, Np);
cudaError_t err = cudaGetLastError();
if (cudaSuccess != err){
printf("CUDA error in ScaLBL_D3Q19_AAodd_Poisson_Potential_BC_Z (kernel): %s \n",cudaGetErrorString(err));
}
}
extern "C" void ScaLBL_D3Q19_AAodd_Poisson(int *neighborList, int *Map, extern "C" void ScaLBL_D3Q19_AAodd_Poisson(int *neighborList, int *Map,
double *dist, double *Den_charge, double *dist, double *Den_charge,
double *Psi, double *ElectricField, double *Psi, double *ElectricField,

40
docs/source/userGuide/models/color/index.rst
View File

@ -176,42 +176,70 @@ The non-zero equilibrium moments are defined as
:nowrap: :nowrap:
$$ $$
m_1^{eq} = (j_x^2+j_y^2+j_z^2) - \alpha |\textbf{C}|, \\ m_1^{eq} = 19\frac{ j_x^2+j_y^2+j_z^2}{\rho_0} - 11\rho - 19 \alpha |\textbf{C}|, \\
$$
.. math::
:nowrap:
$$
m_2^{eq} = 3\rho - \frac{11( j_x^2+j_y^2+j_z^2)}{2\rho_0}, \\
$$
.. math::
:nowrap:
$$
m_4^{eq} = -\frac{2 j_x}{3}, \\
$$
.. math::
:nowrap:
$$
m_6^{eq} = -\frac{2 j_y}{3}, \\
$$
.. math::
:nowrap:
$$
m_8^{eq} = -\frac{2 j_z}{3}, \\
$$ $$
.. math:: .. math::
:nowrap: :nowrap:
$$ $$
m_9^{eq} = (2j_x^2-j_y^2-j_z^2)+ \alpha \frac{|\textbf{C}|}{2}(2n_x^2-n_y^2-n_z^2), \\ m_9^{eq} = \frac{2j_x^2-j_y^2-j_z^2}{\rho_0}+ \alpha \frac{|\textbf{C}|}{2}(2n_x^2-n_y^2-n_z^2), \\
$$ $$
.. math:: .. math::
:nowrap: :nowrap:
$$ $$
m_{11}^{eq} = (j_y^2-j_z^2) + \alpha \frac{|\textbf{C}|}{2}(n_y^2-n_z^2), \\ m_{11}^{eq} = \frac{j_y^2-j_z^2}{\rho_0} + \alpha \frac{|\textbf{C}|}{2}(n_y^2-n_z^2), \\
$$ $$
.. math:: .. math::
:nowrap: :nowrap:
$$ $$
m_{13}^{eq} = j_x j_y + \alpha \frac{|\textbf{C}|}{2} n_x n_y\;, \\ m_{13}^{eq} = \frac{j_x j_y}{\rho_0} + \alpha \frac{|\textbf{C}|}{2} n_x n_y\;, \\
$$ $$
.. math:: .. math::
:nowrap: :nowrap:
$$ $$
m_{14}^{eq} = j_y j_z + \alpha \frac{|\textbf{C}|}{2} n_y n_z\;, \\ m_{14}^{eq} = \frac{j_y j_z}{\rho_0} + \alpha \frac{|\textbf{C}|}{2} n_y n_z\;, \\
$$ $$
.. math:: .. math::
:nowrap: :nowrap:
$$ $$
m_{15}^{eq} = j_x j_z + \alpha \frac{|\textbf{C}|}{2} n_x n_z\;, m_{15}^{eq} = \frac{j_x j_z}{\rho_0} + \alpha \frac{|\textbf{C}|}{2} n_x n_z\;.
$$ $$
where the color gradient is determined from the phase indicator field where the color gradient is determined from the phase indicator field

41
docs/source/userGuide/models/greyscaleColor/greyscaleColor.rst
View File

@ -241,46 +241,75 @@ The relaxation parameters are determined from the relaxation time:
The non-zero equilibrium moments are defined as The non-zero equilibrium moments are defined as
.. math:: .. math::
:nowrap: :nowrap:
$$ $$
m_1^{eq} = (j_x^2+j_y^2+j_z^2) - \alpha |\textbf{C}|, \\ m_1^{eq} = 19\frac{ j_x^2+j_y^2+j_z^2}{\rho_0} - 11\rho - 19 \alpha |\textbf{C}|, \\
$$
.. math::
:nowrap:
$$
m_2^{eq} = 3\rho - \frac{11( j_x^2+j_y^2+j_z^2)}{2\rho_0}, \\
$$
.. math::
:nowrap:
$$
m_4^{eq} = -\frac{2 j_x}{3}, \\
$$
.. math::
:nowrap:
$$
m_6^{eq} = -\frac{2 j_y}{3}, \\
$$
.. math::
:nowrap:
$$
m_8^{eq} = -\frac{2 j_z}{3}, \\
$$ $$
.. math:: .. math::
:nowrap: :nowrap:
$$ $$
m_9^{eq} = (2j_x^2-j_y^2-j_z^2)+ \alpha \frac{|\textbf{C}|}{2}(2n_x^2-n_y^2-n_z^2), \\ m_9^{eq} = \frac{2j_x^2-j_y^2-j_z^2}{\rho_0}+ \alpha \frac{|\textbf{C}|}{2}(2n_x^2-n_y^2-n_z^2), \\
$$ $$
.. math:: .. math::
:nowrap: :nowrap:
$$ $$
m_{11}^{eq} = (j_y^2-j_z^2) + \alpha \frac{|\textbf{C}|}{2}(n_y^2-n_z^2), \\ m_{11}^{eq} = \frac{j_y^2-j_z^2}{\rho_0} + \alpha \frac{|\textbf{C}|}{2}(n_y^2-n_z^2), \\
$$ $$
.. math:: .. math::
:nowrap: :nowrap:
$$ $$
m_{13}^{eq} = j_x j_y + \alpha \frac{|\textbf{C}|}{2} n_x n_y\;, \\ m_{13}^{eq} = \frac{j_x j_y}{\rho_0} + \alpha \frac{|\textbf{C}|}{2} n_x n_y\;, \\
$$ $$
.. math:: .. math::
:nowrap: :nowrap:
$$ $$
m_{14}^{eq} = j_y j_z + \alpha \frac{|\textbf{C}|}{2} n_y n_z\;, \\ m_{14}^{eq} = \frac{j_y j_z}{\rho_0} + \alpha \frac{|\textbf{C}|}{2} n_y n_z\;, \\
$$ $$
.. math:: .. math::
:nowrap: :nowrap:
$$ $$
m_{15}^{eq} = j_x j_z + \alpha \frac{|\textbf{C}|}{2} n_x n_z\;, m_{15}^{eq} = \frac{j_x j_z}{\rho_0} + \alpha \frac{|\textbf{C}|}{2} n_x n_z\;.
$$ $$
where the color gradient is determined from the phase indicator field where the color gradient is determined from the phase indicator field

147
hip/Poisson.hip
View File

@ -301,7 +301,7 @@ __global__ void dvc_ScaLBL_D3Q19_AAodd_Poisson(int *neighborList, int *Map,
//........Get 1-D index for this thread.................... //........Get 1-D index for this thread....................
n = S*blockIdx.x*blockDim.x + s*blockDim.x + threadIdx.x + start; n = S*blockIdx.x*blockDim.x + s*blockDim.x + threadIdx.x + start;
if (n<finish) { if (n<finish) {
//Load data //Load data
//When Helmholtz-Smoluchowski slipping velocity BC is used, the bulk fluid is considered as electroneutral //When Helmholtz-Smoluchowski slipping velocity BC is used, the bulk fluid is considered as electroneutral
//and thus the net space charge density is zero. //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;
@ -451,7 +451,8 @@ __global__ void dvc_ScaLBL_D3Q19_AAodd_Poisson(int *neighborList, int *Map,
dist[nr18] = W2*psi; //f17 * (1.0 - rlx) +W2* (rlx * psi) - (1.0-0.5*rlx)*0.02777777777777778*rho_e; dist[nr18] = W2*psi; //f17 * (1.0 - rlx) +W2* (rlx * psi) - (1.0-0.5*rlx)*0.02777777777777778*rho_e;
// q = 18 // q = 18
dist[nr17] = W2*psi; dist[nr17] = W2*psi; //f18 * (1.0 - rlx) +W2* (rlx * psi) - (1.0-0.5*rlx)*0.02777777777777778*rho_e;
} }
} }
} }
@ -479,7 +480,7 @@ __global__ void dvc_ScaLBL_D3Q19_AAeven_Poisson(int *Map, double *dist,
//........Get 1-D index for this thread.................... //........Get 1-D index for this thread....................
n = S*blockIdx.x*blockDim.x + s*blockDim.x + threadIdx.x + start; n = S*blockIdx.x*blockDim.x + s*blockDim.x + threadIdx.x + start;
if (n<finish) { if (n<finish) {
//Load data //Load data
//When Helmholtz-Smoluchowski slipping velocity BC is used, the bulk fluid is considered as electroneutral //When Helmholtz-Smoluchowski slipping velocity BC is used, the bulk fluid is considered as electroneutral
//and thus the net space charge density is zero. //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;
@ -505,6 +506,12 @@ __global__ void dvc_ScaLBL_D3Q19_AAeven_Poisson(int *Map, double *dist,
f17 = dist[18 * Np + n]; f17 = dist[18 * Np + n];
f18 = dist[17 * Np + n]; f18 = dist[17 * Np + n];
/* Ex = (f1 - f2) * rlx *
4.0; //NOTE the unit of electric field here is V/lu
Ey = (f3 - f4) * rlx *
4.0; //factor 4.0 is D3Q7 lattice squared speed of sound
Ez = (f5 - f6) * rlx * 4.0;
*/
Ex = (f1 - f2 + 0.5*(f7 - f8 + f9 - f10 + f11 - f12 + f13 - f14))*4.0; //NOTE the unit of electric field here is V/lu Ex = (f1 - f2 + 0.5*(f7 - f8 + f9 - f10 + f11 - f12 + f13 - f14))*4.0; //NOTE the unit of electric field here is V/lu
Ey = (f3 - f4 + 0.5*(f7 - f8 - f9 + f10 + f15 - f16 + f17 - f18))*4.0; Ey = (f3 - f4 + 0.5*(f7 - f8 - f9 + f10 + f15 - f16 + f17 - f18))*4.0;
Ez = (f5 - f6 + 0.5*(f11 - f12 - f13 + f14 + f15 - f16 - f17 + f18))*4.0; Ez = (f5 - f6 + 0.5*(f11 - f12 - f13 + f14 + f15 - f16 - f17 + f18))*4.0;
@ -514,12 +521,14 @@ __global__ void dvc_ScaLBL_D3Q19_AAeven_Poisson(int *Map, double *dist,
sum_q = f1+f2+f3+f4+f5+f6+f7+f8+f9+f10+f11+f12+f13+f14+f15+f16+f17+f18; sum_q = f1+f2+f3+f4+f5+f6+f7+f8+f9+f10+f11+f12+f13+f14+f15+f16+f17+f18;
error = 8.0*(sum_q - f0) + rho_e; error = 8.0*(sum_q - f0) + rho_e;
Error[n] = error;
psi = 2.0*(f0*(1.0 - rlx) + rlx*(sum_q + 0.125*rho_e)); psi = 2.0*(f0*(1.0 - rlx) + rlx*(sum_q + 0.125*rho_e));
idx = Map[n]; idx = Map[n];
Psi[idx] = psi; Psi[idx] = psi;
// q = 0 // q = 0
dist[n] = W0*psi;// dist[n] = W0*psi;//
@ -553,7 +562,6 @@ __global__ void dvc_ScaLBL_D3Q19_AAeven_Poisson(int *Map, double *dist,
dist[16 * Np + n] = W2*psi;//f16 * (1.0 - rlx) +W2* (rlx * psi) - (1.0-0.5*rlx)*0.02777777777777778*rho_e; dist[16 * Np + n] = W2*psi;//f16 * (1.0 - rlx) +W2* (rlx * psi) - (1.0-0.5*rlx)*0.02777777777777778*rho_e;
dist[17 * Np + n] = W2*psi;//f17 * (1.0 - rlx) +W2* (rlx * psi) - (1.0-0.5*rlx)*0.02777777777777778*rho_e; dist[17 * Np + n] = W2*psi;//f17 * (1.0 - rlx) +W2* (rlx * psi) - (1.0-0.5*rlx)*0.02777777777777778*rho_e;
dist[18 * Np + n] = W2*psi;//f18 * (1.0 - rlx) +W2* (rlx * psi) - (1.0-0.5*rlx)*0.02777777777777778*rho_e; dist[18 * Np + n] = W2*psi;//f18 * (1.0 - rlx) +W2* (rlx * psi) - (1.0-0.5*rlx)*0.02777777777777778*rho_e;
//........................................................................ //........................................................................
} }
} }
@ -594,6 +602,129 @@ __global__ void dvc_ScaLBL_D3Q19_Poisson_Init(int *Map, double *dist, double *P
} }
} }
} }
__global__ void dvc_ScaLBL_D3Q19_AAeven_Poisson_Potential_BC_z(int *list, double *dist, double Vin, int count, int Np) {
double W1 = 1.0/24.0;
double W2 = 1.0/48.0;
int idx = blockIdx.x*blockDim.x + threadIdx.x;
if (idx < count){
int n = list[idx];
dist[6 * Np + n] = W1*Vin;
dist[12 * Np + n] = W2*Vin;
dist[13 * Np + n] = W2*Vin;
dist[16 * Np + n] = W2*Vin;
dist[17 * Np + n] = W2*Vin;
}
}
__global__ void dvc_ScaLBL_D3Q19_AAeven_Poisson_Potential_BC_Z(int *list, double *dist, double Vout, int count, int Np) {
double W1 = 1.0/24.0;
double W2 = 1.0/48.0;
int idx = blockIdx.x*blockDim.x + threadIdx.x;
if (idx < count){
int n = list[idx];
dist[5 * Np + n] = W1*Vout;
dist[11 * Np + n] = W2*Vout;
dist[14 * Np + n] = W2*Vout;
dist[15 * Np + n] = W2*Vout;
dist[18 * Np + n] = W2*Vout;
}
}
__global__ void dvc_ScaLBL_D3Q19_AAodd_Poisson_Potential_BC_z(int *d_neighborList, int *list, double *dist, double Vin, int count, int Np) {
double W1 = 1.0/24.0;
double W2 = 1.0/48.0;
int nr5, nr11, nr14, nr15, nr18;
int idx = blockIdx.x*blockDim.x + threadIdx.x;
if (idx < count){
int n = list[idx];
// Unknown distributions
nr5 = d_neighborList[n + 4 * Np];
nr11 = d_neighborList[n + 10 * Np];
nr15 = d_neighborList[n + 14 * Np];
nr14 = d_neighborList[n + 13 * Np];
nr18 = d_neighborList[n + 17 * Np];
dist[nr5] = W1*Vin;
dist[nr11] = W2*Vin;
dist[nr15] = W2*Vin;
dist[nr14] = W2*Vin;
dist[nr18] = W2*Vin;
}
}
__global__ void dvc_ScaLBL_D3Q19_AAodd_Poisson_Potential_BC_Z(int *d_neighborList, int *list, double *dist, double Vout, int count, int Np) {
double W1 = 1.0/24.0;
double W2 = 1.0/48.0;
int nr6, nr12, nr13, nr16, nr17;
int idx = blockIdx.x*blockDim.x + threadIdx.x;
if (idx < count){
int n = list[idx];
// unknown distributions
nr6 = d_neighborList[n + 5 * Np];
nr12 = d_neighborList[n + 11 * Np];
nr16 = d_neighborList[n + 15 * Np];
nr17 = d_neighborList[n + 16 * Np];
nr13 = d_neighborList[n + 12 * Np];
dist[nr6] = W1*Vout;
dist[nr12] = W2*Vout;
dist[nr16] = W2*Vout;
dist[nr17] = W2*Vout;
dist[nr13] = W2*Vout;
}
}
/* wrapper functions to launch kernels */
extern "C" void ScaLBL_D3Q19_AAeven_Poisson_Potential_BC_z(int *list, double *dist, double Vin, int count, int Np){
int GRID = count / 512 + 1;
dvc_ScaLBL_D3Q19_AAeven_Poisson_Potential_BC_z<<<GRID,512>>>(list, dist, Vin, count, Np);
cudaError_t err = cudaGetLastError();
if (cudaSuccess != err){
printf("Hip error in ScaLBL_D3Q19_AAeven_Poisson_Potential_BC_z (kernel): %s \n",cudaGetErrorString(err));
}
}
//
extern "C" void ScaLBL_D3Q19_AAeven_Poisson_Potential_BC_Z(int *list, double *dist, double Vout, int count, int Np){
int GRID = count / 512 + 1;
dvc_ScaLBL_D3Q19_AAeven_Poisson_Potential_BC_Z<<<GRID,512>>>(list, dist, Vout, count, Np);
cudaError_t err = cudaGetLastError();
if (cudaSuccess != err){
printf("Hip error in ScaLBL_D3Q19_AAeven_Poisson_Potential_BC_Z (kernel): %s \n",cudaGetErrorString(err));
}
}
extern "C" void ScaLBL_D3Q19_AAodd_Poisson_Potential_BC_z(int *d_neighborList, int *list, double *dist, double Vin, int count,int Np) {
int GRID = count / 512 + 1;
dvc_ScaLBL_D3Q19_AAodd_Poisson_Potential_BC_z<<<GRID,512>>>(d_neighborList, list, dist, Vin, count, Np);
cudaError_t err = cudaGetLastError();
if (cudaSuccess != err){
printf("Hip error in ScaLBL_D3Q19_AAodd_Poisson_Potential_BC_z (kernel): %s \n",cudaGetErrorString(err));
}
}
//
extern "C" void ScaLBL_D3Q19_AAodd_Poisson_Potential_BC_Z(int *d_neighborList, int *list, double *dist, double Vout, int count, int Np) {
int GRID = count / 512 + 1;
dvc_ScaLBL_D3Q19_AAodd_Poisson_Potential_BC_Z<<<GRID,512>>>(d_neighborList, list, dist, Vout, count, Np);
cudaError_t err = cudaGetLastError();
if (cudaSuccess != err){
printf("Hip error in ScaLBL_D3Q19_AAodd_Poisson_Potential_BC_Z (kernel): %s \n",cudaGetErrorString(err));
}
}
extern "C" void ScaLBL_D3Q19_AAodd_Poisson(int *neighborList, int *Map, extern "C" void ScaLBL_D3Q19_AAodd_Poisson(int *neighborList, int *Map,
double *dist, double *Den_charge, double *dist, double *Den_charge,

4
models/IonModel.cpp
View File

@ -300,7 +300,8 @@ void ScaLBL_IonModel::ReadParams(string filename, vector<int> &num_iter) {
void ScaLBL_IonModel::ReadParams(string filename) { void ScaLBL_IonModel::ReadParams(string filename) {
//NOTE: the maximum iteration timesteps for ions are left unspecified //NOTE: the maximum iteration timesteps for ions are left unspecified
// it relies on the multiphys controller to compute the max timestep // it relies on the multiphys controller to compute the max timestep
USE_MEMBRANE = true; USE_MEMBRANE = false;
Restart = false;
// read the input database // read the input database
db = std::make_shared<Database>(filename); db = std::make_shared<Database>(filename);
domain_db = db->getDatabase("Domain"); domain_db = db->getDatabase("Domain");
@ -1549,7 +1550,6 @@ void ScaLBL_IonModel::RunMembrane(double *Velocity, double *ElectricField, doubl
IonMembrane->IonTransport(&fq[ic * Np * 7],&Ci[ic * Np]); IonMembrane->IonTransport(&fq[ic * Np * 7],&Ci[ic * Np]);
/* if (BoundaryConditionSolid == 1) { /* if (BoundaryConditionSolid == 1) {
//TODO IonSolid may also be species-dependent //TODO IonSolid may also be species-dependent
ScaLBL_Comm->SolidDirichletD3Q7(&fq[ic * Np * 7], IonSolid); ScaLBL_Comm->SolidDirichletD3Q7(&fq[ic * Np * 7], IonSolid);

254
models/PoissonSolver.cpp
View File

@ -86,7 +86,7 @@ void ScaLBL_Poisson::ReadParams(string filename){
} }
//'tolerance_method' can be {"MSE","MSE_max"} //'tolerance_method' can be {"MSE","MSE_max"}
tolerance_method = electric_db->getWithDefault<std::string>( "tolerance_method", "MSE" ); tolerance_method = electric_db->getWithDefault<std::string>( "tolerance_method", "MSE" );
lattice_scheme = electric_db->getWithDefault<std::string>( "lattice_scheme", "D3Q7" ); lattice_scheme = electric_db->getWithDefault<std::string>( "lattice_scheme", "D3Q19" );
if (electric_db->keyExists( "epsilonR" )){ if (electric_db->keyExists( "epsilonR" )){
epsilonR = electric_db->getScalar<double>( "epsilonR" ); epsilonR = electric_db->getScalar<double>( "epsilonR" );
} }
@ -726,13 +726,13 @@ void ScaLBL_Poisson::Run(double *ChargeDensity, bool UseSlippingVelBC, int times
// *************ODD TIMESTEP*************// // *************ODD TIMESTEP*************//
timestep++; timestep++;
SolveElectricPotentialAAodd(timestep_from_Study);//update electric potential SolveElectricPotentialAAodd(timestep_from_Study);//update electric potential
SolvePoissonAAodd(ChargeDensity, UseSlippingVelBC);//perform collision SolvePoissonAAodd(ChargeDensity, UseSlippingVelBC,timestep);//perform collision
ScaLBL_Comm->Barrier(); comm.barrier(); ScaLBL_Comm->Barrier(); comm.barrier();
// *************EVEN TIMESTEP*************// // *************EVEN TIMESTEP*************//
timestep++; timestep++;
SolveElectricPotentialAAeven(timestep_from_Study);//update electric potential SolveElectricPotentialAAeven(timestep_from_Study);//update electric potential
SolvePoissonAAeven(ChargeDensity, UseSlippingVelBC);//perform collision SolvePoissonAAeven(ChargeDensity, UseSlippingVelBC,timestep);//perform collision
ScaLBL_Comm->Barrier(); comm.barrier(); ScaLBL_Comm->Barrier(); comm.barrier();
//************************************************************************/ //************************************************************************/
@ -797,18 +797,17 @@ void ScaLBL_Poisson::Run(double *ChargeDensity, bool UseSlippingVelBC, int times
//************************************************************************/ //************************************************************************/
// *************ODD TIMESTEP*************// // *************ODD TIMESTEP*************//
timestep++; timestep++;
//SolveElectricPotentialAAodd(timestep_from_Study,ChargeDensity, UseSlippingVelBC);//update electric potential //SolveElectricPotentialAAodd(timestep_from_Study);//,ChargeDensity, UseSlippingVelBC);//update electric potential
SolvePoissonAAodd(ChargeDensity, UseSlippingVelBC);//perform collision SolvePoissonAAodd(ChargeDensity, UseSlippingVelBC,timestep);//perform collision
ScaLBL_Comm->Barrier(); comm.barrier(); ScaLBL_Comm->Barrier(); comm.barrier();
// *************EVEN TIMESTEP*************// // *************EVEN TIMESTEP*************//
timestep++; timestep++;
//SolveElectricPotentialAAeven(timestep_from_Study,ChargeDensity, UseSlippingVelBC);//update electric potential //SolveElectricPotentialAAeven(timestep_from_Study);//,ChargeDensity, UseSlippingVelBC);//update electric potential
SolvePoissonAAeven(ChargeDensity, UseSlippingVelBC);//perform collision SolvePoissonAAeven(ChargeDensity, UseSlippingVelBC,timestep);//perform collision
ScaLBL_Comm->Barrier(); comm.barrier(); ScaLBL_Comm->Barrier(); comm.barrier();
//************************************************************************/ //************************************************************************/
// Check convergence of steady-state solution // Check convergence of steady-state solution
if (timestep==2){ if (timestep==2){
//save electric potential for convergence check //save electric potential for convergence check
@ -863,6 +862,145 @@ void ScaLBL_Poisson::Run(double *ChargeDensity, bool UseSlippingVelBC, int times
} }
} }
void ScaLBL_Poisson::Run(double *ChargeDensity, DoubleArray MembraneDistance, bool UseSlippingVelBC, int timestep_from_Study){
double error = 1.0;
double threshold = 10000000.0;
bool SET_THRESHOLD = false;
if (electric_db->keyExists( "rescale_at_distance" )){
SET_THRESHOLD = true;
threshold = electric_db->getScalar<double>( "rescale_at_distance" );
}
if (BoundaryConditionInlet > 0) SET_THRESHOLD = false;
if (BoundaryConditionOutlet > 0) SET_THRESHOLD = false;
double *host_Error;
host_Error = new double [Np];
timestep=0;
auto t1 = std::chrono::system_clock::now();
while (timestep < timestepMax && error > tolerance) {
//************************************************************************/
// *************ODD TIMESTEP*************//
timestep++;
//SolveElectricPotentialAAodd(timestep_from_Study,ChargeDensity, UseSlippingVelBC);//update electric potential
SolvePoissonAAodd(ChargeDensity, UseSlippingVelBC,timestep);//perform collision
ScaLBL_Comm->Barrier(); comm.barrier();
// *************EVEN TIMESTEP*************//
timestep++;
//SolveElectricPotentialAAeven(timestep_from_Study,ChargeDensity, UseSlippingVelBC);//update electric potential
SolvePoissonAAeven(ChargeDensity, UseSlippingVelBC,timestep);//perform collision
ScaLBL_Comm->Barrier(); comm.barrier();
//************************************************************************/
// Check convergence of steady-state solution
if (timestep==2){
//save electric potential for convergence check
}
if (timestep%analysis_interval==0){
/* get the elecric potential */
ScaLBL_CopyToHost(Psi_host.data(),Psi,sizeof(double)*Nx*Ny*Nz);
if (rank==0) printf(" ... getting Poisson solver error \n");
double err = 0.0;
double max_error = 0.0;
ScaLBL_CopyToHost(host_Error,ResidualError,sizeof(double)*Np);
for (int idx=0; idx<Np; idx++){
err = host_Error[idx]*host_Error[idx];
if (err > max_error ){
max_error = err;
}
}
error=Dm->Comm.maxReduce(max_error);
if (error > tolerance && SET_THRESHOLD){
/* don't use this with an external BC */
// cpompute the far-field electric potential
double inside_local = 0.0;
double outside_local = 0.0;
double inside_count_local = 0.0;
double outside_count_local = 0.0;
/* global values */
double inside_global = 0.0;
double outside_global = 0.0;
double inside_count_global = 0.0;
double outside_count_global = 0.0;
for (int k=1; k<Nz; k++){
for (int j=1; j<Ny; j++){
for (int i=1; i<Nx; i++){
int n = k*Nx*Ny + j*Nx + i;
double distance = MembraneDistance(i,j,k);
if (distance > threshold && distance < (threshold + 1.0)){
outside_count_local += 1.0;
outside_local += Psi_host(n);
}
else if (distance < (-1.0)*threshold && distance > (-1.0)*(threshold + 1.0)){
inside_count_local += 1.0;
inside_local += Psi_host(n);
}
}
}
}
inside_count_global = Dm->Comm.sumReduce(inside_count_local);
outside_count_global = Dm->Comm.sumReduce(outside_count_local);
outside_global = Dm->Comm.sumReduce(outside_local);
inside_global = Dm->Comm.sumReduce(inside_local);
outside_global /= outside_count_global;
inside_global /= inside_count_global;
if (rank==0) printf(" Rescaling far-field electric potential to value (outside): %f \n",outside_global);
if (rank==0) printf(" Rescaling far-field electric potential to value (inside): %f \n",inside_global);
// rescale the far-field electric potential
for (int k=1; k<Nz; k++){
for (int j=1; j<Ny; j++){
for (int i=1; i<Nx; i++){
int n = k*Nx*Ny + j*Nx + i;
double distance = MembraneDistance(i,j,k);
if ( distance > (threshold + 1.0)){
Psi_host(n) = outside_global;
}
else if ( distance < (-1.0)*(threshold + 1.0)){
Psi_host(n) = inside_global;
}
}
}
}
ScaLBL_CopyToDevice(Psi,Psi_host.data(),sizeof(double)*Nx*Ny*Nz);
}
/* compute the eletric field */
//ScaLBL_D3Q19_Poisson_getElectricField(fq, ElectricField, tau, Np);
}
}
if (rank == 0)
printf("---------------------------------------------------------------"
"----\n");
// Compute the walltime per timestep
auto t2 = std::chrono::system_clock::now();
double cputime = std::chrono::duration<double>(t2 - t1).count() / timestep;
// Performance obtained from each node
double MLUPS = double(Np) / cputime / 1000000;
if (rank == 0)
printf("********************************************************\n");
if (rank == 0)
printf("CPU time = %f \n", cputime);
if (rank == 0)
printf("Lattice update rate (per core)= %f MLUPS \n", MLUPS);
MLUPS *= nprocs;
if (rank == 0)
printf("Lattice update rate (total)= %f MLUPS \n", MLUPS);
if (rank == 0)
printf("********************************************************\n");
delete [] host_Error;
//************************************************************************/
if(WriteLog==true){
getConvergenceLog(timestep,error);
}
}
void ScaLBL_Poisson::getConvergenceLog(int timestep,double error){ void ScaLBL_Poisson::getConvergenceLog(int timestep,double error){
if ( rank == 0 ) { if ( rank == 0 ) {
fprintf(TIMELOG,"%i %.5g\n",timestep,error); fprintf(TIMELOG,"%i %.5g\n",timestep,error);
@ -1008,7 +1146,7 @@ void ScaLBL_Poisson::SolveElectricPotentialAAeven(int timestep_from_Study){
} }
} }
void ScaLBL_Poisson::SolvePoissonAAodd(double *ChargeDensity, bool UseSlippingVelBC){ void ScaLBL_Poisson::SolvePoissonAAodd(double *ChargeDensity, bool UseSlippingVelBC, int timestep){
if (lattice_scheme.compare("D3Q7")==0){ if (lattice_scheme.compare("D3Q7")==0){
ScaLBL_D3Q7_AAodd_Poisson(NeighborList, dvcMap, fq, ChargeDensity, Psi, ElectricField, tau, epsilon_LB, UseSlippingVelBC, ScaLBL_Comm->FirstInterior(), ScaLBL_Comm->LastInterior(), Np); ScaLBL_D3Q7_AAodd_Poisson(NeighborList, dvcMap, fq, ChargeDensity, Psi, ElectricField, tau, epsilon_LB, UseSlippingVelBC, ScaLBL_Comm->FirstInterior(), ScaLBL_Comm->LastInterior(), Np);
@ -1029,6 +1167,30 @@ void ScaLBL_Poisson::SolvePoissonAAodd(double *ChargeDensity, bool UseSlippingVe
ScaLBL_D3Q19_AAodd_Poisson(NeighborList, dvcMap, fq, ChargeDensity, Psi, ElectricField, tau, epsilon_LB, UseSlippingVelBC, ScaLBL_Comm->FirstInterior(), ScaLBL_Comm->LastInterior(), Np); ScaLBL_D3Q19_AAodd_Poisson(NeighborList, dvcMap, fq, ChargeDensity, Psi, ElectricField, tau, epsilon_LB, UseSlippingVelBC, ScaLBL_Comm->FirstInterior(), ScaLBL_Comm->LastInterior(), Np);
//ScaLBL_Comm->RecvD3Q19AA(fq); //WRITE INTO OPPOSITE //ScaLBL_Comm->RecvD3Q19AA(fq); //WRITE INTO OPPOSITE
ScaLBL_Comm->RecvD3Q19AA(fq); //WRITE INTO OPPOSITE ScaLBL_Comm->RecvD3Q19AA(fq); //WRITE INTO OPPOSITE
// Set boundary conditions
if (BoundaryConditionInlet > 0){
switch (BoundaryConditionInlet){
case 1:
ScaLBL_Comm->D3Q7_Poisson_Potential_BC_z(NeighborList, fq, Vin, timestep);
break;
case 2:
Vin = getBoundaryVoltagefromPeriodicBC(Vin0,freqIn,PhaseShift_In,timestep);
ScaLBL_Comm->D3Q7_Poisson_Potential_BC_z(NeighborList, fq, Vin, timestep);
break;
}
}
if (BoundaryConditionOutlet > 0){
switch (BoundaryConditionOutlet){
case 1:
ScaLBL_Comm->D3Q7_Poisson_Potential_BC_Z(NeighborList, fq, Vout, timestep);
break;
case 2:
Vout = getBoundaryVoltagefromPeriodicBC(Vout0,freqOut,PhaseShift_Out,timestep);
ScaLBL_Comm->D3Q7_Poisson_Potential_BC_Z(NeighborList, fq, Vout, timestep);
break;
}
}
ScaLBL_D3Q19_AAodd_Poisson(NeighborList, dvcMap, fq, ChargeDensity, Psi, ElectricField, tau, epsilon_LB, UseSlippingVelBC, 0, ScaLBL_Comm->LastExterior(), Np); ScaLBL_D3Q19_AAodd_Poisson(NeighborList, dvcMap, fq, ChargeDensity, Psi, ElectricField, tau, epsilon_LB, UseSlippingVelBC, 0, ScaLBL_Comm->LastExterior(), Np);
ScaLBL_Comm->Barrier(); ScaLBL_Comm->Barrier();
@ -1038,7 +1200,7 @@ void ScaLBL_Poisson::SolvePoissonAAodd(double *ChargeDensity, bool UseSlippingVe
} }
} }
void ScaLBL_Poisson::SolvePoissonAAeven(double *ChargeDensity, bool UseSlippingVelBC){ void ScaLBL_Poisson::SolvePoissonAAeven(double *ChargeDensity, bool UseSlippingVelBC, int timestep){
if (lattice_scheme.compare("D3Q7")==0){ if (lattice_scheme.compare("D3Q7")==0){
ScaLBL_D3Q7_AAeven_Poisson(dvcMap, fq, ChargeDensity, Psi, ElectricField, tau, epsilon_LB, UseSlippingVelBC, ScaLBL_Comm->FirstInterior(), ScaLBL_Comm->LastInterior(), Np); ScaLBL_D3Q7_AAeven_Poisson(dvcMap, fq, ChargeDensity, Psi, ElectricField, tau, epsilon_LB, UseSlippingVelBC, ScaLBL_Comm->FirstInterior(), ScaLBL_Comm->LastInterior(), Np);
@ -1056,6 +1218,31 @@ void ScaLBL_Poisson::SolvePoissonAAeven(double *ChargeDensity, bool UseSlippingV
ScaLBL_D3Q19_AAeven_Poisson(dvcMap, fq, ChargeDensity, Psi, ElectricField, ResidualError, tau, epsilon_LB, UseSlippingVelBC, ScaLBL_Comm->FirstInterior(), ScaLBL_Comm->LastInterior(), Np); ScaLBL_D3Q19_AAeven_Poisson(dvcMap, fq, ChargeDensity, Psi, ElectricField, ResidualError, tau, epsilon_LB, UseSlippingVelBC, ScaLBL_Comm->FirstInterior(), ScaLBL_Comm->LastInterior(), Np);
ScaLBL_Comm->RecvD3Q19AA(fq); //WRITE INTO OPPOSITE ScaLBL_Comm->RecvD3Q19AA(fq); //WRITE INTO OPPOSITE
// ScaLBL_Comm->RecvD3Q19AA(fq); //WRITE INTO OPPOSITE // ScaLBL_Comm->RecvD3Q19AA(fq); //WRITE INTO OPPOSITE
// Set boundary conditions
if (BoundaryConditionInlet > 0){
switch (BoundaryConditionInlet){
case 1:
ScaLBL_Comm->D3Q7_Poisson_Potential_BC_z(NeighborList, fq, Vin, timestep);
break;
case 2:
Vin = getBoundaryVoltagefromPeriodicBC(Vin0,freqIn,PhaseShift_In,timestep);
ScaLBL_Comm->D3Q7_Poisson_Potential_BC_z(NeighborList, fq, Vin, timestep);
break;
}
}
if (BoundaryConditionOutlet > 0){
switch (BoundaryConditionOutlet){
case 1:
ScaLBL_Comm->D3Q7_Poisson_Potential_BC_Z(NeighborList, fq, Vout, timestep);
break;
case 2:
Vout = getBoundaryVoltagefromPeriodicBC(Vout0,freqOut,PhaseShift_Out,timestep);
ScaLBL_Comm->D3Q7_Poisson_Potential_BC_Z(NeighborList, fq, Vout, timestep);
break;
}
}
ScaLBL_D3Q19_AAeven_Poisson(dvcMap, fq, ChargeDensity, Psi, ElectricField, ResidualError, tau, epsilon_LB, UseSlippingVelBC, 0, ScaLBL_Comm->LastExterior(), Np); ScaLBL_D3Q19_AAeven_Poisson(dvcMap, fq, ChargeDensity, Psi, ElectricField, ResidualError, tau, epsilon_LB, UseSlippingVelBC, 0, ScaLBL_Comm->LastExterior(), Np);
ScaLBL_Comm->Barrier(); ScaLBL_Comm->Barrier();
@ -1182,6 +1369,53 @@ void ScaLBL_Poisson::ElectricField_LB_to_Phys(DoubleArray &Efield_reg){
} }
} }
void ScaLBL_Poisson::WriteVis( int timestep) {
auto vis_db = db->getDatabase("Visualization");
auto format = vis_db->getWithDefault<string>( "format", "hdf5" );
DoubleArray ElectricalPotential(Nx, Ny, Nz);
std::vector<IO::MeshDataStruct> visData;
fillHalo<double> fillData(Dm->Comm, Dm->rank_info,
{Dm->Nx - 2, Dm->Ny - 2, Dm->Nz - 2}, {1, 1, 1},
0, 1);
IO::initialize("",format,"false");
// Create the MeshDataStruct
visData.resize(1);
visData[0].meshName = "domain";
visData[0].mesh =
std::make_shared<IO::DomainMesh>(Dm->rank_info, Dm->Nx - 2, Dm->Ny - 2,
Dm->Nz - 2, Dm->Lx, Dm->Ly, Dm->Lz);
//electric potential
auto ElectricPotentialVar = std::make_shared<IO::Variable>();
//--------------------------------------------------------------------------------------------------------------------
//-------------------------------------Create Names for Variables------------------------------------------------------
if (vis_db->getWithDefault<bool>("save_electric_potential", true)) {
ElectricPotentialVar->name = "ElectricPotential";
ElectricPotentialVar->type = IO::VariableType::VolumeVariable;
ElectricPotentialVar->dim = 1;
ElectricPotentialVar->data.resize(Dm->Nx - 2, Dm->Ny - 2, Dm->Nz - 2);
visData[0].vars.push_back(ElectricPotentialVar);
}
//--------------------------------------------------------------------------------------------------------------------
//------------------------------------Save All Variables--------------------------------------------------------------
if (vis_db->getWithDefault<bool>("save_electric_potential", true)) {
ASSERT(visData[0].vars[0]->name == "ElectricPotential");
getElectricPotential(ElectricalPotential);
Array<double> &ElectricPotentialData = visData[0].vars[0]->data;
fillData.copy(ElectricalPotential, ElectricPotentialData);
}
if (vis_db->getWithDefault<bool>("write_silo", true))
IO::writeData(timestep, visData, Dm->Comm);
//--------------------------------------------------------------------------------------------------------------------
}
//void ScaLBL_Poisson::SolveElectricField(){ //void ScaLBL_Poisson::SolveElectricField(){
// ScaLBL_Comm_Regular->SendHalo(Psi); // ScaLBL_Comm_Regular->SendHalo(Psi);
// ScaLBL_D3Q7_Poisson_ElectricField(NeighborList, dvcMap, dvcID, Psi, ElectricField, BoundaryConditionSolid, // ScaLBL_D3Q7_Poisson_ElectricField(NeighborList, dvcMap, dvcID, Psi, ElectricField, BoundaryConditionSolid,

8
models/PoissonSolver.h
View File

@ -22,6 +22,7 @@
#ifndef ScaLBL_POISSON_INC #ifndef ScaLBL_POISSON_INC
#define ScaLBL_POISSON_INC #define ScaLBL_POISSON_INC
class ScaLBL_Poisson { class ScaLBL_Poisson {
public: public:
ScaLBL_Poisson(int RANK, int NP, const Utilities::MPI &COMM); ScaLBL_Poisson(int RANK, int NP, const Utilities::MPI &COMM);
@ -35,12 +36,15 @@ public:
void Create(); void Create();
void Initialize(double time_conv_from_Study); void Initialize(double time_conv_from_Study);
void Run(double *ChargeDensity, bool UseSlippingVelBC, int timestep_from_Study); void Run(double *ChargeDensity, bool UseSlippingVelBC, int timestep_from_Study);
void Run(double *ChargeDensity, DoubleArray MembraneDistance,
bool UseSlippingVelBC, int timestep_from_Study);
void getElectricPotential(DoubleArray &ReturnValues); void getElectricPotential(DoubleArray &ReturnValues);
void getElectricPotential_debug(int timestep); void getElectricPotential_debug(int timestep);
void getElectricField(DoubleArray &Values_x, DoubleArray &Values_y, void getElectricField(DoubleArray &Values_x, DoubleArray &Values_y,
DoubleArray &Values_z); DoubleArray &Values_z);
void getElectricField_debug(int timestep); void getElectricField_debug(int timestep);
void Checkpoint(); void Checkpoint();
void WriteVis( int timestep);
void DummyChargeDensity(); //for debugging void DummyChargeDensity(); //for debugging
@ -114,8 +118,8 @@ private:
void SolveElectricPotentialAAeven(int timestep_from_Study); void SolveElectricPotentialAAeven(int timestep_from_Study);
void SolveElectricPotentialAAodd(int timestep_from_Study); void SolveElectricPotentialAAodd(int timestep_from_Study);
//void SolveElectricField(); //void SolveElectricField();
void SolvePoissonAAodd(double *ChargeDensity, bool UseSlippingVelBC); void SolvePoissonAAodd(double *ChargeDensity, bool UseSlippingVelBC, int timestep);
void SolvePoissonAAeven(double *ChargeDensity, bool UseSlippingVelBC); void SolvePoissonAAeven(double *ChargeDensity, bool UseSlippingVelBC, int timestep);
void getConvergenceLog(int timestep,double error); void getConvergenceLog(int timestep,double error);
double getBoundaryVoltagefromPeriodicBC(double V0,double freq,double t0,int time_step); double getBoundaryVoltagefromPeriodicBC(double V0,double freq,double t0,int time_step);

1
sample_scripts/configure_arden
View File

@ -13,6 +13,7 @@ cmake -D CMAKE_C_COMPILER:PATH=/opt/arden/openmpi/3.1.2/bin/mpicc \
-D CUDA_HOST_COMPILER="/usr/bin/gcc" \ -D CUDA_HOST_COMPILER="/usr/bin/gcc" \
-D USE_HDF5=1 \ -D USE_HDF5=1 \
-D HDF5_DIRECTORY="/opt/arden/hdf5/1.8.12" \ -D HDF5_DIRECTORY="/opt/arden/hdf5/1.8.12" \
-D USE_DOXYGEN=true \
-D USE_CUDA=0 \ -D USE_CUDA=0 \
-D USE_TIMER=0 \ -D USE_TIMER=0 \
~/Programs/LBPM ~/Programs/LBPM

3
tests/TestPoissonSolver.cpp
View File

@ -76,11 +76,14 @@ int main(int argc, char **argv)
PoissonSolver.getElectricField_debug(timestep); PoissonSolver.getElectricField_debug(timestep);
} }
} }
PoissonSolver.WriteVis(timestep);
} }
else { else {
int timestep = 1;
PoissonSolver.Run(PoissonSolver.ChargeDensityDummy,false,1); PoissonSolver.Run(PoissonSolver.ChargeDensityDummy,false,1);
PoissonSolver.getElectricPotential_debug(1); PoissonSolver.getElectricPotential_debug(1);
PoissonSolver.getElectricField_debug(1); PoissonSolver.getElectricField_debug(1);
PoissonSolver.WriteVis(timestep);
} }
if (rank==0) printf("Maximum timestep is reached and the simulation is completed\n"); if (rank==0) printf("Maximum timestep is reached and the simulation is completed\n");

4
tests/lbpm_nernst_planck_cell_simulator.cpp
View File

@ -125,7 +125,7 @@ int main(int argc, char **argv)
while (timestep < Study.timestepMax){ while (timestep < Study.timestepMax){
timestep++; timestep++;
PoissonSolver.Run(IonModel.ChargeDensity,SlipBC,timestep);//solve Poisson equtaion to get steady-state electrical potental PoissonSolver.Run(IonModel.ChargeDensity,IonModel.MembraneDistance,SlipBC,timestep);//solve Poisson equtaion to get steady-state electrical potental
//comm.barrier(); //comm.barrier();
//if (rank == 0) printf(" Poisson step %i \n",timestep); //if (rank == 0) printf(" Poisson step %i \n",timestep);
//StokesModel.Run_Lite(IonModel.ChargeDensity, PoissonSolver.ElectricField);// Solve the N-S equations to get velocity //StokesModel.Run_Lite(IonModel.ChargeDensity, PoissonSolver.ElectricField);// Solve the N-S equations to get velocity
@ -133,7 +133,7 @@ int main(int argc, char **argv)
IonModel.RunMembrane(IonModel.FluidVelocityDummy,PoissonSolver.ElectricField,PoissonSolver.Psi); //solve for ion transport with membrane IonModel.RunMembrane(IonModel.FluidVelocityDummy,PoissonSolver.ElectricField,PoissonSolver.Psi); //solve for ion transport with membrane
//comm.barrier(); //comm.barrier();
if (rank == 0) printf(" Membrane step %i \n",timestep); //if (rank == 0) printf(" Membrane step %i \n",timestep);
fflush(stdout); fflush(stdout);

1
tests/lbpm_nernst_planck_simulator.cpp
View File

@ -104,7 +104,6 @@ int main(int argc, char **argv)
//StokesModel.Run_Lite(IonModel.ChargeDensity, PoissonSolver.ElectricField);// Solve the N-S equations to get velocity //StokesModel.Run_Lite(IonModel.ChargeDensity, PoissonSolver.ElectricField);// Solve the N-S equations to get velocity
IonModel.Run(IonModel.FluidVelocityDummy,PoissonSolver.ElectricField); //solve for ion transport and electric potential IonModel.Run(IonModel.FluidVelocityDummy,PoissonSolver.ElectricField); //solve for ion transport and electric potential
//if (timestep%Study.analysis_interval==0){ //if (timestep%Study.analysis_interval==0){
// Analysis.Basic(IonModel,PoissonSolver,StokesModel,timestep); // Analysis.Basic(IonModel,PoissonSolver,StokesModel,timestep);
//} //}