cuda build succeeds

cuda/Ion.cu

@@ -6,6 +6,48 @@
 #define NTHREADS 256
 
 
+extern "C" void Membrane_D3Q19_Unpack(int q, int *list, int *links, int start, int linkCount,
+                                    double *recvbuf, double *dist, int N) {
+    //....................................................................................
+    // Unack distribution from the recv buffer
+    // Distribution q matche Cqx, Cqy, Cqz
+    // 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, link;
+    for (link=0; link<linkCount; link++){
+
+    	idx = links[start+link];
+        // Get the value from the list -- note that n is the index is from the send (non-local) process
+        n = list[start + idx];
+        // unpack the distribution to the proper location
+        if (!(n < 0))
+            dist[q * N + n] = recvbuf[start + idx];
+    }
+}
+
+extern "C" void Membrane_D3Q19_Transport(int q, int *list, int *links, double *coef, int start, int offset, 
+		int linkCount, double *recvbuf, double *dist, int N){
+    //....................................................................................
+    // Unack distribution from the recv buffer
+    // Distribution q matche Cqx, Cqy, Cqz
+    // 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, link;
+    double alpha;
+    for (link=offset; link<linkCount; link++){
+
+    	idx = list[start+link];
+        // Get the value from the list -- note that n is the index is from the send (non-local) process
+        n = list[start + idx];
+        alpha = coef[start + idx];
+        // unpack the distribution to the proper location
+        if (!(n < 0))
+            dist[q * N + n] = alpha*recvbuf[start + idx];
+    }
+}
+
 __global__  void dvc_ScaLBL_D3Q7_Membrane_AssignLinkCoef(int *membrane, int *Map, double *Distance, double *Psi, double *coef,
 		double Threshold, double MassFractionIn, double MassFractionOut, double ThresholdMassFractionIn, double ThresholdMassFractionOut,
 		int memLinks, int Nx, int Ny, int Nz, int Np){

hip/Ion.hip

@@ -5,6 +5,48 @@
 #define NBLOCKS 1024
 #define NTHREADS 256
 
+extern "C" void Membrane_D3Q19_Unpack(int q, int *list, int *links, int start, int linkCount,
+                                    double *recvbuf, double *dist, int N) {
+    //....................................................................................
+    // Unack distribution from the recv buffer
+    // Distribution q matche Cqx, Cqy, Cqz
+    // 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, link;
+    for (link=0; link<linkCount; link++){
+
+    	idx = links[start+link];
+        // Get the value from the list -- note that n is the index is from the send (non-local) process
+        n = list[start + idx];
+        // unpack the distribution to the proper location
+        if (!(n < 0))
+            dist[q * N + n] = recvbuf[start + idx];
+    }
+}
+
+extern "C" void Membrane_D3Q19_Transport(int q, int *list, int *links, double *coef, int start, int offset, 
+		int linkCount, double *recvbuf, double *dist, int N){
+    //....................................................................................
+    // Unack distribution from the recv buffer
+    // Distribution q matche Cqx, Cqy, Cqz
+    // 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, link;
+    double alpha;
+    for (link=offset; link<linkCount; link++){
+
+    	idx = list[start+link];
+        // Get the value from the list -- note that n is the index is from the send (non-local) process
+        n = list[start + idx];
+        alpha = coef[start + idx];
+        // unpack the distribution to the proper location
+        if (!(n < 0))
+            dist[q * N + n] = alpha*recvbuf[start + idx];
+    }
+}
+
 __global__  void dvc_ScaLBL_D3Q7_Membrane_AssignLinkCoef(int *membrane, int *Map, double *Distance, double *Psi, double *coef,
 		double Threshold, double MassFractionIn, double MassFractionOut, double ThresholdMassFractionIn, double ThresholdMassFractionOut,
 		int memLinks, int Nx, int Ny, int Nz, int Np){

tests/lbpm_cell_simulator.cpp

@@ -53,7 +53,9 @@ int main(int argc, char **argv)
         ScaLBL_IonModel IonModel(rank,nprocs,comm);
         ScaLBL_Poisson PoissonSolver(rank,nprocs,comm); 
         ScaLBL_Multiphys_Controller Study(rank,nprocs,comm);//multiphysics controller coordinating multi-model coupling
-        
+
+	bool SlipBC = false;
+	
         // Load controller information
         Study.ReadParams(filename);
 
@@ -94,7 +96,7 @@ int main(int argc, char **argv)
         while (timestep < Study.timestepMax){
             
             timestep++;
-            PoissonSolver.Run(IonModel.ChargeDensity,timestep);//solve Poisson equtaion to get steady-state electrical potental
+            PoissonSolver.Run(IonModel.ChargeDensity,SlipBC,timestep);//solve Poisson equtaion to get steady-state electrical potental
             StokesModel.Run_Lite(IonModel.ChargeDensity, PoissonSolver.ElectricField);// Solve the N-S equations to get velocity
             IonModel.RunMembrane(StokesModel.Velocity,PoissonSolver.ElectricField,PoissonSolver.Psi); //solve for ion transport with membrane