Back to home page

EIC code displayed by LXR

 
 

    


File indexing completed on 2026-09-19 09:40:34

0001 /*
0002  * \file GenTrapStruct.h
0003  * \brief Parameter/storage struct and small helper types for a Generic Trapezoid (Arb8).
0004  *
0005  * This header contains:
0006  *  - Low-level helpers in ::arb4helpers for lateral bilinear “arb4” faces:
0007  *      * Arb4Surf     : implicit/bilinear face equation and cheap safety bounds
0008  *      * BoundingPlane: robust plane representation used by the mesh bounder
0009  *      * Arb4Mesh     : two-triangle per side bounding slab per face for fast culling
0010  *  - GenTrapStruct<T> : a POD-like container of all precomputed data for a GenTrap
0011  *
0012  * No heavy kernels live here; kernels are in GenTrapImplementation.h and only consume
0013  * the data populated by GenTrapStruct::Initialize().
0014  *
0015  *  Created on: 17.07.2016
0016  *      Author: mgheata
0017  *  Revised on: 15.09.2025 andrei.gheata@cern.ch
0018  */
0019 
0020 #ifndef VECGEOM_VOLUMES_GENTRAPSTRUCT_H_
0021 #define VECGEOM_VOLUMES_GENTRAPSTRUCT_H_
0022 
0023 #include <VecGeom/base/Global.h>
0024 #include <VecGeom/base/Vector3D.h>
0025 #include <VecGeom/base/Vector2D.h>
0026 #ifndef VECCORE_CUDA
0027 #include <VecGeom/volumes/TessellatedSection.h>
0028 #endif
0029 
0030 // ------------------------------------------------------------------
0031 // Hybrid precompute switch:
0032 // 1 (default): store a few cheap per-face scalars (tiny memory cost)
0033 // 0           : compute everything on-the-fly
0034 #ifndef GENTRAP_USE_HYBRID_COEFFS
0035 #define GENTRAP_USE_HYBRID_COEFFS 1
0036 #endif
0037 
0038 namespace vecgeom {
0039 
0040 inline namespace VECGEOM_IMPL_NAMESPACE {
0041 
0042 namespace arb4helpers {
0043 
0044 /**
0045  * \brief Implicit representation for a lateral face of the GenTrap.
0046  *
0047  * For twisted faces the zero set is a quadratic in (x,y,z):
0048  * \code
0049  *   f(x,y,z) = (A x + B y + C z) z + D x + E y + F z + G
0050  * \endcode
0051  * with outward orientation chosen by construction. For planar faces the
0052  * quadratic part is zero (A=B=C=0) and the equation reduces to the plane
0053  * \( D x + E y + F z + G = 0 \).
0054  *
0055  * The struct also stores a per-face Lipschitz constant (4*k) used to build
0056  * conservative safety underestimates, computed from (A,B,C).
0057  *
0058  * \tparam Real_v scalar or vector floating-point type
0059  */
0060 template <typename Real_v>
0061 struct Arb4Surf {
0062   Real_v f4k{0.}; ///< 4×Lipschitz constant used in SafetyLipschitz()
0063   Real_v A{0.};   ///< Bilinear coefficient (quadratic in z): (A x + B y + C z) z
0064   Real_v B{0.};   ///< Bilinear coefficient (quadratic in z): (A x + B y + C z) z
0065   Real_v C{0.};   ///< Bilinear coefficient (quadratic in z): (A x + B y + C z) z
0066   Real_v D{0.};   ///< Planar term coefficient for x
0067   Real_v E{0.};   ///< Planar term coefficient for y
0068   Real_v F{0.};   ///< Planar term coefficient for z
0069   Real_v G{0.};   ///< Constant term
0070 
0071   /**
0072    * \brief Set all coefficients and precompute Lipschitz bound.
0073    *
0074    * The Lipschitz constant \c k is \f$|C|+\sqrt{A^2+B^2+C^2}\f$, we store 4*k to
0075    * avoid multiplications in SafetyLipschitz().
0076    */
0077   VECCORE_ATT_HOST_DEVICE VECGEOM_FORCE_INLINE void Set(Real_v a, Real_v b, Real_v c, Real_v d, Real_v e, Real_v f,
0078                                                         Real_v g)
0079   {
0080     // Coefficients
0081     A = a;
0082     B = b;
0083     C = c;
0084     D = d;
0085     E = e;
0086     F = f;
0087     G = g;
0088     // Lipschitz constant
0089     Real_v k = Abs(c) + Sqrt(a * a + b * b + c * c);
0090     f4k      = 4. * k;
0091   }
0092 
0093   /// \brief No-op for API symmetry with vector lanes.
0094   VECCORE_ATT_HOST_DEVICE VECGEOM_FORCE_INLINE void SyncLane() {}
0095 
0096   /// \brief Evaluate the full (twisted) implicit function \(f(x,y,z)\).
0097   VECCORE_ATT_HOST_DEVICE VECGEOM_FORCE_INLINE Real_v Evaluate(Vector3D<Real_v> const &p) const
0098   {
0099     return (A * p.x() + B * p.y() + C * p.z()) * p.z() + D * p.x() + E * p.y() + F * p.z() + G;
0100   }
0101 
0102   /// \brief Evaluate the planar part \( D x + E y + F z + G \).
0103   VECCORE_ATT_HOST_DEVICE VECGEOM_FORCE_INLINE Real_v EvaluatePlanar(Vector3D<Real_v> const &p) const
0104   {
0105     return D * p.x() + E * p.y() + F * p.z() + G;
0106   }
0107 
0108   /// \brief Gradient of the full (twisted) implicit function at \p p.
0109   VECCORE_ATT_HOST_DEVICE VECGEOM_FORCE_INLINE Vector3D<Real_v> EvaluateGradient(Vector3D<Real_v> const &p) const
0110   {
0111     return Vector3D<Real_v>(A * p.z() + D, B * p.z() + E, A * p.x() + B * p.y() + 2. * C * p.z() + F);
0112   }
0113 
0114   /// \brief Dot product of the planar normal (D,E,F) with a vector \p v.
0115   VECCORE_ATT_HOST_DEVICE VECGEOM_FORCE_INLINE Real_v DotPlaneNormal(Vector3D<Real_v> const &v) const
0116   {
0117     return D * v.x() + E * v.y() + F * v.z();
0118   }
0119 
0120   /// \brief Return the (outward) plane normal (D,E,F). For twisted faces this is the linear part’s normal.
0121   VECCORE_ATT_HOST_DEVICE VECGEOM_FORCE_INLINE Vector3D<Real_v> GetPlaneNormal() const { return {D, E, F}; }
0122 
0123   /**
0124    * \brief Conservative safety underestimate to this surface via Lipschitz bound.
0125    *
0126    * Computes a tight but conservative \em signed distance-like value based on
0127    * the function value and its gradient norm (see Geant4 MR5261 notes).
0128    * Result is positive when \c Evaluate(p) has outward sign.
0129    */
0130   VECCORE_ATT_HOST_DEVICE VECGEOM_FORCE_INLINE Real_v SafetyLipschitz(Vector3D<Real_v> const &p) const
0131   {
0132     auto gradx    = A * p.z() + D;
0133     auto grady    = B * p.z() + E;
0134     auto Cz       = C * p.z();
0135     auto CzF      = Cz + F;
0136     auto gradz    = A * p.x() + B * p.y() + Cz + CzF;
0137     auto fun      = gradx * p.x() + grady * p.y() + CzF * p.z() + G;
0138     auto grad2    = gradx * gradx + grady * grady + gradz * gradz;
0139     auto divisor  = Sqrt(grad2) + Sqrt(grad2 + f4k * Abs(fun));
0140     auto distSurf = 2. * fun / divisor;
0141     return distSurf;
0142   }
0143 
0144   /**
0145    * \brief Build the quadratic coefficients a,b,c for f(p+t d) along a ray.
0146    *
0147    * Using the implicit form:
0148    *   f(x,y,z) = (A x + B y + C z) z + D x + E y + F z + G
0149    * with p=(px,py,pz), d=(dx,dy,dz), one obtains:
0150    *   S1   = A*dx + B*dy + C*dz
0151    *   S0   = A*px + B*py + C*pz
0152    *   Ldir = D*dx + E*dy + F*dz
0153    *   Lpt  = D*px + E*py + F*pz
0154    *   a = S1*dz
0155    *   b = S1*pz + S0*dz + Ldir
0156    *   c = S0*pz + Lpt + G
0157    */
0158   VECCORE_ATT_HOST_DEVICE
0159   VECGEOM_FORCE_INLINE void RayQuadratic(Vector3D<Real_v> const &p, Vector3D<Real_v> const &d, Real_v &a, Real_v &b,
0160                                          Real_v &c) const
0161   {
0162     const Real_v S1   = A * d.x() + B * d.y() + C * d.z();
0163     const Real_v S0   = A * p.x() + B * p.y() + C * p.z();
0164     const Real_v Ldir = D * d.x() + E * d.y() + F * d.z();
0165     const Real_v Lpt  = D * p.x() + E * p.y() + F * p.z();
0166     a                 = S1 * d.z();
0167     b                 = S1 * p.z() + S0 * d.z() + Ldir;
0168     c                 = S0 * p.z() + Lpt + G;
0169   }
0170 };
0171 
0172 /**
0173  * \brief Specialization for double with 16B-aligned packed coefficients.
0174  *
0175  * This packs coefficients in 16-byte “wide” pairs to help both host SIMD and
0176  * device (CUDA) vector loads (e.g. \c ld.global.v2.f64).
0177  */
0178 template <>
0179 struct alignas(16) Arb4Surf<double> {
0180 
0181   struct alignas(16) CoeffLane {
0182     wide::d2 la, bc, de, fg; // four 16B chunks
0183   } lane;
0184 
0185   /// \copydoc Arb4Surf::Set
0186   VECCORE_ATT_HOST_DEVICE VECGEOM_FORCE_INLINE void Set(double a, double b, double c, double d, double e, double f,
0187                                                         double g)
0188   {
0189     // Lipschitz constant
0190     double k = Abs(c) + Sqrt(a * a + b * b + c * c);
0191     lane.la.set(4. * k, a);
0192     lane.bc.set(b, c);
0193     lane.de.set(d, e);
0194     lane.fg.set(f, g);
0195   }
0196 
0197   /// \copydoc Arb4Surf::Evaluate
0198   VECCORE_ATT_HOST_DEVICE VECGEOM_FORCE_INLINE double Evaluate(const Vector3D<double> &p) const
0199   {
0200     // Each of these temporaries copies a 16B object → NVCC emits ld.global.v2.f64 on device
0201     auto la = lane.la, bc = lane.bc, de = lane.de, fg = lane.fg;
0202     double ax = la.y() * p.x(), by = bc.x() * p.y();
0203     double cz = bc.y() * p.z(), dx = de.x() * p.x();
0204     double ey = de.y() * p.y(), fz = fg.x() * p.z();
0205     return (ax + by + cz) * p.z() + dx + ey + fz + fg.y();
0206   }
0207 
0208   /// \copydoc Arb4Surf::EvaluatePlanar
0209   VECCORE_ATT_HOST_DEVICE VECGEOM_FORCE_INLINE double EvaluatePlanar(Vector3D<double> const &p) const
0210   {
0211     auto de = lane.de, fg = lane.fg;
0212     double dx = de.x() * p.x(), ey = de.y() * p.y(), fz = fg.x() * p.z();
0213     return dx + ey + fz + fg.y();
0214   }
0215 
0216   /// \copydoc Arb4Surf::EvaluateGradient
0217   VECCORE_ATT_HOST_DEVICE VECGEOM_FORCE_INLINE Vector3D<double> EvaluateGradient(const Vector3D<double> &p) const
0218   {
0219     auto la = lane.la, bc = lane.bc, de = lane.de, fg = lane.fg;
0220     return {la.y() * p.z() + de.x(), // A*p.z + D
0221             bc.x() * p.z() + de.y(), // B*p.z + E
0222             la.y() * p.x() + bc.x() * p.y() + 2.0 * bc.y() * p.z() + fg.x()};
0223   }
0224 
0225   /// \copydoc Arb4Surf::DotPlaneNormal
0226   VECCORE_ATT_HOST_DEVICE VECGEOM_FORCE_INLINE double DotPlaneNormal(const Vector3D<double> &v) const
0227   {
0228     auto de = lane.de, fg = lane.fg;
0229     return de.x() * v.x() + de.y() * v.y() + fg.x() * v.z();
0230   }
0231 
0232   /// \copydoc Arb4Surf::GetPlaneNormal
0233   VECCORE_ATT_HOST_DEVICE VECGEOM_FORCE_INLINE Vector3D<double> GetPlaneNormal() const
0234   {
0235     return {lane.de.x(), lane.de.y(), lane.fg.x()};
0236   }
0237 
0238   /// \copydoc Arb4Surf::SafetyLipschitz
0239   VECCORE_ATT_HOST_DEVICE VECGEOM_FORCE_INLINE double SafetyLipschitz(Vector3D<double> const &p) const
0240   {
0241     auto la = lane.la, bc = lane.bc, de = lane.de, fg = lane.fg;
0242     auto gradx    = la.y() * p.z() + de.x();
0243     auto grady    = bc.x() * p.z() + de.y();
0244     auto Cz       = bc.y() * p.z();
0245     auto CzF      = Cz + fg.x();
0246     auto gradz    = la.y() * p.x() + bc.x() * p.y() + Cz + CzF;
0247     auto fun      = gradx * p.x() + grady * p.y() + CzF * p.z() + fg.y();
0248     auto grad2    = gradx * gradx + grady * grady + gradz * gradz;
0249     auto divisor  = Sqrt(grad2) + Sqrt(grad2 + la.x() * Abs(fun));
0250     auto distSurf = 2. * fun / divisor;
0251     return distSurf;
0252   }
0253 
0254   /**
0255    * \brief Build the quadratic coefficients a,b,c for f(p+t d) along a ray (double specialization).
0256    */
0257   VECCORE_ATT_HOST_DEVICE
0258   VECGEOM_FORCE_INLINE void RayQuadratic(const Vector3D<double> &p, const Vector3D<double> &d, double &a, double &b,
0259                                          double &c) const
0260   {
0261     // Unpack coefficients once (compiler emits ld.global.v2.f64 for the pairs)
0262     auto la = lane.la, bc = lane.bc, de = lane.de, fg = lane.fg;
0263     const double A = la.y();
0264     const double B = bc.x();
0265     const double C = bc.y();
0266     const double D = de.x();
0267     const double E = de.y();
0268     const double F = fg.x();
0269     const double G = fg.y();
0270 
0271     const double S1   = A * d.x() + B * d.y() + C * d.z();
0272     const double S0   = A * p.x() + B * p.y() + C * p.z();
0273     const double Ldir = D * d.x() + E * d.y() + F * d.z();
0274     const double Lpt  = D * p.x() + E * p.y() + F * p.z();
0275 
0276     a = S1 * d.z();
0277     b = S1 * p.z() + S0 * d.z() + Ldir;
0278     c = S0 * p.z() + Lpt + G;
0279   }
0280 };
0281 
0282 static_assert(alignof(Arb4Surf<double>) >= 16, "Need 16B alignment");
0283 static_assert(sizeof(Arb4Surf<double>) % 16 == 0, "Size multiple of 16 helps arrays");
0284 
0285 /**
0286  * \brief Robust plane used as bounding slab for a lateral face.
0287  *
0288  * Plane is represented as \( \mathbf{n}\cdot\mathbf{x} + d = 0 \) with \c n unit length.
0289  * Built from three points using a robust triangle normal.
0290  */
0291 template <typename Real_v>
0292 struct BoundingPlane {
0293 
0294   using Vertex_t = Vector3D<Real_v>;
0295 
0296   Vertex_t n;   ///< unit normal
0297   Real_v d{0.}; ///< plane constant so that n·x + d = 0
0298 
0299   BoundingPlane() = default;
0300 
0301   /// \brief Construct from triangle (a,b,c) with robust orientation.
0302   VECCORE_ATT_HOST_DEVICE
0303   BoundingPlane(Vertex_t const &a, Vertex_t const &b, Vertex_t const &c) { Set(a, b, c); }
0304 
0305   /// \brief Reset plane from triangle (a,b,c).
0306   VECCORE_ATT_HOST_DEVICE
0307   void Set(Vertex_t const &a, Vertex_t const &b, Vertex_t const &c) { TriangleUnitNormalRobust(a, b, c, n, d); }
0308 
0309   /// \brief Signed distance from \p point (positive on the normal side).
0310   VECCORE_ATT_HOST_DEVICE VECGEOM_FORCE_INLINE Real_v SignedDistance(const Vertex_t &point) const
0311   {
0312     return n.Dot(point) + d; // because plane is n·x + d = 0
0313   }
0314 };
0315 
0316 /**
0317  * \brief Two-triangle bounding slab per side for a lateral (twisted/planar) face.
0318  *
0319  * Each Arb4 face is conservatively bounded by two oriented triangles. We keep
0320  * the “inbound/outbound” half-spaces to enable very cheap “may hit” culls and
0321  * fast signed safeties.
0322  */
0323 template <typename Real_v>
0324 struct Arb4Mesh {
0325   using Vertex_t    = Vector3D<Real_v>;
0326   using Direction_t = Vertex_t;
0327 
0328   BoundingPlane<Real_v> fBplanes[4]; ///< Two outbound + two inbound oriented planes
0329 
0330   Arb4Mesh() = default;
0331 
0332   /// \brief Build the bounding planes from the four face corners (bottom v1,v2 and top v3,v4).
0333   VECCORE_ATT_HOST_DEVICE
0334   Arb4Mesh(Vertex_t const &v1, Vertex_t const &v2, Vertex_t const &v3, Vertex_t const &v4) { Set(v1, v2, v3, v4); }
0335 
0336   /**
0337    * \brief Set planes from corners; orientation depends on whether v4 is behind (v1,v3)×(v1,v2).
0338    *
0339    * We build: 2 “outbound” planes (fBplanes[0..1]) and 2 “inbound” planes (fBplanes[2..3]).
0340    * The Inside/MaysHit helpers pick the appropriate pair based on whether the start is inside the solid.
0341    */
0342   VECCORE_ATT_HOST_DEVICE
0343   void Set(Vertex_t const &v1, Vertex_t const &v2, Vertex_t const &v3, Vertex_t const &v4)
0344   {
0345     // Check if P4 is behind (P1,P3) X (P1, P2)
0346     bool behind = ((v3 - v1).Cross(v2 - v1)).Dot(v4 - v1) < 0;
0347     if (behind) {
0348       // outbound triangles
0349       fBplanes[0].Set(v1, v3, v2);
0350       fBplanes[1].Set(v2, v3, v4);
0351       // inbound triangles
0352       fBplanes[2].Set(v1, v2, v4);
0353       fBplanes[3].Set(v1, v4, v3);
0354     } else {
0355       // outbound triangles
0356       fBplanes[0].Set(v1, v4, v2);
0357       fBplanes[1].Set(v1, v3, v4);
0358       // inbound triangles
0359       fBplanes[2].Set(v1, v2, v3);
0360       fBplanes[3].Set(v2, v4, v3);
0361     }
0362   }
0363 
0364   /// \brief True if \p point is inside all four planes.
0365   VECCORE_ATT_HOST_DEVICE VECGEOM_FORCE_INLINE bool Inside(Vertex_t const &point) const
0366   {
0367     for (auto i = 0; i < 4; ++i) {
0368       if (fBplanes[i].SignedDistance(point) > Real_v(0.)) return false;
0369     }
0370     return true;
0371   }
0372 
0373   /**
0374    * \brief Optimized inside check knowing whether the point is inside the \em solid.
0375    *
0376    * When \p point_inside is known, only two planes (either inbound or outbound) matter.
0377    */
0378   VECCORE_ATT_HOST_DEVICE VECGEOM_FORCE_INLINE bool Inside(Vertex_t const &point, bool point_inside) const
0379   {
0380     int i = 2 * int(point_inside);
0381     return (fBplanes[i].SignedDistance(point) < Real_v(0.)) && (fBplanes[i + 1].SignedDistance(point) < Real_v(0.));
0382   }
0383 
0384   /**
0385    * \brief Cheap necessary test for a potential hit along direction \p dir.
0386    *
0387    * If the point is outside the mesh, the ray can only hit if it points
0388    * against at least one of the relevant plane normals.
0389    */
0390   VECCORE_ATT_HOST_DEVICE VECGEOM_FORCE_INLINE bool MayHit(Vertex_t const &point, Direction_t const &dir,
0391                                                            bool point_inside) const
0392   {
0393     // If point is inside the mesh, the direction may always hit
0394     if (Inside(point, point_inside)) return true;
0395     // Otherwise: must be opposite to at least one normal
0396     int i = 2 * int(point_inside);
0397     return (dir.Dot(fBplanes[i].n) < Real_v(0.)) || (dir.Dot(fBplanes[i + 1].n) < Real_v(0.));
0398   }
0399 
0400   /**
0401    * \brief Conservative (unsigned) safety to the bounding planes.
0402    *
0403    * If the point is outside the two relevant planes, return the minimum absolute
0404    * distance to either plane, otherwise return 0.
0405    */
0406   VECCORE_ATT_HOST_DEVICE VECGEOM_FORCE_INLINE Real_v Safety(Vertex_t const &point, bool inside) const
0407   {
0408     if (Inside(point)) return Real_v(0.);
0409     int i     = 2 * int(inside);
0410     Real_v s0 = Abs(fBplanes[i].SignedDistance(point));
0411     Real_v s1 = Abs(fBplanes[i + 1].SignedDistance(point));
0412     return Min(s0, s1);
0413   }
0414 
0415   /**
0416    * \brief Very cheap signed safety: positive outside the slab, 0 inside, negative inside-solid mode.
0417    *
0418    * Useful for coarse pruning; not accurate vs. the true triangular patch.
0419    */
0420   VECCORE_ATT_HOST_DEVICE VECGEOM_FORCE_INLINE Real_v FastSignedSafety(Vertex_t const &point, bool inside) const
0421   {
0422     int i         = 2 * int(inside);
0423     Real_v safety = Max(fBplanes[i].SignedDistance(point), fBplanes[i + 1].SignedDistance(point));
0424     safety        = Max(Real_v(0.), safety); // truncate to 0 when inside the two planes
0425     safety *= Real_v(1 - i);                 // outside: +1, inside: -1  (i is 0 or 2)
0426     return safety;
0427   }
0428 };
0429 } // namespace arb4helpers
0430 
0431 /**
0432  * \brief POD-like storage for a Generalized Trapezoid (Arb8) geometry.
0433  *
0434  * This struct owns only precomputed, immutable geometry data (vertices, per-face
0435  * coefficients, cheap bounders, etc.). Navigation kernels are implemented in
0436  * GenTrapImplementation and read from this struct.
0437  *
0438  * \tparam T floating precision
0439  */
0440 template <typename T = double>
0441 struct alignas(16) GenTrapStruct {
0442   using Vertex_t = Vector3D<T>;
0443   using PointXY  = Vector2D<T>;
0444   using UChar_t  = unsigned char;
0445 
0446   Vertex_t fBBdimensions;            ///< Half-lengths of the AABB (x,y,z)
0447   Vertex_t fBBorigin;                ///< Center of the AABB in local coordinates
0448   T fDz{0.};                         ///< Half-height along Z (top/bottom at ±fDz)
0449   T fDz2{0.};                        ///< Precomputed 1/(2*fDz) used for fast v(z) mapping
0450   Vertex_t fVertices[8];             ///< The eight corner points: 0..3 bottom (z=-Dz), 4..7 top (z=+Dz)
0451   Vector2D<T> fT[4];                 ///< Per-face generator slopes in XY (bottom->top) scaled by 1/(2*fDz)
0452   Vector2D<T> fMidXY[4];             ///< Midpoint at z=0 for each generator: v[i].XY() + fT[i]*fDz
0453   arb4helpers::Arb4Surf<T> fSurf[4]; ///< Twisted/planar implicit coefficients for each lateral face
0454   arb4helpers::Arb4Mesh<T> fMesh[4]; ///< Two-triangle bounding slabs per lateral face side
0455   UChar_t fFlags{0};                 ///< Bitfield: [0..3]=twisted flags per face, [4..7]=degenerated flags
0456 
0457 #if GENTRAP_USE_HYBRID_COEFFS
0458   // --- Hybrid precompute (tiny storage, reduces ALU) ---------------------
0459   // For each lateral face i (with j=(i+1)&3), we cache:
0460   //   fNext[i]  = j
0461   //   fTT[i]    = T[i]   × T[j]
0462   //   fMM[i]    = Mid[i] × Mid[j]
0463   //   fMT[i]    = Mid[i] × T[j]  -  Mid[j] × T[i]
0464   int fNext[4]; ///< next index j = (i+1)%4 for face i
0465   T fTT[4];     ///< crossZ(T[i],   T[j])
0466   T fMM[4];     ///< crossZ(Mid[i], Mid[j])
0467   T fMT[4];     ///< crossZ(Mid[i], T[j]) - crossZ(Mid[j], T[i])
0468 
0469   // Extended caches (vectors) used by twisted intersection, XYZ->UV, and UNormal:
0470   Vector2D<T> fDmid[4]; ///< Mid[j] - Mid[i]
0471   Vector2D<T> fDT[4];   ///< T[j]   - T[i]
0472   Vector3D<T> fE0[4];   ///< C - A  (bottom edge, z=0)
0473   Vector3D<T> fE1[4];   ///< D - B  (top edge,    z=0)
0474   Vector3D<T> fG0[4];   ///< B - A  (generator at i,    z=+2Dz)
0475   Vector3D<T> fG1[4];   ///< D - C  (generator at j,    z=+2Dz)
0476 #endif
0477 
0478 #ifndef VECCORE_CUDA
0479   TessellatedSection<T> *fTslHelper = nullptr; ///< SIMD helper using tessellated clusters for the planar case
0480 #endif
0481 
0482   GenTrapStruct() = default;
0483 
0484   /**
0485    * \brief Construct from XY vertex arrays and half-height.
0486    *
0487    * \param verticesx 8 X-coordinates (first 4 bottom, next 4 top)
0488    * \param verticesy 8 Y-coordinates (first 4 bottom, next 4 top)
0489    * \param halfzheight Half height in Z (fDz)
0490    */
0491   VECCORE_ATT_HOST_DEVICE
0492   GenTrapStruct(const Precision verticesx[], const Precision verticesy[], Precision halfzheight)
0493   {
0494     // Constructor
0495     Initialize(verticesx, verticesy, halfzheight);
0496   }
0497 
0498   /**
0499    * \brief Initialize all derived data from input vertices and half-height.
0500    *
0501    * Performs:
0502    *  - Vertex placement at z=±fDz
0503    *  - Degeneracy flags per lateral face
0504    *  - Shear vectors fT and midpoints fMidXY
0505    *  - Orientation fix to ensure clockwise order in XY
0506    *  - Robust validation: no crossing opposite segments; convex quads
0507    *  - Per-face classification: planar vs twisted; plane or twist coefficients
0508    *  - Optional TessellatedSection build for pure planar case (host only)
0509    *  - AABB computation
0510    *
0511    * \return true on success; validation failures are guarded by VECGEOM_VALIDATE
0512    */
0513   VECCORE_ATT_HOST_DEVICE
0514   bool Initialize(const Precision verticesx[], const Precision verticesy[], Precision halfzheight)
0515   {
0516     // Initialization based on vertices and half length
0517     fDz  = halfzheight;
0518     fDz2 = 1. / (2. * fDz);
0519 
0520     // Set vertices in Vector3D form
0521     for (int i = 0; i < 4; ++i) {
0522       fVertices[i].Set(verticesx[i], verticesy[i], -fDz);
0523       fVertices[i + 4].Set(verticesx[i + 4], verticesy[i + 4], fDz);
0524     }
0525 
0526     // Compute degenerated faces, generator connections and midpoints
0527     for (int i = 0; i < 4; ++i) {
0528       int j          = (i + 1) % 4;
0529       const auto &p1 = fVertices[i];
0530       const auto &p2 = fVertices[j];
0531       const auto &p3 = fVertices[i + 4];
0532       const auto &p4 = fVertices[j + 4];
0533       auto lbot      = (p2 - p1).Mag2();
0534       auto ltop      = (p4 - p3).Mag2();
0535       SetDegenerated(i, Max(lbot, ltop) < kToleranceSquared);
0536       fT[i]     = fDz2 * (fVertices[i + 4].XY() - fVertices[i].XY());
0537       fMidXY[i] = fVertices[i].XY() + fT[i] * fDz;
0538     }
0539 
0540 #if GENTRAP_USE_HYBRID_COEFFS
0541     // Precompute light, read-only per-face scalars used in twisted intersection
0542     for (int i = 0; i < 4; ++i) {
0543       int j    = (i + 1) % 4;
0544       fNext[i] = j;
0545       // cross products in XY (z-component of 3D cross)
0546       const auto &Ti = fT[i];
0547       const auto &Tj = fT[j];
0548       const auto &Mi = fMidXY[i];
0549       const auto &Mj = fMidXY[j];
0550       fTT[i]         = Ti.CrossZ(Tj);
0551       fMM[i]         = Mi.CrossZ(Mj);
0552       fMT[i]         = Mi.CrossZ(Tj) - Mj.CrossZ(Ti);
0553 
0554       // Extended vector caches
0555       fDmid[i] = Mj - Mi;
0556       fDT[i]   = Tj - Ti;
0557 
0558       // Face vertices: A=i, C=j on bottom; B=i+4, D=j+4 on top
0559       const auto &A = fVertices[i];
0560       const auto &C = fVertices[j];
0561       const auto &B = fVertices[i + 4];
0562       const auto &D = fVertices[j + 4];
0563       fE0[i]        = C - A; // z=0
0564       fE1[i]        = D - B; // z=0
0565       fG0[i]        = B - A; // z=+2Dz
0566       fG1[i]        = D - C; // z=+2Dz
0567     }
0568 #endif
0569 
0570     // Make sure vertices are defined clockwise on both z-planes
0571     Precision sum1 = 0.;
0572     Precision sum2 = 0.;
0573     for (int i = 0; i < 4; ++i) {
0574       int j = (i + 1) % 4;
0575       sum1 += fVertices[i].CrossZ(fVertices[j]);
0576       sum2 += fVertices[i + 4].CrossZ(fVertices[j + 4]);
0577     }
0578 
0579     // We should generate an exception here
0580     if (sum1 * sum2 < -kTolerance) Print();
0581     VECGEOM_VALIDATE(sum1 * sum2 > -kTolerance, << "Unplaced generic trap defined with opposite clockwise in XY");
0582 
0583     // Revert sequence of vertices to have them clockwise if needed (bottom/top simultaneously)
0584     if (sum1 > kTolerance) {
0585       printf("INFO: Reverting to clockwise vertices of GenTrap shape:\n");
0586       Print();
0587       auto swap = [](auto &a, auto &b) {
0588         auto t = a;
0589         a      = b;
0590         b      = t;
0591       };
0592       swap(fVertices[1], fVertices[3]);
0593       swap(fVertices[5], fVertices[7]);
0594     }
0595 
0596     // Check that opposite segments are not crossing -> fatal exception
0597     bool xing0123 = SegmentsCrossing(fVertices[0], fVertices[1], fVertices[3], fVertices[2]);
0598     bool xing0312 = SegmentsCrossing(fVertices[1], fVertices[2], fVertices[0], fVertices[3]);
0599     bool xing4567 = SegmentsCrossing(fVertices[4], fVertices[5], fVertices[7], fVertices[6]);
0600     bool xing4756 = SegmentsCrossing(fVertices[5], fVertices[6], fVertices[4], fVertices[7]);
0601     if (xing0123 || xing0312 || xing4567 || xing4756) Print();
0602     VECGEOM_VALIDATE(!xing0123, << "Unplaced generic trap defined with crossing opposite segments (01) (23)");
0603     VECGEOM_VALIDATE(!xing0312, << "Unplaced generic trap defined with crossing opposite segments (03) (12)");
0604     VECGEOM_VALIDATE(!xing4567, << "Unplaced generic trap defined with crossing opposite segments (45) (67)");
0605     VECGEOM_VALIDATE(!xing4756, << "Unplaced generic trap defined with crossing opposite segments (47) (56)");
0606 
0607     // Check that top and bottom quadrilaterals are convex
0608     bool convexquads = ComputeIsConvexQuadrilaterals();
0609     if (!convexquads) Print();
0610     VECGEOM_VALIDATE(convexquads, << "Unplaced generic trap defined with top/bottom quadrilaterals not convex");
0611 
0612     // Mark twisted faces and precompute per-face coefficients/mesh
0613     ComputeTwistedFaces();
0614 
0615     // fSurfaceShell.Initialize(fVertices, fDz);
0616 
0617 #ifndef VECCORE_CUDA
0618     //  Create the tessellated helper if the faces are planar (host-side acceleration)
0619     if (IsPlanar()) {
0620       fTslHelper = new TessellatedSection<T>(4, -fDz, fDz);
0621       fTslHelper->AddQuadrilateralFacet(fVertices[0], fVertices[4], fVertices[5], fVertices[1]);
0622       fTslHelper->AddQuadrilateralFacet(fVertices[1], fVertices[5], fVertices[6], fVertices[2]);
0623       fTslHelper->AddQuadrilateralFacet(fVertices[2], fVertices[6], fVertices[7], fVertices[3]);
0624       fTslHelper->AddQuadrilateralFacet(fVertices[3], fVertices[7], fVertices[4], fVertices[0]);
0625     }
0626 #endif
0627     ComputeBoundingBox();
0628     return true;
0629   }
0630 
0631   /// \brief True if all four lateral faces are planar.
0632   VECCORE_ATT_HOST_DEVICE
0633   bool IsPlanar() const { return (fFlags & 0x0F) == 0; }
0634 
0635   /// \brief True if lateral face \p i is twisted (non-planar).
0636   VECCORE_ATT_HOST_DEVICE
0637   bool IsTwisted(int i) const { return (fFlags >> i) & 1u; }
0638 
0639   /// \brief Set twisted flag for face \p i.
0640   VECCORE_ATT_HOST_DEVICE
0641   void SetTwisted(int i, bool flag)
0642   {
0643     UChar_t mask = 1u << i;
0644     if (flag)
0645       fFlags |= mask;
0646     else
0647       fFlags &= ~mask;
0648   }
0649 
0650   /// \brief True if lateral face \p i is degenerate (collapsed segment on top or bottom).
0651   VECCORE_ATT_HOST_DEVICE
0652   bool IsDegenerated(int i) const { return (fFlags >> (i + 4)) & 1u; }
0653 
0654   /// \brief Set degenerate flag for face \p i.
0655   VECCORE_ATT_HOST_DEVICE
0656   void SetDegenerated(int i, bool flag)
0657   {
0658     UChar_t mask = 1u << (i + 4);
0659     if (flag)
0660       fFlags |= mask;
0661     else
0662       fFlags &= ~mask;
0663   }
0664 
0665   /// \brief Compute AABB origin and half-dimensions.
0666   VECCORE_ATT_HOST_DEVICE
0667   void ComputeBoundingBox()
0668   {
0669     // Computes bounding box parameters
0670     Vertex_t aMin, aMax;
0671     Extent(aMin, aMax);
0672     fBBorigin     = 0.5 * (aMin + aMax);
0673     fBBdimensions = 0.5 * (aMax - aMin);
0674   }
0675 
0676   /**
0677    * \brief Compute axis-aligned extent (min/max) of the solid including z=±fDz.
0678    *
0679    * \param[out] aMin minimum x/y/z
0680    * \param[out] aMax maximum x/y/z
0681    */
0682   VECCORE_ATT_HOST_DEVICE
0683   void Extent(Vertex_t &aMin, Vertex_t &aMax) const
0684   {
0685     // Returns the full 3D cartesian extent of the solid.
0686     aMin = aMax = fVertices[0];
0687     aMin[2]     = -fDz;
0688     aMax[2]     = fDz;
0689     for (int i = 0; i < 4; ++i) {
0690       // lower -fDz vertices
0691       if (aMin[0] > fVertices[i].x()) aMin[0] = fVertices[i].x();
0692       if (aMax[0] < fVertices[i].x()) aMax[0] = fVertices[i].x();
0693       if (aMin[1] > fVertices[i].y()) aMin[1] = fVertices[i].y();
0694       if (aMax[1] < fVertices[i].y()) aMax[1] = fVertices[i].y();
0695       // upper fDz vertices
0696       if (aMin[0] > fVertices[i + 4].x()) aMin[0] = fVertices[i + 4].x();
0697       if (aMax[0] < fVertices[i + 4].x()) aMax[0] = fVertices[i + 4].x();
0698       if (aMin[1] > fVertices[i + 4].y()) aMin[1] = fVertices[i + 4].y();
0699       if (aMax[1] < fVertices[i + 4].y()) aMax[1] = fVertices[i + 4].y();
0700     }
0701   }
0702 
0703   /**
0704    * \brief Validate that the bottom and top quadrilaterals are convex.
0705    *
0706    * Assumes clockwise ordering in XY. Returns false if any consecutive edge
0707    * pairs produce a positive z-cross (i.e. non-convex or reversed).
0708    */
0709   VECCORE_ATT_HOST_DEVICE
0710   bool ComputeIsConvexQuadrilaterals()
0711   {
0712     for (int i = 0; i < 4; ++i) {
0713       int j = (i + 1) % 4;
0714       int k = (i + 2) % 4;
0715       // Bottom face
0716       auto vij     = fVertices[j] - fVertices[i];
0717       auto vjk     = fVertices[k] - fVertices[j];
0718       auto crossij = vij.Cross(vjk).z();
0719       if (crossij > kTolerance) return false;
0720       // Top face
0721       vij     = fVertices[j + 4] - fVertices[i + 4];
0722       vjk     = fVertices[k + 4] - fVertices[j + 4];
0723       crossij = vij.Cross(vjk).z();
0724       if (crossij > kTolerance) return false;
0725     }
0726     return true;
0727   }
0728 
0729   /**
0730    * \brief Analyze each lateral face: tag twisted/planar, precompute coefficients and mesh bounders.
0731    *
0732    * For twisted faces we compute the implicit coefficients (A..G) and normalize
0733    * them by the magnitude of the linear normal part (D,E,F) to maintain numeric stability.
0734    * For planar faces we compute the outward plane normal and its offset G.
0735    *
0736    * \return number of twisted faces
0737    */
0738   VECCORE_ATT_HOST_DEVICE
0739   int ComputeTwistedFaces()
0740   {
0741     // Check if the trapezoid is twisted. A lateral face is twisted if the top and
0742     // bottom segments are not parallel (cross product not null)
0743 
0744     int ntwisted = 0;
0745 
0746     for (int i = 0; i < 4; ++i) {
0747       auto j         = (i + 1) % 4;
0748       auto const &p1 = fVertices[i];
0749       auto const &p2 = fVertices[j];
0750       auto const &p3 = fVertices[i + 4];
0751       auto const &p4 = fVertices[j + 4];
0752       auto v12       = p2 - p1;
0753       auto v34       = p4 - p3;
0754       auto lbot      = v12.Length();
0755       auto ltop      = v34.Length();
0756       auto zcross    = v12.CrossZ(v34);
0757       auto eps       = kTolerance * Max(lbot, ltop);
0758       bool planar    = (Abs(zcross) < eps) || (Min(lbot, ltop) < kTolerance);
0759       SetTwisted(i, !planar);
0760       ntwisted += !planar;
0761       if (!planar) {
0762         // Build bounding slab for culling
0763         fMesh[i].Set(fVertices[i], fVertices[j], fVertices[i + 4], fVertices[j + 4]);
0764 
0765         // Implicit coefficients for bilinear patch (outward orientation)
0766         T a, b, c, d, e, f, g;
0767         a = -2. * fDz * (v34.y() - v12.y());
0768         b = 2. * fDz * (v34.x() - v12.x());
0769         c = -(p4 - p2).CrossZ(p3 - p1);
0770         d = -2. * fDz * fDz * (v34.y() + v12.y());
0771         e = 2. * fDz * fDz * (v34.x() + v12.x());
0772         f = 2. * fDz * (p3.CrossZ(p4) - p1.CrossZ(p2));
0773         g = -fDz * fDz * ((p4 + p2).CrossZ(p1 + p3));
0774 
0775         // Normalize by the magnitude of (d,e,f) to reduce scale sensitivity and store planar faces normalized
0776         // coefficients
0777         auto magnitude = Vertex_t(d, e, f).Mag();
0778         VECGEOM_VALIDATE(magnitude > kToleranceDist<T>, << "Wrong twist parameters");
0779         a /= magnitude;
0780         b /= magnitude;
0781         c /= magnitude;
0782         d /= magnitude;
0783         e /= magnitude;
0784         f /= magnitude;
0785         g /= magnitude;
0786         fSurf[i].Set(a, b, c, d, e, f, g);
0787       } else {
0788         // Compute normal
0789         auto normal = Vertex_t::Cross(p3 - p2, p4 - p1);
0790         if (normal.Mag2() < kTolerance) {
0791           normal.Set(0., 0., 1.); // No surface, just a line
0792         }
0793         normal.Normalize();
0794 
0795         T d, e, f, g;
0796         d = normal.x();
0797         e = normal.y();
0798         f = normal.z();
0799         g = -normal.Dot((p1 + p2 + p3 + p4) / 4.); // plane through the face centroid
0800         fSurf[i].Set(0., 0., 0., d, e, f, g);
0801       }
0802     }
0803     return ntwisted;
0804   }
0805 
0806   /**
0807    * \brief Convenience function compatible with Geant4 “twist” reporting.
0808    *
0809    * Returns a signed angle between bottom edge (i->i+1) and top edge ((i+4)->(i+5))
0810    * with sign taken from the z-cross product.
0811    */
0812   VECCORE_ATT_HOST_DEVICE
0813   Precision GetTwist(int i) const
0814   {
0815     if (!IsTwisted(i)) return 0.;
0816     auto j      = (i + 1) % 4;
0817     auto AC     = fVertices[j] - fVertices[i];
0818     auto BD     = fVertices[j + 4] - fVertices[i + 4];
0819     auto lbot   = AC.Length();
0820     auto ltop   = BD.Length();
0821     auto zcross = AC.CrossZ(BD);
0822     auto angle  = vecCore::math::ACos(AC.Dot(BD) / (lbot * ltop));
0823     return CopySign(angle, zcross);
0824   }
0825 
0826   /// \brief Debug print of geometry and per-face classification.
0827   VECCORE_ATT_HOST_DEVICE
0828   void Print() const
0829   {
0830     printf("UnplacedGenTrap: { halfZ: %f mm,  planar: %s }\n", fDz, (IsPlanar() ? "true" : "false"));
0831     printf("bottom:");
0832     for (int i = 0; i < 8; ++i) {
0833       printf(" %d:{%f, %f}", i, fVertices[i].x(), fVertices[i].y());
0834       if (i == 3) printf("\ntop:   ");
0835     }
0836     printf("\nlateral:");
0837     for (int i = 0; i < 4; ++i) {
0838       auto j = (i + 1) % 4;
0839       printf(" %d%d%d%d:", i, j, j + 4, i + 4);
0840       if (IsDegenerated(i))
0841         printf(" degenerated");
0842       else {
0843         if (IsTwisted(i))
0844           printf(" twisted");
0845         else
0846           printf(" planar");
0847       }
0848     }
0849     printf("\n");
0850   }
0851 
0852   /**
0853    * \brief Check whether two XY segments (p1,p2) and (q1,q2) properly cross.
0854    *
0855    * Colinear/parallel/degenerate pairs return false. Crossing is accepted only if
0856    * both parametric coordinates t and u lie in (kTolerance, 1-kTolerance).
0857    */
0858   VECCORE_ATT_HOST_DEVICE
0859   bool SegmentsCrossing(Vertex_t const &p1, Vertex_t const &p2, Vertex_t const &q1, Vertex_t const &q2) const
0860   {
0861     auto r         = p2 - p1;
0862     auto s         = q2 - q1;
0863     auto r_cross_s = r.CrossZ(s);
0864     if (r_cross_s < kTolerance) // parallel, colinear or degenerated - ignore crossing
0865       return false;
0866     // The segments are crossing if:
0867     //   t = ((q-p) × s) / (r × s)   and   u = ((q-p) × r) / (r × s)
0868     // fall into (kTolerance, 1-kTolerance)
0869     auto t = (q1 - p1).CrossZ(s) / r_cross_s;
0870     if (t < kTolerance || t > 1. - kTolerance) return false;
0871     auto u = (q1 - p1).CrossZ(r) / r_cross_s;
0872     if (u < kTolerance || u > 1. - kTolerance) return false;
0873     return true;
0874   }
0875 };
0876 } // namespace VECGEOM_IMPL_NAMESPACE
0877 } // namespace vecgeom
0878 
0879 #endif