From 543c565e5d676914ff93998d0dd878a75e6ff84e Mon Sep 17 00:00:00 2001
From: Miikka Vaisala <mvaisala@asiaa.sinica.edu.tw>
Date: Mon, 23 Nov 2020 15:43:52 +0800
Subject: [PATCH] Dummy at the moment, but now the boundary condition kernel
 caller can see what vertex buffer name is in use.

---
 src/core/device.cc              |  2 +-
 src/core/kernels/boundconds.cuh | 55 ++++++++++++++++++++-------------
 src/core/kernels/kernels.h      |  3 +-
 3 files changed, 36 insertions(+), 24 deletions(-)

diff --git a/src/core/device.cc b/src/core/device.cc
index 9130f34..0d2127e 100644
--- a/src/core/device.cc
+++ b/src/core/device.cc
@@ -440,7 +440,7 @@ acDeviceGeneralBoundcondStep(const Device device, const Stream stream,
 {
     cudaSetDevice(device->id);
     return acKernelGeneralBoundconds(device->streams[stream], start, end,
-                                     device->vba.in[vtxbuf_handle], config, bindex);
+                                     device->vba.in[vtxbuf_handle], vtxbuf_handle, config, bindex);
 }
 
 AcResult
diff --git a/src/core/kernels/boundconds.cuh b/src/core/kernels/boundconds.cuh
index 541b22b..f0518c0 100644
--- a/src/core/kernels/boundconds.cuh
+++ b/src/core/kernels/boundconds.cuh
@@ -191,7 +191,8 @@ acKernelPeriodicBoundconds(const cudaStream_t stream, const int3 start, const in
 
 AcResult 
 acKernelGeneralBoundconds(const cudaStream_t stream, const int3 start, const int3 end,
-                          AcReal* vtxbuf, const AcMeshInfo config, const int3 bindex)
+                          AcReal* vtxbuf, const VertexBufferHandle vtxbuf_handle, 
+                          const AcMeshInfo config, const int3 bindex)   
 {
     const dim3 tpb(8, 2, 8);
     const dim3 bpg((unsigned int)ceil((end.x - start.x) / (float)tpb.x),
@@ -203,28 +204,38 @@ acKernelGeneralBoundconds(const cudaStream_t stream, const int3 start, const int
     int3 bc_bot = {config.int_params[AC_bc_type_bot_x], config.int_params[AC_bc_type_bot_y], 
                    config.int_params[AC_bc_type_bot_z]};
 
-    if (bc_top.x == AC_BOUNDCOND_SYMMETRIC) 
-    {
-        kernel_symmetric_boundconds<<<bpg, tpb, 0, stream>>>(start, end, vtxbuf, bindex,  1);
-        ERRCHK_CUDA_KERNEL();
-    } 
-    else if (bc_top.x == AC_BOUNDCOND_ANTISYMMETRIC) 
-    {
-        kernel_symmetric_boundconds<<<bpg, tpb, 0, stream>>>(start, end, vtxbuf, bindex, -1);
-        ERRCHK_CUDA_KERNEL();
-    } 
-    else if (bc_top.x == AC_BOUNDCOND_PERIODIC) 
-    {
-        kernel_periodic_boundconds<<<bpg, tpb, 0, stream>>>(start, end, vtxbuf);
-        ERRCHK_CUDA_KERNEL();
-    } 
-    else 
-    {
-        printf("ERROR: Boundary condition not recognized!\n");
-        printf("ERROR: bc_top = %i, %i, %i \n", bc_top.x, bc_top.y, bc_top.z);
-        printf("ERROR: bc_bot = %i, %i, %i \n", bc_bot.x, bc_bot.y, bc_bot.z);
+//#if AC_MPI_ENABLED
+//    printf( "WARNING : NON-PERIODIC BOUNDARY CONDITIONS NOT SUPPORTER BY MPI! Only working at node level.\n");
+//    return AC_FAILURE;
+//#endif
+
+    if ( vtxbuf_handle != -1) // This is a dummy to make swithing boundary condition with respect to   more possible later  
+    {
+
+        if (bc_top.x == AC_BOUNDCOND_SYMMETRIC) 
+        {
+            kernel_symmetric_boundconds<<<bpg, tpb, 0, stream>>>(start, end, vtxbuf, bindex,  1);
+            ERRCHK_CUDA_KERNEL();
+        } 
+        else if (bc_top.x == AC_BOUNDCOND_ANTISYMMETRIC) 
+        {
+            kernel_symmetric_boundconds<<<bpg, tpb, 0, stream>>>(start, end, vtxbuf, bindex, -1);
+            ERRCHK_CUDA_KERNEL();
+        } 
+        else if (bc_top.x == AC_BOUNDCOND_PERIODIC) 
+        {
+            kernel_periodic_boundconds<<<bpg, tpb, 0, stream>>>(start, end, vtxbuf);
+            ERRCHK_CUDA_KERNEL();
+        } 
+        else 
+        {
+            printf("ERROR: Boundary condition not recognized!\n");
+            printf("ERROR: bc_top = %i, %i, %i \n", bc_top.x, bc_top.y, bc_top.z);
+            printf("ERROR: bc_bot = %i, %i, %i \n", bc_bot.x, bc_bot.y, bc_bot.z);
+
+            return AC_FAILURE;
+        }
 
-        return AC_FAILURE;
     }
 
     return AC_SUCCESS;
diff --git a/src/core/kernels/kernels.h b/src/core/kernels/kernels.h
index bc94185..a9e05f9 100644
--- a/src/core/kernels/kernels.h
+++ b/src/core/kernels/kernels.h
@@ -45,7 +45,8 @@ AcResult acKernelPeriodicBoundconds(const cudaStream_t stream, const int3 start,
                                     AcReal* vtxbuf);
 /** */
 AcResult acKernelGeneralBoundconds(const cudaStream_t stream, const int3 start, const int3 end,
-                                    AcReal* vtxbuf, const AcMeshInfo config, const int3 bindex);
+                                   AcReal* vtxbuf, const VertexBufferHandle vtxbuf_handle,
+                                   const AcMeshInfo config, const int3 bindex);
 
 
 /** */