From 2820eb8a142b4b0493917ddcb3a31a70a2929713 Mon Sep 17 00:00:00 2001 From: James Le Houx Date: Wed, 6 May 2026 20:47:37 +0000 Subject: [PATCH] move EffDiff setupMatrixEquation/generateActiveMask to public section MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit CUDA build of openimpala-cuda failed in 6ea391c with three nvcc errors: error: The enclosing parent function ("generateActiveMask") for an extended __device__ lambda cannot have private or protected access within its class The pre-existing LoopOnCpu calls compiled fine because they don't use extended __device__ lambdas. My new ParallelFor / ReduceOps::eval calls do, and the CUDA Programming Guide §I.4.1 forbids that inside private or protected member functions: nvcc cannot emit the necessary internal linkage for the lambda's hidden context capture. TortuosityHypre.H already exposes setupMatrixEquation, initializeDiffCoeff, buildTraversableMask, preconditionPhaseFab as public for exactly this reason (line 159+). Match that pattern for the EffDiff equivalents: move setupMatrixEquation and generateActiveMask out of `private:` and into the public section, with a comment explaining the constraint. These methods aren't part of any user-facing API — the EffectiveDiffusivityHypre class is consumed via solve() / value() / getChiSolution() — so widening their visibility has no semantic impact, only relaxes the visibility on internal helpers that user code shouldn't be calling anyway. https://claude.ai/code/session_011dJ5Bwq4Tnr8wxH597XJFf --- src/props/EffectiveDiffusivityHypre.H | 10 +++++++--- 1 file changed, 7 insertions(+), 3 deletions(-) diff --git a/src/props/EffectiveDiffusivityHypre.H b/src/props/EffectiveDiffusivityHypre.H index 26a3255b..ef8ad5be 100644 --- a/src/props/EffectiveDiffusivityHypre.H +++ b/src/props/EffectiveDiffusivityHypre.H @@ -121,12 +121,16 @@ public: return HypreStructSolver::hiV(b); } - -private: - // --- Private Methods --- + // Implementation methods. They are nominally internal but must live in + // the public section because nvcc forbids extended __device__ lambdas in + // private/protected member functions (see CUDA Programming Guide §I.4.1). + // Compare TortuosityHypre.H, which exposes the same methods publicly for + // the same reason. void setupMatrixEquation(); void generateActiveMask(); // Simpler version for D=0/1 based on phase_id + +private: // --- Member Variables --- // Configuration (solver config m_solvertype/m_eps/m_maxiter/m_verbose are in base class) std::string m_resultspath;