From 0800d6b9c3a78ce7d1409ce7962e516c0a786672 Mon Sep 17 00:00:00 2001 From: julianlitz Date: Sun, 3 May 2026 00:50:10 +0200 Subject: [PATCH 1/6] Use Kokkos::Parallel in ExtrapolatedSmootherTake ApplyAscOrtho --- .../applyAscOrtho.inl | 98 ++++++++++++++----- 1 file changed, 72 insertions(+), 26 deletions(-) diff --git a/include/ExtrapolatedSmoother/ExtrapolatedSmootherTake/applyAscOrtho.inl b/include/ExtrapolatedSmoother/ExtrapolatedSmootherTake/applyAscOrtho.inl index 7686504c..3d102b0e 100644 --- a/include/ExtrapolatedSmoother/ExtrapolatedSmootherTake/applyAscOrtho.inl +++ b/include/ExtrapolatedSmoother/ExtrapolatedSmootherTake/applyAscOrtho.inl @@ -481,16 +481,26 @@ void ExtrapolatedSmootherTake::applyAscOrthoBlackCircleSection(C ConstVector detDF = level_cache.detDF(); ConstVector coeff_beta = level_cache.coeff_beta(); + const PolarGrid* grid_ptr = &grid; + /* The outer most circle next to the radial section is defined to be black. */ const int start_black_circles = (grid.numberSmootherCircles() % 2 == 0) ? 1 : 0; - -#pragma omp parallel for num_threads(num_omp_threads) - for (int i_r = start_black_circles; i_r < grid.numberSmootherCircles(); i_r += 2) { - for (int i_theta = 0; i_theta < grid.ntheta(); i_theta++) { - nodeApplyAscOrthoCircleTake(i_r, i_theta, grid, DirBC_Interior, x, rhs, temp, arr, att, art, detDF, + const int num_black_circles = (grid.numberSmootherCircles() - start_black_circles + 1) / 2; + + Kokkos::parallel_for( + "ExtrapolatedSmootherTake: ApplyAscOrtho (Black Circular)", + Kokkos::MDRangePolicy>( // Rank of the index space + {0, 0}, // Starting point of the index space + {num_black_circles, grid.ntheta()} // Ending point of the index space + ), + // Kokkos lambda function to execute for each point in the index space + KOKKOS_LAMBDA(const int circle_task, const int i_theta) { + int i_r = start_black_circles + circle_task * 2; + nodeApplyAscOrthoCircleTake(i_r, i_theta, *grid_ptr, DirBC_Interior, x, rhs, temp, arr, att, art, detDF, coeff_beta); - } - } + }); + + Kokkos::fence(); } template @@ -514,16 +524,26 @@ void ExtrapolatedSmootherTake::applyAscOrthoWhiteCircleSection(C ConstVector detDF = level_cache.detDF(); ConstVector coeff_beta = level_cache.coeff_beta(); + const PolarGrid* grid_ptr = &grid; + /* The outer most circle next to the radial section is defined to be black. */ const int start_white_circles = (grid.numberSmootherCircles() % 2 == 0) ? 0 : 1; - -#pragma omp parallel for num_threads(num_omp_threads) - for (int i_r = start_white_circles; i_r < grid.numberSmootherCircles(); i_r += 2) { - for (int i_theta = 0; i_theta < grid.ntheta(); i_theta++) { - nodeApplyAscOrthoCircleTake(i_r, i_theta, grid, DirBC_Interior, x, rhs, temp, arr, att, art, detDF, + const int num_white_circles = (grid.numberSmootherCircles() - start_white_circles + 1) / 2; + + Kokkos::parallel_for( + "ExtrapolatedSmootherTake: ApplyAscOrtho (White Circular)", + Kokkos::MDRangePolicy>( // Rank of the index space + {0, 0}, // Starting point of the index space + {num_white_circles, grid.ntheta()} // Ending point of the index space + ), + // Kokkos lambda function to execute for each point in the index space + KOKKOS_LAMBDA(const int circle_task, const int i_theta) { + const int i_r = start_white_circles + circle_task * 2; + nodeApplyAscOrthoCircleTake(i_r, i_theta, *grid_ptr, DirBC_Interior, x, rhs, temp, arr, att, art, detDF, coeff_beta); - } - } + }); + + Kokkos::fence(); } template @@ -547,13 +567,26 @@ void ExtrapolatedSmootherTake::applyAscOrthoBlackRadialSection(C ConstVector detDF = level_cache.detDF(); ConstVector coeff_beta = level_cache.coeff_beta(); -#pragma omp parallel for num_threads(num_omp_threads) - for (int i_theta = 0; i_theta < grid.ntheta(); i_theta += 2) { - for (int i_r = grid.numberSmootherCircles(); i_r < grid.nr(); i_r++) { - nodeApplyAscOrthoRadialTake(i_r, i_theta, grid, DirBC_Interior, x, rhs, temp, arr, att, art, detDF, + const PolarGrid* grid_ptr = &grid; + + assert(grid.ntheta() % 2 == 0); + const int start_black_radials = 0; + const int num_black_radial_lines = grid.ntheta() / 2; + + Kokkos::parallel_for( + "ExtrapolatedSmootherTake: ApplyAscOrtho (Black Radial)", + Kokkos::MDRangePolicy>( // Rank of the index space + {0, grid.numberSmootherCircles()}, // Starting point of the index space + {num_black_radial_lines, grid.nr()} // Ending point of the index space + ), + // Kokkos lambda function to execute for each point in the index space + KOKKOS_LAMBDA(const int radial_task, const int i_r) { + const int i_theta = start_black_radials + radial_task * 2; + nodeApplyAscOrthoRadialTake(i_r, i_theta, *grid_ptr, DirBC_Interior, x, rhs, temp, arr, att, art, detDF, coeff_beta); - } - } + }); + + Kokkos::fence(); } template @@ -577,11 +610,24 @@ void ExtrapolatedSmootherTake::applyAscOrthoWhiteRadialSection(C ConstVector detDF = level_cache.detDF(); ConstVector coeff_beta = level_cache.coeff_beta(); -#pragma omp parallel for num_threads(num_omp_threads) - for (int i_theta = 1; i_theta < grid.ntheta(); i_theta += 2) { - for (int i_r = grid.numberSmootherCircles(); i_r < grid.nr(); i_r++) { - nodeApplyAscOrthoRadialTake(i_r, i_theta, grid, DirBC_Interior, x, rhs, temp, arr, att, art, detDF, + const PolarGrid* grid_ptr = &grid; + + assert(grid.ntheta() % 2 == 0); + const int start_white_radials = 1; + const int num_white_radial_lines = grid.ntheta() / 2; + + Kokkos::parallel_for( + "ExtrapolatedSmootherTake: ApplyAscOrtho (White Radial)", + Kokkos::MDRangePolicy>( // Rank of the index space + {0, grid.numberSmootherCircles()}, // Starting point of the index space + {num_white_radial_lines, grid.nr()} // Ending point of the index space + ), + // Kokkos lambda function to execute for each point in the index space + KOKKOS_LAMBDA(const int radial_task, const int i_r) { + const int i_theta = start_white_radials + radial_task * 2; + nodeApplyAscOrthoRadialTake(i_r, i_theta, *grid_ptr, DirBC_Interior, x, rhs, temp, arr, att, art, detDF, coeff_beta); - } - } + }); + + Kokkos::fence(); } \ No newline at end of file From 50075ffe3ba0cc2fb19ac297aa07bdbba637266c Mon Sep 17 00:00:00 2001 From: Julian Litz <91479202+julianlitz@users.noreply.github.com> Date: Wed, 6 May 2026 16:27:34 +0200 Subject: [PATCH 2/6] Update applyAscOrtho.inl --- .../ExtrapolatedSmootherTake/applyAscOrtho.inl | 16 ++++------------ 1 file changed, 4 insertions(+), 12 deletions(-) diff --git a/include/ExtrapolatedSmoother/ExtrapolatedSmootherTake/applyAscOrtho.inl b/include/ExtrapolatedSmoother/ExtrapolatedSmootherTake/applyAscOrtho.inl index 3d102b0e..86e35b3f 100644 --- a/include/ExtrapolatedSmoother/ExtrapolatedSmootherTake/applyAscOrtho.inl +++ b/include/ExtrapolatedSmoother/ExtrapolatedSmootherTake/applyAscOrtho.inl @@ -481,8 +481,6 @@ void ExtrapolatedSmootherTake::applyAscOrthoBlackCircleSection(C ConstVector detDF = level_cache.detDF(); ConstVector coeff_beta = level_cache.coeff_beta(); - const PolarGrid* grid_ptr = &grid; - /* The outer most circle next to the radial section is defined to be black. */ const int start_black_circles = (grid.numberSmootherCircles() % 2 == 0) ? 1 : 0; const int num_black_circles = (grid.numberSmootherCircles() - start_black_circles + 1) / 2; @@ -496,7 +494,7 @@ void ExtrapolatedSmootherTake::applyAscOrthoBlackCircleSection(C // Kokkos lambda function to execute for each point in the index space KOKKOS_LAMBDA(const int circle_task, const int i_theta) { int i_r = start_black_circles + circle_task * 2; - nodeApplyAscOrthoCircleTake(i_r, i_theta, *grid_ptr, DirBC_Interior, x, rhs, temp, arr, att, art, detDF, + nodeApplyAscOrthoCircleTake(i_r, i_theta, grid, DirBC_Interior, x, rhs, temp, arr, att, art, detDF, coeff_beta); }); @@ -524,8 +522,6 @@ void ExtrapolatedSmootherTake::applyAscOrthoWhiteCircleSection(C ConstVector detDF = level_cache.detDF(); ConstVector coeff_beta = level_cache.coeff_beta(); - const PolarGrid* grid_ptr = &grid; - /* The outer most circle next to the radial section is defined to be black. */ const int start_white_circles = (grid.numberSmootherCircles() % 2 == 0) ? 0 : 1; const int num_white_circles = (grid.numberSmootherCircles() - start_white_circles + 1) / 2; @@ -539,7 +535,7 @@ void ExtrapolatedSmootherTake::applyAscOrthoWhiteCircleSection(C // Kokkos lambda function to execute for each point in the index space KOKKOS_LAMBDA(const int circle_task, const int i_theta) { const int i_r = start_white_circles + circle_task * 2; - nodeApplyAscOrthoCircleTake(i_r, i_theta, *grid_ptr, DirBC_Interior, x, rhs, temp, arr, att, art, detDF, + nodeApplyAscOrthoCircleTake(i_r, i_theta, grid, DirBC_Interior, x, rhs, temp, arr, att, art, detDF, coeff_beta); }); @@ -567,8 +563,6 @@ void ExtrapolatedSmootherTake::applyAscOrthoBlackRadialSection(C ConstVector detDF = level_cache.detDF(); ConstVector coeff_beta = level_cache.coeff_beta(); - const PolarGrid* grid_ptr = &grid; - assert(grid.ntheta() % 2 == 0); const int start_black_radials = 0; const int num_black_radial_lines = grid.ntheta() / 2; @@ -582,7 +576,7 @@ void ExtrapolatedSmootherTake::applyAscOrthoBlackRadialSection(C // Kokkos lambda function to execute for each point in the index space KOKKOS_LAMBDA(const int radial_task, const int i_r) { const int i_theta = start_black_radials + radial_task * 2; - nodeApplyAscOrthoRadialTake(i_r, i_theta, *grid_ptr, DirBC_Interior, x, rhs, temp, arr, att, art, detDF, + nodeApplyAscOrthoRadialTake(i_r, i_theta, grid, DirBC_Interior, x, rhs, temp, arr, att, art, detDF, coeff_beta); }); @@ -610,8 +604,6 @@ void ExtrapolatedSmootherTake::applyAscOrthoWhiteRadialSection(C ConstVector detDF = level_cache.detDF(); ConstVector coeff_beta = level_cache.coeff_beta(); - const PolarGrid* grid_ptr = &grid; - assert(grid.ntheta() % 2 == 0); const int start_white_radials = 1; const int num_white_radial_lines = grid.ntheta() / 2; @@ -625,7 +617,7 @@ void ExtrapolatedSmootherTake::applyAscOrthoWhiteRadialSection(C // Kokkos lambda function to execute for each point in the index space KOKKOS_LAMBDA(const int radial_task, const int i_r) { const int i_theta = start_white_radials + radial_task * 2; - nodeApplyAscOrthoRadialTake(i_r, i_theta, *grid_ptr, DirBC_Interior, x, rhs, temp, arr, att, art, detDF, + nodeApplyAscOrthoRadialTake(i_r, i_theta, grid, DirBC_Interior, x, rhs, temp, arr, att, art, detDF, coeff_beta); }); From f2f4aadbb0811c2fd2a88e2b37bf34b0f8937359 Mon Sep 17 00:00:00 2001 From: BOURNE Emily EPFL Date: Tue, 12 May 2026 13:44:20 +0200 Subject: [PATCH 3/6] Missing annotation --- .../ExtrapolatedSmootherGive/applyAscOrtho.inl | 6 +++--- .../ExtrapolatedSmootherTake/applyAscOrtho.inl | 6 +++--- include/Smoother/SmootherGive/applyAscOrtho.inl | 6 +++--- 3 files changed, 9 insertions(+), 9 deletions(-) diff --git a/include/ExtrapolatedSmoother/ExtrapolatedSmootherGive/applyAscOrtho.inl b/include/ExtrapolatedSmoother/ExtrapolatedSmootherGive/applyAscOrtho.inl index 3a845970..e1b1a8cd 100644 --- a/include/ExtrapolatedSmoother/ExtrapolatedSmootherGive/applyAscOrtho.inl +++ b/include/ExtrapolatedSmoother/ExtrapolatedSmootherGive/applyAscOrtho.inl @@ -3,7 +3,7 @@ namespace extrapolated_smoother_give { -static inline void nodeApplyAscOrthoCircleGive(int i_r, int i_theta, const PolarGrid& grid, bool DirBC_Interior, +static KOKKOS_INLINE_FUNCTION void nodeApplyAscOrthoCircleGive(int i_r, int i_theta, const PolarGrid& grid, bool DirBC_Interior, SmootherColor smoother_color, ConstVector& x, ConstVector& rhs, Vector& result, double arr, double att, double art, double detDF, double coeff_beta) @@ -400,7 +400,7 @@ static inline void nodeApplyAscOrthoCircleGive(int i_r, int i_theta, const Polar namespace extrapolated_smoother_give { -static inline void nodeApplyAscOrthoRadialGive(int i_r, int i_theta, const PolarGrid& grid, bool DirBC_Interior, +static KOKKOS_INLINE_FUNCTION void nodeApplyAscOrthoRadialGive(int i_r, int i_theta, const PolarGrid& grid, bool DirBC_Interior, SmootherColor smoother_color, ConstVector& x, ConstVector& rhs, Vector& result, double arr, double att, double art, double detDF, double coeff_beta) @@ -946,4 +946,4 @@ void ExtrapolatedSmootherGive::applyAscOrthoRadialSection(int i_ nodeApplyAscOrthoRadialGive(i_r, i_theta, grid, DirBC_Interior, smoother_color, x, rhs, temp, arr, att, art, detDF, coeff_beta); } -} \ No newline at end of file +} diff --git a/include/ExtrapolatedSmoother/ExtrapolatedSmootherTake/applyAscOrtho.inl b/include/ExtrapolatedSmoother/ExtrapolatedSmootherTake/applyAscOrtho.inl index 1aceed82..4bd15e09 100644 --- a/include/ExtrapolatedSmoother/ExtrapolatedSmootherTake/applyAscOrtho.inl +++ b/include/ExtrapolatedSmoother/ExtrapolatedSmootherTake/applyAscOrtho.inl @@ -3,7 +3,7 @@ namespace extrapolated_smoother_take { -static inline void nodeApplyAscOrthoCircleTake(int i_r, int i_theta, const PolarGrid& grid, bool DirBC_Interior, +static KOKKOS_INLINE_FUNCTION void nodeApplyAscOrthoCircleTake(int i_r, int i_theta, const PolarGrid& grid, bool DirBC_Interior, ConstVector& x, ConstVector& rhs, Vector& result, ConstVector& arr, ConstVector& att, ConstVector& art, ConstVector& detDF, @@ -185,7 +185,7 @@ static inline void nodeApplyAscOrthoCircleTake(int i_r, int i_theta, const Polar } } -static inline void nodeApplyAscOrthoRadialTake(int i_r, int i_theta, const PolarGrid& grid, bool DirBC_Interior, +static KOKKOS_INLINE_FUNCTION void nodeApplyAscOrthoRadialTake(int i_r, int i_theta, const PolarGrid& grid, bool DirBC_Interior, ConstVector& x, ConstVector& rhs, Vector& result, ConstVector& arr, const ConstVector& att, ConstVector& art, const ConstVector& detDF, @@ -622,4 +622,4 @@ void ExtrapolatedSmootherTake::applyAscOrthoWhiteRadialSection(C }); Kokkos::fence(); -} \ No newline at end of file +} diff --git a/include/Smoother/SmootherGive/applyAscOrtho.inl b/include/Smoother/SmootherGive/applyAscOrtho.inl index f91f5366..70b1d5b4 100644 --- a/include/Smoother/SmootherGive/applyAscOrtho.inl +++ b/include/Smoother/SmootherGive/applyAscOrtho.inl @@ -3,7 +3,7 @@ namespace smoother_give { -static inline void nodeApplyAscOrthoCircleGive(int i_r, int i_theta, const PolarGrid& grid, bool DirBC_Interior, +static KOKKOS_INLINE_FUNCTION void nodeApplyAscOrthoCircleGive(int i_r, int i_theta, const PolarGrid& grid, bool DirBC_Interior, SmootherColor smoother_color, ConstVector& x, ConstVector& rhs, Vector& result, double arr, double att, double art, double detDF, double coeff_beta) @@ -187,7 +187,7 @@ static inline void nodeApplyAscOrthoCircleGive(int i_r, int i_theta, const Polar } } -static inline void nodeApplyAscOrthoRadialGive(int i_r, int i_theta, const PolarGrid& grid, bool DirBC_Interior, +static KOKKOS_INLINE_FUNCTION void nodeApplyAscOrthoRadialGive(int i_r, int i_theta, const PolarGrid& grid, bool DirBC_Interior, SmootherColor smoother_color, ConstVector& x, ConstVector& rhs, Vector& result, double arr, double att, double art, double detDF, double coeff_beta) @@ -479,4 +479,4 @@ void SmootherGive::applyAscOrthoRadialSection(const int i_theta, nodeApplyAscOrthoRadialGive(i_r, i_theta, grid, DirBC_Interior, smoother_color, x, rhs, temp, arr, att, art, detDF, coeff_beta); } -} \ No newline at end of file +} From 4231c6c16fdbc85e1698a725809d29b33f8bbf16 Mon Sep 17 00:00:00 2001 From: BOURNE Emily EPFL Date: Tue, 12 May 2026 13:44:28 +0200 Subject: [PATCH 4/6] Functions must be public --- .../ExtrapolatedSmootherTake/extrapolatedSmootherTake.h | 4 +++- 1 file changed, 3 insertions(+), 1 deletion(-) diff --git a/include/ExtrapolatedSmoother/ExtrapolatedSmootherTake/extrapolatedSmootherTake.h b/include/ExtrapolatedSmoother/ExtrapolatedSmootherTake/extrapolatedSmootherTake.h index 8f5e78fe..530aa44a 100644 --- a/include/ExtrapolatedSmoother/ExtrapolatedSmootherTake/extrapolatedSmootherTake.h +++ b/include/ExtrapolatedSmoother/ExtrapolatedSmootherTake/extrapolatedSmootherTake.h @@ -172,6 +172,8 @@ class ExtrapolatedSmootherTake : public ExtrapolatedSmoother /* Orthogonal application */ /* ---------------------- */ + // Functions must be public due to cuda restriction +public: // Compute temp = f_sc − A_sc^ortho * u_sc^ortho (precomputed right-hand side) // where x = u_sc and rhs = f_sc void applyAscOrthoBlackCircleSection(ConstVector x, ConstVector rhs, Vector temp); @@ -182,7 +184,7 @@ class ExtrapolatedSmootherTake : public ExtrapolatedSmoother /* ----------------- */ /* Line-wise solvers */ /* ----------------- */ - +private: // Solve the linear system: // A_sc * u_sc = f_sc − A_sc^ortho * u_sc^ortho // Parameter mapping: From 437a57cf42507348bcca255b6ccbe36e75682ead Mon Sep 17 00:00:00 2001 From: BOURNE Emily EPFL Date: Tue, 12 May 2026 13:49:13 +0200 Subject: [PATCH 5/6] Run on host --- .../ExtrapolatedSmootherTake/applyAscOrtho.inl | 8 ++++---- 1 file changed, 4 insertions(+), 4 deletions(-) diff --git a/include/ExtrapolatedSmoother/ExtrapolatedSmootherTake/applyAscOrtho.inl b/include/ExtrapolatedSmoother/ExtrapolatedSmootherTake/applyAscOrtho.inl index 4bd15e09..aa9f3f20 100644 --- a/include/ExtrapolatedSmoother/ExtrapolatedSmootherTake/applyAscOrtho.inl +++ b/include/ExtrapolatedSmoother/ExtrapolatedSmootherTake/applyAscOrtho.inl @@ -487,7 +487,7 @@ void ExtrapolatedSmootherTake::applyAscOrthoBlackCircleSection(C Kokkos::parallel_for( "ExtrapolatedSmootherTake: ApplyAscOrtho (Black Circular)", - Kokkos::MDRangePolicy>( // Rank of the index space + Kokkos::MDRangePolicy>( // Rank of the index space {0, 0}, // Starting point of the index space {num_black_circles, grid.ntheta()} // Ending point of the index space ), @@ -528,7 +528,7 @@ void ExtrapolatedSmootherTake::applyAscOrthoWhiteCircleSection(C Kokkos::parallel_for( "ExtrapolatedSmootherTake: ApplyAscOrtho (White Circular)", - Kokkos::MDRangePolicy>( // Rank of the index space + Kokkos::MDRangePolicy>( // Rank of the index space {0, 0}, // Starting point of the index space {num_white_circles, grid.ntheta()} // Ending point of the index space ), @@ -569,7 +569,7 @@ void ExtrapolatedSmootherTake::applyAscOrthoBlackRadialSection(C Kokkos::parallel_for( "ExtrapolatedSmootherTake: ApplyAscOrtho (Black Radial)", - Kokkos::MDRangePolicy>( // Rank of the index space + Kokkos::MDRangePolicy>( // Rank of the index space {0, grid.numberSmootherCircles()}, // Starting point of the index space {num_black_radial_lines, grid.nr()} // Ending point of the index space ), @@ -610,7 +610,7 @@ void ExtrapolatedSmootherTake::applyAscOrthoWhiteRadialSection(C Kokkos::parallel_for( "ExtrapolatedSmootherTake: ApplyAscOrtho (White Radial)", - Kokkos::MDRangePolicy>( // Rank of the index space + Kokkos::MDRangePolicy>( // Rank of the index space {0, grid.numberSmootherCircles()}, // Starting point of the index space {num_white_radial_lines, grid.nr()} // Ending point of the index space ), From 029128ca2ccac4a2cce2d948d98a1ed811de91b3 Mon Sep 17 00:00:00 2001 From: Emily Bourne Date: Tue, 12 May 2026 13:51:09 +0200 Subject: [PATCH 6/6] Clang formatting --- .../DirectSolverGive/buildSolverMatrix.inl | 4 ++-- .../DirectSolverTake/buildSolverMatrix.inl | 4 ++-- .../applyAscOrtho.inl | 18 +++++++++------- .../buildInnerBoundaryAsc.inl | 4 ++-- .../applyAscOrtho.inl | 21 ++++++++++--------- .../buildInnerBoundaryAsc.inl | 4 ++-- .../extrapolatedSmootherTake.h | 2 +- .../Smoother/SmootherGive/applyAscOrtho.inl | 18 +++++++++------- .../SmootherGive/buildInnerBoundaryAsc.inl | 4 ++-- .../SmootherTake/buildInnerBoundaryAsc.inl | 6 +++--- 10 files changed, 45 insertions(+), 40 deletions(-) diff --git a/include/DirectSolver/DirectSolverGive/buildSolverMatrix.inl b/include/DirectSolver/DirectSolverGive/buildSolverMatrix.inl index 59637462..48b2877a 100644 --- a/include/DirectSolver/DirectSolverGive/buildSolverMatrix.inl +++ b/include/DirectSolver/DirectSolverGive/buildSolverMatrix.inl @@ -5,8 +5,8 @@ namespace direct_solver_coo_mumps_give #ifdef GMGPOLAR_USE_MUMPS // When using the MUMPS solver, the matrix is assembled in COO format. -static inline void updateMatrixElement(SparseMatrixCOO& matrix, int ptr, int offset, int row, int column, - double value) +static inline void updateMatrixElement(SparseMatrixCOO& matrix, int ptr, int offset, int row, + int column, double value) { matrix.set_row_index(ptr + offset, row); matrix.set_col_index(ptr + offset, column); diff --git a/include/DirectSolver/DirectSolverTake/buildSolverMatrix.inl b/include/DirectSolver/DirectSolverTake/buildSolverMatrix.inl index 4cdb4487..a144a9c6 100644 --- a/include/DirectSolver/DirectSolverTake/buildSolverMatrix.inl +++ b/include/DirectSolver/DirectSolverTake/buildSolverMatrix.inl @@ -5,8 +5,8 @@ namespace direct_solver_coo_mumps_take #ifdef GMGPOLAR_USE_MUMPS // When using the MUMPS solver, the matrix is assembled in COO format. -static inline void updateMatrixElement(SparseMatrixCOO& matrix, int ptr, int offset, int row, int column, - double value) +static inline void updateMatrixElement(SparseMatrixCOO& matrix, int ptr, int offset, int row, + int column, double value) { matrix.set_row_index(ptr + offset, row); matrix.set_col_index(ptr + offset, column); diff --git a/include/ExtrapolatedSmoother/ExtrapolatedSmootherGive/applyAscOrtho.inl b/include/ExtrapolatedSmoother/ExtrapolatedSmootherGive/applyAscOrtho.inl index e1b1a8cd..7371c2c5 100644 --- a/include/ExtrapolatedSmoother/ExtrapolatedSmootherGive/applyAscOrtho.inl +++ b/include/ExtrapolatedSmoother/ExtrapolatedSmootherGive/applyAscOrtho.inl @@ -3,10 +3,11 @@ namespace extrapolated_smoother_give { -static KOKKOS_INLINE_FUNCTION void nodeApplyAscOrthoCircleGive(int i_r, int i_theta, const PolarGrid& grid, bool DirBC_Interior, - SmootherColor smoother_color, ConstVector& x, - ConstVector& rhs, Vector& result, double arr, double att, - double art, double detDF, double coeff_beta) +static KOKKOS_INLINE_FUNCTION void nodeApplyAscOrthoCircleGive(int i_r, int i_theta, const PolarGrid& grid, + bool DirBC_Interior, SmootherColor smoother_color, + ConstVector& x, ConstVector& rhs, + Vector& result, double arr, double att, + double art, double detDF, double coeff_beta) { assert(i_r >= 0 && i_r <= grid.numberSmootherCircles()); @@ -400,10 +401,11 @@ static KOKKOS_INLINE_FUNCTION void nodeApplyAscOrthoCircleGive(int i_r, int i_th namespace extrapolated_smoother_give { -static KOKKOS_INLINE_FUNCTION void nodeApplyAscOrthoRadialGive(int i_r, int i_theta, const PolarGrid& grid, bool DirBC_Interior, - SmootherColor smoother_color, ConstVector& x, - ConstVector& rhs, Vector& result, double arr, double att, - double art, double detDF, double coeff_beta) +static KOKKOS_INLINE_FUNCTION void nodeApplyAscOrthoRadialGive(int i_r, int i_theta, const PolarGrid& grid, + bool DirBC_Interior, SmootherColor smoother_color, + ConstVector& x, ConstVector& rhs, + Vector& result, double arr, double att, + double art, double detDF, double coeff_beta) { assert(i_r >= grid.numberSmootherCircles() - 1 && i_r < grid.nr()); diff --git a/include/ExtrapolatedSmoother/ExtrapolatedSmootherGive/buildInnerBoundaryAsc.inl b/include/ExtrapolatedSmoother/ExtrapolatedSmootherGive/buildInnerBoundaryAsc.inl index 0ff8b1b2..ad6b3362 100644 --- a/include/ExtrapolatedSmoother/ExtrapolatedSmootherGive/buildInnerBoundaryAsc.inl +++ b/include/ExtrapolatedSmoother/ExtrapolatedSmootherGive/buildInnerBoundaryAsc.inl @@ -7,8 +7,8 @@ namespace extrapolated_smoother_give #ifdef GMGPOLAR_USE_MUMPS // When using the MUMPS solver, the matrix is assembled in COO format. -static inline void update_CSR_COO_MatrixElement(SparseMatrixCOO& matrix, int ptr, int offset, int row, - int column, double value) +static inline void update_CSR_COO_MatrixElement(SparseMatrixCOO& matrix, int ptr, int offset, + int row, int column, double value) { matrix.set_row_index(ptr + offset, row); matrix.set_col_index(ptr + offset, column); diff --git a/include/ExtrapolatedSmoother/ExtrapolatedSmootherTake/applyAscOrtho.inl b/include/ExtrapolatedSmoother/ExtrapolatedSmootherTake/applyAscOrtho.inl index aa9f3f20..32b4f420 100644 --- a/include/ExtrapolatedSmoother/ExtrapolatedSmootherTake/applyAscOrtho.inl +++ b/include/ExtrapolatedSmoother/ExtrapolatedSmootherTake/applyAscOrtho.inl @@ -3,11 +3,12 @@ namespace extrapolated_smoother_take { -static KOKKOS_INLINE_FUNCTION void nodeApplyAscOrthoCircleTake(int i_r, int i_theta, const PolarGrid& grid, bool DirBC_Interior, - ConstVector& x, ConstVector& rhs, Vector& result, - ConstVector& arr, ConstVector& att, - ConstVector& art, ConstVector& detDF, - ConstVector& coeff_beta) +static KOKKOS_INLINE_FUNCTION void nodeApplyAscOrthoCircleTake(int i_r, int i_theta, const PolarGrid& grid, + bool DirBC_Interior, ConstVector& x, + ConstVector& rhs, Vector& result, + ConstVector& arr, ConstVector& att, + ConstVector& art, ConstVector& detDF, + ConstVector& coeff_beta) { assert(i_r >= 0 && i_r <= grid.numberSmootherCircles()); @@ -185,11 +186,11 @@ static KOKKOS_INLINE_FUNCTION void nodeApplyAscOrthoCircleTake(int i_r, int i_th } } -static KOKKOS_INLINE_FUNCTION void nodeApplyAscOrthoRadialTake(int i_r, int i_theta, const PolarGrid& grid, bool DirBC_Interior, - ConstVector& x, ConstVector& rhs, Vector& result, - ConstVector& arr, const ConstVector& att, - ConstVector& art, const ConstVector& detDF, - ConstVector& coeff_beta) +static KOKKOS_INLINE_FUNCTION void +nodeApplyAscOrthoRadialTake(int i_r, int i_theta, const PolarGrid& grid, bool DirBC_Interior, ConstVector& x, + ConstVector& rhs, Vector& result, ConstVector& arr, + const ConstVector& att, ConstVector& art, const ConstVector& detDF, + ConstVector& coeff_beta) { assert(i_r >= grid.numberSmootherCircles() - 1 && i_r < grid.nr()); diff --git a/include/ExtrapolatedSmoother/ExtrapolatedSmootherTake/buildInnerBoundaryAsc.inl b/include/ExtrapolatedSmoother/ExtrapolatedSmootherTake/buildInnerBoundaryAsc.inl index aae19b48..ebf2eb4d 100644 --- a/include/ExtrapolatedSmoother/ExtrapolatedSmootherTake/buildInnerBoundaryAsc.inl +++ b/include/ExtrapolatedSmoother/ExtrapolatedSmootherTake/buildInnerBoundaryAsc.inl @@ -5,8 +5,8 @@ namespace extrapolated_smoother_take #ifdef GMGPOLAR_USE_MUMPS // When using the MUMPS solver, the matrix is assembled in COO format. -static inline void update_CSR_COO_MatrixElement(SparseMatrixCOO& matrix, int ptr, int offset, int row, - int column, double value) +static inline void update_CSR_COO_MatrixElement(SparseMatrixCOO& matrix, int ptr, int offset, + int row, int column, double value) { matrix.set_row_index(ptr + offset, row); matrix.set_col_index(ptr + offset, column); diff --git a/include/ExtrapolatedSmoother/ExtrapolatedSmootherTake/extrapolatedSmootherTake.h b/include/ExtrapolatedSmoother/ExtrapolatedSmootherTake/extrapolatedSmootherTake.h index 530aa44a..099a790e 100644 --- a/include/ExtrapolatedSmoother/ExtrapolatedSmootherTake/extrapolatedSmootherTake.h +++ b/include/ExtrapolatedSmoother/ExtrapolatedSmootherTake/extrapolatedSmootherTake.h @@ -172,7 +172,7 @@ class ExtrapolatedSmootherTake : public ExtrapolatedSmoother /* Orthogonal application */ /* ---------------------- */ - // Functions must be public due to cuda restriction + // Functions must be public due to cuda restriction public: // Compute temp = f_sc − A_sc^ortho * u_sc^ortho (precomputed right-hand side) // where x = u_sc and rhs = f_sc diff --git a/include/Smoother/SmootherGive/applyAscOrtho.inl b/include/Smoother/SmootherGive/applyAscOrtho.inl index 70b1d5b4..c197e188 100644 --- a/include/Smoother/SmootherGive/applyAscOrtho.inl +++ b/include/Smoother/SmootherGive/applyAscOrtho.inl @@ -3,10 +3,11 @@ namespace smoother_give { -static KOKKOS_INLINE_FUNCTION void nodeApplyAscOrthoCircleGive(int i_r, int i_theta, const PolarGrid& grid, bool DirBC_Interior, - SmootherColor smoother_color, ConstVector& x, - ConstVector& rhs, Vector& result, double arr, double att, - double art, double detDF, double coeff_beta) +static KOKKOS_INLINE_FUNCTION void nodeApplyAscOrthoCircleGive(int i_r, int i_theta, const PolarGrid& grid, + bool DirBC_Interior, SmootherColor smoother_color, + ConstVector& x, ConstVector& rhs, + Vector& result, double arr, double att, + double art, double detDF, double coeff_beta) { assert(i_r >= 0 && i_r <= grid.numberSmootherCircles()); @@ -187,10 +188,11 @@ static KOKKOS_INLINE_FUNCTION void nodeApplyAscOrthoCircleGive(int i_r, int i_th } } -static KOKKOS_INLINE_FUNCTION void nodeApplyAscOrthoRadialGive(int i_r, int i_theta, const PolarGrid& grid, bool DirBC_Interior, - SmootherColor smoother_color, ConstVector& x, - ConstVector& rhs, Vector& result, double arr, double att, - double art, double detDF, double coeff_beta) +static KOKKOS_INLINE_FUNCTION void nodeApplyAscOrthoRadialGive(int i_r, int i_theta, const PolarGrid& grid, + bool DirBC_Interior, SmootherColor smoother_color, + ConstVector& x, ConstVector& rhs, + Vector& result, double arr, double att, + double art, double detDF, double coeff_beta) { assert(i_r >= grid.numberSmootherCircles() - 1 && i_r < grid.nr()); diff --git a/include/Smoother/SmootherGive/buildInnerBoundaryAsc.inl b/include/Smoother/SmootherGive/buildInnerBoundaryAsc.inl index dd92f089..c59c0359 100644 --- a/include/Smoother/SmootherGive/buildInnerBoundaryAsc.inl +++ b/include/Smoother/SmootherGive/buildInnerBoundaryAsc.inl @@ -5,8 +5,8 @@ namespace smoother_give #ifdef GMGPOLAR_USE_MUMPS // When using the MUMPS solver, the matrix is assembled in COO format. -static inline void update_CSR_COO_MatrixElement(SparseMatrixCOO& matrix, int ptr, int offset, int row, - int column, double value) +static inline void update_CSR_COO_MatrixElement(SparseMatrixCOO& matrix, int ptr, int offset, + int row, int column, double value) { matrix.set_row_index(ptr + offset, row); matrix.set_col_index(ptr + offset, column); diff --git a/include/Smoother/SmootherTake/buildInnerBoundaryAsc.inl b/include/Smoother/SmootherTake/buildInnerBoundaryAsc.inl index b2f3cfd3..243962c9 100644 --- a/include/Smoother/SmootherTake/buildInnerBoundaryAsc.inl +++ b/include/Smoother/SmootherTake/buildInnerBoundaryAsc.inl @@ -5,12 +5,12 @@ namespace smoother_take #ifdef GMGPOLAR_USE_MUMPS // When using the MUMPS solver, the matrix is assembled in COO format. -static inline void update_CSR_COO_MatrixElement(SparseMatrixCOO& matrix, int ptr, int offset, int row, - int column, double value) +static inline void update_CSR_COO_MatrixElement(SparseMatrixCOO& matrix, int ptr, int offset, + int row, int column, double value) { matrix.set_row_index(ptr + offset, row); matrix.set_col_index(ptr + offset, column); - matrix.set_value(ptr + offset , value); + matrix.set_value(ptr + offset, value); } #else // When using the in-house solver, the matrix is stored in CSR format.