embedded3 R6-1: the body's φ and the cut geometry on the device (RTX_E3_GEOM_DEVICE=1)

- e3_geom.cu: corner φ on the narrow band (the flag test's circle + capsule
  around the step's centreline polyline + span cuts + fillet union), face
  apertures / face-centre φ by the Kuhn triangles, cell volumes by the six
  Kuhn tets and wall vectors by closure — CutGeometry::build_from in fp64,
  host operation order, FMA contraction off.
- Body::with_device_sdf / DeviceSdf: the device form of φ (the host passes the
  polyline per step); the flag wake test attaches it (same arithmetic as its
  closure; the polylines factored out unchanged).
- step/device/geom.rs DeviceGeom: persistent φ / bound / volume / wall buffers;
  the face tables written straight into DeviceCut's predictor apertures and
  distances (update skips their scatters); the host mirror = the band's entries
  gathered compactly onto recycled arrays of retired device generations
  (GeomPool; band-list history) — Mask::from_cut classifies it.
- RTX_E3_BAND_CHECK=1 with the knob: mirror AND device tables against the host
  build_from bit for bit on every refresh.
- Knob off: byte-identical (slab ny 62 one period CSV = main's); host suite 21/21.

Co-Authored-By: Claude Opus 5.5 (1M context) <[email protected]>

crates/specialized/rtx-cfd/src/kernels/cuda/e3_geom.cu

@@ -0,0 +1,277 @@
+/**
+ * embedded3 R6-1: the body's φ and the cut geometry on the device — the
+ * host `cut.rs::CutGeometry::build_from` expression for expression in fp64
+ * (compiled with FMA contraction off; `/` and `sqrt` are IEEE round-to-
+ * nearest in double on the device, as on the host), and the flag test's
+ * φ (`tests/embedded3_flag_wake.rs`: the circle, the capsule around the
+ * centreline polyline, the span cut with rounded edges, the fillet union).
+ *
+ * Kernels:
+ *   e3_geom_phi    per corner: the narrow band's keep-or-evaluate (bound −
+ *                  motion > band keeps φ), φ from the polyline otherwise;
+ *   e3_geom_faces  per face of one component: the aperture by the two
+ *                  Kuhn triangles and the face-centre φ, where a corner was
+ *                  evaluated (every face on a full pass);
+ *   e3_geom_cells  per cell: the fluid volume by the six Kuhn tetrahedra
+ *                  and the wall vector by closure of the face apertures;
+ *   e3_geom_gather compact copies for the host mirror.
+ */
+
+struct GeomSdf {
+    double cx, cy, rc;      /* the circle */
+    double zc, span, r_edge;/* the span cut */
+    double half;            /* capsule half-thickness */
+    double fillet;          /* root fillet radius (0: min) */
+    int cyl_cut, flag_cut;  /* cut to the span */
+    int npts;               /* polyline points (x, y interleaved) */
+};
+
+struct GeomGrid {
+    int nx, ny, nz;
+    double dx, dy, dz;
+};
+
+__device__ __forceinline__ double rs_min(double a, double b) { return b < a ? b : a; }
+__device__ __forceinline__ double rs_max(double a, double b) { return b > a ? b : a; }
+
+/* The span cut with rounded edges (flag_3d / cylinder_3d). */
+__device__ __forceinline__ double span_cut(double d2, double z, const GeomSdf& s)
+{
+    double r = s.r_edge;
+    double q1 = d2 + r;
+    double q2 = fabs(z - s.zc) - 0.5 * s.span + r;
+    double m1 = rs_max(q1, 0.0), m2 = rs_max(q2, 0.0);
+    double outside = sqrt(m1 * m1 + m2 * m2);
+    return outside + rs_min(rs_max(q1, q2), 0.0) - r;
+}
+
+__device__ double geom_phi_at(double x, double y, double z, const GeomSdf& s, const double* __restrict__ poly)
+{
+    /* the circle */
+    double ex0 = x - s.cx, ey0 = y - s.cy;
+    double dc = sqrt(ex0 * ex0 + ey0 * ey0) - s.rc;
+    if (s.cyl_cut) dc = span_cut(dc, z, s);
+    /* the capsule: distance to the polyline */
+    double best = 1.0 / 0.0;
+    for (int m = 0; m + 1 < s.npts; ++m) {
+        double ax = poly[2 * m], ay = poly[2 * m + 1];
+        double bx = poly[2 * m + 2], by = poly[2 * m + 3];
+        double ex = bx - ax, ey = by - ay;
+        double l2 = ex * ex + ey * ey;
+        double u = ((x - ax) * ex + (y - ay) * ey) / l2;
+        if (u < 0.0) u = 0.0;
+        if (u > 1.0) u = 1.0;
+        double px = ax + u * ex, py = ay + u * ey;
+        double qx = x - px, qy = y - py;
+        double d = sqrt(qx * qx + qy * qy);
+        if (d < best) best = d;
+    }
+    double df = best - s.half;
+    if (s.flag_cut) df = span_cut(df, z, s);
+    double r = s.fillet;
+    if (r > 0.0 && dc < r && df < r) {
+        double a = r - dc, b = r - df;
+        return r - sqrt(a * a + b * b);
+    }
+    return rs_min(dc, df);
+}
+
+/* Per corner. `has_prev`: keep φ where the decayed bound stays beyond the band. */
+extern "C" __global__ void e3_geom_phi(
+    GeomGrid g, GeomSdf s, const double* __restrict__ poly,
+    int has_prev, double band, double motion,
+    double* __restrict__ phi, double* __restrict__ bound, unsigned char* __restrict__ touched)
+{
+    long long n = (long long) blockIdx.x * blockDim.x + threadIdx.x;
+    long long nn = (long long) (g.nx + 1) * (g.ny + 1) * (g.nz + 1);
+    if (n >= nn) return;
+    long long nxy = (long long) (g.nx + 1) * (g.ny + 1);
+    int k = (int) (n / nxy);
+    int j = (int) ((n / (g.nx + 1)) % (g.ny + 1));
+    int i = (int) (n % (g.nx + 1));
+    if (has_prev) {
+        double b = bound[n] - motion;
+        if (b > band) {
+            bound[n] = b;
+            touched[n] = 0;
+            return;
+        }
+    }
+    double v = geom_phi_at((double) i * g.dx, (double) j * g.dy, (double) k * g.dz, s, poly);
+    phi[n] = v;
+    bound[n] = fabs(v);
+    touched[n] = 1;
+}
+
+__device__ __forceinline__ long long gnode(const GeomGrid& g, int k, int j, int i)
+{
+    return ((long long) k * (g.ny + 1) + j) * (g.nx + 1) + i;
+}
+
+/* cut.rs tri_area_fraction */
+__device__ double tri_frac(double p0, double p1, double p2)
+{
+    double v[3] = {p0, p1, p2};
+    int pos = (p0 >= 0.0) + (p1 >= 0.0) + (p2 >= 0.0);
+    if (pos == 0) return 0.0;
+    if (pos == 3) return 1.0;
+    if (pos == 1) {
+        int a = v[0] >= 0.0 ? 0 : (v[1] >= 0.0 ? 1 : 2);
+        int b = (a + 1) % 3, c = (a + 2) % 3;
+        return (v[a] / (v[a] - v[b])) * (v[a] / (v[a] - v[c]));
+    }
+    int a = v[0] < 0.0 ? 0 : (v[1] < 0.0 ? 1 : 2);
+    int b = (a + 1) % 3, c = (a + 2) % 3;
+    return 1.0 - (v[a] / (v[a] - v[b])) * (v[a] / (v[a] - v[c]));
+}
+
+__device__ __forceinline__ double quad_frac(double q00, double q10, double q01, double q11)
+{
+    return 0.5 * (tri_frac(q00, q10, q11) + tri_frac(q00, q11, q01));
+}
+
+/* Per face of component c (0 u, 1 v, 2 w): aperture and face-centre φ. */
+extern "C" __global__ void e3_geom_faces(
+    GeomGrid g, int c, int full, const double* __restrict__ phi, const unsigned char* __restrict__ touched,
+    double* __restrict__ a, double* __restrict__ d)
+{
+    long long f = (long long) blockIdx.x * blockDim.x + threadIdx.x;
+    int ni = g.nx + (c == 0), nj = g.ny + (c == 1), nk = g.nz + (c == 2);
+    if (f >= (long long) ni * nj * nk) return;
+    int k = (int) (f / ((long long) nj * ni));
+    int j = (int) ((f / ni) % nj);
+    int i = (int) (f % ni);
+    long long n00, n10, n01, n11;
+    if (c == 0) {        /* x-face: (j, k) */
+        n00 = gnode(g, k, j, i); n10 = gnode(g, k, j + 1, i);
+        n01 = gnode(g, k + 1, j, i); n11 = gnode(g, k + 1, j + 1, i);
+    } else if (c == 1) { /* y-face: (i, k) */
+        n00 = gnode(g, k, j, i); n10 = gnode(g, k, j, i + 1);
+        n01 = gnode(g, k + 1, j, i); n11 = gnode(g, k + 1, j, i + 1);
+    } else {             /* z-face: (i, j) */
+        n00 = gnode(g, k, j, i); n10 = gnode(g, k, j, i + 1);
+        n01 = gnode(g, k, j + 1, i); n11 = gnode(g, k, j + 1, i + 1);
+    }
+    if (!full && !(touched[n00] | touched[n10] | touched[n01] | touched[n11])) return;
+    double q00 = phi[n00], q10 = phi[n10], q01 = phi[n01], q11 = phi[n11];
+    a[f] = quad_frac(q00, q10, q01, q11);
+    d[f] = 0.25 * (q00 + q10 + q01 + q11);
+}
+
+__device__ __forceinline__ double det3(const double* a, const double* b, const double* c)
+{
+    return a[0] * (b[1] * c[2] - b[2] * c[1]) - a[1] * (b[0] * c[2] - b[2] * c[0])
+        + a[2] * (b[0] * c[1] - b[1] * c[0]);
+}
+
+__device__ double tet_vol(const double* p0, const double* p1, const double* p2, const double* p3)
+{
+    double e1[3] = {p1[0] - p0[0], p1[1] - p0[1], p1[2] - p0[2]};
+    double e2[3] = {p2[0] - p0[0], p2[1] - p0[1], p2[2] - p0[2]};
+    double e3[3] = {p3[0] - p0[0], p3[1] - p0[1], p3[2] - p0[2]};
+    return fabs(det3(e1, e2, e3)) / 6.0;
+}
+
+__device__ __forceinline__ void lerp3(const double* a, const double* b, double t, double* out)
+{
+    out[0] = a[0] + t * (b[0] - a[0]);
+    out[1] = a[1] + t * (b[1] - a[1]);
+    out[2] = a[2] + t * (b[2] - a[2]);
+}
+
+/* cut.rs tet_fluid_volume */
+__device__ double tet_fluid(const double pts[4][3], const double* v)
+{
+    double total = tet_vol(pts[0], pts[1], pts[2], pts[3]);
+    int pos[4], neg[4], np = 0, nn = 0;
+    for (int q = 0; q < 4; ++q) {
+        if (v[q] >= 0.0) pos[np++] = q;
+    }
+    for (int q = 0; q < 4; ++q) {
+        if (v[q] < 0.0) neg[nn++] = q;
+    }
+    if (np == 0) return 0.0;
+    if (np == 4) return total;
+    if (np == 1 || np == 3) {
+        int a = np == 1 ? pos[0] : neg[0];
+        const int* o = np == 1 ? neg : pos;
+        double pb[3], pc[3], pd[3];
+        lerp3(pts[a], pts[o[0]], v[a] / (v[a] - v[o[0]]), pb);
+        lerp3(pts[a], pts[o[1]], v[a] / (v[a] - v[o[1]]), pc);
+        lerp3(pts[a], pts[o[2]], v[a] / (v[a] - v[o[2]]), pd);
+        double t = tet_vol(pts[a], pb, pc, pd);
+        return np == 1 ? t : total - t;
+    }
+    int a = pos[0], b = pos[1], c = neg[0], d = neg[1];
+    double pac[3], pad[3], pbc[3], pbd[3];
+    lerp3(pts[a], pts[c], v[a] / (v[a] - v[c]), pac);
+    lerp3(pts[a], pts[d], v[a] / (v[a] - v[d]), pad);
+    lerp3(pts[b], pts[c], v[b] / (v[b] - v[c]), pbc);
+    lerp3(pts[b], pts[d], v[b] / (v[b] - v[d]), pbd);
+    return tet_vol(pts[a], pts[b], pbc, pbd) + tet_vol(pts[a], pac, pbc, pbd)
+        + tet_vol(pts[a], pac, pad, pbd);
+}
+
+/* The Kuhn split around (0,0,0)–(1,1,1), corners as (x, y, z). */
+__constant__ int KUHN[6][4][3] = {
+    {{0, 0, 0}, {1, 0, 0}, {1, 1, 0}, {1, 1, 1}},
+    {{0, 0, 0}, {1, 0, 0}, {1, 0, 1}, {1, 1, 1}},
+    {{0, 0, 0}, {0, 1, 0}, {1, 1, 0}, {1, 1, 1}},
+    {{0, 0, 0}, {0, 1, 0}, {0, 1, 1}, {1, 1, 1}},
+    {{0, 0, 0}, {0, 0, 1}, {1, 0, 1}, {1, 1, 1}},
+    {{0, 0, 0}, {0, 0, 1}, {0, 1, 1}, {1, 1, 1}},
+};
+
+/* Per cell: fluid volume fraction and the wall vector (3 interleaved). */
+extern "C" __global__ void e3_geom_cells(
+    GeomGrid g, int full, const double* __restrict__ phi, const unsigned char* __restrict__ touched,
+    const double* __restrict__ a_u, const double* __restrict__ a_v, const double* __restrict__ a_w,
+    double* __restrict__ vol, double* __restrict__ wall, unsigned char* __restrict__ touched_cell)
+{
+    long long idx = (long long) blockIdx.x * blockDim.x + threadIdx.x;
+    long long nxy = (long long) g.nx * g.ny;
+    if (idx >= nxy * g.nz) return;
+    int k = (int) (idx / nxy);
+    int j = (int) ((idx / g.nx) % g.ny);
+    int i = (int) (idx % g.nx);
+    unsigned char t = 0;
+    for (int dk = 0; dk < 2; ++dk)
+        for (int dj = 0; dj < 2; ++dj)
+            for (int di = 0; di < 2; ++di)
+                t |= touched[gnode(g, k + dk, j + dj, i + di)];
+    touched_cell[idx] = t;
+    if (!full && !t) return;
+    double fluid = 0.0;
+    for (int q = 0; q < 6; ++q) {
+        double pts[4][3];
+        double vals[4];
+        for (int m = 0; m < 4; ++m) {
+            pts[m][0] = (double) KUHN[q][m][0];
+            pts[m][1] = (double) KUHN[q][m][1];
+            pts[m][2] = (double) KUHN[q][m][2];
+            vals[m] = phi[gnode(g, k + KUHN[q][m][2], j + KUHN[q][m][1], i + KUHN[q][m][0])];
+        }
+        fluid += tet_fluid(pts, vals);
+    }
+    vol[idx] = fluid;
+    double ax = g.dy * g.dz, ay = g.dx * g.dz, az = g.dx * g.dy;
+    long long fu0 = ((long long) k * g.ny + j) * (g.nx + 1) + i;
+    long long fv0 = ((long long) k * (g.ny + 1) + j) * g.nx + i;
+    long long fw0 = ((long long) k * g.ny + j) * g.nx + i;
+    double sx = (a_u[fu0 + 1] - a_u[fu0]) * ax;
+    double sy = (a_v[fv0 + g.nx] - a_v[fv0]) * ay;
+    double sz = (a_w[fw0 + nxy] - a_w[fw0]) * az;
+    wall[3 * idx] = -sx;
+    wall[3 * idx + 1] = -sy;
+    wall[3 * idx + 2] = -sz;
+}
+
+/* out[w t + c] = src[w idx[t] + c], c < w (compact copies for the host mirror). */
+extern "C" __global__ void e3_geom_gather(
+    int n, int w, const unsigned int* __restrict__ idx, const double* __restrict__ src, double* __restrict__ out)
+{
+    int t = blockIdx.x * blockDim.x + threadIdx.x;
+    if (t >= n) return;
+    long long s = (long long) w * idx[t];
+    for (int c = 0; c < w; ++c) out[(long long) w * t + c] = src[s + c];
+}