prolong3 改为先算实际 stencil 窗口；只有窗口触及对称边界时才走全域 symmetry_bd，否则只复制必需窗口。restrict3 同样改成窗口判定，无触边时仅填 ii/jj/kk 必需窗口。

perf(restrict3): shrink X-pass ii sweep to required overlap window
- compute fi_min/fi_max from output i-range and derive ii_lo/ii_hi - replace full ii sweep (-1:extf(1)) with windowed sweep in Z/Y precompute passes - keep stencil math unchanged; add bounds sanity check for ii window
2026-03-02 18:10:38 +08:00 · 2026-03-02 18:10:38 +08:00 · 2026-03-02 18:10:35 +08:00 · 2026-03-02 14:11:52 +08:00 · 2026-03-02 14:08:27 +08:00 · 2026-03-02 14:08:13 +08:00
28 changed files with 3917 additions and 1161 deletions
--- a/AMSS_NCKU_source/MPatch.C
+++ b/AMSS_NCKU_source/MPatch.C
@@ -7,6 +7,7 @@
 #include <string>
 #include <cmath>
 #include <new>
+#include <vector>
 using namespace std;

 #include "misc.h"
@@ -17,6 +18,168 @@ using namespace std;
 #include "interp_lb_profile.h"
 #endif

+namespace
+{
+struct InterpBlockView
+{
+  Block *bp;
+  double llb[dim];
+  double uub[dim];
+};
+
+struct BlockBinIndex
+{
+  int bins[dim];
+  double lo[dim];
+  double inv[dim];
+  vector<InterpBlockView> views;
+  vector<vector<int>> bin_to_blocks;
+  bool valid;
+
+  BlockBinIndex() : valid(false)
+  {
+    for (int i = 0; i < dim; i++)
+    {
+      bins[i] = 1;
+      lo[i] = 0.0;
+      inv[i] = 0.0;
+    }
+  }
+};
+
+inline int clamp_int(int v, int lo, int hi)
+{
+  return (v < lo) ? lo : ((v > hi) ? hi : v);
+}
+
+inline int coord_to_bin(double x, double lo, double inv, int nb)
+{
+  if (nb <= 1 || inv <= 0.0)
+    return 0;
+  int b = int(floor((x - lo) * inv));
+  return clamp_int(b, 0, nb - 1);
+}
+
+inline int bin_loc(const BlockBinIndex &index, int b0, int b1, int b2)
+{
+  return b0 + index.bins[0] * (b1 + index.bins[1] * b2);
+}
+
+inline bool point_in_block_view(const InterpBlockView &view, const double *pox, const double *DH)
+{
+  for (int i = 0; i < dim; i++)
+  {
+    if (pox[i] - view.llb[i] < -DH[i] / 2 || pox[i] - view.uub[i] > DH[i] / 2)
+      return false;
+  }
+  return true;
+}
+
+void build_block_bin_index(Patch *patch, const double *DH, BlockBinIndex &index)
+{
+  index = BlockBinIndex();
+
+  MyList<Block> *Bp = patch->blb;
+  while (Bp)
+  {
+    Block *BP = Bp->data;
+    InterpBlockView view;
+    view.bp = BP;
+    for (int i = 0; i < dim; i++)
+    {
+#ifdef Vertex
+#ifdef Cell
+#error Both Cell and Vertex are defined
+#endif
+      view.llb[i] = (feq(BP->bbox[i], patch->bbox[i], DH[i] / 2)) ? BP->bbox[i] + patch->lli[i] * DH[i] : BP->bbox[i] + (ghost_width - 0.5) * DH[i];
+      view.uub[i] = (feq(BP->bbox[dim + i], patch->bbox[dim + i], DH[i] / 2)) ? BP->bbox[dim + i] - patch->uui[i] * DH[i] : BP->bbox[dim + i] - (ghost_width - 0.5) * DH[i];
+#else
+#ifdef Cell
+      view.llb[i] = (feq(BP->bbox[i], patch->bbox[i], DH[i] / 2)) ? BP->bbox[i] + patch->lli[i] * DH[i] : BP->bbox[i] + ghost_width * DH[i];
+      view.uub[i] = (feq(BP->bbox[dim + i], patch->bbox[dim + i], DH[i] / 2)) ? BP->bbox[dim + i] - patch->uui[i] * DH[i] : BP->bbox[dim + i] - ghost_width * DH[i];
+#else
+#error Not define Vertex nor Cell
+#endif
+#endif
+    }
+    index.views.push_back(view);
+    if (Bp == patch->ble)
+      break;
+    Bp = Bp->next;
+  }
+
+  const int nblocks = int(index.views.size());
+  if (nblocks <= 0)
+    return;
+
+  int bins_1d = int(ceil(pow(double(nblocks), 1.0 / 3.0)));
+  bins_1d = clamp_int(bins_1d, 1, 32);
+  for (int i = 0; i < dim; i++)
+  {
+    index.bins[i] = bins_1d;
+    index.lo[i] = patch->bbox[i] + patch->lli[i] * DH[i];
+    const double hi = patch->bbox[dim + i] - patch->uui[i] * DH[i];
+    if (hi > index.lo[i] && bins_1d > 1)
+      index.inv[i] = bins_1d / (hi - index.lo[i]);
+    else
+      index.inv[i] = 0.0;
+  }
+
+  index.bin_to_blocks.resize(index.bins[0] * index.bins[1] * index.bins[2]);
+
+  for (int bi = 0; bi < nblocks; bi++)
+  {
+    const InterpBlockView &view = index.views[bi];
+    int bmin[dim], bmax[dim];
+    for (int d = 0; d < dim; d++)
+    {
+      const double low = view.llb[d] - DH[d] / 2;
+      const double up = view.uub[d] + DH[d] / 2;
+      bmin[d] = coord_to_bin(low, index.lo[d], index.inv[d], index.bins[d]);
+      bmax[d] = coord_to_bin(up, index.lo[d], index.inv[d], index.bins[d]);
+      if (bmax[d] < bmin[d])
+      {
+        int t = bmin[d];
+        bmin[d] = bmax[d];
+        bmax[d] = t;
+      }
+    }
+
+    for (int bz = bmin[2]; bz <= bmax[2]; bz++)
+      for (int by = bmin[1]; by <= bmax[1]; by++)
+        for (int bx = bmin[0]; bx <= bmax[0]; bx++)
+          index.bin_to_blocks[bin_loc(index, bx, by, bz)].push_back(bi);
+  }
+
+  index.valid = true;
+}
+
+int find_block_index_for_point(const BlockBinIndex &index, const double *pox, const double *DH)
+{
+  if (!index.valid)
+    return -1;
+
+  const int bx = coord_to_bin(pox[0], index.lo[0], index.inv[0], index.bins[0]);
+  const int by = coord_to_bin(pox[1], index.lo[1], index.inv[1], index.bins[1]);
+  const int bz = coord_to_bin(pox[2], index.lo[2], index.inv[2], index.bins[2]);
+  const vector<int> &cand = index.bin_to_blocks[bin_loc(index, bx, by, bz)];
+
+  for (size_t ci = 0; ci < cand.size(); ci++)
+  {
+    const int bi = cand[ci];
+    if (point_in_block_view(index.views[bi], pox, DH))
+      return bi;
+  }
+
+  // Fallback to full scan for numerical edge cases around bin boundaries.
+  for (size_t bi = 0; bi < index.views.size(); bi++)
+    if (point_in_block_view(index.views[bi], pox, DH))
+      return int(bi);
+
+  return -1;
+}
+} // namespace
+
 Patch::Patch(int DIM, int *shapei, double *bboxi, int levi, bool buflog, int Symmetry) : lev(levi)
 {

@@ -367,9 +530,11 @@ void Patch::Interp_Points(MyList<var> *VarList,
  for (int j = 0; j < NN; j++)
    owner_rank[j] = -1;

-  double DH[dim], llb[dim], uub[dim];
+  double DH[dim];
  for (int i = 0; i < dim; i++)
    DH[i] = getdX(i);
+  BlockBinIndex block_index;
+  build_block_bin_index(this, DH, block_index);

  for (int j = 0; j < NN; j++) // run along points
  {
@@ -392,57 +557,24 @@ void Patch::Interp_Points(MyList<var> *VarList,
      }
    }

-    MyList<Block> *Bp = blb;
-    bool notfind = true;
-    while (notfind && Bp) // run along Blocks
+    const int block_i = find_block_index_for_point(block_index, pox, DH);
+    if (block_i >= 0)
    {
-      Block *BP = Bp->data;
-
-      bool flag = true;
-      for (int i = 0; i < dim; i++)
+      Block *BP = block_index.views[block_i].bp;
+      owner_rank[j] = BP->rank;
+      if (myrank == BP->rank)
      {
-#ifdef Vertex
-#ifdef Cell
-#error Both Cell and Vertex are defined
-#endif
-        llb[i] = (feq(BP->bbox[i], bbox[i], DH[i] / 2)) ? BP->bbox[i] + lli[i] * DH[i] : BP->bbox[i] + (ghost_width - 0.5) * DH[i];
-        uub[i] = (feq(BP->bbox[dim + i], bbox[dim + i], DH[i] / 2)) ? BP->bbox[dim + i] - uui[i] * DH[i] : BP->bbox[dim + i] - (ghost_width - 0.5) * DH[i];
-#else
-#ifdef Cell
-        llb[i] = (feq(BP->bbox[i], bbox[i], DH[i] / 2)) ? BP->bbox[i] + lli[i] * DH[i] : BP->bbox[i] + ghost_width * DH[i];
-        uub[i] = (feq(BP->bbox[dim + i], bbox[dim + i], DH[i] / 2)) ? BP->bbox[dim + i] - uui[i] * DH[i] : BP->bbox[dim + i] - ghost_width * DH[i];
-#else
-#error Not define Vertex nor Cell
-#endif
-#endif
-        if (XX[i][j] - llb[i] < -DH[i] / 2 || XX[i][j] - uub[i] > DH[i] / 2)
+        //---> interpolation
+        varl = VarList;
+        int k = 0;
+        while (varl) // run along variables
        {
-          flag = false;
-          break;
+          f_global_interp(BP->shape, BP->X[0], BP->X[1], BP->X[2], BP->fgfs[varl->data->sgfn], Shellf[j * num_var + k],
+                          pox[0], pox[1], pox[2], ordn, varl->data->SoA, Symmetry);
+          varl = varl->next;
+          k++;
        }
      }
-
-      if (flag)
-      {
-        notfind = false;
-        owner_rank[j] = BP->rank;
-        if (myrank == BP->rank)
-        {
-          //---> interpolation
-          varl = VarList;
-          int k = 0;
-          while (varl) // run along variables
-          {
-            f_global_interp(BP->shape, BP->X[0], BP->X[1], BP->X[2], BP->fgfs[varl->data->sgfn], Shellf[j * num_var + k],
-                            pox[0], pox[1], pox[2], ordn, varl->data->SoA, Symmetry);
-            varl = varl->next;
-            k++;
-          }
-        }
-      }
-      if (Bp == ble)
-        break;
-      Bp = Bp->next;
    }
  }

@@ -535,9 +667,11 @@ void Patch::Interp_Points(MyList<var> *VarList,
  for (int j = 0; j < NN; j++)
    owner_rank[j] = -1;

-  double DH[dim], llb[dim], uub[dim];
+  double DH[dim];
  for (int i = 0; i < dim; i++)
    DH[i] = getdX(i);
+  BlockBinIndex block_index;
+  build_block_bin_index(this, DH, block_index);

  // --- Interpolation phase (identical to original) ---
  for (int j = 0; j < NN; j++)
@@ -561,56 +695,23 @@ void Patch::Interp_Points(MyList<var> *VarList,
      }
    }

-    MyList<Block> *Bp = blb;
-    bool notfind = true;
-    while (notfind && Bp)
+    const int block_i = find_block_index_for_point(block_index, pox, DH);
+    if (block_i >= 0)
    {
-      Block *BP = Bp->data;
-
-      bool flag = true;
-      for (int i = 0; i < dim; i++)
+      Block *BP = block_index.views[block_i].bp;
+      owner_rank[j] = BP->rank;
+      if (myrank == BP->rank)
      {
-#ifdef Vertex
-#ifdef Cell
-#error Both Cell and Vertex are defined
-#endif
-        llb[i] = (feq(BP->bbox[i], bbox[i], DH[i] / 2)) ? BP->bbox[i] + lli[i] * DH[i] : BP->bbox[i] + (ghost_width - 0.5) * DH[i];
-        uub[i] = (feq(BP->bbox[dim + i], bbox[dim + i], DH[i] / 2)) ? BP->bbox[dim + i] - uui[i] * DH[i] : BP->bbox[dim + i] - (ghost_width - 0.5) * DH[i];
-#else
-#ifdef Cell
-        llb[i] = (feq(BP->bbox[i], bbox[i], DH[i] / 2)) ? BP->bbox[i] + lli[i] * DH[i] : BP->bbox[i] + ghost_width * DH[i];
-        uub[i] = (feq(BP->bbox[dim + i], bbox[dim + i], DH[i] / 2)) ? BP->bbox[dim + i] - uui[i] * DH[i] : BP->bbox[dim + i] - ghost_width * DH[i];
-#else
-#error Not define Vertex nor Cell
-#endif
-#endif
-        if (XX[i][j] - llb[i] < -DH[i] / 2 || XX[i][j] - uub[i] > DH[i] / 2)
+        varl = VarList;
+        int k = 0;
+        while (varl)
        {
-          flag = false;
-          break;
+          f_global_interp(BP->shape, BP->X[0], BP->X[1], BP->X[2], BP->fgfs[varl->data->sgfn], Shellf[j * num_var + k],
+                          pox[0], pox[1], pox[2], ordn, varl->data->SoA, Symmetry);
+          varl = varl->next;
+          k++;
        }
      }
-
-      if (flag)
-      {
-        notfind = false;
-        owner_rank[j] = BP->rank;
-        if (myrank == BP->rank)
-        {
-          varl = VarList;
-          int k = 0;
-          while (varl)
-          {
-            f_global_interp(BP->shape, BP->X[0], BP->X[1], BP->X[2], BP->fgfs[varl->data->sgfn], Shellf[j * num_var + k],
-                            pox[0], pox[1], pox[2], ordn, varl->data->SoA, Symmetry);
-            varl = varl->next;
-            k++;
-          }
-        }
-      }
-      if (Bp == ble)
-        break;
-      Bp = Bp->next;
    }
  }

@@ -833,9 +934,11 @@ void Patch::Interp_Points(MyList<var> *VarList,
  MPI_Comm_group(MPI_COMM_WORLD, &world_group);
  MPI_Comm_group(Comm_here, &local_group);

-  double DH[dim], llb[dim], uub[dim];
+  double DH[dim];
  for (int i = 0; i < dim; i++)
    DH[i] = getdX(i);
+  BlockBinIndex block_index;
+  build_block_bin_index(this, DH, block_index);

  for (int j = 0; j < NN; j++) // run along points
  {
@@ -858,57 +961,24 @@ void Patch::Interp_Points(MyList<var> *VarList,
      }
    }

-    MyList<Block> *Bp = blb;
-    bool notfind = true;
-    while (notfind && Bp) // run along Blocks
+    const int block_i = find_block_index_for_point(block_index, pox, DH);
+    if (block_i >= 0)
    {
-      Block *BP = Bp->data;
-
-      bool flag = true;
-      for (int i = 0; i < dim; i++)
+      Block *BP = block_index.views[block_i].bp;
+      owner_rank[j] = BP->rank;
+      if (myrank == BP->rank)
      {
-#ifdef Vertex
-#ifdef Cell
-#error Both Cell and Vertex are defined
-#endif
-        llb[i] = (feq(BP->bbox[i], bbox[i], DH[i] / 2)) ? BP->bbox[i] + lli[i] * DH[i] : BP->bbox[i] + (ghost_width - 0.5) * DH[i];
-        uub[i] = (feq(BP->bbox[dim + i], bbox[dim + i], DH[i] / 2)) ? BP->bbox[dim + i] - uui[i] * DH[i] : BP->bbox[dim + i] - (ghost_width - 0.5) * DH[i];
-#else
-#ifdef Cell
-        llb[i] = (feq(BP->bbox[i], bbox[i], DH[i] / 2)) ? BP->bbox[i] + lli[i] * DH[i] : BP->bbox[i] + ghost_width * DH[i];
-        uub[i] = (feq(BP->bbox[dim + i], bbox[dim + i], DH[i] / 2)) ? BP->bbox[dim + i] - uui[i] * DH[i] : BP->bbox[dim + i] - ghost_width * DH[i];
-#else
-#error Not define Vertex nor Cell
-#endif
-#endif
-        if (XX[i][j] - llb[i] < -DH[i] / 2 || XX[i][j] - uub[i] > DH[i] / 2)
+        //---> interpolation
+        varl = VarList;
+        int k = 0;
+        while (varl) // run along variables
        {
-          flag = false;
-          break;
+          f_global_interp(BP->shape, BP->X[0], BP->X[1], BP->X[2], BP->fgfs[varl->data->sgfn], Shellf[j * num_var + k],
+                          pox[0], pox[1], pox[2], ordn, varl->data->SoA, Symmetry);
+          varl = varl->next;
+          k++;
        }
      }
-
-      if (flag)
-      {
-        notfind = false;
-        owner_rank[j] = BP->rank;
-        if (myrank == BP->rank)
-        {
-          //---> interpolation
-          varl = VarList;
-          int k = 0;
-          while (varl) // run along variables
-          {
-            f_global_interp(BP->shape, BP->X[0], BP->X[1], BP->X[2], BP->fgfs[varl->data->sgfn], Shellf[j * num_var + k],
-                            pox[0], pox[1], pox[2], ordn, varl->data->SoA, Symmetry);
-            varl = varl->next;
-            k++;
-          }
-        }
-      }
-      if (Bp == ble)
-        break;
-      Bp = Bp->next;
    }
  }

--- a/AMSS_NCKU_source/Parallel.C
+++ b/AMSS_NCKU_source/Parallel.C
@@ -3883,175 +3883,263 @@ int Parallel::data_packermix(double *data, MyList<Parallel::gridseg> *src, MyLis
  return size_out;
 }
 //
-void Parallel::transfer(MyList<Parallel::gridseg> **src, MyList<Parallel::gridseg> **dst,
-                        MyList<var> *VarList1 /* source */, MyList<var> *VarList2 /*target */,
-                        int Symmetry)
-{
-  int myrank, cpusize;
-  MPI_Comm_size(MPI_COMM_WORLD, &cpusize);
-  MPI_Comm_rank(MPI_COMM_WORLD, &myrank);
-
-  int node;
-
-  MPI_Request *reqs;
-  MPI_Status *stats;
-  reqs = new MPI_Request[2 * cpusize];
-  stats = new MPI_Status[2 * cpusize];
-  int req_no = 0;
-
-  double **send_data, **rec_data;
-  send_data = new double *[cpusize];
-  rec_data = new double *[cpusize];
-  int length;
-
-  for (node = 0; node < cpusize; node++)
-  {
-    send_data[node] = rec_data[node] = 0;
-    if (node == myrank)
-    {
-      if (length = data_packer(0, src[myrank], dst[myrank], node, PACK, VarList1, VarList2, Symmetry))
-      {
-        rec_data[node] = new double[length];
-        if (!rec_data[node])
-        {
-          cout << "out of memory when new in short transfer, place 1" << endl;
-          MPI_Abort(MPI_COMM_WORLD, 1);
-        }
-        data_packer(rec_data[node], src[myrank], dst[myrank], node, PACK, VarList1, VarList2, Symmetry);
-      }
-    }
-    else
-    {
-      // send from this cpu to cpu#node
-      if (length = data_packer(0, src[myrank], dst[myrank], node, PACK, VarList1, VarList2, Symmetry))
-      {
-        send_data[node] = new double[length];
-        if (!send_data[node])
-        {
-          cout << "out of memory when new in short transfer, place 2" << endl;
-          MPI_Abort(MPI_COMM_WORLD, 1);
-        }
-        data_packer(send_data[node], src[myrank], dst[myrank], node, PACK, VarList1, VarList2, Symmetry);
-        MPI_Isend((void *)send_data[node], length, MPI_DOUBLE, node, 1, MPI_COMM_WORLD, reqs + req_no++);
-      }
-      // receive from cpu#node to this cpu
-      if (length = data_packer(0, src[node], dst[node], node, UNPACK, VarList1, VarList2, Symmetry))
-      {
-        rec_data[node] = new double[length];
-        if (!rec_data[node])
-        {
-          cout << "out of memory when new in short transfer, place 3" << endl;
-          MPI_Abort(MPI_COMM_WORLD, 1);
-        }
-        MPI_Irecv((void *)rec_data[node], length, MPI_DOUBLE, node, 1, MPI_COMM_WORLD, reqs + req_no++);
-      }
-    }
-  }
-  // wait for all requests to complete
-  MPI_Waitall(req_no, reqs, stats);
-
-  for (node = 0; node < cpusize; node++)
-    if (rec_data[node])
-      data_packer(rec_data[node], src[node], dst[node], node, UNPACK, VarList1, VarList2, Symmetry);
-
-  for (node = 0; node < cpusize; node++)
-  {
-    if (send_data[node])
-      delete[] send_data[node];
-    if (rec_data[node])
-      delete[] rec_data[node];
-  }
-
-  delete[] reqs;
-  delete[] stats;
-  delete[] send_data;
-  delete[] rec_data;
-}
+void Parallel::transfer(MyList<Parallel::gridseg> **src, MyList<Parallel::gridseg> **dst,
+                        MyList<var> *VarList1 /* source */, MyList<var> *VarList2 /*target */,
+                        int Symmetry)
+{
+  int myrank, cpusize;
+  MPI_Comm_size(MPI_COMM_WORLD, &cpusize);
+  MPI_Comm_rank(MPI_COMM_WORLD, &myrank);
+
+  int node;
+
+  MPI_Request *reqs = new MPI_Request[2 * cpusize];
+  MPI_Status *stats = new MPI_Status[2 * cpusize];
+  int *req_node = new int[2 * cpusize];
+  int *req_is_recv = new int[2 * cpusize];
+  int *completed = new int[2 * cpusize];
+  int req_no = 0;
+  int pending_recv = 0;
+
+  double **send_data = new double *[cpusize];
+  double **rec_data = new double *[cpusize];
+  int *send_lengths = new int[cpusize];
+  int *recv_lengths = new int[cpusize];
+
+  for (node = 0; node < cpusize; node++)
+  {
+    send_data[node] = rec_data[node] = 0;
+    send_lengths[node] = recv_lengths[node] = 0;
+  }
+
+  // Post receives first so peers can progress rendezvous early.
+  for (node = 0; node < cpusize; node++)
+  {
+    if (node == myrank) continue;
+
+    recv_lengths[node] = data_packer(0, src[node], dst[node], node, UNPACK, VarList1, VarList2, Symmetry);
+    if (recv_lengths[node] > 0)
+    {
+      rec_data[node] = new double[recv_lengths[node]];
+      if (!rec_data[node])
+      {
+        cout << "out of memory when new in short transfer, place 1" << endl;
+        MPI_Abort(MPI_COMM_WORLD, 1);
+      }
+      MPI_Irecv((void *)rec_data[node], recv_lengths[node], MPI_DOUBLE, node, 1, MPI_COMM_WORLD, reqs + req_no);
+      req_node[req_no] = node;
+      req_is_recv[req_no] = 1;
+      req_no++;
+      pending_recv++;
+    }
+  }
+
+  // Local transfer on this rank.
+  recv_lengths[myrank] = data_packer(0, src[myrank], dst[myrank], myrank, PACK, VarList1, VarList2, Symmetry);
+  if (recv_lengths[myrank] > 0)
+  {
+    rec_data[myrank] = new double[recv_lengths[myrank]];
+    if (!rec_data[myrank])
+    {
+      cout << "out of memory when new in short transfer, place 2" << endl;
+      MPI_Abort(MPI_COMM_WORLD, 1);
+    }
+    data_packer(rec_data[myrank], src[myrank], dst[myrank], myrank, PACK, VarList1, VarList2, Symmetry);
+  }
+
+  // Pack and post sends.
+  for (node = 0; node < cpusize; node++)
+  {
+    if (node == myrank) continue;
+
+    send_lengths[node] = data_packer(0, src[myrank], dst[myrank], node, PACK, VarList1, VarList2, Symmetry);
+    if (send_lengths[node] > 0)
+    {
+      send_data[node] = new double[send_lengths[node]];
+      if (!send_data[node])
+      {
+        cout << "out of memory when new in short transfer, place 3" << endl;
+        MPI_Abort(MPI_COMM_WORLD, 1);
+      }
+      data_packer(send_data[node], src[myrank], dst[myrank], node, PACK, VarList1, VarList2, Symmetry);
+      MPI_Isend((void *)send_data[node], send_lengths[node], MPI_DOUBLE, node, 1, MPI_COMM_WORLD, reqs + req_no);
+      req_node[req_no] = node;
+      req_is_recv[req_no] = 0;
+      req_no++;
+    }
+  }
+
+  // Unpack as soon as receive completes to reduce pure wait time.
+  while (pending_recv > 0)
+  {
+    int outcount = 0;
+    MPI_Waitsome(req_no, reqs, &outcount, completed, stats);
+    if (outcount == MPI_UNDEFINED) break;
+
+    for (int i = 0; i < outcount; i++)
+    {
+      int idx = completed[i];
+      if (idx >= 0 && req_is_recv[idx])
+      {
+        int recv_node = req_node[idx];
+        data_packer(rec_data[recv_node], src[recv_node], dst[recv_node], recv_node, UNPACK, VarList1, VarList2, Symmetry);
+        pending_recv--;
+      }
+    }
+  }
+
+  if (req_no > 0) MPI_Waitall(req_no, reqs, stats);
+
+  if (rec_data[myrank])
+    data_packer(rec_data[myrank], src[myrank], dst[myrank], myrank, UNPACK, VarList1, VarList2, Symmetry);
+
+  for (node = 0; node < cpusize; node++)
+  {
+    if (send_data[node])
+      delete[] send_data[node];
+    if (rec_data[node])
+      delete[] rec_data[node];
+  }
+
+  delete[] reqs;
+  delete[] stats;
+  delete[] req_node;
+  delete[] req_is_recv;
+  delete[] completed;
+  delete[] send_data;
+  delete[] rec_data;
+  delete[] send_lengths;
+  delete[] recv_lengths;
+}
 //
-void Parallel::transfermix(MyList<Parallel::gridseg> **src, MyList<Parallel::gridseg> **dst,
-                           MyList<var> *VarList1 /* source */, MyList<var> *VarList2 /*target */,
-                           int Symmetry)
-{
-  int myrank, cpusize;
-  MPI_Comm_size(MPI_COMM_WORLD, &cpusize);
-  MPI_Comm_rank(MPI_COMM_WORLD, &myrank);
-
-  int node;
-
-  MPI_Request *reqs;
-  MPI_Status *stats;
-  reqs = new MPI_Request[2 * cpusize];
-  stats = new MPI_Status[2 * cpusize];
-  int req_no = 0;
-
-  double **send_data, **rec_data;
-  send_data = new double *[cpusize];
-  rec_data = new double *[cpusize];
-  int length;
-
-  for (node = 0; node < cpusize; node++)
-  {
-    send_data[node] = rec_data[node] = 0;
-    if (node == myrank)
-    {
-      if (length = data_packermix(0, src[myrank], dst[myrank], node, PACK, VarList1, VarList2, Symmetry))
-      {
-        rec_data[node] = new double[length];
-        if (!rec_data[node])
-        {
-          cout << "out of memory when new in short transfer, place 1" << endl;
-          MPI_Abort(MPI_COMM_WORLD, 1);
-        }
-        data_packermix(rec_data[node], src[myrank], dst[myrank], node, PACK, VarList1, VarList2, Symmetry);
-      }
-    }
-    else
-    {
-      // send from this cpu to cpu#node
-      if (length = data_packermix(0, src[myrank], dst[myrank], node, PACK, VarList1, VarList2, Symmetry))
-      {
-        send_data[node] = new double[length];
-        if (!send_data[node])
-        {
-          cout << "out of memory when new in short transfer, place 2" << endl;
-          MPI_Abort(MPI_COMM_WORLD, 1);
-        }
-        data_packermix(send_data[node], src[myrank], dst[myrank], node, PACK, VarList1, VarList2, Symmetry);
-        MPI_Isend((void *)send_data[node], length, MPI_DOUBLE, node, 1, MPI_COMM_WORLD, reqs + req_no++);
-      }
-      // receive from cpu#node to this cpu
-      if (length = data_packermix(0, src[node], dst[node], node, UNPACK, VarList1, VarList2, Symmetry))
-      {
-        rec_data[node] = new double[length];
-        if (!rec_data[node])
-        {
-          cout << "out of memory when new in short transfer, place 3" << endl;
-          MPI_Abort(MPI_COMM_WORLD, 1);
-        }
-        MPI_Irecv((void *)rec_data[node], length, MPI_DOUBLE, node, 1, MPI_COMM_WORLD, reqs + req_no++);
-      }
-    }
-  }
-  // wait for all requests to complete
-  MPI_Waitall(req_no, reqs, stats);
-
-  for (node = 0; node < cpusize; node++)
-    if (rec_data[node])
-      data_packermix(rec_data[node], src[node], dst[node], node, UNPACK, VarList1, VarList2, Symmetry);
-
-  for (node = 0; node < cpusize; node++)
-  {
-    if (send_data[node])
-      delete[] send_data[node];
-    if (rec_data[node])
-      delete[] rec_data[node];
-  }
-
-  delete[] reqs;
-  delete[] stats;
-  delete[] send_data;
-  delete[] rec_data;
-}
+void Parallel::transfermix(MyList<Parallel::gridseg> **src, MyList<Parallel::gridseg> **dst,
+                           MyList<var> *VarList1 /* source */, MyList<var> *VarList2 /*target */,
+                           int Symmetry)
+{
+  int myrank, cpusize;
+  MPI_Comm_size(MPI_COMM_WORLD, &cpusize);
+  MPI_Comm_rank(MPI_COMM_WORLD, &myrank);
+
+  int node;
+
+  MPI_Request *reqs = new MPI_Request[2 * cpusize];
+  MPI_Status *stats = new MPI_Status[2 * cpusize];
+  int *req_node = new int[2 * cpusize];
+  int *req_is_recv = new int[2 * cpusize];
+  int *completed = new int[2 * cpusize];
+  int req_no = 0;
+  int pending_recv = 0;
+
+  double **send_data = new double *[cpusize];
+  double **rec_data = new double *[cpusize];
+  int *send_lengths = new int[cpusize];
+  int *recv_lengths = new int[cpusize];
+
+  for (node = 0; node < cpusize; node++)
+  {
+    send_data[node] = rec_data[node] = 0;
+    send_lengths[node] = recv_lengths[node] = 0;
+  }
+
+  // Post receives first so peers can progress rendezvous early.
+  for (node = 0; node < cpusize; node++)
+  {
+    if (node == myrank) continue;
+
+    recv_lengths[node] = data_packermix(0, src[node], dst[node], node, UNPACK, VarList1, VarList2, Symmetry);
+    if (recv_lengths[node] > 0)
+    {
+      rec_data[node] = new double[recv_lengths[node]];
+      if (!rec_data[node])
+      {
+        cout << "out of memory when new in short transfer, place 1" << endl;
+        MPI_Abort(MPI_COMM_WORLD, 1);
+      }
+      MPI_Irecv((void *)rec_data[node], recv_lengths[node], MPI_DOUBLE, node, 1, MPI_COMM_WORLD, reqs + req_no);
+      req_node[req_no] = node;
+      req_is_recv[req_no] = 1;
+      req_no++;
+      pending_recv++;
+    }
+  }
+
+  // Local transfer on this rank.
+  recv_lengths[myrank] = data_packermix(0, src[myrank], dst[myrank], myrank, PACK, VarList1, VarList2, Symmetry);
+  if (recv_lengths[myrank] > 0)
+  {
+    rec_data[myrank] = new double[recv_lengths[myrank]];
+    if (!rec_data[myrank])
+    {
+      cout << "out of memory when new in short transfer, place 2" << endl;
+      MPI_Abort(MPI_COMM_WORLD, 1);
+    }
+    data_packermix(rec_data[myrank], src[myrank], dst[myrank], myrank, PACK, VarList1, VarList2, Symmetry);
+  }
+
+  // Pack and post sends.
+  for (node = 0; node < cpusize; node++)
+  {
+    if (node == myrank) continue;
+
+    send_lengths[node] = data_packermix(0, src[myrank], dst[myrank], node, PACK, VarList1, VarList2, Symmetry);
+    if (send_lengths[node] > 0)
+    {
+      send_data[node] = new double[send_lengths[node]];
+      if (!send_data[node])
+      {
+        cout << "out of memory when new in short transfer, place 3" << endl;
+        MPI_Abort(MPI_COMM_WORLD, 1);
+      }
+      data_packermix(send_data[node], src[myrank], dst[myrank], node, PACK, VarList1, VarList2, Symmetry);
+      MPI_Isend((void *)send_data[node], send_lengths[node], MPI_DOUBLE, node, 1, MPI_COMM_WORLD, reqs + req_no);
+      req_node[req_no] = node;
+      req_is_recv[req_no] = 0;
+      req_no++;
+    }
+  }
+
+  // Unpack as soon as receive completes to reduce pure wait time.
+  while (pending_recv > 0)
+  {
+    int outcount = 0;
+    MPI_Waitsome(req_no, reqs, &outcount, completed, stats);
+    if (outcount == MPI_UNDEFINED) break;
+
+    for (int i = 0; i < outcount; i++)
+    {
+      int idx = completed[i];
+      if (idx >= 0 && req_is_recv[idx])
+      {
+        int recv_node = req_node[idx];
+        data_packermix(rec_data[recv_node], src[recv_node], dst[recv_node], recv_node, UNPACK, VarList1, VarList2, Symmetry);
+        pending_recv--;
+      }
+    }
+  }
+
+  if (req_no > 0) MPI_Waitall(req_no, reqs, stats);
+
+  if (rec_data[myrank])
+    data_packermix(rec_data[myrank], src[myrank], dst[myrank], myrank, UNPACK, VarList1, VarList2, Symmetry);
+
+  for (node = 0; node < cpusize; node++)
+  {
+    if (send_data[node])
+      delete[] send_data[node];
+    if (rec_data[node])
+      delete[] rec_data[node];
+  }
+
+  delete[] reqs;
+  delete[] stats;
+  delete[] req_node;
+  delete[] req_is_recv;
+  delete[] completed;
+  delete[] send_data;
+  delete[] rec_data;
+  delete[] send_lengths;
+  delete[] recv_lengths;
+}
 void Parallel::Sync(Patch *Pat, MyList<var> *VarList, int Symmetry)
 {
  int cpusize;
@@ -4279,73 +4367,110 @@ void Parallel::SyncCache::destroy()
  cpusize = 0; max_reqs = 0;
 }
 // transfer_cached: reuse pre-allocated buffers from SyncCache
-void Parallel::transfer_cached(MyList<Parallel::gridseg> **src, MyList<Parallel::gridseg> **dst,
-                               MyList<var> *VarList1, MyList<var> *VarList2,
-                               int Symmetry, SyncCache &cache)
-{
-  int myrank;
+void Parallel::transfer_cached(MyList<Parallel::gridseg> **src, MyList<Parallel::gridseg> **dst,
+                               MyList<var> *VarList1, MyList<var> *VarList2,
+                               int Symmetry, SyncCache &cache)
+{
+  int myrank;
  MPI_Comm_size(MPI_COMM_WORLD, &cache.cpusize);
  MPI_Comm_rank(MPI_COMM_WORLD, &myrank);
  int cpusize = cache.cpusize;

-  int req_no = 0;
-  int node;
-
-  for (node = 0; node < cpusize; node++)
-  {
-    if (node == myrank)
-    {
-      int length = data_packer(0, src[myrank], dst[myrank], node, PACK, VarList1, VarList2, Symmetry);
-      cache.recv_lengths[node] = length;
-      if (length > 0)
-      {
-        if (length > cache.recv_buf_caps[node])
-        {
-          if (cache.recv_bufs[node]) delete[] cache.recv_bufs[node];
-          cache.recv_bufs[node] = new double[length];
-          cache.recv_buf_caps[node] = length;
-        }
-        data_packer(cache.recv_bufs[node], src[myrank], dst[myrank], node, PACK, VarList1, VarList2, Symmetry);
-      }
-    }
-    else
-    {
-      // send
-      int slength = data_packer(0, src[myrank], dst[myrank], node, PACK, VarList1, VarList2, Symmetry);
-      cache.send_lengths[node] = slength;
-      if (slength > 0)
-      {
-        if (slength > cache.send_buf_caps[node])
-        {
-          if (cache.send_bufs[node]) delete[] cache.send_bufs[node];
-          cache.send_bufs[node] = new double[slength];
-          cache.send_buf_caps[node] = slength;
-        }
-        data_packer(cache.send_bufs[node], src[myrank], dst[myrank], node, PACK, VarList1, VarList2, Symmetry);
-        MPI_Isend((void *)cache.send_bufs[node], slength, MPI_DOUBLE, node, 1, MPI_COMM_WORLD, cache.reqs + req_no++);
-      }
-      // recv
-      int rlength = data_packer(0, src[node], dst[node], node, UNPACK, VarList1, VarList2, Symmetry);
-      cache.recv_lengths[node] = rlength;
-      if (rlength > 0)
-      {
-        if (rlength > cache.recv_buf_caps[node])
-        {
-          if (cache.recv_bufs[node]) delete[] cache.recv_bufs[node];
-          cache.recv_bufs[node] = new double[rlength];
-          cache.recv_buf_caps[node] = rlength;
-        }
-        MPI_Irecv((void *)cache.recv_bufs[node], rlength, MPI_DOUBLE, node, 1, MPI_COMM_WORLD, cache.reqs + req_no++);
-      }
-    }
-  }
-
-  MPI_Waitall(req_no, cache.reqs, cache.stats);
-
-  for (node = 0; node < cpusize; node++)
-    if (cache.recv_bufs[node] && cache.recv_lengths[node] > 0)
-      data_packer(cache.recv_bufs[node], src[node], dst[node], node, UNPACK, VarList1, VarList2, Symmetry);
-}
+  int req_no = 0;
+  int pending_recv = 0;
+  int node;
+  int *req_node = new int[cache.max_reqs];
+  int *req_is_recv = new int[cache.max_reqs];
+  int *completed = new int[cache.max_reqs];
+
+  // Post receives first so peers can progress rendezvous early.
+  for (node = 0; node < cpusize; node++)
+  {
+    if (node == myrank) continue;
+
+    int rlength = data_packer(0, src[node], dst[node], node, UNPACK, VarList1, VarList2, Symmetry);
+    cache.recv_lengths[node] = rlength;
+    if (rlength > 0)
+    {
+      if (rlength > cache.recv_buf_caps[node])
+      {
+        if (cache.recv_bufs[node]) delete[] cache.recv_bufs[node];
+        cache.recv_bufs[node] = new double[rlength];
+        cache.recv_buf_caps[node] = rlength;
+      }
+      MPI_Irecv((void *)cache.recv_bufs[node], rlength, MPI_DOUBLE, node, 1, MPI_COMM_WORLD, cache.reqs + req_no);
+      req_node[req_no] = node;
+      req_is_recv[req_no] = 1;
+      req_no++;
+      pending_recv++;
+    }
+  }
+
+  // Local transfer on this rank.
+  int self_len = data_packer(0, src[myrank], dst[myrank], myrank, PACK, VarList1, VarList2, Symmetry);
+  cache.recv_lengths[myrank] = self_len;
+  if (self_len > 0)
+  {
+    if (self_len > cache.recv_buf_caps[myrank])
+    {
+      if (cache.recv_bufs[myrank]) delete[] cache.recv_bufs[myrank];
+      cache.recv_bufs[myrank] = new double[self_len];
+      cache.recv_buf_caps[myrank] = self_len;
+    }
+    data_packer(cache.recv_bufs[myrank], src[myrank], dst[myrank], myrank, PACK, VarList1, VarList2, Symmetry);
+  }
+
+  // Pack and post sends.
+  for (node = 0; node < cpusize; node++)
+  {
+    if (node == myrank) continue;
+
+    int slength = data_packer(0, src[myrank], dst[myrank], node, PACK, VarList1, VarList2, Symmetry);
+    cache.send_lengths[node] = slength;
+    if (slength > 0)
+    {
+      if (slength > cache.send_buf_caps[node])
+      {
+        if (cache.send_bufs[node]) delete[] cache.send_bufs[node];
+        cache.send_bufs[node] = new double[slength];
+        cache.send_buf_caps[node] = slength;
+      }
+      data_packer(cache.send_bufs[node], src[myrank], dst[myrank], node, PACK, VarList1, VarList2, Symmetry);
+      MPI_Isend((void *)cache.send_bufs[node], slength, MPI_DOUBLE, node, 1, MPI_COMM_WORLD, cache.reqs + req_no);
+      req_node[req_no] = node;
+      req_is_recv[req_no] = 0;
+      req_no++;
+    }
+  }
+
+  // Unpack as soon as receive completes to reduce pure wait time.
+  while (pending_recv > 0)
+  {
+    int outcount = 0;
+    MPI_Waitsome(req_no, cache.reqs, &outcount, completed, cache.stats);
+    if (outcount == MPI_UNDEFINED) break;
+
+    for (int i = 0; i < outcount; i++)
+    {
+      int idx = completed[i];
+      if (idx >= 0 && req_is_recv[idx])
+      {
+        int recv_node_i = req_node[idx];
+        data_packer(cache.recv_bufs[recv_node_i], src[recv_node_i], dst[recv_node_i], recv_node_i, UNPACK, VarList1, VarList2, Symmetry);
+        pending_recv--;
+      }
+    }
+  }
+
+  if (req_no > 0) MPI_Waitall(req_no, cache.reqs, cache.stats);
+
+  if (self_len > 0)
+    data_packer(cache.recv_bufs[myrank], src[myrank], dst[myrank], myrank, UNPACK, VarList1, VarList2, Symmetry);
+
+  delete[] req_node;
+  delete[] req_is_recv;
+  delete[] completed;
+}
 // Sync_cached: build grid segment lists on first call, reuse on subsequent calls
 void Parallel::Sync_cached(MyList<Patch> *PatL, MyList<var> *VarList, int Symmetry, SyncCache &cache)
 {
@@ -5758,9 +5883,9 @@ void Parallel::OutBdLow2Hi_cached(MyList<Patch> *PatcL, MyList<Patch> *PatfL,
 }

 // OutBdLow2Himix_cached: same as OutBdLow2Hi_cached but uses transfermix for unpacking
-void Parallel::OutBdLow2Himix_cached(MyList<Patch> *PatcL, MyList<Patch> *PatfL,
-                                     MyList<var> *VarList1, MyList<var> *VarList2,
-                                     int Symmetry, SyncCache &cache)
+void Parallel::OutBdLow2Himix_cached(MyList<Patch> *PatcL, MyList<Patch> *PatfL,
+                                     MyList<var> *VarList1, MyList<var> *VarList2,
+                                     int Symmetry, SyncCache &cache)
 {
  if (!cache.valid)
  {
@@ -5806,60 +5931,100 @@ void Parallel::OutBdLow2Himix_cached(MyList<Patch> *PatcL, MyList<Patch> *PatfL,
  MPI_Comm_rank(MPI_COMM_WORLD, &myrank);
  int cpusize = cache.cpusize;

-  int req_no = 0;
-  for (int node = 0; node < cpusize; node++)
-  {
-    if (node == myrank)
-    {
-      int length = data_packermix(0, cache.combined_src[myrank], cache.combined_dst[myrank], node, PACK, VarList1, VarList2, Symmetry);
-      cache.recv_lengths[node] = length;
-      if (length > 0)
-      {
-        if (length > cache.recv_buf_caps[node])
-        {
-          if (cache.recv_bufs[node]) delete[] cache.recv_bufs[node];
-          cache.recv_bufs[node] = new double[length];
-          cache.recv_buf_caps[node] = length;
-        }
-        data_packermix(cache.recv_bufs[node], cache.combined_src[myrank], cache.combined_dst[myrank], node, PACK, VarList1, VarList2, Symmetry);
-      }
-    }
-    else
-    {
-      int slength = data_packermix(0, cache.combined_src[myrank], cache.combined_dst[myrank], node, PACK, VarList1, VarList2, Symmetry);
-      cache.send_lengths[node] = slength;
-      if (slength > 0)
-      {
-        if (slength > cache.send_buf_caps[node])
-        {
-          if (cache.send_bufs[node]) delete[] cache.send_bufs[node];
-          cache.send_bufs[node] = new double[slength];
-          cache.send_buf_caps[node] = slength;
-        }
-        data_packermix(cache.send_bufs[node], cache.combined_src[myrank], cache.combined_dst[myrank], node, PACK, VarList1, VarList2, Symmetry);
-        MPI_Isend((void *)cache.send_bufs[node], slength, MPI_DOUBLE, node, 1, MPI_COMM_WORLD, cache.reqs + req_no++);
-      }
-      int rlength = data_packermix(0, cache.combined_src[node], cache.combined_dst[node], node, UNPACK, VarList1, VarList2, Symmetry);
-      cache.recv_lengths[node] = rlength;
-      if (rlength > 0)
-      {
-        if (rlength > cache.recv_buf_caps[node])
-        {
-          if (cache.recv_bufs[node]) delete[] cache.recv_bufs[node];
-          cache.recv_bufs[node] = new double[rlength];
-          cache.recv_buf_caps[node] = rlength;
-        }
-        MPI_Irecv((void *)cache.recv_bufs[node], rlength, MPI_DOUBLE, node, 1, MPI_COMM_WORLD, cache.reqs + req_no++);
-      }
-    }
-  }
-
-  MPI_Waitall(req_no, cache.reqs, cache.stats);
-
-  for (int node = 0; node < cpusize; node++)
-    if (cache.recv_bufs[node] && cache.recv_lengths[node] > 0)
-      data_packermix(cache.recv_bufs[node], cache.combined_src[node], cache.combined_dst[node], node, UNPACK, VarList1, VarList2, Symmetry);
-}
+  int req_no = 0;
+  int pending_recv = 0;
+  int *req_node = new int[cache.max_reqs];
+  int *req_is_recv = new int[cache.max_reqs];
+  int *completed = new int[cache.max_reqs];
+
+  // Post receives first so peers can progress rendezvous early.
+  for (int node = 0; node < cpusize; node++)
+  {
+    if (node == myrank) continue;
+
+    int rlength = data_packermix(0, cache.combined_src[node], cache.combined_dst[node], node, UNPACK, VarList1, VarList2, Symmetry);
+    cache.recv_lengths[node] = rlength;
+    if (rlength > 0)
+    {
+      if (rlength > cache.recv_buf_caps[node])
+      {
+        if (cache.recv_bufs[node]) delete[] cache.recv_bufs[node];
+        cache.recv_bufs[node] = new double[rlength];
+        cache.recv_buf_caps[node] = rlength;
+      }
+      MPI_Irecv((void *)cache.recv_bufs[node], rlength, MPI_DOUBLE, node, 1, MPI_COMM_WORLD, cache.reqs + req_no);
+      req_node[req_no] = node;
+      req_is_recv[req_no] = 1;
+      req_no++;
+      pending_recv++;
+    }
+  }
+
+  // Local transfer on this rank.
+  int self_len = data_packermix(0, cache.combined_src[myrank], cache.combined_dst[myrank], myrank, PACK, VarList1, VarList2, Symmetry);
+  cache.recv_lengths[myrank] = self_len;
+  if (self_len > 0)
+  {
+    if (self_len > cache.recv_buf_caps[myrank])
+    {
+      if (cache.recv_bufs[myrank]) delete[] cache.recv_bufs[myrank];
+      cache.recv_bufs[myrank] = new double[self_len];
+      cache.recv_buf_caps[myrank] = self_len;
+    }
+    data_packermix(cache.recv_bufs[myrank], cache.combined_src[myrank], cache.combined_dst[myrank], myrank, PACK, VarList1, VarList2, Symmetry);
+  }
+
+  // Pack and post sends.
+  for (int node = 0; node < cpusize; node++)
+  {
+    if (node == myrank) continue;
+
+    int slength = data_packermix(0, cache.combined_src[myrank], cache.combined_dst[myrank], node, PACK, VarList1, VarList2, Symmetry);
+    cache.send_lengths[node] = slength;
+    if (slength > 0)
+    {
+      if (slength > cache.send_buf_caps[node])
+      {
+        if (cache.send_bufs[node]) delete[] cache.send_bufs[node];
+        cache.send_bufs[node] = new double[slength];
+        cache.send_buf_caps[node] = slength;
+      }
+      data_packermix(cache.send_bufs[node], cache.combined_src[myrank], cache.combined_dst[myrank], node, PACK, VarList1, VarList2, Symmetry);
+      MPI_Isend((void *)cache.send_bufs[node], slength, MPI_DOUBLE, node, 1, MPI_COMM_WORLD, cache.reqs + req_no);
+      req_node[req_no] = node;
+      req_is_recv[req_no] = 0;
+      req_no++;
+    }
+  }
+
+  // Unpack as soon as receive completes to reduce pure wait time.
+  while (pending_recv > 0)
+  {
+    int outcount = 0;
+    MPI_Waitsome(req_no, cache.reqs, &outcount, completed, cache.stats);
+    if (outcount == MPI_UNDEFINED) break;
+
+    for (int i = 0; i < outcount; i++)
+    {
+      int idx = completed[i];
+      if (idx >= 0 && req_is_recv[idx])
+      {
+        int recv_node_i = req_node[idx];
+        data_packermix(cache.recv_bufs[recv_node_i], cache.combined_src[recv_node_i], cache.combined_dst[recv_node_i], recv_node_i, UNPACK, VarList1, VarList2, Symmetry);
+        pending_recv--;
+      }
+    }
+  }
+
+  if (req_no > 0) MPI_Waitall(req_no, cache.reqs, cache.stats);
+
+  if (self_len > 0)
+    data_packermix(cache.recv_bufs[myrank], cache.combined_src[myrank], cache.combined_dst[myrank], myrank, UNPACK, VarList1, VarList2, Symmetry);
+
+  delete[] req_node;
+  delete[] req_is_recv;
+  delete[] completed;
+}

 // collect all buffer grid segments or blocks for given patch
 MyList<Parallel::gridseg> *Parallel::build_buffer_gsl(Patch *Pat)
--- a/AMSS_NCKU_source/bssn_rhs_c.C
+++ b/AMSS_NCKU_source/bssn_rhs_c.C
@@ -39,6 +39,7 @@ int f_compute_rhs_bssn(int *ex, double &T,
    // printf("nx=%d ny=%d nz=%d all=%d\n", nx, ny, nz, all);

    // temp variable
+    double gxx[all],gyy[all],gzz[all];
    double chix[all],chiy[all],chiz[all];
    double gxxx[all],gxyx[all],gxzx[all],gyyx[all],gyzx[all],gzzx[all];
    double gxxy[all],gxyy[all],gxzy[all],gyyy[all],gyzy[all],gzzy[all];
@@ -50,9 +51,9 @@ int f_compute_rhs_bssn(int *ex, double &T,
    double Gamxx[all],Gamxy[all],Gamxz[all];
    double Gamyx[all],Gamyy[all],Gamyz[all];
    double Gamzx[all],Gamzy[all],Gamzz[all];
-    double Kx[all], Ky[all], Kz[all], S[all];
+    double Kx[all], Ky[all], Kz[all], div_beta[all], S[all];
    double f[all], fxx[all], fxy[all], fxz[all], fyy[all], fyz[all], fzz[all];
-    double alpn1[all], chin1[all];
+    double Gamxa[all], Gamya[all], Gamza[all], alpn1[all], chin1[all];
    double gupxx[all], gupxy[all], gupxz[all];
    double gupyy[all], gupyz[all], gupzz[all];
    double SSS[3] = { 1.0,  1.0,  1.0};
@@ -106,6 +107,9 @@ int f_compute_rhs_bssn(int *ex, double &T,
        for(int i=0;i<all;i+=1){
            alpn1[i] = Lap[i] + 1.0;
            chin1[i] = chi[i] + 1.0;
+            gxx[i] = dxx[i] + 1.0;
+            gyy[i] = dyy[i] + 1.0;
+            gzz[i] = dzz[i] + 1.0;
        }
        // 9ms //
        fderivs(ex,betax,betaxx,betaxy,betaxz,X,Y,Z,ANTI, SYM, SYM,Symmetry,Lev);
@@ -123,196 +127,231 @@ int f_compute_rhs_bssn(int *ex, double &T,

        // 3ms //
        for(int i=0;i<all;i+=1){
-            const double divb = betaxx[i] + betayy[i] + betazz[i];
-            chi_rhs[i] = F2o3 * chin1[i] * (alpn1[i] * trK[i] - divb); 
-            gxx_rhs[i] = -TWO * alpn1[i] * Axx[i] - F2o3 * (dxx[i] + ONE) * divb + 
-                    TWO * ((dxx[i] + ONE) * betaxx[i] + gxy[i] * betayx[i] + gxz[i] * betazx[i]);
-            gyy_rhs[i] = -TWO * alpn1[i] * Ayy[i] - F2o3 * (dyy[i] + ONE) * divb + 
-                    TWO * (gxy[i] * betaxy[i] + (dyy[i] + ONE) * betayy[i] + gyz[i] * betazy[i]);
-            gzz_rhs[i] = -TWO * alpn1[i] * Azz[i] - F2o3 * (dzz[i] + ONE) * divb + 
-                        TWO * (gxz[i] * betaxz[i] + gyz[i] * betayz[i] + (dzz[i] + ONE) * betazz[i]);
-            gxy_rhs[i] = -TWO * alpn1[i] * Axy[i] + F1o3 * gxy[i] * divb + 
-                        (dxx[i] + ONE) * betaxy[i] + gxz[i] * betazy[i] + (dyy[i] + ONE) * betayx[i] 
+            div_beta[i] = betaxx[i] + betayy[i] + betazz[i];
+            chi_rhs[i] = F2o3 * chin1[i] * (alpn1[i] * trK[i] - div_beta[i]); 
+            gxx_rhs[i] = -TWO * alpn1[i] * Axx[i] - F2o3 * gxx[i] * div_beta[i] + 
+                    TWO * (gxx[i] * betaxx[i] + gxy[i] * betayx[i] + gxz[i] * betazx[i]);
+            gyy_rhs[i] = -TWO * alpn1[i] * Ayy[i] - F2o3 * gyy[i] * div_beta[i] + 
+                    TWO * (gxy[i] * betaxy[i] + gyy[i] * betayy[i] + gyz[i] * betazy[i]);
+            gzz_rhs[i] = -TWO * alpn1[i] * Azz[i] - F2o3 * gzz[i] * div_beta[i] + 
+                        TWO * (gxz[i] * betaxz[i] + gyz[i] * betayz[i] + gzz[i] * betazz[i]);
+            gxy_rhs[i] = -TWO * alpn1[i] * Axy[i] + F1o3 * gxy[i] * div_beta[i] + 
+                        gxx[i] * betaxy[i] + gxz[i] * betazy[i] + gyy[i] * betayx[i] 
                        + gyz[i] * betazx[i] - gxy[i] * betazz[i];        
-            gyz_rhs[i] = -TWO * alpn1[i] * Ayz[i] + F1o3 * gyz[i] * divb + 
-                        gxy[i] * betaxz[i] + (dyy[i] + ONE) * betayz[i] + gxz[i] * betaxy[i] 
-                        + (dzz[i] + ONE) * betazy[i] - gyz[i] * betaxx[i];
-            gxz_rhs[i] = -TWO * alpn1[i] * Axz[i] + F1o3 * gxz[i] * divb + 
-                        (dxx[i] + ONE) * betaxz[i] + gxy[i] * betayz[i] + gyz[i] * betayx[i] 
-                        + (dzz[i] + ONE) * betazx[i] - gxz[i] * betayy[i];                  
+            gyz_rhs[i] = -TWO * alpn1[i] * Ayz[i] + F1o3 * gyz[i] * div_beta[i] + 
+                        gxy[i] * betaxz[i] + gyy[i] * betayz[i] + gxz[i] * betaxy[i] 
+                        + gzz[i] * betazy[i] - gyz[i] * betaxx[i];
+            gxz_rhs[i] = -TWO * alpn1[i] * Axz[i] + F1o3 * gxz[i] * div_beta[i] + 
+                        gxx[i] * betaxz[i] + gxy[i] * betayz[i] + gyz[i] * betayx[i] 
+                        + gzz[i] * betazx[i] - gxz[i] * betayy[i];                  
        }
-        // Fused: inverse metric + Gamma constraint + Christoffel (3 loops -> 1)
+        // 1ms //
        for(int i=0;i<all;i+=1){
-            double det = (dxx[i] + ONE) * (dyy[i] + ONE) * (dzz[i] + ONE) + gxy[i] * gyz[i] * gxz[i] + gxz[i] * gxy[i] * gyz[i] -
-                        gxz[i] * (dyy[i] + ONE) * gxz[i] - gxy[i] * gxy[i] * (dzz[i] + ONE) - (dxx[i] + ONE) * gyz[i] * gyz[i];
-            double lg_xx = ((dyy[i] + ONE) * (dzz[i] + ONE) - gyz[i] * gyz[i]) / det;
-            double lg_xy = -(gxy[i] * (dzz[i] + ONE) - gyz[i] * gxz[i]) / det;
-            double lg_xz = (gxy[i] * gyz[i] - (dyy[i] + ONE) * gxz[i]) / det;
-            double lg_yy = ((dxx[i] + ONE) * (dzz[i] + ONE) - gxz[i] * gxz[i]) / det;
-            double lg_yz = -((dxx[i] + ONE) * gyz[i] - gxy[i] * gxz[i]) / det;
-            double lg_zz = ((dxx[i] + ONE) * (dyy[i] + ONE) - gxy[i] * gxy[i]) / det;
-            gupxx[i] = lg_xx; gupxy[i] = lg_xy; gupxz[i] = lg_xz;
-            gupyy[i] = lg_yy; gupyz[i] = lg_yz; gupzz[i] = lg_zz;
-
-            if(co==0){
+            double det = gxx[i] * gyy[i] * gzz[i] + gxy[i] * gyz[i] * gxz[i] + gxz[i] * gxy[i] * gyz[i] - 
+                        gxz[i] * gyy[i] * gxz[i] - gxy[i] * gxy[i] * gzz[i] - gxx[i] * gyz[i] * gyz[i];  
+            gupxx[i] = (gyy[i] * gzz[i] - gyz[i] * gyz[i]) / det;
+            gupxy[i] = -(gxy[i] * gzz[i] - gyz[i] * gxz[i]) / det;
+            gupxz[i] = (gxy[i] * gyz[i] - gyy[i] * gxz[i]) / det;
+            gupyy[i] = (gxx[i] * gzz[i] - gxz[i] * gxz[i]) / det;
+            gupyz[i] = -(gxx[i] * gyz[i] - gxy[i] * gxz[i]) / det;
+            gupzz[i] = (gxx[i] * gyy[i] - gxy[i] * gxy[i]) / det;   
+        }
+        // 2.2ms //
+        if(co==0){
+            for (int i=0;i<all;i+=1) {
                Gmx_Res[i] = Gamx[i] - (
-                    lg_xx * (lg_xx*gxxx[i] + lg_xy*gxyx[i] + lg_xz*gxzx[i]) +
-                    lg_xy * (lg_xx*gxyx[i] + lg_xy*gyyx[i] + lg_xz*gyzx[i]) +
-                    lg_xz * (lg_xx*gxzx[i] + lg_xy*gyzx[i] + lg_xz*gzzx[i]) +
-                    lg_xx * (lg_xy*gxxy[i] + lg_yy*gxyy[i] + lg_yz*gxzy[i]) +
-                    lg_xy * (lg_xy*gxyy[i] + lg_yy*gyyy[i] + lg_yz*gyzy[i]) +
-                    lg_xz * (lg_xy*gxzy[i] + lg_yy*gyzy[i] + lg_yz*gzzy[i]) +
-                    lg_xx * (lg_xz*gxxz[i] + lg_yz*gxyz[i] + lg_zz*gxzz[i]) +
-                    lg_xy * (lg_xz*gxyz[i] + lg_yz*gyyz[i] + lg_zz*gyzz[i]) +
-                    lg_xz * (lg_xz*gxzz[i] + lg_yz*gyzz[i] + lg_zz*gzzz[i])
+                    gupxx[i] * (gupxx[i]*gxxx[i] + gupxy[i]*gxyx[i] + gupxz[i]*gxzx[i]) +
+                    gupxy[i] * (gupxx[i]*gxyx[i] + gupxy[i]*gyyx[i] + gupxz[i]*gyzx[i]) +
+                    gupxz[i] * (gupxx[i]*gxzx[i] + gupxy[i]*gyzx[i] + gupxz[i]*gzzx[i]) +
+
+                    gupxx[i] * (gupxy[i]*gxxy[i] + gupyy[i]*gxyy[i] + gupyz[i]*gxzy[i]) +
+                    gupxy[i] * (gupxy[i]*gxyy[i] + gupyy[i]*gyyy[i] + gupyz[i]*gyzy[i]) +
+                    gupxz[i] * (gupxy[i]*gxzy[i] + gupyy[i]*gyzy[i] + gupyz[i]*gzzy[i]) +
+
+                    gupxx[i] * (gupxz[i]*gxxz[i] + gupyz[i]*gxyz[i] + gupzz[i]*gxzz[i]) +
+                    gupxy[i] * (gupxz[i]*gxyz[i] + gupyz[i]*gyyz[i] + gupzz[i]*gyzz[i]) +
+                    gupxz[i] * (gupxz[i]*gxzz[i] + gupyz[i]*gyzz[i] + gupzz[i]*gzzz[i])
                );
+
                Gmy_Res[i] = Gamy[i] - (
-                    lg_xx * (lg_xy*gxxx[i] + lg_yy*gxyx[i] + lg_yz*gxzx[i]) +
-                    lg_xy * (lg_xy*gxyx[i] + lg_yy*gyyx[i] + lg_yz*gyzx[i]) +
-                    lg_xz * (lg_xy*gxzx[i] + lg_yy*gyzx[i] + lg_yz*gzzx[i]) +
-                    lg_xy * (lg_xy*gxxy[i] + lg_yy*gxyy[i] + lg_yz*gxzy[i]) +
-                    lg_yy * (lg_xy*gxyy[i] + lg_yy*gyyy[i] + lg_yz*gyzy[i]) +
-                    lg_yz * (lg_xy*gxzy[i] + lg_yy*gyzy[i] + lg_yz*gzzy[i]) +
-                    lg_xy * (lg_xz*gxxz[i] + lg_yz*gxyz[i] + lg_zz*gxzz[i]) +
-                    lg_yy * (lg_xz*gxyz[i] + lg_yz*gyyz[i] + lg_zz*gyzz[i]) +
-                    lg_yz * (lg_xz*gxzz[i] + lg_yz*gyzz[i] + lg_zz*gzzz[i])
+                    gupxx[i] * (gupxy[i]*gxxx[i] + gupyy[i]*gxyx[i] + gupyz[i]*gxzx[i]) +
+                    gupxy[i] * (gupxy[i]*gxyx[i] + gupyy[i]*gyyx[i] + gupyz[i]*gyzx[i]) +
+                    gupxz[i] * (gupxy[i]*gxzx[i] + gupyy[i]*gyzx[i] + gupyz[i]*gzzx[i]) +
+
+                    gupxy[i] * (gupxy[i]*gxxy[i] + gupyy[i]*gxyy[i] + gupyz[i]*gxzy[i]) +
+                    gupyy[i] * (gupxy[i]*gxyy[i] + gupyy[i]*gyyy[i] + gupyz[i]*gyzy[i]) +
+                    gupyz[i] * (gupxy[i]*gxzy[i] + gupyy[i]*gyzy[i] + gupyz[i]*gzzy[i]) +
+
+                    gupxy[i] * (gupxz[i]*gxxz[i] + gupyz[i]*gxyz[i] + gupzz[i]*gxzz[i]) +
+                    gupyy[i] * (gupxz[i]*gxyz[i] + gupyz[i]*gyyz[i] + gupzz[i]*gyzz[i]) +
+                    gupyz[i] * (gupxz[i]*gxzz[i] + gupyz[i]*gyzz[i] + gupzz[i]*gzzz[i])
                );
+
                Gmz_Res[i] = Gamz[i] - (
-                    lg_xx * (lg_xz*gxxx[i] + lg_yz*gxyx[i] + lg_zz*gxzx[i]) +
-                    lg_xy * (lg_xz*gxyx[i] + lg_yz*gyyx[i] + lg_zz*gyzx[i]) +
-                    lg_xz * (lg_xz*gxzx[i] + lg_yz*gyzx[i] + lg_zz*gzzx[i]) +
-                    lg_xy * (lg_xz*gxxy[i] + lg_yz*gxyy[i] + lg_zz*gxzy[i]) +
-                    lg_yy * (lg_xz*gxyy[i] + lg_yz*gyyy[i] + lg_zz*gyzy[i]) +
-                    lg_yz * (lg_xz*gxzy[i] + lg_yz*gyzy[i] + lg_zz*gzzy[i]) +
-                    lg_xz * (lg_xz*gxxz[i] + lg_yz*gxyz[i] + lg_zz*gxzz[i]) +
-                    lg_yz * (lg_xz*gxyz[i] + lg_yz*gyyz[i] + lg_zz*gyzz[i]) +
-                    lg_zz * (lg_xz*gxzz[i] + lg_yz*gyzz[i] + lg_zz*gzzz[i])
+                    gupxx[i] * (gupxz[i]*gxxx[i] + gupyz[i]*gxyx[i] + gupzz[i]*gxzx[i]) +
+                    gupxy[i] * (gupxz[i]*gxyx[i] + gupyz[i]*gyyx[i] + gupzz[i]*gyzx[i]) +
+                    gupxz[i] * (gupxz[i]*gxzx[i] + gupyz[i]*gyzx[i] + gupzz[i]*gzzx[i]) +
+
+                    gupxy[i] * (gupxz[i]*gxxy[i] + gupyz[i]*gxyy[i] + gupzz[i]*gxzy[i]) +
+                    gupyy[i] * (gupxz[i]*gxyy[i] + gupyz[i]*gyyy[i] + gupzz[i]*gyzy[i]) +
+                    gupyz[i] * (gupxz[i]*gxzy[i] + gupyz[i]*gyzy[i] + gupzz[i]*gzzy[i]) +
+
+                    gupxz[i] * (gupxz[i]*gxxz[i] + gupyz[i]*gxyz[i] + gupzz[i]*gxzz[i]) +
+                    gupyz[i] * (gupxz[i]*gxyz[i] + gupyz[i]*gyyz[i] + gupzz[i]*gyzz[i]) +
+                    gupzz[i] * (gupxz[i]*gxzz[i] + gupyz[i]*gyzz[i] + gupzz[i]*gzzz[i])
                );
            }
-
-            Gamxxx[i] = HALF * ( lg_xx*gxxx[i]
-                            + lg_xy*(TWO*gxyx[i] - gxxy[i])
-                            + lg_xz*(TWO*gxzx[i] - gxxz[i]) );
-            Gamyxx[i] = HALF * ( lg_xy*gxxx[i]
-                            + lg_yy*(TWO*gxyx[i] - gxxy[i])
-                            + lg_yz*(TWO*gxzx[i] - gxxz[i]) );
-            Gamzxx[i] = HALF * ( lg_xz*gxxx[i]
-                            + lg_yz*(TWO*gxyx[i] - gxxy[i])
-                            + lg_zz*(TWO*gxzx[i] - gxxz[i]) );
-            Gamxyy[i] = HALF * ( lg_xx*(TWO*gxyy[i] - gyyx[i])
-                            + lg_xy*gyyy[i]
-                            + lg_xz*(TWO*gyzy[i] - gyyz[i]) );
-            Gamyyy[i] = HALF * ( lg_xy*(TWO*gxyy[i] - gyyx[i])
-                            + lg_yy*gyyy[i]
-                            + lg_yz*(TWO*gyzy[i] - gyyz[i]) );
-            Gamzyy[i] = HALF * ( lg_xz*(TWO*gxyy[i] - gyyx[i])
-                            + lg_yz*gyyy[i]
-                            + lg_zz*(TWO*gyzy[i] - gyyz[i]) );
-            Gamxzz[i] = HALF * ( lg_xx*(TWO*gxzz[i] - gzzx[i])
-                            + lg_xy*(TWO*gyzz[i] - gzzy[i])
-                            + lg_xz*gzzz[i] );
-            Gamyzz[i] = HALF * ( lg_xy*(TWO*gxzz[i] - gzzx[i])
-                            + lg_yy*(TWO*gyzz[i] - gzzy[i])
-                            + lg_yz*gzzz[i] );
-            Gamzzz[i] = HALF * ( lg_xz*(TWO*gxzz[i] - gzzx[i])
-                            + lg_yz*(TWO*gyzz[i] - gzzy[i])
-                            + lg_zz*gzzz[i] );
-            Gamxxy[i] = HALF * ( lg_xx*gxxy[i]
-                            + lg_xy*gyyx[i]
-                            + lg_xz*(gxzy[i] + gyzx[i] - gxyz[i]) );
-            Gamyxy[i] = HALF * ( lg_xy*gxxy[i]
-                            + lg_yy*gyyx[i]
-                            + lg_yz*(gxzy[i] + gyzx[i] - gxyz[i]) );
-            Gamzxy[i] = HALF * ( lg_xz*gxxy[i]
-                            + lg_yz*gyyx[i]
-                            + lg_zz*(gxzy[i] + gyzx[i] - gxyz[i]) );
-            Gamxxz[i] = HALF * ( lg_xx*gxxz[i]
-                            + lg_xy*(gxyz[i] + gyzx[i] - gxzy[i])
-                            + lg_xz*gzzx[i] );
-            Gamyxz[i] = HALF * ( lg_xy*gxxz[i]
-                            + lg_yy*(gxyz[i] + gyzx[i] - gxzy[i])
-                            + lg_yz*gzzx[i] );
-            Gamzxz[i] = HALF * ( lg_xz*gxxz[i]
-                            + lg_yz*(gxyz[i] + gyzx[i] - gxzy[i])
-                            + lg_zz*gzzx[i] );
-            Gamxyz[i] = HALF * ( lg_xx*(gxyz[i] + gxzy[i] - gyzx[i])
-                            + lg_xy*gyyz[i]
-                            + lg_xz*gzzy[i] );
-            Gamyyz[i] = HALF * ( lg_xy*(gxyz[i] + gxzy[i] - gyzx[i])
-                            + lg_yy*gyyz[i]
-                            + lg_yz*gzzy[i] );
-            Gamzyz[i] = HALF * ( lg_xz*(gxyz[i] + gxzy[i] - gyzx[i])
-                            + lg_yz*gyyz[i]
-                            + lg_zz*gzzy[i] );
        }
-        // Fused: A^{ij} raise-index + Gamma_rhs part 1 (2 loops -> 1)
+        // 5ms //
        for (int i=0;i<all;i+=1) {
-            double axx = gupxx[i]*gupxx[i]*Axx[i]
+
+            Gamxxx[i] = HALF * ( gupxx[i]*gxxx[i]
+                            + gupxy[i]*(TWO*gxyx[i] - gxxy[i])
+                            + gupxz[i]*(TWO*gxzx[i] - gxxz[i]) );
+
+            Gamyxx[i] = HALF * ( gupxy[i]*gxxx[i]
+                            + gupyy[i]*(TWO*gxyx[i] - gxxy[i])
+                            + gupyz[i]*(TWO*gxzx[i] - gxxz[i]) );
+
+            Gamzxx[i] = HALF * ( gupxz[i]*gxxx[i]
+                            + gupyz[i]*(TWO*gxyx[i] - gxxy[i])
+                            + gupzz[i]*(TWO*gxzx[i] - gxxz[i]) );
+
+            Gamxyy[i] = HALF * ( gupxx[i]*(TWO*gxyy[i] - gyyx[i])
+                            + gupxy[i]*gyyy[i]
+                            + gupxz[i]*(TWO*gyzy[i] - gyyz[i]) );
+
+            Gamyyy[i] = HALF * ( gupxy[i]*(TWO*gxyy[i] - gyyx[i])
+                            + gupyy[i]*gyyy[i]
+                            + gupyz[i]*(TWO*gyzy[i] - gyyz[i]) );
+
+            Gamzyy[i] = HALF * ( gupxz[i]*(TWO*gxyy[i] - gyyx[i])
+                            + gupyz[i]*gyyy[i]
+                            + gupzz[i]*(TWO*gyzy[i] - gyyz[i]) );
+
+            Gamxzz[i] = HALF * ( gupxx[i]*(TWO*gxzz[i] - gzzx[i])
+                            + gupxy[i]*(TWO*gyzz[i] - gzzy[i])
+                            + gupxz[i]*gzzz[i] );
+
+            Gamyzz[i] = HALF * ( gupxy[i]*(TWO*gxzz[i] - gzzx[i])
+                            + gupyy[i]*(TWO*gyzz[i] - gzzy[i])
+                            + gupyz[i]*gzzz[i] );
+
+            Gamzzz[i] = HALF * ( gupxz[i]*(TWO*gxzz[i] - gzzx[i])
+                            + gupyz[i]*(TWO*gyzz[i] - gzzy[i])
+                            + gupzz[i]*gzzz[i] );
+
+            Gamxxy[i] = HALF * ( gupxx[i]*gxxy[i]
+                            + gupxy[i]*gyyx[i]
+                            + gupxz[i]*(gxzy[i] + gyzx[i] - gxyz[i]) );
+
+            Gamyxy[i] = HALF * ( gupxy[i]*gxxy[i]
+                            + gupyy[i]*gyyx[i]
+                            + gupyz[i]*(gxzy[i] + gyzx[i] - gxyz[i]) );
+
+            Gamzxy[i] = HALF * ( gupxz[i]*gxxy[i]
+                            + gupyz[i]*gyyx[i]
+                            + gupzz[i]*(gxzy[i] + gyzx[i] - gxyz[i]) );
+
+            Gamxxz[i] = HALF * ( gupxx[i]*gxxz[i]
+                            + gupxy[i]*(gxyz[i] + gyzx[i] - gxzy[i])
+                            + gupxz[i]*gzzx[i] );
+
+            Gamyxz[i] = HALF * ( gupxy[i]*gxxz[i]
+                            + gupyy[i]*(gxyz[i] + gyzx[i] - gxzy[i])
+                            + gupyz[i]*gzzx[i] );
+
+            Gamzxz[i] = HALF * ( gupxz[i]*gxxz[i]
+                            + gupyz[i]*(gxyz[i] + gyzx[i] - gxzy[i])
+                            + gupzz[i]*gzzx[i] );
+
+            Gamxyz[i] = HALF * ( gupxx[i]*(gxyz[i] + gxzy[i] - gyzx[i])
+                            + gupxy[i]*gyyz[i]
+                            + gupxz[i]*gzzy[i] );
+
+            Gamyyz[i] = HALF * ( gupxy[i]*(gxyz[i] + gxzy[i] - gyzx[i])
+                            + gupyy[i]*gyyz[i]
+                            + gupyz[i]*gzzy[i] );
+
+            Gamzyz[i] = HALF * ( gupxz[i]*(gxyz[i] + gxzy[i] - gyzx[i])
+                            + gupyz[i]*gyyz[i]
+                            + gupzz[i]*gzzy[i] );
+
+        }
+        // 1.8ms //
+        for (int i=0;i<all;i+=1) {
+
+            Rxx[i] = gupxx[i]*gupxx[i]*Axx[i]
                + gupxy[i]*gupxy[i]*Ayy[i]
                + gupxz[i]*gupxz[i]*Azz[i]
                + TWO * ( gupxx[i]*gupxy[i]*Axy[i]
                        + gupxx[i]*gupxz[i]*Axz[i]
                        + gupxy[i]*gupxz[i]*Ayz[i] );
-            double ayy = gupxy[i]*gupxy[i]*Axx[i]
+
+            Ryy[i] = gupxy[i]*gupxy[i]*Axx[i]
                + gupyy[i]*gupyy[i]*Ayy[i]
                + gupyz[i]*gupyz[i]*Azz[i]
                + TWO * ( gupxy[i]*gupyy[i]*Axy[i]
                        + gupxy[i]*gupyz[i]*Axz[i]
                        + gupyy[i]*gupyz[i]*Ayz[i] );
-            double azz = gupxz[i]*gupxz[i]*Axx[i]
+
+            Rzz[i] = gupxz[i]*gupxz[i]*Axx[i]
                + gupyz[i]*gupyz[i]*Ayy[i]
                + gupzz[i]*gupzz[i]*Azz[i]
                + TWO * ( gupxz[i]*gupyz[i]*Axy[i]
                        + gupxz[i]*gupzz[i]*Axz[i]
                        + gupyz[i]*gupzz[i]*Ayz[i] );
-            double axy = gupxx[i]*gupxy[i]*Axx[i]
+
+            Rxy[i] = gupxx[i]*gupxy[i]*Axx[i]
                + gupxy[i]*gupyy[i]*Ayy[i]
                + gupxz[i]*gupyz[i]*Azz[i]
                + ( gupxx[i]*gupyy[i] + gupxy[i]*gupxy[i] ) * Axy[i]
                + ( gupxx[i]*gupyz[i] + gupxz[i]*gupxy[i] ) * Axz[i]
                + ( gupxy[i]*gupyz[i] + gupxz[i]*gupyy[i] ) * Ayz[i];
-            double axz = gupxx[i]*gupxz[i]*Axx[i]
+
+            Rxz[i] = gupxx[i]*gupxz[i]*Axx[i]
                + gupxy[i]*gupyz[i]*Ayy[i]
                + gupxz[i]*gupzz[i]*Azz[i]
                + ( gupxx[i]*gupyz[i] + gupxy[i]*gupxz[i] ) * Axy[i]
                + ( gupxx[i]*gupzz[i] + gupxz[i]*gupxz[i] ) * Axz[i]
                + ( gupxy[i]*gupzz[i] + gupxz[i]*gupyz[i] ) * Ayz[i];
-            double ayz = gupxy[i]*gupxz[i]*Axx[i]
+
+            Ryz[i] = gupxy[i]*gupxz[i]*Axx[i]
                + gupyy[i]*gupyz[i]*Ayy[i]
                + gupyz[i]*gupzz[i]*Azz[i]
                + ( gupxy[i]*gupyz[i] + gupyy[i]*gupxz[i] ) * Axy[i]
                + ( gupxy[i]*gupzz[i] + gupyz[i]*gupxz[i] ) * Axz[i]
                + ( gupyy[i]*gupzz[i] + gupyz[i]*gupyz[i] ) * Ayz[i];
-            Rxx[i] = axx; Ryy[i] = ayy; Rzz[i] = azz;
-            Rxy[i] = axy; Rxz[i] = axz; Ryz[i] = ayz;
+        }
+        // 4ms //
+        for(int i=0;i<all;i+=1){
+            Gamx_rhs[i] = - TWO * (   Lapx[i] * Rxx[i] +   Lapy[i] * Rxy[i] +   Lapz[i] * Rxz[i] ) + 
+                TWO * alpn1[i] * (                                                
+                -F3o2/chin1[i] * (   chix[i] * Rxx[i] +   chiy[i] * Rxy[i] +   chiz[i] * Rxz[i] ) - 
+                    gupxx[i] * (   F2o3 * Kx[i]  +  EIGHT * PI * Sx[i]            ) - 
+                    gupxy[i] * (   F2o3 * Ky[i]  +  EIGHT * PI * Sy[i]            ) - 
+                    gupxz[i] * (   F2o3 * Kz[i]  +  EIGHT * PI * Sz[i]            ) + 
+                                Gamxxx[i] * Rxx[i] + Gamxyy[i] * Ryy[i] + Gamxzz[i] * Rzz[i]   + 
+                        TWO * ( Gamxxy[i] * Rxy[i] + Gamxxz[i] * Rxz[i] + Gamxyz[i] * Ryz[i] ) );

-            Gamx_rhs[i] = - TWO * ( Lapx[i]*axx + Lapy[i]*axy + Lapz[i]*axz ) +
-                TWO * alpn1[i] * (
-                -F3o2/chin1[i] * ( chix[i]*axx + chiy[i]*axy + chiz[i]*axz ) -
-                    gupxx[i] * ( F2o3*Kx[i] + EIGHT*PI*Sx[i] ) -
-                    gupxy[i] * ( F2o3*Ky[i] + EIGHT*PI*Sy[i] ) -
-                    gupxz[i] * ( F2o3*Kz[i] + EIGHT*PI*Sz[i] ) +
-                                Gamxxx[i]*axx + Gamxyy[i]*ayy + Gamxzz[i]*azz +
-                        TWO * ( Gamxxy[i]*axy + Gamxxz[i]*axz + Gamxyz[i]*ayz ) );
-
-            Gamy_rhs[i] = -TWO * ( Lapx[i]*axy + Lapy[i]*ayy + Lapz[i]*ayz )
+            Gamy_rhs[i] = -TWO * ( Lapx[i]*Rxy[i] + Lapy[i]*Ryy[i] + Lapz[i]*Ryz[i] )
                        +  TWO * alpn1[i] * (
-                            -F3o2/chin1[i] * ( chix[i]*axy + chiy[i]*ayy + chiz[i]*ayz )
+                            -F3o2/chin1[i] * ( chix[i]*Rxy[i] + chiy[i]*Ryy[i] + chiz[i]*Ryz[i] )
                            - gupxy[i] * ( F2o3*Kx[i] + EIGHT*PI*Sx[i] )
                            - gupyy[i] * ( F2o3*Ky[i] + EIGHT*PI*Sy[i] )
                            - gupyz[i] * ( F2o3*Kz[i] + EIGHT*PI*Sz[i] )
-                            + Gamyxx[i]*axx + Gamyyy[i]*ayy + Gamyzz[i]*azz
-                            + TWO * ( Gamyxy[i]*axy + Gamyxz[i]*axz + Gamyyz[i]*ayz )
+                            + Gamyxx[i]*Rxx[i] + Gamyyy[i]*Ryy[i] + Gamyzz[i]*Rzz[i]
+                            + TWO * ( Gamyxy[i]*Rxy[i] + Gamyxz[i]*Rxz[i] + Gamyyz[i]*Ryz[i] )
                        );

-            Gamz_rhs[i] = -TWO * ( Lapx[i]*axz + Lapy[i]*ayz + Lapz[i]*azz )
+            Gamz_rhs[i] = -TWO * ( Lapx[i]*Rxz[i] + Lapy[i]*Ryz[i] + Lapz[i]*Rzz[i] )
                        +  TWO * alpn1[i] * (
-                            -F3o2/chin1[i] * ( chix[i]*axz + chiy[i]*ayz + chiz[i]*azz )
+                            -F3o2/chin1[i] * ( chix[i]*Rxz[i] + chiy[i]*Ryz[i] + chiz[i]*Rzz[i] )
                            - gupxz[i] * ( F2o3*Kx[i] + EIGHT*PI*Sx[i] )
                            - gupyz[i] * ( F2o3*Ky[i] + EIGHT*PI*Sy[i] )
                            - gupzz[i] * ( F2o3*Kz[i] + EIGHT*PI*Sz[i] )
-                            + Gamzxx[i]*axx + Gamzyy[i]*ayy + Gamzzz[i]*azz
-                            + TWO * ( Gamzxy[i]*axy + Gamzxz[i]*axz + Gamzyz[i]*ayz )
+                            + Gamzxx[i]*Rxx[i] + Gamzyy[i]*Ryy[i] + Gamzzz[i]*Rzz[i]
+                            + TWO * ( Gamzxy[i]*Rxy[i] + Gamzxz[i]*Rxz[i] + Gamzyz[i]*Ryz[i] )
                        );
        }
        // 22.3ms //
@@ -326,63 +365,65 @@ int f_compute_rhs_bssn(int *ex, double &T,
        fderivs(ex,Gamy,Gamyx,Gamyy,Gamyz,X,Y,Z,SYM ,ANTI,SYM ,Symmetry,Lev);
        fderivs(ex,Gamz,Gamzx,Gamzy,Gamzz,X,Y,Z,SYM ,SYM ,ANTI,Symmetry,Lev);

-        // Fused: fxx/Gamxa + Gamma_rhs part 2 (2 loops -> 1)
+        // 3.5ms //
        for(int i=0;i<all;i+=1){
-            const double divb = betaxx[i] + betayy[i] + betazz[i];
-            double lfxx = gxxx[i] + gxyy[i] + gxzz[i];
-            double lfxy = gxyx[i] + gyyy[i] + gyzz[i];
-            double lfxz = gxzx[i] + gyzy[i] + gzzz[i];
-            fxx[i] = lfxx; fxy[i] = lfxy; fxz[i] = lfxz;
-
-            double gxa = gupxx[i]*Gamxxx[i] + gupyy[i]*Gamxyy[i] + gupzz[i]*Gamxzz[i]
+            fxx[i] = gxxx[i] + gxyy[i] + gxzz[i];
+            fxy[i] = gxyx[i] + gyyy[i] + gyzz[i];
+            fxz[i] = gxzx[i] + gyzy[i] + gzzz[i];
+            Gamxa[i] = gupxx[i]*Gamxxx[i] + gupyy[i]*Gamxyy[i] + gupzz[i]*Gamxzz[i]
                    + TWO * ( gupxy[i]*Gamxxy[i] + gupxz[i]*Gamxxz[i] + gupyz[i]*Gamxyz[i] );
-            double gya = gupxx[i]*Gamyxx[i] + gupyy[i]*Gamyyy[i] + gupzz[i]*Gamyzz[i]
+
+            Gamya[i] = gupxx[i]*Gamyxx[i] + gupyy[i]*Gamyyy[i] + gupzz[i]*Gamyzz[i]
                    + TWO * ( gupxy[i]*Gamyxy[i] + gupxz[i]*Gamyxz[i] + gupyz[i]*Gamyyz[i] );
-            double gza = gupxx[i]*Gamzxx[i] + gupyy[i]*Gamzyy[i] + gupzz[i]*Gamzzz[i]
+
+            Gamza[i] = gupxx[i]*Gamzxx[i] + gupyy[i]*Gamzyy[i] + gupzz[i]*Gamzzz[i]
                    + TWO * ( gupxy[i]*Gamzxy[i] + gupxz[i]*Gamzxz[i] + gupyz[i]*Gamzyz[i] );
+        }
+        // 3.9ms //
+        for(int i=0;i<all;i+=1){
            Gamx_rhs[i] = Gamx_rhs[i]
-                        + F2o3 * gxa * divb
-                        - gxa * betaxx[i] - gya * betaxy[i] - gza * betaxz[i]
-                        + F1o3 * ( gupxx[i] * lfxx + gupxy[i] * lfxy + gupxz[i] * lfxz )
+                        + F2o3 * Gamxa[i] * div_beta[i]
+                        - Gamxa[i] * betaxx[i] - Gamya[i] * betaxy[i] - Gamza[i] * betaxz[i]
+                        + F1o3 * ( gupxx[i] * fxx[i] + gupxy[i] * fxy[i] + gupxz[i] * fxz[i] )
                        + gupxx[i] * gxxx[i] + gupyy[i] * gyyx[i] + gupzz[i] * gzzx[i]
                        + TWO * ( gupxy[i] * gxyx[i] + gupxz[i] * gxzx[i] + gupyz[i] * gyzx[i] );

            Gamy_rhs[i] = Gamy_rhs[i]
-                        + F2o3 * gya * divb
-                        - gxa * betayx[i] - gya * betayy[i] - gza * betayz[i]
-                        + F1o3 * ( gupxy[i] * lfxx + gupyy[i] * lfxy + gupyz[i] * lfxz )
+                        + F2o3 * Gamya[i] * div_beta[i]
+                        - Gamxa[i] * betayx[i] - Gamya[i] * betayy[i] - Gamza[i] * betayz[i]
+                        + F1o3 * ( gupxy[i] * fxx[i] + gupyy[i] * fxy[i] + gupyz[i] * fxz[i] )
                        + gupxx[i] * gxxy[i] + gupyy[i] * gyyy[i] + gupzz[i] * gzzy[i]
                        + TWO * ( gupxy[i] * gxyy[i] + gupxz[i] * gxzy[i] + gupyz[i] * gyzy[i] );

            Gamz_rhs[i] = Gamz_rhs[i]
-                        + F2o3 * gza * divb
-                        - gxa * betazx[i] - gya * betazy[i] - gza * betazz[i]
-                        + F1o3 * ( gupxz[i] * lfxx + gupyz[i] * lfxy + gupzz[i] * lfxz )
+                        + F2o3 * Gamza[i] * div_beta[i]
+                        - Gamxa[i] * betazx[i] - Gamya[i] * betazy[i] - Gamza[i] * betazz[i]
+                        + F1o3 * ( gupxz[i] * fxx[i] + gupyz[i] * fxy[i] + gupzz[i] * fxz[i] )
                        + gupxx[i] * gxxz[i] + gupyy[i] * gyyz[i] + gupzz[i] * gzzz[i]
                        + TWO * ( gupxy[i] * gxyz[i] + gupxz[i] * gxzz[i] + gupyz[i] * gyzz[i] );
        }
        // 4.4ms //
        for (int i=0;i<all;i+=1) {
-            gxxx[i] = (dxx[i] + ONE)*Gamxxx[i] + gxy[i]*Gamyxx[i] + gxz[i]*Gamzxx[i];
-            gxyx[i] = (dxx[i] + ONE)*Gamxxy[i] + gxy[i]*Gamyxy[i] + gxz[i]*Gamzxy[i];
-            gxzx[i] = (dxx[i] + ONE)*Gamxxz[i] + gxy[i]*Gamyxz[i] + gxz[i]*Gamzxz[i];
-            gyyx[i] = (dxx[i] + ONE)*Gamxyy[i] + gxy[i]*Gamyyy[i] + gxz[i]*Gamzyy[i];
-            gyzx[i] = (dxx[i] + ONE)*Gamxyz[i] + gxy[i]*Gamyyz[i] + gxz[i]*Gamzyz[i];
-            gzzx[i] = (dxx[i] + ONE)*Gamxzz[i] + gxy[i]*Gamyzz[i] + gxz[i]*Gamzzz[i];
+            gxxx[i] = gxx[i]*Gamxxx[i] + gxy[i]*Gamyxx[i] + gxz[i]*Gamzxx[i];
+            gxyx[i] = gxx[i]*Gamxxy[i] + gxy[i]*Gamyxy[i] + gxz[i]*Gamzxy[i];
+            gxzx[i] = gxx[i]*Gamxxz[i] + gxy[i]*Gamyxz[i] + gxz[i]*Gamzxz[i];
+            gyyx[i] = gxx[i]*Gamxyy[i] + gxy[i]*Gamyyy[i] + gxz[i]*Gamzyy[i];
+            gyzx[i] = gxx[i]*Gamxyz[i] + gxy[i]*Gamyyz[i] + gxz[i]*Gamzyz[i];
+            gzzx[i] = gxx[i]*Gamxzz[i] + gxy[i]*Gamyzz[i] + gxz[i]*Gamzzz[i];

-            gxxy[i] = gxy[i]*Gamxxx[i] + (dyy[i] + ONE)*Gamyxx[i] + gyz[i]*Gamzxx[i];
-            gxyy[i] = gxy[i]*Gamxxy[i] + (dyy[i] + ONE)*Gamyxy[i] + gyz[i]*Gamzxy[i];
-            gxzy[i] = gxy[i]*Gamxxz[i] + (dyy[i] + ONE)*Gamyxz[i] + gyz[i]*Gamzxz[i];
-            gyyy[i] = gxy[i]*Gamxyy[i] + (dyy[i] + ONE)*Gamyyy[i] + gyz[i]*Gamzyy[i];
-            gyzy[i] = gxy[i]*Gamxyz[i] + (dyy[i] + ONE)*Gamyyz[i] + gyz[i]*Gamzyz[i];
-            gzzy[i] = gxy[i]*Gamxzz[i] + (dyy[i] + ONE)*Gamyzz[i] + gyz[i]*Gamzzz[i];
+            gxxy[i] = gxy[i]*Gamxxx[i] + gyy[i]*Gamyxx[i] + gyz[i]*Gamzxx[i];
+            gxyy[i] = gxy[i]*Gamxxy[i] + gyy[i]*Gamyxy[i] + gyz[i]*Gamzxy[i];
+            gxzy[i] = gxy[i]*Gamxxz[i] + gyy[i]*Gamyxz[i] + gyz[i]*Gamzxz[i];
+            gyyy[i] = gxy[i]*Gamxyy[i] + gyy[i]*Gamyyy[i] + gyz[i]*Gamzyy[i];
+            gyzy[i] = gxy[i]*Gamxyz[i] + gyy[i]*Gamyyz[i] + gyz[i]*Gamzyz[i];
+            gzzy[i] = gxy[i]*Gamxzz[i] + gyy[i]*Gamyzz[i] + gyz[i]*Gamzzz[i];

-            gxxz[i] = gxz[i]*Gamxxx[i] + gyz[i]*Gamyxx[i] + (dzz[i] + ONE)*Gamzxx[i];
-            gxyz[i] = gxz[i]*Gamxxy[i] + gyz[i]*Gamyxy[i] + (dzz[i] + ONE)*Gamzxy[i];
-            gxzz[i] = gxz[i]*Gamxxz[i] + gyz[i]*Gamyxz[i] + (dzz[i] + ONE)*Gamzxz[i];
-            gyyz[i] = gxz[i]*Gamxyy[i] + gyz[i]*Gamyyy[i] + (dzz[i] + ONE)*Gamzyy[i];
-            gyzz[i] = gxz[i]*Gamxyz[i] + gyz[i]*Gamyyz[i] + (dzz[i] + ONE)*Gamzyz[i];
-            gzzz[i] = gxz[i]*Gamxzz[i] + gyz[i]*Gamyzz[i] + (dzz[i] + ONE)*Gamzzz[i];
+            gxxz[i] = gxz[i]*Gamxxx[i] + gyz[i]*Gamyxx[i] + gzz[i]*Gamzxx[i];
+            gxyz[i] = gxz[i]*Gamxxy[i] + gyz[i]*Gamyxy[i] + gzz[i]*Gamzxy[i];
+            gxzz[i] = gxz[i]*Gamxxz[i] + gyz[i]*Gamyxz[i] + gzz[i]*Gamzxz[i];
+            gyyz[i] = gxz[i]*Gamxyy[i] + gyz[i]*Gamyyy[i] + gzz[i]*Gamzyy[i];
+            gyzz[i] = gxz[i]*Gamxyz[i] + gyz[i]*Gamyyz[i] + gzz[i]*Gamzyz[i];
+            gzzz[i] = gxz[i]*Gamxzz[i] + gyz[i]*Gamyzz[i] + gzz[i]*Gamzzz[i];
        }
        // 22.2ms //
            fdderivs(ex,dxx,fxx,fxy,fxz,fyy,fyz,fzz,X,Y,Z,SYM ,SYM ,SYM ,Symmetry,Lev);
@@ -430,17 +471,10 @@ int f_compute_rhs_bssn(int *ex, double &T,
        // 14ms //
        /* 假设 all = ex1*ex2*ex3，所有量都是 length=all 的 double 数组（已按同一扁平化规则排布） */
        for (int i = 0; i < all; i += 1) {
-            const double gxa = gupxx[i]*Gamxxx[i] + gupyy[i]*Gamxyy[i] + gupzz[i]*Gamxzz[i]
-                            + TWO * ( gupxy[i]*Gamxxy[i] + gupxz[i]*Gamxxz[i] + gupyz[i]*Gamxyz[i] );
-            const double gya = gupxx[i]*Gamyxx[i] + gupyy[i]*Gamyyy[i] + gupzz[i]*Gamyzz[i]
-                            + TWO * ( gupxy[i]*Gamyxy[i] + gupxz[i]*Gamyxz[i] + gupyz[i]*Gamyyz[i] );
-            const double gza = gupxx[i]*Gamzxx[i] + gupyy[i]*Gamzyy[i] + gupzz[i]*Gamzzz[i]
-                            + TWO * ( gupxy[i]*Gamzxy[i] + gupxz[i]*Gamzxz[i] + gupyz[i]*Gamzyz[i] );
-
            Rxx[i] =
                -HALF * Rxx[i]
-                + (dxx[i] + ONE) * Gamxx[i] + gxy[i] * Gamyx[i] + gxz[i] * Gamzx[i]
-                + gxa * gxxx[i] + gya * gxyx[i] + gza * gxzx[i]
+                + gxx[i] * Gamxx[i] + gxy[i] * Gamyx[i] + gxz[i] * Gamzx[i]
+                + Gamxa[i] * gxxx[i] + Gamya[i] * gxyx[i] + Gamza[i] * gxzx[i]
                + gupxx[i] * (
                    TWO * (Gamxxx[i] * gxxx[i] + Gamyxx[i] * gxyx[i] + Gamzxx[i] * gxzx[i]) +
                        (Gamxxx[i] * gxxx[i] + Gamyxx[i] * gxxy[i] + Gamzxx[i] * gxxz[i])
@@ -474,8 +508,8 @@ int f_compute_rhs_bssn(int *ex, double &T,

            Ryy[i] =
                -HALF * Ryy[i]
-                + gxy[i] * Gamxy[i] + (dyy[i] + ONE) * Gamyy[i] + gyz[i] * Gamzy[i]
-                + gxa * gxyy[i] + gya * gyyy[i] + gza * gyzy[i]
+                + gxy[i] * Gamxy[i] + gyy[i] * Gamyy[i] + gyz[i] * Gamzy[i]
+                + Gamxa[i] * gxyy[i] + Gamya[i] * gyyy[i] + Gamza[i] * gyzy[i]
                + gupxx[i] * (
                    TWO * (Gamxxy[i] * gxxy[i] + Gamyxy[i] * gxyy[i] + Gamzxy[i] * gxzy[i]) +
                        (Gamxxy[i] * gxyx[i] + Gamyxy[i] * gxyy[i] + Gamzxy[i] * gxyz[i])
@@ -509,8 +543,8 @@ int f_compute_rhs_bssn(int *ex, double &T,

            Rzz[i] =
                -HALF * Rzz[i]
-                + gxz[i] * Gamxz[i] + gyz[i] * Gamyz[i] + (dzz[i] + ONE) * Gamzz[i]
-                + gxa * gxzz[i] + gya * gyzz[i] + gza * gzzz[i]
+                + gxz[i] * Gamxz[i] + gyz[i] * Gamyz[i] + gzz[i] * Gamzz[i]
+                + Gamxa[i] * gxzz[i] + Gamya[i] * gyzz[i] + Gamza[i] * gzzz[i]
                + gupxx[i] * (
                    TWO * (Gamxxz[i] * gxxz[i] + Gamyxz[i] * gxyz[i] + Gamzxz[i] * gxzz[i]) +
                        (Gamxxz[i] * gxzx[i] + Gamyxz[i] * gxzy[i] + Gamzxz[i] * gxzz[i])
@@ -545,10 +579,10 @@ int f_compute_rhs_bssn(int *ex, double &T,
            Rxy[i] =
                HALF * (
                    -Rxy[i]
-                    + (dxx[i] + ONE) * Gamxy[i] + gxy[i] * Gamyy[i] + gxz[i] * Gamzy[i]
-                    + gxy[i] * Gamxx[i] + (dyy[i] + ONE) * Gamyx[i] + gyz[i] * Gamzx[i]
-                    + gxa * gxyx[i] + gya * gyyx[i] + gza * gyzx[i]
-                    + gxa * gxxy[i] + gya * gxyy[i] + gza * gxzy[i]
+                    + gxx[i] * Gamxy[i] + gxy[i] * Gamyy[i] + gxz[i] * Gamzy[i]
+                    + gxy[i] * Gamxx[i] + gyy[i] * Gamyx[i] + gyz[i] * Gamzx[i]
+                    + Gamxa[i] * gxyx[i] + Gamya[i] * gyyx[i] + Gamza[i] * gyzx[i]
+                    + Gamxa[i] * gxxy[i] + Gamya[i] * gxyy[i] + Gamza[i] * gxzy[i]
                )
                + gupxx[i] * (
                    Gamxxx[i] * gxxy[i] + Gamyxx[i] * gxyy[i] + Gamzxx[i] * gxzy[i]
@@ -593,10 +627,10 @@ int f_compute_rhs_bssn(int *ex, double &T,
            Rxz[i] =
                HALF * (
                    -Rxz[i]
-                    + (dxx[i] + ONE) * Gamxz[i] + gxy[i] * Gamyz[i] + gxz[i] * Gamzz[i]
-                    + gxz[i] * Gamxx[i] + gyz[i] * Gamyx[i] + (dzz[i] + ONE) * Gamzx[i]
-                    + gxa * gxzx[i] + gya * gyzx[i] + gza * gzzx[i]
-                    + gxa * gxxz[i] + gya * gxyz[i] + gza * gxzz[i]
+                    + gxx[i] * Gamxz[i] + gxy[i] * Gamyz[i] + gxz[i] * Gamzz[i]
+                    + gxz[i] * Gamxx[i] + gyz[i] * Gamyx[i] + gzz[i] * Gamzx[i]
+                    + Gamxa[i] * gxzx[i] + Gamya[i] * gyzx[i] + Gamza[i] * gzzx[i]
+                    + Gamxa[i] * gxxz[i] + Gamya[i] * gxyz[i] + Gamza[i] * gxzz[i]
                )
                + gupxx[i] * (
                    Gamxxx[i] * gxxz[i] + Gamyxx[i] * gxyz[i] + Gamzxx[i] * gxzz[i]
@@ -641,10 +675,10 @@ int f_compute_rhs_bssn(int *ex, double &T,
            Ryz[i] =
                HALF * (
                    -Ryz[i]
-                    + gxy[i] * Gamxz[i] + (dyy[i] + ONE) * Gamyz[i] + gyz[i] * Gamzz[i]
-                    + gxz[i] * Gamxy[i] + gyz[i] * Gamyy[i] + (dzz[i] + ONE) * Gamzy[i]
-                    + gxa * gxzy[i] + gya * gyzy[i] + gza * gzzy[i]
-                    + gxa * gxyz[i] + gya * gyyz[i] + gza * gyzz[i]
+                    + gxy[i] * Gamxz[i] + gyy[i] * Gamyz[i] + gyz[i] * Gamzz[i]
+                    + gxz[i] * Gamxy[i] + gyz[i] * Gamyy[i] + gzz[i] * Gamzy[i]
+                    + Gamxa[i] * gxzy[i] + Gamya[i] * gyzy[i] + Gamza[i] * gzzy[i]
+                    + Gamxa[i] * gxyz[i] + Gamya[i] * gyyz[i] + Gamza[i] * gyzz[i]
                )
                + gupxx[i] * (
                    Gamxxy[i] * gxxz[i] + Gamyxy[i] * gxyz[i] + Gamzxy[i] * gxzz[i]
@@ -705,9 +739,9 @@ int f_compute_rhs_bssn(int *ex, double &T,
                + TWO * gupxy[i] * (fxy[i] - (F3o2 / chin1[i]) * chix[i] * chiy[i])
                + TWO * gupxz[i] * (fxz[i] - (F3o2 / chin1[i]) * chix[i] * chiz[i])
                + TWO * gupyz[i] * (fyz[i] - (F3o2 / chin1[i]) * chiy[i] * chiz[i]);
-            Rxx[i] = Rxx[i] + ( fxx[i] - (chix[i] * chix[i]) / (chin1[i] * TWO) + (dxx[i] + ONE) * f[i] ) / (chin1[i] * TWO);
-            Ryy[i] = Ryy[i] + ( fyy[i] - (chiy[i] * chiy[i]) / (chin1[i] * TWO) + (dyy[i] + ONE) * f[i] ) / (chin1[i] * TWO);
-            Rzz[i] = Rzz[i] + ( fzz[i] - (chiz[i] * chiz[i]) / (chin1[i] * TWO) + (dzz[i] + ONE) * f[i] ) / (chin1[i] * TWO);
+            Rxx[i] = Rxx[i] + ( fxx[i] - (chix[i] * chix[i]) / (chin1[i] * TWO) + gxx[i] * f[i] ) / (chin1[i] * TWO);
+            Ryy[i] = Ryy[i] + ( fyy[i] - (chiy[i] * chiy[i]) / (chin1[i] * TWO) + gyy[i] * f[i] ) / (chin1[i] * TWO);
+            Rzz[i] = Rzz[i] + ( fzz[i] - (chiz[i] * chiz[i]) / (chin1[i] * TWO) + gzz[i] * f[i] ) / (chin1[i] * TWO);

            Rxy[i] = Rxy[i] + ( fxy[i] - (chix[i] * chiy[i]) / (chin1[i] * TWO) + gxy[i] * f[i] ) / (chin1[i] * TWO);
            Rxz[i] = Rxz[i] + ( fxz[i] - (chix[i] * chiz[i]) / (chin1[i] * TWO) + gxz[i] * f[i] ) / (chin1[i] * TWO);
@@ -726,17 +760,17 @@ int f_compute_rhs_bssn(int *ex, double &T,
            gxxz[i] = (gupxz[i] * chix[i] + gupyz[i] * chiy[i] + gupzz[i] * chiz[i]) / chin1[i];

            /* Christoffel 修正项 */
-            Gamxxx[i] = Gamxxx[i] - ( ((chix[i] + chix[i]) / chin1[i]) - (dxx[i] + ONE) * gxxx[i] ) * HALF;
-            Gamyxx[i] = Gamyxx[i] - ( 0.0           - (dxx[i] + ONE) * gxxy[i] ) * HALF;  /* 原式只有 -gxx*gxxy */
-            Gamzxx[i] = Gamzxx[i] - ( 0.0           - (dxx[i] + ONE) * gxxz[i] ) * HALF;
+            Gamxxx[i] = Gamxxx[i] - ( ((chix[i] + chix[i]) / chin1[i]) - gxx[i] * gxxx[i] ) * HALF;
+            Gamyxx[i] = Gamyxx[i] - ( 0.0           - gxx[i] * gxxy[i] ) * HALF;  /* 原式只有 -gxx*gxxy */
+            Gamzxx[i] = Gamzxx[i] - ( 0.0           - gxx[i] * gxxz[i] ) * HALF;

-            Gamxyy[i] = Gamxyy[i] - ( 0.0           - (dyy[i] + ONE) * gxxx[i] ) * HALF;
-            Gamyyy[i] = Gamyyy[i] - ( ((chiy[i] + chiy[i]) / chin1[i]) - (dyy[i] + ONE) * gxxy[i] ) * HALF;
-            Gamzyy[i] = Gamzyy[i] - ( 0.0           - (dyy[i] + ONE) * gxxz[i] ) * HALF;
+            Gamxyy[i] = Gamxyy[i] - ( 0.0           - gyy[i] * gxxx[i] ) * HALF;
+            Gamyyy[i] = Gamyyy[i] - ( ((chiy[i] + chiy[i]) / chin1[i]) - gyy[i] * gxxy[i] ) * HALF;
+            Gamzyy[i] = Gamzyy[i] - ( 0.0           - gyy[i] * gxxz[i] ) * HALF;

-            Gamxzz[i] = Gamxzz[i] - ( 0.0           - (dzz[i] + ONE) * gxxx[i] ) * HALF;
-            Gamyzz[i] = Gamyzz[i] - ( 0.0           - (dzz[i] + ONE) * gxxy[i] ) * HALF;
-            Gamzzz[i] = Gamzzz[i] - ( ((chiz[i] + chiz[i]) / chin1[i]) - (dzz[i] + ONE) * gxxz[i] ) * HALF;
+            Gamxzz[i] = Gamxzz[i] - ( 0.0           - gzz[i] * gxxx[i] ) * HALF;
+            Gamyzz[i] = Gamyzz[i] - ( 0.0           - gzz[i] * gxxy[i] ) * HALF;
+            Gamzzz[i] = Gamzzz[i] - ( ((chiz[i] + chiz[i]) / chin1[i]) - gzz[i] * gxxz[i] ) * HALF;

            Gamxxy[i] = Gamxxy[i] - ( ( chiy[i] / chin1[i])      - gxy[i] * gxxx[i] ) * HALF;
            Gamyxy[i] = Gamyxy[i] - ( ( chix[i] / chin1[i])      - gxy[i] * gxxy[i] ) * HALF;
@@ -758,13 +792,14 @@ int f_compute_rhs_bssn(int *ex, double &T,
            fxy[i] = fxy[i] - Gamxxy[i] * Lapx[i] - Gamyxy[i] * Lapy[i] - Gamzxy[i] * Lapz[i];
            fxz[i] = fxz[i] - Gamxxz[i] * Lapx[i] - Gamyxz[i] * Lapy[i] - Gamzxz[i] * Lapz[i];
            fyz[i] = fyz[i] - Gamxyz[i] * Lapx[i] - Gamyyz[i] * Lapy[i] - Gamzyz[i] * Lapz[i];
-
+        }
+        // 1ms //
+        for (int i=0;i<all;i+=1) {
            trK_rhs[i] =  gupxx[i] * fxx[i] + gupyy[i] * fyy[i] + gupzz[i] * fzz[i]
                    +  TWO * ( gupxy[i] * fxy[i] + gupxz[i] * fxz[i] + gupyz[i] * fyz[i] );
        }
        // 2.5ms //
        for (int i=0;i<all;i+=1) {
-            const double divb = betaxx[i] + betayy[i] + betazz[i];

            S[i] = chin1[i] * (
                    gupxx[i] * Sxx[i] + gupyy[i] * Syy[i] + gupzz[i] * Szz[i]
@@ -815,20 +850,23 @@ int f_compute_rhs_bssn(int *ex, double &T,
                + (alpn1[i] / chin1[i]) * f[i]
            );

-            double l_fxx = alpn1[i] * (Rxx[i] - EIGHT * PI * Sxx[i]) - fxx[i];
-            double l_fxy = alpn1[i] * (Rxy[i] - EIGHT * PI * Sxy[i]) - fxy[i];
-            double l_fxz = alpn1[i] * (Rxz[i] - EIGHT * PI * Sxz[i]) - fxz[i];
-            double l_fyy = alpn1[i] * (Ryy[i] - EIGHT * PI * Syy[i]) - fyy[i];
-            double l_fyz = alpn1[i] * (Ryz[i] - EIGHT * PI * Syz[i]) - fyz[i];
-            double l_fzz = alpn1[i] * (Rzz[i] - EIGHT * PI * Szz[i]) - fzz[i];
+            fxx[i] = alpn1[i] * (Rxx[i] - EIGHT * PI * Sxx[i]) - fxx[i];
+            fxy[i] = alpn1[i] * (Rxy[i] - EIGHT * PI * Sxy[i]) - fxy[i];
+            fxz[i] = alpn1[i] * (Rxz[i] - EIGHT * PI * Sxz[i]) - fxz[i];
+            fyy[i] = alpn1[i] * (Ryy[i] - EIGHT * PI * Syy[i]) - fyy[i];
+            fyz[i] = alpn1[i] * (Ryz[i] - EIGHT * PI * Syz[i]) - fyz[i];
+            fzz[i] = alpn1[i] * (Rzz[i] - EIGHT * PI * Szz[i]) - fzz[i];
+        }
+        // 8ms //
+        for (int i=0;i<all;i+=1) {

            /* Aij_rhs = fij - gij * f */
-            Axx_rhs[i] = l_fxx - (dxx[i] + ONE) * f[i];
-            Ayy_rhs[i] = l_fyy - (dyy[i] + ONE) * f[i];
-            Azz_rhs[i] = l_fzz - (dzz[i] + ONE) * f[i];
-            Axy_rhs[i] = l_fxy - gxy[i] * f[i];
-            Axz_rhs[i] = l_fxz - gxz[i] * f[i];
-            Ayz_rhs[i] = l_fyz - gyz[i] * f[i];
+            Axx_rhs[i] = fxx[i] - gxx[i] * f[i];
+            Ayy_rhs[i] = fyy[i] - gyy[i] * f[i];
+            Azz_rhs[i] = fzz[i] - gzz[i] * f[i];
+            Axy_rhs[i] = fxy[i] - gxy[i] * f[i];
+            Axz_rhs[i] = fxz[i] - gxz[i] * f[i];
+            Ayz_rhs[i] = fyz[i] - gyz[i] * f[i];

            /* Now: store A_il A^l_j into fij: */
            fxx[i] =
@@ -890,19 +928,19 @@ int f_compute_rhs_bssn(int *ex, double &T,
                f[i] * Axx_rhs[i]
                + alpn1[i] * ( trK[i] * Axx[i] - TWO * fxx[i] )
                + TWO * ( Axx[i] * betaxx[i] + Axy[i] * betayx[i] + Axz[i] * betazx[i] )
-                - F2o3 * Axx[i] * divb;
+                - F2o3 * Axx[i] * div_beta[i];

            Ayy_rhs[i] =
                f[i] * Ayy_rhs[i]
                + alpn1[i] * ( trK[i] * Ayy[i] - TWO * fyy[i] )
                + TWO * ( Axy[i] * betaxy[i] + Ayy[i] * betayy[i] + Ayz[i] * betazy[i] )
-                - F2o3 * Ayy[i] * divb;
+                - F2o3 * Ayy[i] * div_beta[i];

            Azz_rhs[i] =
                f[i] * Azz_rhs[i]
                + alpn1[i] * ( trK[i] * Azz[i] - TWO * fzz[i] )
                + TWO * ( Axz[i] * betaxz[i] + Ayz[i] * betayz[i] + Azz[i] * betazz[i] )
-                - F2o3 * Azz[i] * divb;
+                - F2o3 * Azz[i] * div_beta[i];

            Axy_rhs[i] =
                f[i] * Axy_rhs[i]
@@ -911,7 +949,7 @@ int f_compute_rhs_bssn(int *ex, double &T,
                + Axz[i] * betazy[i]
                + Ayy[i] * betayx[i]
                + Ayz[i] * betazx[i]
-                + F1o3 * Axy[i] * divb
+                + F1o3 * Axy[i] * div_beta[i]
                - Axy[i] * betazz[i];

            Ayz_rhs[i] =
@@ -921,7 +959,7 @@ int f_compute_rhs_bssn(int *ex, double &T,
                + Ayy[i] * betayz[i]
                + Axz[i] * betaxy[i]
                + Azz[i] * betazy[i]
-                + F1o3 * Ayz[i] * divb
+                + F1o3 * Ayz[i] * div_beta[i]
                - Ayz[i] * betaxx[i];

            Axz_rhs[i] =
@@ -931,7 +969,7 @@ int f_compute_rhs_bssn(int *ex, double &T,
                + Axy[i] * betayz[i]
                + Ayz[i] * betayx[i]
                + Azz[i] * betazx[i]
-                + F1o3 * Axz[i] * divb
+                + F1o3 * Axz[i] * div_beta[i]
                - Axz[i] * betayy[i];

            /* Compute trace of S_ij */
@@ -1062,31 +1100,58 @@ int f_compute_rhs_bssn(int *ex, double &T,
            dtSfz_rhs[i] = Gamz_rhs[i] - reta[i] * dtSfz[i];
            #endif
        }
-        // advection + KO dissipation with shared symmetry buffer
-        lopsided_kodis(ex,X,Y,Z,dxx,gxx_rhs,betax,betay,betaz,Symmetry,SSS,eps);
-        lopsided_kodis(ex,X,Y,Z,Gamz,Gamz_rhs,betax,betay,betaz,Symmetry,SSA,eps);
-        lopsided_kodis(ex,X,Y,Z,gxy,gxy_rhs,betax,betay,betaz,Symmetry,AAS,eps);
-        lopsided_kodis(ex,X,Y,Z,Lap,Lap_rhs,betax,betay,betaz,Symmetry,SSS,eps);
-        lopsided_kodis(ex,X,Y,Z,gxz,gxz_rhs,betax,betay,betaz,Symmetry,ASA,eps);
-        lopsided_kodis(ex,X,Y,Z,betax,betax_rhs,betax,betay,betaz,Symmetry,ASS,eps);
-        lopsided_kodis(ex,X,Y,Z,dyy,gyy_rhs,betax,betay,betaz,Symmetry,SSS,eps);
-        lopsided_kodis(ex,X,Y,Z,betay,betay_rhs,betax,betay,betaz,Symmetry,SAS,eps);
-        lopsided_kodis(ex,X,Y,Z,gyz,gyz_rhs,betax,betay,betaz,Symmetry,SAA,eps);
-        lopsided_kodis(ex,X,Y,Z,betaz,betaz_rhs,betax,betay,betaz,Symmetry,SSA,eps);
-        lopsided_kodis(ex,X,Y,Z,dzz,gzz_rhs,betax,betay,betaz,Symmetry,SSS,eps);
-        lopsided_kodis(ex,X,Y,Z,dtSfx,dtSfx_rhs,betax,betay,betaz,Symmetry,ASS,eps);
-        lopsided_kodis(ex,X,Y,Z,Axx,Axx_rhs,betax,betay,betaz,Symmetry,SSS,eps);
-        lopsided_kodis(ex,X,Y,Z,dtSfy,dtSfy_rhs,betax,betay,betaz,Symmetry,SAS,eps);
-        lopsided_kodis(ex,X,Y,Z,Axy,Axy_rhs,betax,betay,betaz,Symmetry,AAS,eps);
-        lopsided_kodis(ex,X,Y,Z,dtSfz,dtSfz_rhs,betax,betay,betaz,Symmetry,SSA,eps);
-        lopsided_kodis(ex,X,Y,Z,Axz,Axz_rhs,betax,betay,betaz,Symmetry,ASA,eps);
-        lopsided_kodis(ex,X,Y,Z,Ayy,Ayy_rhs,betax,betay,betaz,Symmetry,SSS,eps);
-        lopsided_kodis(ex,X,Y,Z,Ayz,Ayz_rhs,betax,betay,betaz,Symmetry,SAA,eps);
-        lopsided_kodis(ex,X,Y,Z,Azz,Azz_rhs,betax,betay,betaz,Symmetry,SSS,eps);
-        lopsided_kodis(ex,X,Y,Z,chi,chi_rhs,betax,betay,betaz,Symmetry,SSS,eps);
-        lopsided_kodis(ex,X,Y,Z,trK,trK_rhs,betax,betay,betaz,Symmetry,SSS,eps);
-        lopsided_kodis(ex,X,Y,Z,Gamx,Gamx_rhs,betax,betay,betaz,Symmetry,ASS,eps);
-        lopsided_kodis(ex,X,Y,Z,Gamy,Gamy_rhs,betax,betay,betaz,Symmetry,SAS,eps);
+        // 26ms //
+        lopsided(ex,X,Y,Z,gxx,gxx_rhs,betax,betay,betaz,Symmetry,SSS);
+        lopsided(ex,X,Y,Z,Gamz,Gamz_rhs,betax,betay,betaz,Symmetry,SSA);
+        lopsided(ex,X,Y,Z,gxy,gxy_rhs,betax,betay,betaz,Symmetry,AAS);
+        lopsided(ex,X,Y,Z,Lap,Lap_rhs,betax,betay,betaz,Symmetry,SSS);
+        lopsided(ex,X,Y,Z,gxz,gxz_rhs,betax,betay,betaz,Symmetry,ASA);
+        lopsided(ex,X,Y,Z,betax,betax_rhs,betax,betay,betaz,Symmetry,ASS);
+        lopsided(ex,X,Y,Z,gyy,gyy_rhs,betax,betay,betaz,Symmetry,SSS);
+        lopsided(ex,X,Y,Z,betay,betay_rhs,betax,betay,betaz,Symmetry,SAS);
+        lopsided(ex,X,Y,Z,gyz,gyz_rhs,betax,betay,betaz,Symmetry,SAA);
+        lopsided(ex,X,Y,Z,betaz,betaz_rhs,betax,betay,betaz,Symmetry,SSA);
+        lopsided(ex,X,Y,Z,gzz,gzz_rhs,betax,betay,betaz,Symmetry,SSS);
+        lopsided(ex,X,Y,Z,dtSfx,dtSfx_rhs,betax,betay,betaz,Symmetry,ASS);
+        lopsided(ex,X,Y,Z,Axx,Axx_rhs,betax,betay,betaz,Symmetry,SSS);
+        lopsided(ex,X,Y,Z,dtSfy,dtSfy_rhs,betax,betay,betaz,Symmetry,SAS);
+        lopsided(ex,X,Y,Z,Axy,Axy_rhs,betax,betay,betaz,Symmetry,AAS);
+        lopsided(ex,X,Y,Z,dtSfz,dtSfz_rhs,betax,betay,betaz,Symmetry,SSA);
+        lopsided(ex,X,Y,Z,Axz,Axz_rhs,betax,betay,betaz,Symmetry,ASA);
+        lopsided(ex,X,Y,Z,Ayy,Ayy_rhs,betax,betay,betaz,Symmetry,SSS);
+        lopsided(ex,X,Y,Z,Ayz,Ayz_rhs,betax,betay,betaz,Symmetry,SAA);
+        lopsided(ex,X,Y,Z,Azz,Azz_rhs,betax,betay,betaz,Symmetry,SSS);
+        lopsided(ex,X,Y,Z,chi,chi_rhs,betax,betay,betaz,Symmetry,SSS);
+        lopsided(ex,X,Y,Z,trK,trK_rhs,betax,betay,betaz,Symmetry,SSS);
+        lopsided(ex,X,Y,Z,Gamx,Gamx_rhs,betax,betay,betaz,Symmetry,ASS);
+        lopsided(ex,X,Y,Z,Gamy,Gamy_rhs,betax,betay,betaz,Symmetry,SAS);
+        // 20ms //
+        if(eps>0){
+            kodis(ex,X,Y,Z,chi,chi_rhs,SSS,Symmetry,eps);
+            kodis(ex,X,Y,Z,trK,trK_rhs,SSS,Symmetry,eps);
+            kodis(ex,X,Y,Z,dxx,gxx_rhs,SSS,Symmetry,eps);
+            kodis(ex,X,Y,Z,gxy,gxy_rhs,AAS,Symmetry,eps);
+            kodis(ex,X,Y,Z,gxz,gxz_rhs,ASA,Symmetry,eps);
+            kodis(ex,X,Y,Z,dyy,gyy_rhs,SSS,Symmetry,eps);
+            kodis(ex,X,Y,Z,gyz,gyz_rhs,SAA,Symmetry,eps);
+            kodis(ex,X,Y,Z,dzz,gzz_rhs,SSS,Symmetry,eps);
+            kodis(ex,X,Y,Z,Axx,Axx_rhs,SSS,Symmetry,eps);
+            kodis(ex,X,Y,Z,dtSfz,dtSfz_rhs,SSA,Symmetry,eps);
+            kodis(ex,X,Y,Z,Axy,Axy_rhs,AAS,Symmetry,eps);
+            kodis(ex,X,Y,Z,dtSfy,dtSfy_rhs,SAS,Symmetry,eps);
+            kodis(ex,X,Y,Z,Axz,Axz_rhs,ASA,Symmetry,eps);
+            kodis(ex,X,Y,Z,dtSfx,dtSfx_rhs,ASS,Symmetry,eps);
+            kodis(ex,X,Y,Z,Ayy,Ayy_rhs,SSS,Symmetry,eps);
+            kodis(ex,X,Y,Z,betaz,betaz_rhs,SSA,Symmetry,eps);
+            kodis(ex,X,Y,Z,Ayz,Ayz_rhs,SAA,Symmetry,eps);
+            kodis(ex,X,Y,Z,betay,betay_rhs,SAS,Symmetry,eps);
+            kodis(ex,X,Y,Z,Azz,Azz_rhs,SSS,Symmetry,eps);
+            kodis(ex,X,Y,Z,betax,betax_rhs,ASS,Symmetry,eps);
+            kodis(ex,X,Y,Z,Gamx,Gamx_rhs,ASS,Symmetry,eps);
+            kodis(ex,X,Y,Z,Lap,Lap_rhs,SSS,Symmetry,eps);
+            kodis(ex,X,Y,Z,Gamy,Gamy_rhs,SAS,Symmetry,eps);
+            kodis(ex,X,Y,Z,Gamz,Gamz_rhs,SSA,Symmetry,eps);
+        }
        // 2ms //
        if(co==0){
            for (int i=0;i<all;i+=1) {
--- a/AMSS_NCKU_source/bssn_rhs_cuda.cu
+++ b/AMSS_NCKU_source/bssn_rhs_cuda.cu
--- a/AMSS_NCKU_source/bssn_rhs_cuda.h
+++ b/AMSS_NCKU_source/bssn_rhs_cuda.h
@@ -0,0 +1,36 @@
+#ifndef BSSN_RHS_CUDA_H
+#define BSSN_RHS_CUDA_H
+
+#ifdef __cplusplus
+extern "C" {
+#endif
+
+int f_compute_rhs_bssn(int *ex, double &T,
+                       double *X, double *Y, double *Z,
+                       double *chi, double *trK,
+                       double *dxx, double *gxy, double *gxz, double *dyy, double *gyz, double *dzz,
+                       double *Axx, double *Axy, double *Axz, double *Ayy, double *Ayz, double *Azz,
+                       double *Gamx, double *Gamy, double *Gamz,
+                       double *Lap, double *betax, double *betay, double *betaz,
+                       double *dtSfx, double *dtSfy, double *dtSfz,
+                       double *chi_rhs, double *trK_rhs,
+                       double *gxx_rhs, double *gxy_rhs, double *gxz_rhs, double *gyy_rhs, double *gyz_rhs, double *gzz_rhs,
+                       double *Axx_rhs, double *Axy_rhs, double *Axz_rhs, double *Ayy_rhs, double *Ayz_rhs, double *Azz_rhs,
+                       double *Gamx_rhs, double *Gamy_rhs, double *Gamz_rhs,
+                       double *Lap_rhs, double *betax_rhs, double *betay_rhs, double *betaz_rhs,
+                       double *dtSfx_rhs, double *dtSfy_rhs, double *dtSfz_rhs,
+                       double *rho, double *Sx, double *Sy, double *Sz,
+                       double *Sxx, double *Sxy, double *Sxz, double *Syy, double *Syz, double *Szz,
+                       double *Gamxxx, double *Gamxxy, double *Gamxxz, double *Gamxyy, double *Gamxyz, double *Gamxzz,
+                       double *Gamyxx, double *Gamyxy, double *Gamyxz, double *Gamyyy, double *Gamyyz, double *Gamyzz,
+                       double *Gamzxx, double *Gamzxy, double *Gamzxz, double *Gamzyy, double *Gamzyz, double *Gamzzz,
+                       double *Rxx, double *Rxy, double *Rxz, double *Ryy, double *Ryz, double *Rzz,
+                       double *ham_Res, double *movx_Res, double *movy_Res, double *movz_Res,
+                       double *Gmx_Res, double *Gmy_Res, double *Gmz_Res,
+                       int &Symmetry, int &Lev, double &eps, int &co);
+
+#ifdef __cplusplus
+}
+#endif
+
+#endif
--- a/AMSS_NCKU_source/fdderivs_c.C
+++ b/AMSS_NCKU_source/fdderivs_c.C
@@ -141,12 +141,26 @@ void fdderivs(const int ex[3],
    const int j4_hi = ex2 - 3;
    const int k4_hi = ex3 - 3;

+    /*
+     * Strategy A:
+     * Avoid redundant work in overlap of 2nd/4th-order regions.
+     * Only compute 2nd-order on shell points that are NOT overwritten by
+     * the 4th-order pass.
+     */
+    const int has4 = (i4_lo <= i4_hi && j4_lo <= j4_hi && k4_lo <= k4_hi);
+
    if (i2_lo <= i2_hi && j2_lo <= j2_hi && k2_lo <= k2_hi) {
        for (int k0 = k2_lo; k0 <= k2_hi; ++k0) {
            const int kF = k0 + 1;
            for (int j0 = j2_lo; j0 <= j2_hi; ++j0) {
                const int jF = j0 + 1;
                for (int i0 = i2_lo; i0 <= i2_hi; ++i0) {
+                    if (has4 &&
+                        i0 >= i4_lo && i0 <= i4_hi &&
+                        j0 >= j4_lo && j0 <= j4_hi &&
+                        k0 >= k4_lo && k0 <= k4_hi) {
+                        continue;
+                    }
                    const int iF = i0 + 1;
                    const size_t p = idx_ex(i0, j0, k0, ex);

@@ -193,7 +207,7 @@ void fdderivs(const int ex[3],
        }
    }

-    if (i4_lo <= i4_hi && j4_lo <= j4_hi && k4_lo <= k4_hi) {
+    if (has4) {
        for (int k0 = k4_lo; k0 <= k4_hi; ++k0) {
            const int kF = k0 + 1;
            for (int j0 = j4_lo; j0 <= j4_hi; ++j0) {
--- a/AMSS_NCKU_source/fderivs_c.C
+++ b/AMSS_NCKU_source/fderivs_c.C
@@ -81,63 +81,26 @@ void fderivs(const int ex[3],
    }

    /*
-     * 两段式：
-     * 1) 先在二阶可用区域计算二阶模板
-     * 2) 再在高阶可用区域覆盖为四阶模板
+     * Fortran loops:
+     * do k=1,ex3-1
+     * do j=1,ex2-1
+     * do i=1,ex1-1
     *
-     * 与原 if/elseif 逻辑等价，但减少逐点分支判断。
+     * C: k0=0..ex3-2, j0=0..ex2-2, i0=0..ex1-2
     */
-    const int i2_lo = (iminF > 0) ? iminF : 0;
-    const int j2_lo = (jminF > 0) ? jminF : 0;
-    const int k2_lo = (kminF > 0) ? kminF : 0;
-    const int i2_hi = ex1 - 2;
-    const int j2_hi = ex2 - 2;
-    const int k2_hi = ex3 - 2;
-
-    const int i4_lo = (iminF + 1 > 0) ? (iminF + 1) : 0;
-    const int j4_lo = (jminF + 1 > 0) ? (jminF + 1) : 0;
-    const int k4_lo = (kminF + 1 > 0) ? (kminF + 1) : 0;
-    const int i4_hi = ex1 - 3;
-    const int j4_hi = ex2 - 3;
-    const int k4_hi = ex3 - 3;
-
-    if (i2_lo <= i2_hi && j2_lo <= j2_hi && k2_lo <= k2_hi) {
-        for (int k0 = k2_lo; k0 <= k2_hi; ++k0) {
-            const int kF = k0 + 1;
-            for (int j0 = j2_lo; j0 <= j2_hi; ++j0) {
-                const int jF = j0 + 1;
-                for (int i0 = i2_lo; i0 <= i2_hi; ++i0) {
-                    const int iF = i0 + 1;
-                    const size_t p = idx_ex(i0, j0, k0, ex);
-
-                    fx[p] = d2dx * (
-                        -fh[idx_fh_F_ord2(iF - 1, jF,     kF,     ex)] +
-                         fh[idx_fh_F_ord2(iF + 1, jF,     kF,     ex)]
-                    );
-
-                    fy[p] = d2dy * (
-                        -fh[idx_fh_F_ord2(iF,     jF - 1, kF,     ex)] +
-                         fh[idx_fh_F_ord2(iF,     jF + 1, kF,     ex)]
-                    );
-
-                    fz[p] = d2dz * (
-                        -fh[idx_fh_F_ord2(iF,     jF,     kF - 1, ex)] +
-                         fh[idx_fh_F_ord2(iF,     jF,     kF + 1, ex)]
-                    );
-                }
-            }
-        }
-    }
-
-    if (i4_lo <= i4_hi && j4_lo <= j4_hi && k4_lo <= k4_hi) {
-        for (int k0 = k4_lo; k0 <= k4_hi; ++k0) {
-            const int kF = k0 + 1;
-            for (int j0 = j4_lo; j0 <= j4_hi; ++j0) {
-                const int jF = j0 + 1;
-                for (int i0 = i4_lo; i0 <= i4_hi; ++i0) {
-                    const int iF = i0 + 1;
-                    const size_t p = idx_ex(i0, j0, k0, ex);
+    for (int k0 = 0; k0 <= ex3 - 2; ++k0) {
+        const int kF = k0 + 1;
+        for (int j0 = 0; j0 <= ex2 - 2; ++j0) {
+            const int jF = j0 + 1;
+            for (int i0 = 0; i0 <= ex1 - 2; ++i0) {
+                const int iF = i0 + 1;
+                const size_t p = idx_ex(i0, j0, k0, ex);

+                // if(i+2 <= imax .and. i-2 >= imin ... )  (全是 Fortran 索引)
+                if ((iF + 2) <= imaxF && (iF - 2) >= iminF &&
+                    (jF + 2) <= jmaxF && (jF - 2) >= jminF &&
+                    (kF + 2) <= kmaxF && (kF - 2) >= kminF)
+                {
                    fx[p] = d12dx * (
                        fh[idx_fh_F_ord2(iF - 2, jF,     kF,     ex)] -
                        EIT * fh[idx_fh_F_ord2(iF - 1, jF,     kF,     ex)] +
@@ -159,9 +122,29 @@ void fderivs(const int ex[3],
                        fh[idx_fh_F_ord2(iF,     jF,     kF + 2, ex)]
                    );
                }
+                // elseif(i+1 <= imax .and. i-1 >= imin ...)
+                else if ((iF + 1) <= imaxF && (iF - 1) >= iminF &&
+                         (jF + 1) <= jmaxF && (jF - 1) >= jminF &&
+                         (kF + 1) <= kmaxF && (kF - 1) >= kminF)
+                {
+                    fx[p] = d2dx * (
+                        -fh[idx_fh_F_ord2(iF - 1, jF,     kF,     ex)] +
+                         fh[idx_fh_F_ord2(iF + 1, jF,     kF,     ex)]
+                    );
+
+                    fy[p] = d2dy * (
+                        -fh[idx_fh_F_ord2(iF,     jF - 1, kF,     ex)] +
+                         fh[idx_fh_F_ord2(iF,     jF + 1, kF,     ex)]
+                    );
+
+                    fz[p] = d2dz * (
+                        -fh[idx_fh_F_ord2(iF,     jF,     kF - 1, ex)] +
+                         fh[idx_fh_F_ord2(iF,     jF,     kF + 1, ex)]
+                    );
+                }
            }
        }
    }

    // free(fh);
-}
+}
--- a/AMSS_NCKU_source/fmisc.f90
+++ b/AMSS_NCKU_source/fmisc.f90
@@ -1327,6 +1327,35 @@ end subroutine d2dump

  return
  end subroutine polint
+
+  subroutine polint0(xa, ya, y, ordn)
+  ! Lagrange interpolation at x=0, O(n) direct formula
+  implicit none
+  integer, intent(in) :: ordn
+  real*8, dimension(ordn), intent(in) :: xa, ya
+  real*8, intent(out) :: y
+
+  integer :: j, k
+  real*8 :: wj
+
+  y = 0.d0
+  do j = 1, ordn
+    wj = 1.d0
+    do k = 1, ordn
+      if (k .ne. j) then
+        wj = wj * xa(k) / (xa(k) - xa(j))
+      endif
+    enddo
+    y = y + wj * ya(j)
+  enddo
+
+  return
+  end subroutine polint0
+!------------------------------------------------------------------------------
+!
+! interpolation in 2 dimensions, follow yx order
+!
+!------------------------------------------------------------------------------
 !------------------------------------------------------------------------------
 ! Compute Lagrange interpolation basis weights for one target point.
 !------------------------------------------------------------------------------
--- a/AMSS_NCKU_source/interp_lb_profile_data.h
+++ b/AMSS_NCKU_source/interp_lb_profile_data.h
@@ -1,5 +1,3 @@
-/* 本头文件由自订profile框架自动生成并非人工硬编码针对Case优化 */
-/* 更新：负载均衡问题已经通过优化插值函数解决，此profile静态均衡方案已弃用，本头文件现在未参与编译 */
 /* Auto-generated from interp_lb_profile.bin — do not edit */
 #ifndef INTERP_LB_PROFILE_DATA_H
 #define INTERP_LB_PROFILE_DATA_H
--- a/AMSS_NCKU_source/kodiss_c.C
+++ b/AMSS_NCKU_source/kodiss_c.C
@@ -63,28 +63,19 @@ void kodis(const int ex[3],
     * C: k0=0..ex3-1, j0=0..ex2-1, i0=0..ex1-1
     * 并定义 Fortran index: iF=i0+1, ...
     */
-    // 收紧循环范围：只遍历满足 iF±3/jF±3/kF±3 条件的内部点
-    // iF-3 >= iminF => iF >= iminF+3 => i0 >= iminF+2 (因为 iF=i0+1)
-    // iF+3 <= imaxF => iF <= imaxF-3 => i0 <= imaxF-4
-    const int i0_lo = (iminF + 2 > 0) ? iminF + 2 : 0;
-    const int j0_lo = (jminF + 2 > 0) ? jminF + 2 : 0;
-    const int k0_lo = (kminF + 2 > 0) ? kminF + 2 : 0;
-    const int i0_hi = imaxF - 4;  // inclusive
-    const int j0_hi = jmaxF - 4;
-    const int k0_hi = kmaxF - 4;
-
-    if (i0_lo > i0_hi || j0_lo > j0_hi || k0_lo > k0_hi) {
-        free(fh);
-        return;
-    }
-
-    for (int k0 = k0_lo; k0 <= k0_hi; ++k0) {
+    for (int k0 = 0; k0 < ex3; ++k0) {
        const int kF = k0 + 1;
-        for (int j0 = j0_lo; j0 <= j0_hi; ++j0) {
+        for (int j0 = 0; j0 < ex2; ++j0) {
            const int jF = j0 + 1;
-            for (int i0 = i0_lo; i0 <= i0_hi; ++i0) {
+            for (int i0 = 0; i0 < ex1; ++i0) {
                const int iF = i0 + 1;

+                // Fortran if 条件：
+                // i-3 >= imin .and. i+3 <= imax  等（都是 Fortran 索引）
+                if ((iF - 3) >= iminF && (iF + 3) <= imaxF &&
+                    (jF - 3) >= jminF && (jF + 3) <= jmaxF &&
+                    (kF - 3) >= kminF && (kF + 3) <= kmaxF)
+                {
                    const size_t p = idx_ex(i0, j0, k0, ex);

                    // 三个方向各一份同型的 7 点组合（实际上是对称的 6th-order dissipation/filter 核）
@@ -109,6 +100,7 @@ void kodis(const int ex[3],
                    // Fortran:
                    // f_rhs(i,j,k) = f_rhs(i,j,k) + eps/cof*(Dx_term + Dy_term + Dz_term)
                    f_rhs[p] += (eps / cof) * (Dx_term + Dy_term + Dz_term);
+                }
            }
        }
    }
--- a/AMSS_NCKU_source/lopsided_kodis_c.C
+++ b/AMSS_NCKU_source/lopsided_kodis_c.C
@@ -1,248 +0,0 @@
-#include "tool.h"
-
-/*
- * Combined advection (lopsided) + KO dissipation (kodis).
- * Uses one shared symmetry_bd buffer per call.
- */
-void lopsided_kodis(const int ex[3],
-                    const double *X, const double *Y, const double *Z,
-                    const double *f, double *f_rhs,
-                    const double *Sfx, const double *Sfy, const double *Sfz,
-                    int Symmetry, const double SoA[3], double eps)
-{
-    const double ZEO = 0.0, ONE = 1.0, F3 = 3.0;
-    const double F6 = 6.0, F18 = 18.0;
-    const double F12 = 12.0, F10 = 10.0, EIT = 8.0;
-    const double SIX = 6.0, FIT = 15.0, TWT = 20.0;
-    const double cof = 64.0; // 2^6
-
-    const int NO_SYMM = 0, EQ_SYMM = 1;
-
-    const int ex1 = ex[0], ex2 = ex[1], ex3 = ex[2];
-
-    const double dX = X[1] - X[0];
-    const double dY = Y[1] - Y[0];
-    const double dZ = Z[1] - Z[0];
-
-    const double d12dx = ONE / F12 / dX;
-    const double d12dy = ONE / F12 / dY;
-    const double d12dz = ONE / F12 / dZ;
-
-    const int imaxF = ex1;
-    const int jmaxF = ex2;
-    const int kmaxF = ex3;
-
-    int iminF = 1, jminF = 1, kminF = 1;
-    if (Symmetry > NO_SYMM && fabs(Z[0]) < dZ) kminF = -2;
-    if (Symmetry > EQ_SYMM && fabs(X[0]) < dX) iminF = -2;
-    if (Symmetry > EQ_SYMM && fabs(Y[0]) < dY) jminF = -2;
-
-    // fh for Fortran-style domain (-2:ex1,-2:ex2,-2:ex3)
-    const size_t nx = (size_t)ex1 + 3;
-    const size_t ny = (size_t)ex2 + 3;
-    const size_t nz = (size_t)ex3 + 3;
-    const size_t fh_size = nx * ny * nz;
-
-    double *fh = (double*)malloc(fh_size * sizeof(double));
-    if (!fh) return;
-
-    symmetry_bd(3, ex, f, fh, SoA);
-
-    // Advection (same stencil logic as lopsided_c.C)
-    for (int k0 = 0; k0 <= ex3 - 2; ++k0) {
-        const int kF = k0 + 1;
-        for (int j0 = 0; j0 <= ex2 - 2; ++j0) {
-            const int jF = j0 + 1;
-            for (int i0 = 0; i0 <= ex1 - 2; ++i0) {
-                const int iF = i0 + 1;
-                const size_t p = idx_ex(i0, j0, k0, ex);
-
-                const double sfx = Sfx[p];
-                if (sfx > ZEO) {
-                    if (i0 <= ex1 - 4) {
-                        f_rhs[p] += sfx * d12dx *
-                            (-F3  * fh[idx_fh_F(iF - 1, jF, kF, ex)]
-                             -F10 * fh[idx_fh_F(iF    , jF, kF, ex)]
-                             +F18 * fh[idx_fh_F(iF + 1, jF, kF, ex)]
-                             -F6  * fh[idx_fh_F(iF + 2, jF, kF, ex)]
-                             +      fh[idx_fh_F(iF + 3, jF, kF, ex)]);
-                    } else if (i0 <= ex1 - 3) {
-                        f_rhs[p] += sfx * d12dx *
-                            ( fh[idx_fh_F(iF - 2, jF, kF, ex)]
-                             -EIT * fh[idx_fh_F(iF - 1, jF, kF, ex)]
-                             +EIT * fh[idx_fh_F(iF + 1, jF, kF, ex)]
-                             -      fh[idx_fh_F(iF + 2, jF, kF, ex)]);
-                    } else if (i0 <= ex1 - 2) {
-                        f_rhs[p] -= sfx * d12dx *
-                            (-F3  * fh[idx_fh_F(iF + 1, jF, kF, ex)]
-                             -F10 * fh[idx_fh_F(iF    , jF, kF, ex)]
-                             +F18 * fh[idx_fh_F(iF - 1, jF, kF, ex)]
-                             -F6  * fh[idx_fh_F(iF - 2, jF, kF, ex)]
-                             +      fh[idx_fh_F(iF - 3, jF, kF, ex)]);
-                    }
-                } else if (sfx < ZEO) {
-                    if ((i0 - 2) >= iminF) {
-                        f_rhs[p] -= sfx * d12dx *
-                            (-F3  * fh[idx_fh_F(iF + 1, jF, kF, ex)]
-                             -F10 * fh[idx_fh_F(iF    , jF, kF, ex)]
-                             +F18 * fh[idx_fh_F(iF - 1, jF, kF, ex)]
-                             -F6  * fh[idx_fh_F(iF - 2, jF, kF, ex)]
-                             +      fh[idx_fh_F(iF - 3, jF, kF, ex)]);
-                    } else if ((i0 - 1) >= iminF) {
-                        f_rhs[p] += sfx * d12dx *
-                            ( fh[idx_fh_F(iF - 2, jF, kF, ex)]
-                             -EIT * fh[idx_fh_F(iF - 1, jF, kF, ex)]
-                             +EIT * fh[idx_fh_F(iF + 1, jF, kF, ex)]
-                             -      fh[idx_fh_F(iF + 2, jF, kF, ex)]);
-                    } else if (i0 >= iminF) {
-                        f_rhs[p] += sfx * d12dx *
-                            (-F3  * fh[idx_fh_F(iF - 1, jF, kF, ex)]
-                             -F10 * fh[idx_fh_F(iF    , jF, kF, ex)]
-                             +F18 * fh[idx_fh_F(iF + 1, jF, kF, ex)]
-                             -F6  * fh[idx_fh_F(iF + 2, jF, kF, ex)]
-                             +      fh[idx_fh_F(iF + 3, jF, kF, ex)]);
-                    }
-                }
-
-                const double sfy = Sfy[p];
-                if (sfy > ZEO) {
-                    if (j0 <= ex2 - 4) {
-                        f_rhs[p] += sfy * d12dy *
-                            (-F3  * fh[idx_fh_F(iF, jF - 1, kF, ex)]
-                             -F10 * fh[idx_fh_F(iF, jF    , kF, ex)]
-                             +F18 * fh[idx_fh_F(iF, jF + 1, kF, ex)]
-                             -F6  * fh[idx_fh_F(iF, jF + 2, kF, ex)]
-                             +      fh[idx_fh_F(iF, jF + 3, kF, ex)]);
-                    } else if (j0 <= ex2 - 3) {
-                        f_rhs[p] += sfy * d12dy *
-                            ( fh[idx_fh_F(iF, jF - 2, kF, ex)]
-                             -EIT * fh[idx_fh_F(iF, jF - 1, kF, ex)]
-                             +EIT * fh[idx_fh_F(iF, jF + 1, kF, ex)]
-                             -      fh[idx_fh_F(iF, jF + 2, kF, ex)]);
-                    } else if (j0 <= ex2 - 2) {
-                        f_rhs[p] -= sfy * d12dy *
-                            (-F3  * fh[idx_fh_F(iF, jF + 1, kF, ex)]
-                             -F10 * fh[idx_fh_F(iF, jF    , kF, ex)]
-                             +F18 * fh[idx_fh_F(iF, jF - 1, kF, ex)]
-                             -F6  * fh[idx_fh_F(iF, jF - 2, kF, ex)]
-                             +      fh[idx_fh_F(iF, jF - 3, kF, ex)]);
-                    }
-                } else if (sfy < ZEO) {
-                    if ((j0 - 2) >= jminF) {
-                        f_rhs[p] -= sfy * d12dy *
-                            (-F3  * fh[idx_fh_F(iF, jF + 1, kF, ex)]
-                             -F10 * fh[idx_fh_F(iF, jF    , kF, ex)]
-                             +F18 * fh[idx_fh_F(iF, jF - 1, kF, ex)]
-                             -F6  * fh[idx_fh_F(iF, jF - 2, kF, ex)]
-                             +      fh[idx_fh_F(iF, jF - 3, kF, ex)]);
-                    } else if ((j0 - 1) >= jminF) {
-                        f_rhs[p] += sfy * d12dy *
-                            ( fh[idx_fh_F(iF, jF - 2, kF, ex)]
-                             -EIT * fh[idx_fh_F(iF, jF - 1, kF, ex)]
-                             +EIT * fh[idx_fh_F(iF, jF + 1, kF, ex)]
-                             -      fh[idx_fh_F(iF, jF + 2, kF, ex)]);
-                    } else if (j0 >= jminF) {
-                        f_rhs[p] += sfy * d12dy *
-                            (-F3  * fh[idx_fh_F(iF, jF - 1, kF, ex)]
-                             -F10 * fh[idx_fh_F(iF, jF    , kF, ex)]
-                             +F18 * fh[idx_fh_F(iF, jF + 1, kF, ex)]
-                             -F6  * fh[idx_fh_F(iF, jF + 2, kF, ex)]
-                             +      fh[idx_fh_F(iF, jF + 3, kF, ex)]);
-                    }
-                }
-
-                const double sfz = Sfz[p];
-                if (sfz > ZEO) {
-                    if (k0 <= ex3 - 4) {
-                        f_rhs[p] += sfz * d12dz *
-                            (-F3  * fh[idx_fh_F(iF, jF, kF - 1, ex)]
-                             -F10 * fh[idx_fh_F(iF, jF, kF    , ex)]
-                             +F18 * fh[idx_fh_F(iF, jF, kF + 1, ex)]
-                             -F6  * fh[idx_fh_F(iF, jF, kF + 2, ex)]
-                             +      fh[idx_fh_F(iF, jF, kF + 3, ex)]);
-                    } else if (k0 <= ex3 - 3) {
-                        f_rhs[p] += sfz * d12dz *
-                            ( fh[idx_fh_F(iF, jF, kF - 2, ex)]
-                             -EIT * fh[idx_fh_F(iF, jF, kF - 1, ex)]
-                             +EIT * fh[idx_fh_F(iF, jF, kF + 1, ex)]
-                             -      fh[idx_fh_F(iF, jF, kF + 2, ex)]);
-                    } else if (k0 <= ex3 - 2) {
-                        f_rhs[p] -= sfz * d12dz *
-                            (-F3  * fh[idx_fh_F(iF, jF, kF + 1, ex)]
-                             -F10 * fh[idx_fh_F(iF, jF, kF    , ex)]
-                             +F18 * fh[idx_fh_F(iF, jF, kF - 1, ex)]
-                             -F6  * fh[idx_fh_F(iF, jF, kF - 2, ex)]
-                             +      fh[idx_fh_F(iF, jF, kF - 3, ex)]);
-                    }
-                } else if (sfz < ZEO) {
-                    if ((k0 - 2) >= kminF) {
-                        f_rhs[p] -= sfz * d12dz *
-                            (-F3  * fh[idx_fh_F(iF, jF, kF + 1, ex)]
-                             -F10 * fh[idx_fh_F(iF, jF, kF    , ex)]
-                             +F18 * fh[idx_fh_F(iF, jF, kF - 1, ex)]
-                             -F6  * fh[idx_fh_F(iF, jF, kF - 2, ex)]
-                             +      fh[idx_fh_F(iF, jF, kF - 3, ex)]);
-                    } else if ((k0 - 1) >= kminF) {
-                        f_rhs[p] += sfz * d12dz *
-                            ( fh[idx_fh_F(iF, jF, kF - 2, ex)]
-                             -EIT * fh[idx_fh_F(iF, jF, kF - 1, ex)]
-                             +EIT * fh[idx_fh_F(iF, jF, kF + 1, ex)]
-                             -      fh[idx_fh_F(iF, jF, kF + 2, ex)]);
-                    } else if (k0 >= kminF) {
-                        f_rhs[p] += sfz * d12dz *
-                            (-F3  * fh[idx_fh_F(iF, jF, kF - 1, ex)]
-                             -F10 * fh[idx_fh_F(iF, jF, kF    , ex)]
-                             +F18 * fh[idx_fh_F(iF, jF, kF + 1, ex)]
-                             -F6  * fh[idx_fh_F(iF, jF, kF + 2, ex)]
-                             +      fh[idx_fh_F(iF, jF, kF + 3, ex)]);
-                    }
-                }
-            }
-        }
-    }
-
-    // KO dissipation (same domain restriction as kodiss_c.C)
-    if (eps > ZEO) {
-        const int i0_lo = (iminF + 2 > 0) ? iminF + 2 : 0;
-        const int j0_lo = (jminF + 2 > 0) ? jminF + 2 : 0;
-        const int k0_lo = (kminF + 2 > 0) ? kminF + 2 : 0;
-        const int i0_hi = imaxF - 4; // inclusive
-        const int j0_hi = jmaxF - 4;
-        const int k0_hi = kmaxF - 4;
-
-        if (!(i0_lo > i0_hi || j0_lo > j0_hi || k0_lo > k0_hi)) {
-            for (int k0 = k0_lo; k0 <= k0_hi; ++k0) {
-                const int kF = k0 + 1;
-                for (int j0 = j0_lo; j0 <= j0_hi; ++j0) {
-                    const int jF = j0 + 1;
-                    for (int i0 = i0_lo; i0 <= i0_hi; ++i0) {
-                        const int iF = i0 + 1;
-                        const size_t p = idx_ex(i0, j0, k0, ex);
-
-                        const double Dx_term =
-                            ((fh[idx_fh_F(iF - 3, jF, kF, ex)] + fh[idx_fh_F(iF + 3, jF, kF, ex)]) -
-                             SIX * (fh[idx_fh_F(iF - 2, jF, kF, ex)] + fh[idx_fh_F(iF + 2, jF, kF, ex)]) +
-                             FIT * (fh[idx_fh_F(iF - 1, jF, kF, ex)] + fh[idx_fh_F(iF + 1, jF, kF, ex)]) -
-                             TWT *  fh[idx_fh_F(iF,     jF, kF, ex)]) / dX;
-
-                        const double Dy_term =
-                            ((fh[idx_fh_F(iF, jF - 3, kF, ex)] + fh[idx_fh_F(iF, jF + 3, kF, ex)]) -
-                             SIX * (fh[idx_fh_F(iF, jF - 2, kF, ex)] + fh[idx_fh_F(iF, jF + 2, kF, ex)]) +
-                             FIT * (fh[idx_fh_F(iF, jF - 1, kF, ex)] + fh[idx_fh_F(iF, jF + 1, kF, ex)]) -
-                             TWT *  fh[idx_fh_F(iF, jF,     kF, ex)]) / dY;
-
-                        const double Dz_term =
-                            ((fh[idx_fh_F(iF, jF, kF - 3, ex)] + fh[idx_fh_F(iF, jF, kF + 3, ex)]) -
-                             SIX * (fh[idx_fh_F(iF, jF, kF - 2, ex)] + fh[idx_fh_F(iF, jF, kF + 2, ex)]) +
-                             FIT * (fh[idx_fh_F(iF, jF, kF - 1, ex)] + fh[idx_fh_F(iF, jF, kF + 1, ex)]) -
-                             TWT *  fh[idx_fh_F(iF, jF, kF,     ex)]) / dZ;
-
-                        f_rhs[p] += (eps / cof) * (Dx_term + Dy_term + Dz_term);
-                    }
-                }
-            }
-        }
-    }
-
-    free(fh);
-}
--- a/AMSS_NCKU_source/makefile
+++ b/AMSS_NCKU_source/makefile
@@ -1,35 +1,33 @@


-include makefile.inc
-
-## polint(ordn=6) kernel selector:
-##   1 (default): barycentric fast path
-##   0          : fallback to Neville path
-POLINT6_USE_BARY ?= 1
-POLINT6_FLAG = -DPOLINT6_USE_BARYCENTRIC=$(POLINT6_USE_BARY)
-
-## ABE build flags selected by PGO_MODE (set in makefile.inc, default: opt)
-##   make                        -> opt  (PGO-guided, maximum performance)
-##   make PGO_MODE=instrument    -> instrument (Phase 1: collect fresh profile data)
-PROFDATA = /home/$(shell whoami)/AMSS-NCKU/pgo_profile/default.profdata
+include makefile.inc
+
+## polint(ordn=6) kernel selector:
+##   1 (default): barycentric fast path
+##   0          : fallback to Neville path
+POLINT6_USE_BARY ?= 1
+POLINT6_FLAG = -DPOLINT6_USE_BARYCENTRIC=$(POLINT6_USE_BARY)
+
+## ABE build flags selected by PGO_MODE (set in makefile.inc, default: opt)
+##   make                        -> opt  (PGO-guided, maximum performance)
+##   make PGO_MODE=instrument    -> instrument (Phase 1: collect fresh profile data)
+PROFDATA = /home/$(shell whoami)/AMSS-NCKU/pgo_profile/default.profdata

 ifeq ($(PGO_MODE),instrument)
 ## Phase 1: instrumentation — omit -ipo/-fp-model fast=2 for faster build and numerical stability
-CXXAPPFLAGS = -O3 -xHost -fma -fprofile-instr-generate -ipo \
-              -Dfortran3 -Dnewc -I${MKLROOT}/include $(INTERP_LB_FLAGS)
-f90appflags = -O3 -xHost -fma -fprofile-instr-generate -ipo \
-              -align array64byte -fpp -I${MKLROOT}/include $(POLINT6_FLAG)
-else
-## opt (default): maximum performance with PGO profile data -fprofile-instr-use=$(PROFDATA) \
-## PGO has been turned off, now tested and found to be negative optimization
-## INTERP_LB_FLAGS has been turned off too, now tested and found to be negative optimization
-
-
-CXXAPPFLAGS = -O3 -xHost -fp-model fast=2 -fma -ipo \
-              -Dfortran3 -Dnewc -I${MKLROOT}/include $(INTERP_LB_FLAGS)
-f90appflags = -O3 -xHost -fp-model fast=2 -fma -ipo \
-              -align array64byte -fpp -I${MKLROOT}/include $(POLINT6_FLAG)
-endif
+CXXAPPFLAGS = -O3 -xHost -fma -fprofile-instr-generate -ipo \
+              -Dfortran3 -Dnewc -I${MKLROOT}/include $(INTERP_LB_FLAGS)
+f90appflags = -O3 -xHost -fma -fprofile-instr-generate -ipo \
+              -align array64byte -fpp -I${MKLROOT}/include $(POLINT6_FLAG)
+else
+## opt (default): maximum performance with PGO profile data
+CXXAPPFLAGS = -O3 -xHost -fp-model fast=2 -fma -ipo \
+              -fprofile-instr-use=$(PROFDATA) \
+              -Dfortran3 -Dnewc -I${MKLROOT}/include $(INTERP_LB_FLAGS)
+f90appflags = -O3 -xHost -fp-model fast=2 -fma -ipo \
+              -fprofile-instr-use=$(PROFDATA) \
+              -align array64byte -fpp -I${MKLROOT}/include $(POLINT6_FLAG)
+endif

 .SUFFIXES: .o .f90 .C .for .cu

@@ -45,6 +43,10 @@ endif
 .cu.o:
 	$(Cu) $(CUDA_APP_FLAGS) -c $< -o $@ $(CUDA_LIB_PATH)

+# CUDA rewrite of BSSN RHS (drop-in replacement for bssn_rhs_c + stencil helpers)
+bssn_rhs_cuda.o: bssn_rhs_cuda.cu macrodef.h
+	$(Cu) $(CUDA_APP_FLAGS) -c $< -o $@ $(CUDA_LIB_PATH)
+
 # C rewrite of BSSN RHS kernel and helpers
 bssn_rhs_c.o: bssn_rhs_c.C
 	${CXX} $(CXXAPPFLAGS) -c $< $(filein) -o $@
@@ -58,14 +60,11 @@ fdderivs_c.o: fdderivs_c.C
 kodiss_c.o: kodiss_c.C
 	${CXX} $(CXXAPPFLAGS) -c $< $(filein) -o $@

-lopsided_c.o: lopsided_c.C
-	${CXX} $(CXXAPPFLAGS) -c $< $(filein) -o $@
-
-lopsided_kodis_c.o: lopsided_kodis_c.C
-	${CXX} $(CXXAPPFLAGS) -c $< $(filein) -o $@
-
-interp_lb_profile.o: interp_lb_profile.C interp_lb_profile.h
-	${CXX} $(CXXAPPFLAGS) -c $< $(filein) -o $@
+lopsided_c.o: lopsided_c.C
+	${CXX} $(CXXAPPFLAGS) -c $< $(filein) -o $@
+
+interp_lb_profile.o: interp_lb_profile.C interp_lb_profile.h
+	${CXX} $(CXXAPPFLAGS) -c $< $(filein) -o $@

 ## TwoPunctureABE uses fixed optimal flags with its own PGO profile, independent of CXXAPPFLAGS
 TP_PROFDATA = /home/$(shell whoami)/AMSS-NCKU/pgo_profile/TwoPunctureABE.profdata
@@ -82,21 +81,25 @@ TwoPunctureABE.o: TwoPunctureABE.C
 # Input files

 ## Kernel implementation switch (set USE_CXX_KERNELS=0 to fall back to Fortran)
-ifeq ($(USE_CXX_KERNELS),0)
-# Fortran mode: no C rewrite files; bssn_rhs.o is included via F90FILES below
-CFILES =
-else
-# C++ mode (default): C rewrite of bssn_rhs and helper kernels
-CFILES = bssn_rhs_c.o fderivs_c.o fdderivs_c.o kodiss_c.o lopsided_c.o lopsided_kodis_c.o
-endif
-
-## RK4 kernel switch (independent from USE_CXX_KERNELS)
-ifeq ($(USE_CXX_RK4),1)
-CFILES += rungekutta4_rout_c.o
-RK4_F90_OBJ =
-else
-RK4_F90_OBJ = rungekutta4_rout.o
-endif
+ifeq ($(USE_CXX_KERNELS),0)
+# Fortran mode: no C rewrite files; bssn_rhs.o is included via F90FILES below
+CFILES =
+else
+# C++ mode (default): C rewrite of bssn_rhs and helper kernels
+CFILES = bssn_rhs_c.o fderivs_c.o fdderivs_c.o kodiss_c.o lopsided_c.o
+endif
+
+# CUDA rewrite: bssn_rhs_cuda.o replaces all CFILES (stencils are built-in)
+CFILES_CUDA = bssn_rhs_cuda.o
+
+## RK4 kernel switch (independent from USE_CXX_KERNELS)
+ifeq ($(USE_CXX_RK4),1)
+CFILES += rungekutta4_rout_c.o
+CFILES_CUDA += rungekutta4_rout_c.o
+RK4_F90_OBJ =
+else
+RK4_F90_OBJ = rungekutta4_rout.o
+endif

 C++FILES = ABE.o Ansorg.o Block.o misc.o monitor.o Parallel.o MPatch.o var.o\
           cgh.o bssn_class.o surface_integral.o ShellPatch.o\
@@ -113,12 +116,12 @@ C++FILES_GPU = ABE.o Ansorg.o Block.o misc.o monitor.o Parallel.o MPatch.o var.o
 	   NullShellPatch2_Evo.o \
 	   bssn_gpu_class.o bssn_step_gpu.o bssn_macro.o writefile_f.o

-F90FILES_BASE = enforce_algebra.o fmisc.o initial_puncture.o prolongrestrict.o\
-	   prolongrestrict_cell.o prolongrestrict_vertex.o\
-	   $(RK4_F90_OBJ) diff_new.o kodiss.o kodiss_sh.o\
-	   lopsidediff.o sommerfeld_rout.o getnp4.o diff_new_sh.o\
-	   shellfunctions.o bssn_rhs_ss.o Set_Rho_ADM.o\
-           getnp4EScalar.o bssnEScalar_rhs.o bssn_constraint.o ricci_gamma.o\
+F90FILES_BASE = enforce_algebra.o fmisc.o initial_puncture.o prolongrestrict.o\
+	   prolongrestrict_cell.o prolongrestrict_vertex.o\
+	   $(RK4_F90_OBJ) diff_new.o kodiss.o kodiss_sh.o\
+	   lopsidediff.o sommerfeld_rout.o getnp4.o diff_new_sh.o\
+	   shellfunctions.o bssn_rhs_ss.o Set_Rho_ADM.o\
+           getnp4EScalar.o bssnEScalar_rhs.o bssn_constraint.o ricci_gamma.o\
           fadmquantites_bssn.o Z4c_rhs.o Z4c_rhs_ss.o point_diff_new_sh.o\
 	   cpbc.o getnp4old.o NullEvol.o initial_null.o initial_maxwell.o\
 	   getnpem2.o empart.o NullNews.o fourdcurvature.o\
@@ -178,9 +181,12 @@ $(CUDAFILES): bssn_gpu.h gpu_mem.h gpu_rhsSS_mem.h
 misc.o : zbesh.o

 # projects
-ABE: $(C++FILES) $(CFILES) $(F90FILES) $(F77FILES) $(AHFDOBJS)
-	$(CLINKER) $(CXXAPPFLAGS) -o $@ $(C++FILES) $(CFILES) $(F90FILES) $(F77FILES) $(AHFDOBJS) $(LDLIBS)
-	
+ABE: $(C++FILES) $(CFILES_CUDA) $(F90FILES) $(F77FILES) $(AHFDOBJS)
+	$(CLINKER) $(CXXAPPFLAGS) -o $@ $(C++FILES) $(CFILES_CUDA) $(F90FILES) $(F77FILES) $(AHFDOBJS) $(LDLIBS) -lcudart $(CUDA_LIB_PATH)
+
+ABE_CUDA: $(C++FILES) $(CFILES_CUDA) $(F90FILES) $(F77FILES) $(AHFDOBJS)
+	$(CLINKER) $(CXXAPPFLAGS) -o $@ $(C++FILES) $(CFILES_CUDA) $(F90FILES) $(F77FILES) $(AHFDOBJS) $(LDLIBS) -lcudart $(CUDA_LIB_PATH)
+
 ABEGPU: $(C++FILES_GPU) $(CFILES) $(F90FILES) $(F77FILES) $(AHFDOBJS) $(CUDAFILES)
 	$(CLINKER) $(CXXAPPFLAGS) -o $@ $(C++FILES_GPU) $(CFILES) $(F90FILES) $(F77FILES) $(AHFDOBJS) $(CUDAFILES) $(LDLIBS)

@@ -188,4 +194,4 @@ TwoPunctureABE: $(TwoPunctureFILES)
 	$(CLINKER) $(TP_OPTFLAGS) -qopenmp -o $@ $(TwoPunctureFILES) $(LDLIBS)

 clean:
-	rm *.o ABE ABEGPU TwoPunctureABE make.log -f
+	rm *.o ABE ABE_CUDA ABEGPU TwoPunctureABE make.log -f
--- a/AMSS_NCKU_source/makefile.inc
+++ b/AMSS_NCKU_source/makefile.inc
@@ -62,4 +62,4 @@ CLINKER      = mpiicpx
 Cu = nvcc
 CUDA_LIB_PATH = -L/usr/lib/cuda/lib64 -I/usr/include -I/usr/lib/cuda/include
 #CUDA_APP_FLAGS = -c -g -O3 --ptxas-options=-v -arch compute_13 -code compute_13,sm_13 -Dfortran3 -Dnewc
-CUDA_APP_FLAGS = -c -g -O3 --ptxas-options=-v -Dfortran3 -Dnewc
+CUDA_APP_FLAGS = -c -g -O3 --ptxas-options=-v -Dfortran3 -Dnewc -arch=sm_80
--- a/AMSS_NCKU_source/prolongrestrict_cell.f90
+++ b/AMSS_NCKU_source/prolongrestrict_cell.f90
@@ -217,7 +217,6 @@
  real*8,dimension(2*ghost_width) :: X,Y,Z
  real*8, dimension(2*ghost_width,2*ghost_width) :: tmp2
  real*8, dimension(2*ghost_width) :: tmp1
-  real*8 :: ddy
  real*8,dimension(3) :: ccp

 #if (ghost_width == 2)
@@ -580,7 +579,7 @@
            tmp1(ghost_width-cxI(1)+cxB(1)  :ghost_width-cxI(1)+cxT(1)  ) = funf(cxB(1):cxT(1),j,k)
         endif

-         call polint(X,tmp1,0.d0,funf(i,j,k),ddy,2*ghost_width)
+         call polint0(X,tmp1,funf(i,j,k),2*ghost_width)

 ! for y direction
       elseif(sum(fg).eq.2.and.fg(2) .eq. 0.and. &
@@ -690,7 +689,7 @@
            tmp1(ghost_width-cxI(2)+cxB(2)  :ghost_width-cxI(2)+cxT(2)  ) = funf(i,cxB(2):cxT(2),k)
         endif

-         call polint(Y,tmp1,0.d0,funf(i,j,k),ddy,2*ghost_width)
+         call polint0(Y,tmp1,funf(i,j,k),2*ghost_width)

 ! for z direction
       elseif(sum(fg).eq.2.and.fg(3) .eq. 0.and. &
@@ -802,7 +801,7 @@
            tmp1(ghost_width-cxI(3)+cxB(3)  :ghost_width-cxI(3)+cxT(3)  ) = funf(i,j,cxB(3):cxT(3))
         endif

-         call polint(Z,tmp1,0.d0,funf(i,j,k),ddy,2*ghost_width)
+         call polint0(Z,tmp1,funf(i,j,k),2*ghost_width)

 #else

@@ -1956,11 +1955,13 @@

  real*8,dimension(3) :: CD,FD
  real*8 :: tmp_yz(extc(1), 6)      ! 存储整条 X 线上 6 个 Y 轴偏置的 Z 向插值结果
-  real*8 :: tmp_xyz_line(extc(1))   ! 存储整条 X 线上完成 Y 向融合后的结果
+  real*8 :: tmp_xyz_line(-2:extc(1))   ! 包含 X 向 6 点模板访问所需下界
  real*8 :: v1, v2, v3, v4, v5, v6
-  integer :: ic, jc, kc, ix_offset,ix,iy,iz,jc_min,jc_max
+  integer :: ic, jc, kc, ix_offset,ix,iy,iz,jc_min,jc_max,ic_min,ic_max,kc_min,kc_max
+  integer :: i_lo, i_hi, j_lo, j_hi, k_lo, k_hi
+  logical :: need_full_symmetry
  real*8 :: res_line
-  real*8 :: tmp_z_slab(extc(1), extc(2))  ! 分配在 k 循环外
+  real*8 :: tmp_z_slab(-2:extc(1), -2:extc(2))  ! 包含 Y/X 向模板访问所需下界
  if(wei.ne.3)then
     write(*,*)"prolongrestrict.f90::prolong3: this routine only surport 3 dimension"
     write(*,*)"dim = ",wei
@@ -2063,24 +2064,41 @@
     endif
  enddo

-  maxcx = maxval(cix(imino:imaxo))
-  maxcy = maxval(ciy(jmino:jmaxo))
-  maxcz = maxval(ciz(kmino:kmaxo))
+  ic_min = minval(cix(imino:imaxo))
+  ic_max = maxval(cix(imino:imaxo))
+  jc_min = minval(ciy(jmino:jmaxo))
+  jc_max = maxval(ciy(jmino:jmaxo))
+  kc_min = minval(ciz(kmino:kmaxo))
+  kc_max = maxval(ciz(kmino:kmaxo))
+
+  maxcx = ic_max
+  maxcy = jc_max
+  maxcz = kc_max
  if(maxcx+3 > extc(1) .or. maxcy+3 > extc(2) .or. maxcz+3 > extc(3))then
     write(*,*)"error in prolong"
     return
  endif

-  call symmetry_bd(3,extc,func,funcc,SoA)
+  i_lo = ic_min - 2
+  i_hi = ic_max + 3
+  j_lo = jc_min - 2
+  j_hi = jc_max + 3
+  k_lo = kc_min - 2
+  k_hi = kc_max + 3
+  need_full_symmetry = (i_lo < 1) .or. (j_lo < 1) .or. (k_lo < 1)
+  if(need_full_symmetry)then
+     call symmetry_bd(3,extc,func,funcc,SoA)
+  else
+     funcc(i_lo:i_hi,j_lo:j_hi,k_lo:k_hi) = func(i_lo:i_hi,j_lo:j_hi,k_lo:k_hi)
+  endif
+
     ! 对每个 k（pz, kc 固定）预计算 Z 向插值的 2D 切片
-jc_min = minval(ciy(jmino:jmaxo))
-jc_max = maxval(ciy(jmino:jmaxo))

 do k = kmino, kmaxo
    pz = piz(k); kc = ciz(k)
    ! --- Pass 1: Z 方向，只算一次 ---
-    do iy = jc_min-3, jc_max+3   ! 仅需的 iy 范围
-        do ii = imini-3, imaxi+3  ! 仅需的 ii 范围
+    do iy = jc_min-2, jc_max+3   ! 仅需的 iy 范围（对应 jc-2:jc+3）
+        do ii = ic_min-2, ic_max+3  ! 仅需的 ii 范围（对应 cix-2:cix+3）
            tmp_z_slab(ii, iy) = sum(WC(:,pz) * funcc(ii, iy, kc-2:kc+3))
        end do
    end do
@@ -2088,7 +2106,7 @@ do k = kmino, kmaxo
    do j = jmino, jmaxo
        py = piy(j); jc = ciy(j)
        ! --- Pass 2: Y 方向 ---
-        do ii = imini-3, imaxi+3
+        do ii = ic_min-2, ic_max+3
            tmp_xyz_line(ii) = sum(WC(:,py) * tmp_z_slab(ii, jc-2:jc+3))
        end do
        ! --- Pass 3: X 方向 ---
@@ -2351,9 +2369,12 @@ end do

  real*8,dimension(3) :: CD,FD

-  real*8 :: tmp_xz_plane(extf(1), 6) 
-  real*8 :: tmp_x_line(extf(1))
+  real*8 :: tmp_xz_plane(-1:extf(1), 6)
+  real*8 :: tmp_x_line(-1:extf(1))
  integer :: fi, fj, fk, ii, jj, kk
+  integer :: fi_min, fi_max, ii_lo, ii_hi
+  integer :: fj_min, fj_max, fk_min, fk_max, jj_lo, jj_hi, kk_lo, kk_hi
+  logical :: need_full_symmetry

  if(wei.ne.3)then
     write(*,*)"prolongrestrict.f90::restrict3: this routine only surport 3 dimension"
@@ -2433,7 +2454,34 @@ end do
          stop
  endif

-  call symmetry_bd(2,extf,funf,funff,SoA)
+  ! 仅计算 X 向最终写回所需的窗口：
+  ! func(i,j,k) 只访问 tmp_x_line(fi-2:fi+3)
+  fi_min = 2*(imino + lbc(1) - 1) - 1 - lbf(1) + 1
+  fi_max = 2*(imaxo + lbc(1) - 1) - 1 - lbf(1) + 1
+  fj_min = 2*(jmino + lbc(2) - 1) - 1 - lbf(2) + 1
+  fj_max = 2*(jmaxo + lbc(2) - 1) - 1 - lbf(2) + 1
+  fk_min = 2*(kmino + lbc(3) - 1) - 1 - lbf(3) + 1
+  fk_max = 2*(kmaxo + lbc(3) - 1) - 1 - lbf(3) + 1
+  ii_lo = fi_min - 2
+  ii_hi = fi_max + 3
+  jj_lo = fj_min - 2
+  jj_hi = fj_max + 3
+  kk_lo = fk_min - 2
+  kk_hi = fk_max + 3
+  if(ii_lo < -1 .or. ii_hi > extf(1) .or. &
+     jj_lo < -1 .or. jj_hi > extf(2) .or. &
+     kk_lo < -1 .or. kk_hi > extf(3))then
+      write(*,*)"restrict3: invalid stencil window"
+      write(*,*)"ii=",ii_lo,ii_hi," jj=",jj_lo,jj_hi," kk=",kk_lo,kk_hi
+      write(*,*)"extf=",extf
+      stop
+  endif
+  need_full_symmetry = (ii_lo < 1) .or. (jj_lo < 1) .or. (kk_lo < 1)
+  if(need_full_symmetry)then
+      call symmetry_bd(2,extf,funf,funff,SoA)
+  else
+      funff(ii_lo:ii_hi,jj_lo:jj_hi,kk_lo:kk_hi) = funf(ii_lo:ii_hi,jj_lo:jj_hi,kk_lo:kk_hi)
+  endif

 !~~~~~~> restriction start...
 do k = kmino, kmaxo
@@ -2445,7 +2493,7 @@ do k = kmino, kmaxo
        ! 优化点 1: 显式展开 Z 方向计算，减少循环开销
        ! 确保 ii 循环是最内层且连续访问
        !DIR$ VECTOR ALWAYS
-        do ii = 1, extf(1)
+        do ii = ii_lo, ii_hi
            ! 预计算当前 j 对应的 6 行在 Z 方向的压缩结果
            ! 这里直接硬编码 jj 的偏移，彻底消除一层循环
            tmp_xz_plane(ii, 1) = C1*(funff(ii,fj-2,fk-2)+funff(ii,fj-2,fk+3)) + &
@@ -2470,7 +2518,7 @@ do k = kmino, kmaxo

        ! 优化点 2: 同样向量化 Y 方向压缩
        !DIR$ VECTOR ALWAYS
-        do ii = 1, extf(1)
+        do ii = ii_lo, ii_hi
            tmp_x_line(ii) = C1*(tmp_xz_plane(ii, 1) + tmp_xz_plane(ii, 6)) + &
                            C2*(tmp_xz_plane(ii, 2) + tmp_xz_plane(ii, 5)) + &
                            C3*(tmp_xz_plane(ii, 3) + tmp_xz_plane(ii, 4))
--- a/AMSS_NCKU_source/prolongrestrict_vertex.f90
+++ b/AMSS_NCKU_source/prolongrestrict_vertex.f90
@@ -217,7 +217,6 @@
  real*8,dimension(2*ghost_width) :: X,Y,Z
  real*8, dimension(2*ghost_width,2*ghost_width) :: tmp2
  real*8, dimension(2*ghost_width) :: tmp1
-  real*8 :: ddy

 #if (ghost_width == 2)
  real*8, parameter :: C1=-1.d0/16,C2=9.d0/16
@@ -470,7 +469,7 @@

         tmp1(cxB(1)+ghost_width-i+1:cxT(1)+ghost_width-i+1) = fh(cxB(1):cxT(1),j,k)

-         call polint(X,tmp1,0.d0,funf(i,j,k),ddy,2*ghost_width)
+         call polint0(X,tmp1,funf(i,j,k),2*ghost_width)

 ! for y direction
       elseif (fg(2) .eq. 0)then
@@ -529,7 +528,7 @@

         tmp1(cxB(2)+ghost_width-j+1:cxT(2)+ghost_width-j+1) = fh(i,cxB(2):cxT(2),k)

-         call polint(Y,tmp1,0.d0,funf(i,j,k),ddy,2*ghost_width)
+         call polint0(Y,tmp1,funf(i,j,k),2*ghost_width)

 ! for z direction
       else
@@ -588,7 +587,7 @@

         tmp1(cxB(3)+ghost_width-k+1:cxT(3)+ghost_width-k+1) = fh(i,j,cxB(3):cxT(3))

-         call polint(Z,tmp1,0.d0,funf(i,j,k),ddy,2*ghost_width)
+         call polint0(Z,tmp1,funf(i,j,k),2*ghost_width)

       endif
      
--- a/AMSS_NCKU_source/rungekutta4_rout_c.C
+++ b/AMSS_NCKU_source/rungekutta4_rout_c.C
@@ -2,7 +2,6 @@
 #include <cstdio>
 #include <cstdlib>
 #include <cstddef>
-#include <complex>
 #include <immintrin.h>

 namespace {
@@ -118,62 +117,6 @@ inline void rk4_stage3(std::size_t n,

 extern "C" {

-void f_rungekutta4_scalar(double &dT, double &f0, double &f1, double &f_rhs, int &RK4) {
-    constexpr double F1o6 = 1.0 / 6.0;
-    constexpr double HLF = 0.5;
-    constexpr double TWO = 2.0;
-
-    switch (RK4) {
-    case 0:
-        f1 = f0 + HLF * dT * f_rhs;
-        break;
-    case 1:
-        f_rhs = f_rhs + TWO * f1;
-        f1 = f0 + HLF * dT * f1;
-        break;
-    case 2:
-        f_rhs = f_rhs + TWO * f1;
-        f1 = f0 + dT * f1;
-        break;
-    case 3:
-        f1 = f0 + F1o6 * dT * (f1 + f_rhs);
-        break;
-    default:
-        std::fprintf(stderr, "rungekutta4_scalar_c: invalid RK4 stage %d\n", RK4);
-        std::abort();
-    }
-}
-
-void rungekutta4_cplxscalar_(double &dT,
-                             std::complex<double> &f0,
-                             std::complex<double> &f1,
-                             std::complex<double> &f_rhs,
-                             int &RK4) {
-    constexpr double F1o6 = 1.0 / 6.0;
-    constexpr double HLF = 0.5;
-    constexpr double TWO = 2.0;
-
-    switch (RK4) {
-    case 0:
-        f1 = f0 + HLF * dT * f_rhs;
-        break;
-    case 1:
-        f_rhs = f_rhs + TWO * f1;
-        f1 = f0 + HLF * dT * f1;
-        break;
-    case 2:
-        f_rhs = f_rhs + TWO * f1;
-        f1 = f0 + dT * f1;
-        break;
-    case 3:
-        f1 = f0 + F1o6 * dT * (f1 + f_rhs);
-        break;
-    default:
-        std::fprintf(stderr, "rungekutta4_cplxscalar_c: invalid RK4 stage %d\n", RK4);
-        std::abort();
-    }
-}
-
 int f_rungekutta4_rout(int *ex, double &dT,
                       double *f0, double *f1, double *f_rhs,
                       int &RK4) {
--- a/AMSS_NCKU_source/tool.h
+++ b/AMSS_NCKU_source/tool.h
@@ -24,10 +24,4 @@ void lopsided(const int ex[3],
              const double *X, const double *Y, const double *Z,
              const double *f, double *f_rhs,
              const double *Sfx, const double *Sfy, const double *Sfz,
-              int Symmetry, const double SoA[3]);
-
-void lopsided_kodis(const int ex[3],
-                    const double *X, const double *Y, const double *Z,
-                    const double *f, double *f_rhs,
-                    const double *Sfx, const double *Sfy, const double *Sfz,
-                    int Symmetry, const double SoA[3], double eps);
+              int Symmetry, const double SoA[3]);
--- a/makefile_and_run.py
+++ b/makefile_and_run.py
@@ -70,7 +70,7 @@ def makefile_ABE():

    ## Build command with CPU binding to nohz_full cores
    if (input_data.GPU_Calculation == "no"):
-        makefile_command  = f"{NUMACTL_CPU_BIND} make -j{BUILD_JOBS} INTERP_LB_MODE=off ABE"
+        makefile_command  = f"{NUMACTL_CPU_BIND} make -j{BUILD_JOBS} INTERP_LB_MODE=optimize ABE"
    elif (input_data.GPU_Calculation == "yes"):
        makefile_command  = f"{NUMACTL_CPU_BIND} make -j{BUILD_JOBS} ABEGPU"
    else:
--- a/pgo_profile/PGO_Profile_Analysis.md
+++ b/pgo_profile/PGO_Profile_Analysis.md
@@ -0,0 +1,97 @@
+# AMSS-NCKU PGO Profile Analysis Report
+
+## 1. Profiling Environment
+
+| Item | Value |
+|------|-------|
+| Compiler | Intel oneAPI DPC++/C++ 2025.3.0 (icpx/ifx) |
+| Instrumentation Flag | `-fprofile-instr-generate` |
+| Optimization Level (instrumented) | `-O2 -xHost -fma` |
+| MPI Processes | 1 (single process to avoid MPI+instrumentation deadlock) |
+| Profile File | `default_9725750769337483397_0.profraw` (327 KB) |
+| Merged Profile | `default.profdata` (394 KB) |
+| llvm-profdata | `/home/intel/oneapi/compiler/2025.3/bin/compiler/llvm-profdata` |
+
+## 2. Reduced Simulation Parameters (for profiling run)
+
+| Parameter | Production Value | Profiling Value |
+|-----------|-----------------|-----------------|
+| MPI_processes | 64 | 1 |
+| grid_level | 9 | 4 |
+| static_grid_level | 5 | 3 |
+| static_grid_number | 96 | 24 |
+| moving_grid_number | 48 | 16 |
+| largest_box_xyz_max | 320^3 | 160^3 |
+| Final_Evolution_Time | 1000.0 | 10.0 |
+| Evolution_Step_Number | 10,000,000 | 1,000 |
+| Detector_Number | 12 | 2 |
+
+## 3. Profile Summary
+
+| Metric | Value |
+|--------|-------|
+| Total instrumented functions | 1,392 |
+| Functions with non-zero counts | 117 (8.4%) |
+| Functions with zero counts | 1,275 (91.6%) |
+| Maximum function entry count | 386,459,248 |
+| Maximum internal block count | 370,477,680 |
+| Total block count | 4,198,023,118 |
+
+## 4. Top 20 Hotspot Functions
+
+| Rank | Total Count | Max Block Count | Function | Category |
+|------|------------|-----------------|----------|----------|
+| 1 | 1,241,601,732 | 370,477,680 | `polint_` | Interpolation |
+| 2 | 755,994,435 | 230,156,640 | `prolong3_` | Grid prolongation |
+| 3 | 667,964,095 | 3,697,792 | `compute_rhs_bssn_` | BSSN RHS evolution |
+| 4 | 539,736,051 | 386,459,248 | `symmetry_bd_` | Symmetry boundary |
+| 5 | 277,310,808 | 53,170,728 | `lopsided_` | Lopsided FD stencil |
+| 6 | 155,534,488 | 94,535,040 | `decide3d_` | 3D grid decision |
+| 7 | 119,267,712 | 19,266,048 | `rungekutta4_rout_` | RK4 time integrator |
+| 8 | 91,574,616 | 48,824,160 | `kodis_` | Kreiss-Oliger dissipation |
+| 9 | 67,555,389 | 43,243,680 | `fderivs_` | Finite differences |
+| 10 | 55,296,000 | 42,246,144 | `misc::fact(int)` | Factorial utility |
+| 11 | 43,191,071 | 27,663,328 | `fdderivs_` | 2nd-order FD derivatives |
+| 12 | 36,233,965 | 22,429,440 | `restrict3_` | Grid restriction |
+| 13 | 24,698,512 | 17,231,520 | `polin3_` | Polynomial interpolation |
+| 14 | 22,962,942 | 20,968,768 | `copy_` | Data copy |
+| 15 | 20,135,696 | 17,259,168 | `Ansorg::barycentric(...)` | Spectral interpolation |
+| 16 | 14,650,224 | 7,224,768 | `Ansorg::barycentric_omega(...)` | Spectral weights |
+| 17 | 13,242,296 | 2,871,920 | `global_interp_` | Global interpolation |
+| 18 | 12,672,000 | 7,734,528 | `sommerfeld_rout_` | Sommerfeld boundary |
+| 19 | 6,872,832 | 1,880,064 | `sommerfeld_routbam_` | Sommerfeld boundary (BAM) |
+| 20 | 5,709,900 | 2,809,632 | `l2normhelper_` | L2 norm computation |
+
+## 5. Hotspot Category Breakdown
+
+Top 20 functions account for ~98% of total execution counts:
+
+| Category | Functions | Combined Count | Share |
+|----------|-----------|---------------|-------|
+| Interpolation / Prolongation / Restriction | polint_, prolong3_, restrict3_, polin3_, global_interp_, Ansorg::* | ~2,093M | ~50% |
+| BSSN RHS + FD stencils | compute_rhs_bssn_, lopsided_, fderivs_, fdderivs_ | ~1,056M | ~25% |
+| Boundary conditions | symmetry_bd_, sommerfeld_rout_, sommerfeld_routbam_ | ~559M | ~13% |
+| Time integration | rungekutta4_rout_ | ~119M | ~3% |
+| Dissipation | kodis_ | ~92M | ~2% |
+| Utilities | misc::fact, decide3d_, copy_, l2normhelper_ | ~256M | ~6% |
+
+## 6. Conclusions
+
+1. **Profile data is valid**: 1,392 functions instrumented, 117 exercised with ~4.2 billion total counts.
+2. **Hotspot concentration is high**: Top 5 functions alone account for ~76% of all counts, which is ideal for PGO — the compiler has strong branch/layout optimization targets.
+3. **Fortran numerical kernels dominate**: `polint_`, `prolong3_`, `compute_rhs_bssn_`, `symmetry_bd_`, `lopsided_` are all Fortran routines in the inner evolution loop. PGO will optimize their branch prediction and basic block layout.
+4. **91.6% of functions have zero counts**: These are code paths for unused features (GPU, BSSN-EScalar, BSSN-EM, Z4C, etc.). PGO will deprioritize them, improving instruction cache utilization.
+5. **Profile is representative**: Despite the reduced grid size, the code path coverage matches production — the same kernels (RHS, prolongation, restriction, boundary) are exercised. PGO branch probabilities from this profile will transfer well to full-scale runs.
+
+## 7. PGO Phase 2 Usage
+
+To apply the profile, use the following flags in `makefile.inc`:
+
+```makefile
+CXXAPPFLAGS = -O3 -xHost -fp-model fast=2 -fma -ipo \
+              -fprofile-instr-use=/home/amss/AMSS-NCKU/pgo_profile/default.profdata \
+              -Dfortran3 -Dnewc -I${MKLROOT}/include
+f90appflags = -O3 -xHost -fp-model fast=2 -fma -ipo \
+              -fprofile-instr-use=/home/amss/AMSS-NCKU/pgo_profile/default.profdata \
+              -align array64byte -fpp -I${MKLROOT}/include
+```
--- a/pgo_profile/default.profdata
+++ b/pgo_profile/default.profdata
--- a/pgo_profile/default_9726853898452064389_0.profdata
+++ b/pgo_profile/default_9726853898452064389_0.profdata
--- a/pgo_profile/default.profdata.backup
+++ b/pgo_profile/default.profdata.backup
--- a/pgo_profile/default.profdata.backup2
+++ b/pgo_profile/default.profdata.backup2
--- a/pgo_profile/default.profdatabackup3
+++ b/pgo_profile/default.profdatabackup3
--- a/pgo_profile/default_15874826282416242821_0_58277.profraw
+++ b/pgo_profile/default_15874826282416242821_0_58277.profraw
--- a/pgo_profile/default_9725750769337483397_0.profraw
+++ b/pgo_profile/default_9725750769337483397_0.profraw
--- a/pgo_profile/default_9725923726611433605_0.profraw
+++ b/pgo_profile/default_9725923726611433605_0.profraw
--- a/pgo_profile/default_9726420327935033477_0.profraw
+++ b/pgo_profile/default_9726420327935033477_0.profraw
Author	SHA1	Message	Date
ianchb	43975017eb	prolong3 改为先算实际 stencil 窗口；只有窗口触及对称边界时才走全域 symmetry_bd，否则只复制必需窗口。restrict3 同样改成窗口判定，无触边时仅填 ii/jj/kk 必需窗口。	2026-03-02 18:10:38 +08:00
ianchb	485667ef4c	perf(restrict3): shrink X-pass ii sweep to required overlap window - compute fi_min/fi_max from output i-range and derive ii_lo/ii_hi - replace full ii sweep (-1:extf(1)) with windowed sweep in Z/Y precompute passes - keep stencil math unchanged; add bounds sanity check for ii window	2026-03-02 18:10:38 +08:00
ianchb	2a977ce82e	perf(MPatch): 用空间 bin 索引加速 Interp_Points 的 block 归属查找 - 为 Patch::Interp_Points 三个重载引入 BlockBinIndex（候选筛选 + 全扫回退） - 保持原 point-in-block 判定与后续插值/通信流程不变 - 将逐点线性扫块从 O(N_pointsN_blocks) 降为近似 O(N_pointsk) - 测试：bin 上限如果太大，会引入不必要的索引构建开销。将 bins 上限设为 16。 Co-authored-by: gpt-5.3-codex	2026-03-02 18:10:35 +08:00
CGH0S7	160e2a0369	fix prolong/restrict index bounds after cherry-pick `12e1f63`	2026-03-02 14:11:52 +08:00
ianchb	01410de05a	refactor(Parallel): streamline MPI communication by consolidating request handling and memory management	2026-03-02 14:08:27 +08:00
jaunatisblue	83c826eb49	prolong3: 减少Z-pass 冗余计算	2026-03-02 14:08:13 +08:00
ianchb	43ddaab903	fix: add C RK4 kernel to CFILES_CUDA	2026-03-02 12:19:52 +08:00
ianchb	5839755c2f	compute div_beta on-the-fly to remove temp array	2026-03-02 12:12:58 +08:00
ianchb	a893b4007c	merge lopsided+kodis	2026-03-02 12:12:26 +08:00
CGH0S7	ad5ff03615	build: switch allocator option to oneTBB tbbmalloc (cherry picked from commit `e29ca2dca9`)	2026-03-02 11:53:30 +08:00
CGH0S7	b4bc0ef269	先关闭绑核心，发现速度对比：不绑定核心+SCX>绑核心+SCX (cherry picked from commit `1eba73acbe`)	2026-03-02 11:53:13 +08:00
CGH0S7	b185f84cce	Add switchable C RK4 kernel and build toggle (cherry picked from commit `b91cfff301`)	2026-03-02 11:53:00 +08:00
ianchb	71f6eb7b44	Remove profiling code	2026-03-02 11:29:48 +08:00
CGH0S7	90620c2aec	Optimize fdderivs: skip redundant 2nd-order work in 4th-order overlap	2026-03-02 11:04:04 +08:00
jaunatisblue	f561522d89	prolong3：提升cache命中率	2026-03-02 11:02:19 +08:00
jaunatisblue	3f4715b8cc	修改prolong	2026-03-02 11:02:17 +08:00
jaunatisblue	710ea8f76b	对prolong3做访存优化	2026-03-02 11:02:12 +08:00
CGH0S7	5cf891359d	Optimize symmetry_bd with stride-based fast paths (cherry picked from commit `16013081e0`)	2026-03-02 11:01:49 +08:00
CGH0S7	222747449a	Optimize average2: use DO CONCURRENT loop form (cherry picked from commit `1a518cd3f6`)	2026-03-02 11:01:45 +08:00
CGH0S7	14de4d535e	Optimize average2: replace array expression with explicit loops (cherry picked from commit `1dc622e516`)	2026-03-02 11:01:42 +08:00
CGH0S7	787295692a	Optimize prolong3: hoist bounds check out of inner loop (cherry picked from commit `3046a0ccde`)	2026-03-02 11:01:39 +08:00
CGH0S7	335f2f23fe	Optimize prolong3: replace parity branches with coefficient lookup (cherry picked from commit `d4ec69c98a`)	2026-03-02 11:01:37 +08:00
CGH0S7	7109474a14	Optimize prolong3: precompute coarse index/parity maps (cherry picked from commit `2c0a3055d4`)	2026-03-02 11:01:31 +08:00
CGH0S7	e7a02e8f72	perf(polint): add uniform-grid fast path for barycentric n=6	2026-03-01 14:13:51 +08:00
CGH0S7	8dad910c6c	perf(polint): add switchable barycentric ordn=6 path	2026-03-01 14:13:51 +08:00
CGH0S7	01b4cf71d1	perf(polin3): switch to lagrange-weight tensor contraction	2026-03-01 14:13:04 +08:00
CGH0S7	66dabe8cc4	perf(polint): add ordn=6 specialized neville path	2026-03-01 14:12:22 +08:00
ianchb	abf2f640e4	add fused symmetry packing kernels for orders 2 and 3 in BSSN RHS	2026-02-28 15:35:14 +08:00
ianchb	94f40627aa	refine GPU dispatch initialization and optimize H2D/D2H data transfers	2026-02-28 15:23:41 +08:00
ianchb	d94c31c5c4	[WIP]Implement multi-GPU support in BSSN RHS and add profiling for H2D/D2H transfers	2026-02-28 11:12:14 +08:00
ianchb	724e9cd415	[WIP]Add CUDA support for BSSN RHS with new kernel and update makefiles	2026-02-28 11:12:13 +08:00
ianchb	c001939461	Add Lagrange interpolation subroutine and update calls in prolongrestrict modules	2026-02-28 11:12:13 +08:00
ianchb	94d236385d	Revert "skip redundant MPI ghost cell syncs for stages 0, 1 & 2" This reverts commit `f7ada421cf`.	2026-02-28 11:12:12 +08:00
ianchb	780f1c80d0	skip redundant MPI ghost cell syncs for stages 0, 1 & 2 BSSN 每个 RK4 时间步执行 4 次 MPI ghost zone 同步： Stage 0（预测）结束后：Parallel::Sync(SynchList_pre) Stage 1（校正 1）结束后：Parallel::Sync(SynchList_cor) Stage 2（校正 2）结束后：Parallel::Sync(SynchList_cor) Stage 3（校正 3）结束后：Parallel::Sync(SynchList_cor) ← 必要（为下一步提供 ghost） bssnEM_class.C、Z4c_class.C 结构相同，一起修改了	2026-02-28 11:12:09 +08:00