diff --git a/acc/mhd_solver/stencil_kernel.ac b/acc/mhd_solver/stencil_kernel.ac
index e0efa94..57097f8 100644
--- a/acc/mhd_solver/stencil_kernel.ac
+++ b/acc/mhd_solver/stencil_kernel.ac
@@ -20,9 +20,12 @@
 uniform int AC_max_steps;
 uniform int AC_save_steps;
 uniform int AC_bin_steps;
-uniform int AC_bc_type;
 uniform int AC_start_step;
 
+// In3 params 
+uniform int3 AC_bc_type_bot;
+uniform int3 AC_bc_type_top;
+
 // Real params
 uniform Scalar AC_dt;
 uniform Scalar AC_max_time;
diff --git a/include/astaroth.h b/include/astaroth.h
index 8db7608..ace32c0 100644
--- a/include/astaroth.h
+++ b/include/astaroth.h
@@ -304,7 +304,13 @@ AcResult acGridStoreMesh(const Stream stream, AcMesh* host_mesh);
 AcResult acGridIntegrate(const Stream stream, const AcReal dt);
 
 /** */
+/*   MV: Commented out for a while, but save for the future when standalone_MPI
+         works with periodic boundary conditions. 
+AcResult
+acGridIntegrateNonperiodic(const Stream stream, const AcReal dt)
+
 AcResult acGridIntegrateNonperiodic(const Stream stream, const AcReal dt);
+*/
 
 /** */
 AcResult acGridPeriodicBoundconds(const Stream stream);
diff --git a/src/core/device.cc b/src/core/device.cc
index 9f5a8e1..ad46148 100644
--- a/src/core/device.cc
+++ b/src/core/device.cc
@@ -1907,6 +1907,8 @@ acGridIntegrate(const Stream stream, const AcReal dt)
     return AC_SUCCESS;
 }
 
+/*   MV: Commented out for a while, but save for the future when standalone_MPI
+         works with periodic boundary conditions. 
 AcResult
 acGridIntegrateNonperiodic(const Stream stream, const AcReal dt)
 {
@@ -2130,6 +2132,8 @@ acGridIntegrateNonperiodic(const Stream stream, const AcReal dt)
     return AC_SUCCESS;
 }
 
+*/
+
 
 AcResult
 acGridPeriodicBoundconds(const Stream stream)
diff --git a/src/core/kernels/boundconds.cuh b/src/core/kernels/boundconds.cuh
index 4b35065..01e89fa 100644
--- a/src/core/kernels/boundconds.cuh
+++ b/src/core/kernels/boundconds.cuh
@@ -1,5 +1,9 @@
 #pragma once
 
+// Temporary defines untils we can figure out smarter swithches.
+#define BOUNDCOND_SYM  1
+#define BOUNDCOND_ASYM 2
+
 static __global__ void
 kernel_symmetric_boundconds(const int3 start, const int3 end, AcReal* vtxbuf, const int3 bindex, const int sign)
 {
@@ -134,19 +138,19 @@ 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 int bindex);
+                          AcReal* vtxbuf, const int3 bindex)
 {
     const dim3 tpb(8, 2, 8);
     const dim3 bpg((unsigned int)ceil((end.x - start.x) / (float)tpb.x),
                    (unsigned int)ceil((end.y - start.y) / (float)tpb.y),
                    (unsigned int)ceil((end.z - start.z) / (float)tpb.z));
 
-    if (DCONST(AC_bype) == BOUNDCOND_SYM) 
+    if (DCONST(AC_bc_type) == BOUNDCOND_SYM) 
     {
         kernel_symmetric_boundconds<<<bpg, tpb, 0, stream>>>(start, end, vtxbuf, bindex,  1);
         ERRCHK_CUDA_KERNEL();
     } 
-    else if (DCONST(AC_bype) == BOUNDCOND_ASYM) 
+    else if (DCONST(AC_bc_type) == BOUNDCOND_ASYM) 
     {
         kernel_symmetric_boundconds<<<bpg, tpb, 0, stream>>>(start, end, vtxbuf, bindex, -1);
         ERRCHK_CUDA_KERNEL();