From 9e74f49812ac4fae1cdef888a76be1f4e7df4d53 Mon Sep 17 00:00:00 2001 From: James McClure Date: Sun, 8 May 2022 05:49:23 -0400 Subject: [PATCH] add new ion code for cuda --- cuda/Ion.cu | 258 ++++++++++++++++++++++++++++++++++++++++++++++++++++ 1 file changed, 258 insertions(+) diff --git a/cuda/Ion.cu b/cuda/Ion.cu index 1650449a..acbd0c9a 100644 --- a/cuda/Ion.cu +++ b/cuda/Ion.cu @@ -593,6 +593,264 @@ __global__ void dvc_ScaLBL_D3Q7_Ion_ChargeDensity(double *Den, double *ChargeDe } } } +__global__ void dvc_ScaLBL_D3Q7_AAodd_Ion_v0(int *neighborList, double *dist, + double *Den, double *FluxDiffusive, + double *FluxAdvective, + double *FluxElectrical, double *Velocity, + double *ElectricField, double Di, int zi, + double rlx, double Vt, int start, + int finish, int Np) { + int n; + double Ci; + double ux, uy, uz; + double uEPx, uEPy, uEPz; //electrochemical induced velocity + double Ex, Ey, Ez; //electrical field + double flux_diffusive_x, flux_diffusive_y, flux_diffusive_z; + double f0, f1, f2, f3, f4, f5, f6; + //double X,Y,Z,factor_x, factor_y, factor_z; + int nr1, nr2, nr3, nr4, nr5, nr6; + + int S = Np/NBLOCKS/NTHREADS + 1; + for (int s=0; s 10Np => odd part of dist) + f1 = dist[nr1]; // reading the f1 data into register fq + // q=2 + nr2 = neighborList[n + Np]; // neighbor 1 ( < 10Np => even part of dist) + f2 = dist[nr2]; // reading the f2 data into register fq + // q=3 + nr3 = neighborList[n + 2 * Np]; // neighbor 4 + f3 = dist[nr3]; + // q=4 + nr4 = neighborList[n + 3 * Np]; // neighbor 3 + f4 = dist[nr4]; + // q=5 + nr5 = neighborList[n + 4 * Np]; + f5 = dist[nr5]; + // q=6 + nr6 = neighborList[n + 5 * Np]; + f6 = dist[nr6]; + + // compute diffusive flux + //Ci = f0 + f1 + f2 + f3 + f4 + f5 + f6; + flux_diffusive_x = (1.0 - 0.5 * rlx) * ((f1 - f2) - ux * Ci); + flux_diffusive_y = (1.0 - 0.5 * rlx) * ((f3 - f4) - uy * Ci); + flux_diffusive_z = (1.0 - 0.5 * rlx) * ((f5 - f6) - uz * Ci); + FluxDiffusive[n + 0 * Np] = flux_diffusive_x; + FluxDiffusive[n + 1 * Np] = flux_diffusive_y; + FluxDiffusive[n + 2 * Np] = flux_diffusive_z; + FluxAdvective[n + 0 * Np] = ux * Ci; + FluxAdvective[n + 1 * Np] = uy * Ci; + FluxAdvective[n + 2 * Np] = uz * Ci; + FluxElectrical[n + 0 * Np] = uEPx * Ci; + FluxElectrical[n + 1 * Np] = uEPy * Ci; + FluxElectrical[n + 2 * Np] = uEPz * Ci; + + //Den[n] = Ci; + + /* use logistic function to prevent negative distributions*/ + //X = 4.0 * (ux + uEPx); + //Y = 4.0 * (uy + uEPy); + //Z = 4.0 * (uz + uEPz); + //factor_x = X / sqrt(1 + X*X); + //factor_y = Y / sqrt(1 + Y*Y); + //factor_z = Z / sqrt(1 + Z*Z); + + // q=0 + dist[n] = f0 * (1.0 - rlx) + rlx * 0.25 * Ci; + + // q = 1 + dist[nr2] = + f1 * (1.0 - rlx) + rlx * 0.125 * Ci * (1.0 + 4.0 * (ux + uEPx)); + // f1 * (1.0 - rlx) + rlx * 0.125 * Ci * (1.0 + factor_x); + + + // q=2 + dist[nr1] = + f2 * (1.0 - rlx) + rlx * 0.125 * Ci * (1.0 - 4.0 * (ux + uEPx)); + // f2 * (1.0 - rlx) + rlx * 0.125 * Ci * (1.0 - factor_x); + + // q = 3 + dist[nr4] = + f3 * (1.0 - rlx) + rlx * 0.125 * Ci * (1.0 + 4.0 * (uy + uEPy)); + // f3 * (1.0 - rlx) + rlx * 0.125 * Ci * (1.0 + factor_y ); + + // q = 4 + dist[nr3] = + f4 * (1.0 - rlx) + rlx * 0.125 * Ci * (1.0 - 4.0 * (uy + uEPy)); + // f4 * (1.0 - rlx) + rlx * 0.125 * Ci * (1.0 - factor_y); + + // q = 5 + dist[nr6] = + f5 * (1.0 - rlx) + rlx * 0.125 * Ci * (1.0 + 4.0 * (uz + uEPz)); + // f5 * (1.0 - rlx) + rlx * 0.125 * Ci * (1.0 + factor_z); + + // q = 6 + dist[nr5] = + f6 * (1.0 - rlx) + rlx * 0.125 * Ci * (1.0 - 4.0 * (uz + uEPz)); + // f6 * (1.0 - rlx) + rlx * 0.125 * Ci * (1.0 - factor_z); + + } + } +} + +__global__ void dvc_ScaLBL_D3Q7_AAeven_Ion_v0( + double *dist, double *Den, double *FluxDiffusive, double *FluxAdvective, + double *FluxElectrical, double *Velocity, double *ElectricField, double Di, + int zi, double rlx, double Vt, int start, int finish, int Np) { + int n; + double Ci; + double ux, uy, uz; + double uEPx, uEPy, uEPz; //electrochemical induced velocity + double Ex, Ey, Ez; //electrical field + double flux_diffusive_x, flux_diffusive_y, flux_diffusive_z; + double f0, f1, f2, f3, f4, f5, f6; + //double X,Y,Z, factor_x, factor_y, factor_z; + + int S = Np/NBLOCKS/NTHREADS + 1; + for (int s=0; s>>(dist, + Den, FluxDiffusive, FluxAdvective, + FluxElectrical, Velocity, + ElectricField, Di, zi, + rlx, Vt, start, finish, Np); + + cudaError_t err = cudaGetLastError(); + if (cudaSuccess != err){ + printf("cuda error in dvc_ScaLBL_D3Q7_AAeven_Ion_v0: %s \n",cudaGetErrorString(err)); + } +} + +extern "C" void ScaLBL_D3Q7_AAodd_Ion_v0(int *neighborList, double *dist, + double *Den, double *FluxDiffusive, + double *FluxAdvective, + double *FluxElectrical, double *Velocity, + double *ElectricField, double Di, int zi, + double rlx, double Vt, int start, + int finish, int Np) { + + dvc_ScaLBL_D3Q7_AAodd_Ion_v0<<>>(neighborList, dist, + Den, FluxDiffusive, FluxAdvective, + FluxElectrical, Velocity, + ElectricField, Di, zi, + rlx, Vt, start, + finish, Np); + + cudaError_t err = cudaGetLastError(); + if (cudaSuccess != err){ + printf("cuda error in dvc_ScaLBL_D3Q7_AAodd_Ion_v0: %s \n",cudaGetErrorString(err)); + } +} + extern "C" void ScaLBL_D3Q7_AAodd_IonConcentration(int *neighborList, double *dist, double *Den, int start, int finish, int Np){