From cb2487f81fa6df1d2bed548f8f54eae1f2c3e2e7 Mon Sep 17 00:00:00 2001 From: julianlitz Date: Sat, 2 May 2026 23:15:57 +0200 Subject: [PATCH 01/11] Use Kokkos::Parallel in Smoother (moveUpdatedValues) --- .../solveAscSystem.inl | 64 ++++++++++++------- .../solveAscSystem.inl | 64 ++++++++++++------- .../Smoother/SmootherGive/solveAscSystem.inl | 64 ++++++++++++------- .../Smoother/SmootherTake/solveAscSystem.inl | 64 ++++++++++++------- 4 files changed, 160 insertions(+), 96 deletions(-) diff --git a/include/ExtrapolatedSmoother/ExtrapolatedSmootherGive/solveAscSystem.inl b/include/ExtrapolatedSmoother/ExtrapolatedSmootherGive/solveAscSystem.inl index e17e6278..7941c60e 100644 --- a/include/ExtrapolatedSmoother/ExtrapolatedSmootherGive/solveAscSystem.inl +++ b/include/ExtrapolatedSmoother/ExtrapolatedSmootherGive/solveAscSystem.inl @@ -38,13 +38,16 @@ void ExtrapolatedSmootherGive::solveBlackCircleSection(Vector>({0, 0}, {num_black_circles, grid.ntheta()}), + KOKKOS_LAMBDA(const int circle_task, const int i_theta) { + const int i_r = start_black_circles + circle_task * 2; x[grid.index(i_r, i_theta)] = temp[grid.index(i_r, i_theta)]; - } - } + }); + Kokkos::fence(); } template @@ -85,13 +88,16 @@ void ExtrapolatedSmootherGive::solveWhiteCircleSection(Vector>({0, 0}, {num_white_circles, grid.ntheta()}), + KOKKOS_LAMBDA(const int circle_task, const int i_theta) { + const int i_r = start_white_circles + circle_task * 2; x[grid.index(i_r, i_theta)] = temp[grid.index(i_r, i_theta)]; - } - } + }); + Kokkos::fence(); } template @@ -108,13 +114,18 @@ void ExtrapolatedSmootherGive::solveBlackRadialSection(Vector>({0, grid.numberSmootherCircles()}, {num_black_radial_lines, grid.nr()}), + KOKKOS_LAMBDA(const int radial_task, const int i_r) { + const int i_theta = start_black_radials + radial_task * 2; x[grid.index(i_r, i_theta)] = temp[grid.index(i_r, i_theta)]; - } - } + }); + Kokkos::fence(); } template @@ -131,11 +142,16 @@ void ExtrapolatedSmootherGive::solveWhiteRadialSection(Vector>({0, grid.numberSmootherCircles()}, {num_white_radial_lines, grid.nr()}), + KOKKOS_LAMBDA(const int radial_task, const int i_r) { + const int i_theta = start_white_radials + radial_task * 2; x[grid.index(i_r, i_theta)] = temp[grid.index(i_r, i_theta)]; - } - } + }); + Kokkos::fence(); } diff --git a/include/ExtrapolatedSmoother/ExtrapolatedSmootherTake/solveAscSystem.inl b/include/ExtrapolatedSmoother/ExtrapolatedSmootherTake/solveAscSystem.inl index 5b677e18..879f1645 100644 --- a/include/ExtrapolatedSmoother/ExtrapolatedSmootherTake/solveAscSystem.inl +++ b/include/ExtrapolatedSmoother/ExtrapolatedSmootherTake/solveAscSystem.inl @@ -27,13 +27,16 @@ void ExtrapolatedSmootherTake::solveBlackCircleSection(Vector>({0, 0}, {num_black_circles, grid.ntheta()}), + KOKKOS_LAMBDA(const int circle_task, const int i_theta) { + const int i_r = start_black_circles + circle_task * 2; x[grid.index(i_r, i_theta)] = temp[grid.index(i_r, i_theta)]; - } - } + }); + Kokkos::fence(); } template @@ -63,13 +66,16 @@ void ExtrapolatedSmootherTake::solveWhiteCircleSection(Vector>({0, 0}, {num_white_circles, grid.ntheta()}), + KOKKOS_LAMBDA(const int circle_task, const int i_theta) { + const int i_r = start_white_circles + circle_task * 2; x[grid.index(i_r, i_theta)] = temp[grid.index(i_r, i_theta)]; - } - } + }); + Kokkos::fence(); } template @@ -86,13 +92,18 @@ void ExtrapolatedSmootherTake::solveBlackRadialSection(Vector>({0, grid.numberSmootherCircles()}, {num_black_radial_lines, grid.nr()}), + KOKKOS_LAMBDA(const int radial_task, const int i_r) { + const int i_theta = start_black_radials + radial_task * 2; x[grid.index(i_r, i_theta)] = temp[grid.index(i_r, i_theta)]; - } - } + }); + Kokkos::fence(); } template @@ -109,11 +120,16 @@ void ExtrapolatedSmootherTake::solveWhiteRadialSection(Vector>({0, grid.numberSmootherCircles()}, {num_white_radial_lines, grid.nr()}), + KOKKOS_LAMBDA(const int radial_task, const int i_r) { + const int i_theta = start_white_radials + radial_task * 2; x[grid.index(i_r, i_theta)] = temp[grid.index(i_r, i_theta)]; - } - } + }); + Kokkos::fence(); } diff --git a/include/Smoother/SmootherGive/solveAscSystem.inl b/include/Smoother/SmootherGive/solveAscSystem.inl index 37c7a2d3..8a183957 100644 --- a/include/Smoother/SmootherGive/solveAscSystem.inl +++ b/include/Smoother/SmootherGive/solveAscSystem.inl @@ -33,13 +33,16 @@ void SmootherGive::solveBlackCircleSection(Vector x, Vec } // Move updated values to x - int start_black_circles = is_inner_circle_black ? 0 : 1; -#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++) { + const int start_black_circles = is_inner_circle_black ? 0 : 1; + const int num_black_circles = (grid.numberSmootherCircles() - start_black_circles + 1) / 2; + Kokkos::parallel_for( + "SmootherGive: moveUpdatedValues (Black Circular)", + Kokkos::MDRangePolicy>({0, 0}, {num_black_circles, grid.ntheta()}), + KOKKOS_LAMBDA(const int circle_task, const int i_theta) { + const int i_r = start_black_circles + circle_task * 2; x[grid.index(i_r, i_theta)] = temp[grid.index(i_r, i_theta)]; - } - } + }); + Kokkos::fence(); } template @@ -75,13 +78,16 @@ void SmootherGive::solveWhiteCircleSection(Vector x, Vec } // Move updated values to x - int start_white_circles = is_inner_circle_white ? 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++) { + const int start_white_circles = is_inner_circle_white ? 0 : 1; + const int num_white_circles = (grid.numberSmootherCircles() - start_white_circles + 1) / 2; + Kokkos::parallel_for( + "SmootherGive: moveUpdatedValues (White Circular)", + Kokkos::MDRangePolicy>({0, 0}, {num_white_circles, grid.ntheta()}), + KOKKOS_LAMBDA(const int circle_task, const int i_theta) { + const int i_r = start_white_circles + circle_task * 2; x[grid.index(i_r, i_theta)] = temp[grid.index(i_r, i_theta)]; - } - } + }); + Kokkos::fence(); } template @@ -98,13 +104,18 @@ void SmootherGive::solveBlackRadialSection(Vector x, Vec int batch_stride = 2; radial_tridiagonal_solver_.solve(radial_section, batch_offset, batch_stride); -// Move updated values to x -#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++) { + // Move updated values to x + assert(grid.ntheta() % 2 == 0); + const int start_black_radials = 0; + const int num_black_radial_lines = grid.ntheta() / 2; + Kokkos::parallel_for( + "SmootherGive: moveUpdatedValues (Black Radial)", + Kokkos::MDRangePolicy>({0, grid.numberSmootherCircles()}, {num_black_radial_lines, grid.nr()}), + KOKKOS_LAMBDA(const int radial_task, const int i_r) { + const int i_theta = start_black_radials + radial_task * 2; x[grid.index(i_r, i_theta)] = temp[grid.index(i_r, i_theta)]; - } - } + }); + Kokkos::fence(); } template @@ -121,11 +132,16 @@ void SmootherGive::solveWhiteRadialSection(Vector x, Vec int batch_stride = 2; radial_tridiagonal_solver_.solve(radial_section, batch_offset, batch_stride); -// Move updated values to x -#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++) { + // Move updated values to x + assert(grid.ntheta() % 2 == 0); + const int start_white_radials = 1; + const int num_white_radial_lines = grid.ntheta() / 2; + Kokkos::parallel_for( + "SmootherGive: moveUpdatedValues (White Radial)", + Kokkos::MDRangePolicy>({0, grid.numberSmootherCircles()}, {num_white_radial_lines, grid.nr()}), + KOKKOS_LAMBDA(const int radial_task, const int i_r) { + const int i_theta = start_white_radials + radial_task * 2; x[grid.index(i_r, i_theta)] = temp[grid.index(i_r, i_theta)]; - } - } + }); + Kokkos::fence(); } diff --git a/include/Smoother/SmootherTake/solveAscSystem.inl b/include/Smoother/SmootherTake/solveAscSystem.inl index 6d149b15..6fd08df5 100644 --- a/include/Smoother/SmootherTake/solveAscSystem.inl +++ b/include/Smoother/SmootherTake/solveAscSystem.inl @@ -22,13 +22,16 @@ void SmootherTake::solveBlackCircleSection(Vector x, Vec } // Move updated values to x - int start_black_circles = is_inner_circle_black ? 0 : 1; -#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++) { + const int start_black_circles = is_inner_circle_black ? 0 : 1; + const int num_black_circles = (grid.numberSmootherCircles() - start_black_circles + 1) / 2; + Kokkos::parallel_for( + "SmootherTake: moveUpdatedValues (Black Circular)", + Kokkos::MDRangePolicy>({0, 0}, {num_black_circles, grid.ntheta()}), + KOKKOS_LAMBDA(const int circle_task, const int i_theta) { + const int i_r = start_black_circles + circle_task * 2; x[grid.index(i_r, i_theta)] = temp[grid.index(i_r, i_theta)]; - } - } + }); + Kokkos::fence(); } template @@ -53,13 +56,16 @@ void SmootherTake::solveWhiteCircleSection(Vector x, Vec } // Move updated values to x - int start_white_circles = is_inner_circle_white ? 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++) { + const int start_white_circles = is_inner_circle_white ? 0 : 1; + const int num_white_circles = (grid.numberSmootherCircles() - start_white_circles + 1) / 2; + Kokkos::parallel_for( + "SmootherTake: moveUpdatedValues (White Circular)", + Kokkos::MDRangePolicy>({0, 0}, {num_white_circles, grid.ntheta()}), + KOKKOS_LAMBDA(const int circle_task, const int i_theta) { + const int i_r = start_white_circles + circle_task * 2; x[grid.index(i_r, i_theta)] = temp[grid.index(i_r, i_theta)]; - } - } + }); + Kokkos::fence(); } template @@ -76,13 +82,18 @@ void SmootherTake::solveBlackRadialSection(Vector x, Vec int batch_stride = 2; radial_tridiagonal_solver_.solve(radial_section, batch_offset, batch_stride); -// Move updated values to x -#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++) { + // Move updated values to x + assert(grid.ntheta() % 2 == 0); + const int start_black_radials = 0; + const int num_black_radial_lines = grid.ntheta() / 2; + Kokkos::parallel_for( + "SmootherTake: moveUpdatedValues (Black Radial)", + Kokkos::MDRangePolicy>({0, grid.numberSmootherCircles()}, {num_black_radial_lines, grid.nr()}), + KOKKOS_LAMBDA(const int radial_task, const int i_r) { + const int i_theta = start_black_radials + radial_task * 2; x[grid.index(i_r, i_theta)] = temp[grid.index(i_r, i_theta)]; - } - } + }); + Kokkos::fence(); } template @@ -99,11 +110,16 @@ void SmootherTake::solveWhiteRadialSection(Vector x, Vec int batch_stride = 2; radial_tridiagonal_solver_.solve(radial_section, batch_offset, batch_stride); -// Move updated values to x -#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++) { + // Move updated values to x + assert(grid.ntheta() % 2 == 0); + const int start_white_radials = 1; + const int num_white_radial_lines = grid.ntheta() / 2; + Kokkos::parallel_for( + "SmootherTake: moveUpdatedValues (White Radial)", + Kokkos::MDRangePolicy>({0, grid.numberSmootherCircles()}, {num_white_radial_lines, grid.nr()}), + KOKKOS_LAMBDA(const int radial_task, const int i_r) { + const int i_theta = start_white_radials + radial_task * 2; x[grid.index(i_r, i_theta)] = temp[grid.index(i_r, i_theta)]; - } - } + }); + Kokkos::fence(); } From 46dde6e43817e1d0d3a3b142f9a5ff49c84bf846 Mon Sep 17 00:00:00 2001 From: julianlitz Date: Sun, 3 May 2026 00:02:30 +0200 Subject: [PATCH 02/11] Avoid PolarGrid copyy --- .../solveAscSystem.inl | 24 ++++++++++------ .../solveAscSystem.inl | 28 +++++++++++++------ .../Smoother/SmootherGive/solveAscSystem.inl | 28 +++++++++++++------ .../Smoother/SmootherTake/solveAscSystem.inl | 28 +++++++++++++------ 4 files changed, 76 insertions(+), 32 deletions(-) diff --git a/include/ExtrapolatedSmoother/ExtrapolatedSmootherGive/solveAscSystem.inl b/include/ExtrapolatedSmoother/ExtrapolatedSmootherGive/solveAscSystem.inl index 7941c60e..6ccef597 100644 --- a/include/ExtrapolatedSmoother/ExtrapolatedSmootherGive/solveAscSystem.inl +++ b/include/ExtrapolatedSmoother/ExtrapolatedSmootherGive/solveAscSystem.inl @@ -44,8 +44,9 @@ void ExtrapolatedSmootherGive::solveBlackCircleSection(Vector>({0, 0}, {num_black_circles, grid.ntheta()}), KOKKOS_LAMBDA(const int circle_task, const int i_theta) { - const int i_r = start_black_circles + circle_task * 2; - x[grid.index(i_r, i_theta)] = temp[grid.index(i_r, i_theta)]; + const int i_r = start_black_circles + circle_task * 2; + const int index = grid_ptr->index(i_r, i_theta); + x[index] = temp[index]; }); Kokkos::fence(); } @@ -94,8 +95,9 @@ void ExtrapolatedSmootherGive::solveWhiteCircleSection(Vector>({0, 0}, {num_white_circles, grid.ntheta()}), KOKKOS_LAMBDA(const int circle_task, const int i_theta) { - const int i_r = start_white_circles + circle_task * 2; - x[grid.index(i_r, i_theta)] = temp[grid.index(i_r, i_theta)]; + const int i_r = start_white_circles + circle_task * 2; + const int index = grid_ptr->index(i_r, i_theta); + x[index] = temp[index]; }); Kokkos::fence(); } @@ -114,6 +116,8 @@ void ExtrapolatedSmootherGive::solveBlackRadialSection(Vector::solveBlackRadialSection(Vector>({0, grid.numberSmootherCircles()}, {num_black_radial_lines, grid.nr()}), KOKKOS_LAMBDA(const int radial_task, const int i_r) { - const int i_theta = start_black_radials + radial_task * 2; - x[grid.index(i_r, i_theta)] = temp[grid.index(i_r, i_theta)]; + const int i_theta = start_black_radials + radial_task * 2; + const int index = grid_ptr->index(i_r, i_theta); + x[index] = temp[index]; }); Kokkos::fence(); } @@ -142,6 +147,8 @@ void ExtrapolatedSmootherGive::solveWhiteRadialSection(Vector::solveWhiteRadialSection(Vector>({0, grid.numberSmootherCircles()}, {num_white_radial_lines, grid.nr()}), KOKKOS_LAMBDA(const int radial_task, const int i_r) { - const int i_theta = start_white_radials + radial_task * 2; - x[grid.index(i_r, i_theta)] = temp[grid.index(i_r, i_theta)]; + const int i_theta = start_white_radials + radial_task * 2; + const int index = grid_ptr->index(i_r, i_theta); + x[index] = temp[index]; }); Kokkos::fence(); } diff --git a/include/ExtrapolatedSmoother/ExtrapolatedSmootherTake/solveAscSystem.inl b/include/ExtrapolatedSmoother/ExtrapolatedSmootherTake/solveAscSystem.inl index 879f1645..bb1b43d1 100644 --- a/include/ExtrapolatedSmoother/ExtrapolatedSmootherTake/solveAscSystem.inl +++ b/include/ExtrapolatedSmoother/ExtrapolatedSmootherTake/solveAscSystem.inl @@ -26,6 +26,8 @@ void ExtrapolatedSmootherTake::solveBlackCircleSection(Vector::solveBlackCircleSection(Vector>({0, 0}, {num_black_circles, grid.ntheta()}), KOKKOS_LAMBDA(const int circle_task, const int i_theta) { - const int i_r = start_black_circles + circle_task * 2; - x[grid.index(i_r, i_theta)] = temp[grid.index(i_r, i_theta)]; + const int i_r = start_black_circles + circle_task * 2; + const int index = grid_ptr->index(i_r, i_theta); + x[index] = temp[index]; }); Kokkos::fence(); } @@ -65,6 +68,8 @@ void ExtrapolatedSmootherTake::solveWhiteCircleSection(Vector::solveWhiteCircleSection(Vector>({0, 0}, {num_white_circles, grid.ntheta()}), KOKKOS_LAMBDA(const int circle_task, const int i_theta) { - const int i_r = start_white_circles + circle_task * 2; - x[grid.index(i_r, i_theta)] = temp[grid.index(i_r, i_theta)]; + const int i_r = start_white_circles + circle_task * 2; + const int index = grid_ptr->index(i_r, i_theta); + x[index] = temp[index]; }); Kokkos::fence(); } @@ -92,6 +98,8 @@ void ExtrapolatedSmootherTake::solveBlackRadialSection(Vector::solveBlackRadialSection(Vector>({0, grid.numberSmootherCircles()}, {num_black_radial_lines, grid.nr()}), KOKKOS_LAMBDA(const int radial_task, const int i_r) { - const int i_theta = start_black_radials + radial_task * 2; - x[grid.index(i_r, i_theta)] = temp[grid.index(i_r, i_theta)]; + const int i_theta = start_black_radials + radial_task * 2; + const int index = grid_ptr->index(i_r, i_theta); + x[index] = temp[index]; }); Kokkos::fence(); } @@ -120,6 +129,8 @@ void ExtrapolatedSmootherTake::solveWhiteRadialSection(Vector::solveWhiteRadialSection(Vector>({0, grid.numberSmootherCircles()}, {num_white_radial_lines, grid.nr()}), KOKKOS_LAMBDA(const int radial_task, const int i_r) { - const int i_theta = start_white_radials + radial_task * 2; - x[grid.index(i_r, i_theta)] = temp[grid.index(i_r, i_theta)]; + const int i_theta = start_white_radials + radial_task * 2; + const int index = grid_ptr->index(i_r, i_theta); + x[index] = temp[index]; }); Kokkos::fence(); } diff --git a/include/Smoother/SmootherGive/solveAscSystem.inl b/include/Smoother/SmootherGive/solveAscSystem.inl index 8a183957..40b39478 100644 --- a/include/Smoother/SmootherGive/solveAscSystem.inl +++ b/include/Smoother/SmootherGive/solveAscSystem.inl @@ -32,6 +32,8 @@ void SmootherGive::solveBlackCircleSection(Vector x, Vec #endif } + const PolarGrid* grid_ptr = &grid; + // Move updated values to x const int start_black_circles = is_inner_circle_black ? 0 : 1; const int num_black_circles = (grid.numberSmootherCircles() - start_black_circles + 1) / 2; @@ -39,8 +41,9 @@ void SmootherGive::solveBlackCircleSection(Vector x, Vec "SmootherGive: moveUpdatedValues (Black Circular)", Kokkos::MDRangePolicy>({0, 0}, {num_black_circles, grid.ntheta()}), KOKKOS_LAMBDA(const int circle_task, const int i_theta) { - const int i_r = start_black_circles + circle_task * 2; - x[grid.index(i_r, i_theta)] = temp[grid.index(i_r, i_theta)]; + const int i_r = start_black_circles + circle_task * 2; + const int index = grid_ptr->index(i_r, i_theta); + x[index] = temp[index]; }); Kokkos::fence(); } @@ -77,6 +80,8 @@ void SmootherGive::solveWhiteCircleSection(Vector x, Vec #endif } + const PolarGrid* grid_ptr = &grid; + // Move updated values to x const int start_white_circles = is_inner_circle_white ? 0 : 1; const int num_white_circles = (grid.numberSmootherCircles() - start_white_circles + 1) / 2; @@ -84,8 +89,9 @@ void SmootherGive::solveWhiteCircleSection(Vector x, Vec "SmootherGive: moveUpdatedValues (White Circular)", Kokkos::MDRangePolicy>({0, 0}, {num_white_circles, grid.ntheta()}), KOKKOS_LAMBDA(const int circle_task, const int i_theta) { - const int i_r = start_white_circles + circle_task * 2; - x[grid.index(i_r, i_theta)] = temp[grid.index(i_r, i_theta)]; + const int i_r = start_white_circles + circle_task * 2; + const int index = grid_ptr->index(i_r, i_theta); + x[index] = temp[index]; }); Kokkos::fence(); } @@ -104,6 +110,8 @@ void SmootherGive::solveBlackRadialSection(Vector x, Vec int batch_stride = 2; radial_tridiagonal_solver_.solve(radial_section, batch_offset, batch_stride); + const PolarGrid* grid_ptr = &grid; + // Move updated values to x assert(grid.ntheta() % 2 == 0); const int start_black_radials = 0; @@ -112,8 +120,9 @@ void SmootherGive::solveBlackRadialSection(Vector x, Vec "SmootherGive: moveUpdatedValues (Black Radial)", Kokkos::MDRangePolicy>({0, grid.numberSmootherCircles()}, {num_black_radial_lines, grid.nr()}), KOKKOS_LAMBDA(const int radial_task, const int i_r) { - const int i_theta = start_black_radials + radial_task * 2; - x[grid.index(i_r, i_theta)] = temp[grid.index(i_r, i_theta)]; + const int i_theta = start_black_radials + radial_task * 2; + const int index = grid_ptr->index(i_r, i_theta); + x[index] = temp[index]; }); Kokkos::fence(); } @@ -132,6 +141,8 @@ void SmootherGive::solveWhiteRadialSection(Vector x, Vec int batch_stride = 2; radial_tridiagonal_solver_.solve(radial_section, batch_offset, batch_stride); + const PolarGrid* grid_ptr = &grid; + // Move updated values to x assert(grid.ntheta() % 2 == 0); const int start_white_radials = 1; @@ -140,8 +151,9 @@ void SmootherGive::solveWhiteRadialSection(Vector x, Vec "SmootherGive: moveUpdatedValues (White Radial)", Kokkos::MDRangePolicy>({0, grid.numberSmootherCircles()}, {num_white_radial_lines, grid.nr()}), KOKKOS_LAMBDA(const int radial_task, const int i_r) { - const int i_theta = start_white_radials + radial_task * 2; - x[grid.index(i_r, i_theta)] = temp[grid.index(i_r, i_theta)]; + const int i_theta = start_white_radials + radial_task * 2; + const int index = grid_ptr->index(i_r, i_theta); + x[index] = temp[index]; }); Kokkos::fence(); } diff --git a/include/Smoother/SmootherTake/solveAscSystem.inl b/include/Smoother/SmootherTake/solveAscSystem.inl index 6fd08df5..e639f286 100644 --- a/include/Smoother/SmootherTake/solveAscSystem.inl +++ b/include/Smoother/SmootherTake/solveAscSystem.inl @@ -21,6 +21,8 @@ void SmootherTake::solveBlackCircleSection(Vector x, Vec inner_boundary_solver_.solveInPlace(inner_boundary); } + const PolarGrid* grid_ptr = &grid; + // Move updated values to x const int start_black_circles = is_inner_circle_black ? 0 : 1; const int num_black_circles = (grid.numberSmootherCircles() - start_black_circles + 1) / 2; @@ -28,8 +30,9 @@ void SmootherTake::solveBlackCircleSection(Vector x, Vec "SmootherTake: moveUpdatedValues (Black Circular)", Kokkos::MDRangePolicy>({0, 0}, {num_black_circles, grid.ntheta()}), KOKKOS_LAMBDA(const int circle_task, const int i_theta) { - const int i_r = start_black_circles + circle_task * 2; - x[grid.index(i_r, i_theta)] = temp[grid.index(i_r, i_theta)]; + const int i_r = start_black_circles + circle_task * 2; + const int index = grid_ptr->index(i_r, i_theta); + x[index] = temp[index]; }); Kokkos::fence(); } @@ -55,6 +58,8 @@ void SmootherTake::solveWhiteCircleSection(Vector x, Vec inner_boundary_solver_.solveInPlace(inner_boundary); } + const PolarGrid* grid_ptr = &grid; + // Move updated values to x const int start_white_circles = is_inner_circle_white ? 0 : 1; const int num_white_circles = (grid.numberSmootherCircles() - start_white_circles + 1) / 2; @@ -62,8 +67,9 @@ void SmootherTake::solveWhiteCircleSection(Vector x, Vec "SmootherTake: moveUpdatedValues (White Circular)", Kokkos::MDRangePolicy>({0, 0}, {num_white_circles, grid.ntheta()}), KOKKOS_LAMBDA(const int circle_task, const int i_theta) { - const int i_r = start_white_circles + circle_task * 2; - x[grid.index(i_r, i_theta)] = temp[grid.index(i_r, i_theta)]; + const int i_r = start_white_circles + circle_task * 2; + const int index = grid_ptr->index(i_r, i_theta); + x[index] = temp[index]; }); Kokkos::fence(); } @@ -82,6 +88,8 @@ void SmootherTake::solveBlackRadialSection(Vector x, Vec int batch_stride = 2; radial_tridiagonal_solver_.solve(radial_section, batch_offset, batch_stride); + const PolarGrid* grid_ptr = &grid; + // Move updated values to x assert(grid.ntheta() % 2 == 0); const int start_black_radials = 0; @@ -90,8 +98,9 @@ void SmootherTake::solveBlackRadialSection(Vector x, Vec "SmootherTake: moveUpdatedValues (Black Radial)", Kokkos::MDRangePolicy>({0, grid.numberSmootherCircles()}, {num_black_radial_lines, grid.nr()}), KOKKOS_LAMBDA(const int radial_task, const int i_r) { - const int i_theta = start_black_radials + radial_task * 2; - x[grid.index(i_r, i_theta)] = temp[grid.index(i_r, i_theta)]; + const int i_theta = start_black_radials + radial_task * 2; + const int index = grid_ptr->index(i_r, i_theta); + x[index] = temp[index]; }); Kokkos::fence(); } @@ -110,6 +119,8 @@ void SmootherTake::solveWhiteRadialSection(Vector x, Vec int batch_stride = 2; radial_tridiagonal_solver_.solve(radial_section, batch_offset, batch_stride); + const PolarGrid* grid_ptr = &grid; + // Move updated values to x assert(grid.ntheta() % 2 == 0); const int start_white_radials = 1; @@ -118,8 +129,9 @@ void SmootherTake::solveWhiteRadialSection(Vector x, Vec "SmootherTake: moveUpdatedValues (White Radial)", Kokkos::MDRangePolicy>({0, grid.numberSmootherCircles()}, {num_white_radial_lines, grid.nr()}), KOKKOS_LAMBDA(const int radial_task, const int i_r) { - const int i_theta = start_white_radials + radial_task * 2; - x[grid.index(i_r, i_theta)] = temp[grid.index(i_r, i_theta)]; + const int i_theta = start_white_radials + radial_task * 2; + const int index = grid_ptr->index(i_r, i_theta); + x[index] = temp[index]; }); Kokkos::fence(); } From 382bab7da0326410ef7c840c761da99982e30ddc Mon Sep 17 00:00:00 2001 From: julianlitz Date: Sun, 3 May 2026 00:03:31 +0200 Subject: [PATCH 03/11] Fix --- .../ExtrapolatedSmootherGive/solveAscSystem.inl | 4 ++++ 1 file changed, 4 insertions(+) diff --git a/include/ExtrapolatedSmoother/ExtrapolatedSmootherGive/solveAscSystem.inl b/include/ExtrapolatedSmoother/ExtrapolatedSmootherGive/solveAscSystem.inl index 6ccef597..762c1144 100644 --- a/include/ExtrapolatedSmoother/ExtrapolatedSmootherGive/solveAscSystem.inl +++ b/include/ExtrapolatedSmoother/ExtrapolatedSmootherGive/solveAscSystem.inl @@ -37,6 +37,8 @@ void ExtrapolatedSmootherGive::solveBlackCircleSection(Vector::solveWhiteCircleSection(Vector Date: Wed, 6 May 2026 21:33:02 +0200 Subject: [PATCH 04/11] Update solveAscSystem.inl --- include/Smoother/SmootherTake/solveAscSystem.inl | 16 ++++------------ 1 file changed, 4 insertions(+), 12 deletions(-) diff --git a/include/Smoother/SmootherTake/solveAscSystem.inl b/include/Smoother/SmootherTake/solveAscSystem.inl index e639f286..aae9b3c0 100644 --- a/include/Smoother/SmootherTake/solveAscSystem.inl +++ b/include/Smoother/SmootherTake/solveAscSystem.inl @@ -21,8 +21,6 @@ void SmootherTake::solveBlackCircleSection(Vector x, Vec inner_boundary_solver_.solveInPlace(inner_boundary); } - const PolarGrid* grid_ptr = &grid; - // Move updated values to x const int start_black_circles = is_inner_circle_black ? 0 : 1; const int num_black_circles = (grid.numberSmootherCircles() - start_black_circles + 1) / 2; @@ -31,7 +29,7 @@ void SmootherTake::solveBlackCircleSection(Vector x, Vec Kokkos::MDRangePolicy>({0, 0}, {num_black_circles, grid.ntheta()}), KOKKOS_LAMBDA(const int circle_task, const int i_theta) { const int i_r = start_black_circles + circle_task * 2; - const int index = grid_ptr->index(i_r, i_theta); + const int index = grid.index(i_r, i_theta); x[index] = temp[index]; }); Kokkos::fence(); @@ -58,8 +56,6 @@ void SmootherTake::solveWhiteCircleSection(Vector x, Vec inner_boundary_solver_.solveInPlace(inner_boundary); } - const PolarGrid* grid_ptr = &grid; - // Move updated values to x const int start_white_circles = is_inner_circle_white ? 0 : 1; const int num_white_circles = (grid.numberSmootherCircles() - start_white_circles + 1) / 2; @@ -68,7 +64,7 @@ void SmootherTake::solveWhiteCircleSection(Vector x, Vec Kokkos::MDRangePolicy>({0, 0}, {num_white_circles, grid.ntheta()}), KOKKOS_LAMBDA(const int circle_task, const int i_theta) { const int i_r = start_white_circles + circle_task * 2; - const int index = grid_ptr->index(i_r, i_theta); + const int index = grid.index(i_r, i_theta); x[index] = temp[index]; }); Kokkos::fence(); @@ -88,8 +84,6 @@ void SmootherTake::solveBlackRadialSection(Vector x, Vec int batch_stride = 2; radial_tridiagonal_solver_.solve(radial_section, batch_offset, batch_stride); - const PolarGrid* grid_ptr = &grid; - // Move updated values to x assert(grid.ntheta() % 2 == 0); const int start_black_radials = 0; @@ -99,7 +93,7 @@ void SmootherTake::solveBlackRadialSection(Vector x, Vec Kokkos::MDRangePolicy>({0, grid.numberSmootherCircles()}, {num_black_radial_lines, grid.nr()}), KOKKOS_LAMBDA(const int radial_task, const int i_r) { const int i_theta = start_black_radials + radial_task * 2; - const int index = grid_ptr->index(i_r, i_theta); + const int index = grid.index(i_r, i_theta); x[index] = temp[index]; }); Kokkos::fence(); @@ -119,8 +113,6 @@ void SmootherTake::solveWhiteRadialSection(Vector x, Vec int batch_stride = 2; radial_tridiagonal_solver_.solve(radial_section, batch_offset, batch_stride); - const PolarGrid* grid_ptr = &grid; - // Move updated values to x assert(grid.ntheta() % 2 == 0); const int start_white_radials = 1; @@ -130,7 +122,7 @@ void SmootherTake::solveWhiteRadialSection(Vector x, Vec Kokkos::MDRangePolicy>({0, grid.numberSmootherCircles()}, {num_white_radial_lines, grid.nr()}), KOKKOS_LAMBDA(const int radial_task, const int i_r) { const int i_theta = start_white_radials + radial_task * 2; - const int index = grid_ptr->index(i_r, i_theta); + const int index = grid.index(i_r, i_theta); x[index] = temp[index]; }); Kokkos::fence(); From 47b8f50d797ae176014a87e9ec80bc3f80516bf7 Mon Sep 17 00:00:00 2001 From: Julian Litz <91479202+julianlitz@users.noreply.github.com> Date: Wed, 6 May 2026 21:34:06 +0200 Subject: [PATCH 05/11] Update solveAscSystem.inl --- include/Smoother/SmootherGive/solveAscSystem.inl | 16 ++++------------ 1 file changed, 4 insertions(+), 12 deletions(-) diff --git a/include/Smoother/SmootherGive/solveAscSystem.inl b/include/Smoother/SmootherGive/solveAscSystem.inl index 40b39478..391b03ed 100644 --- a/include/Smoother/SmootherGive/solveAscSystem.inl +++ b/include/Smoother/SmootherGive/solveAscSystem.inl @@ -32,8 +32,6 @@ void SmootherGive::solveBlackCircleSection(Vector x, Vec #endif } - const PolarGrid* grid_ptr = &grid; - // Move updated values to x const int start_black_circles = is_inner_circle_black ? 0 : 1; const int num_black_circles = (grid.numberSmootherCircles() - start_black_circles + 1) / 2; @@ -42,7 +40,7 @@ void SmootherGive::solveBlackCircleSection(Vector x, Vec Kokkos::MDRangePolicy>({0, 0}, {num_black_circles, grid.ntheta()}), KOKKOS_LAMBDA(const int circle_task, const int i_theta) { const int i_r = start_black_circles + circle_task * 2; - const int index = grid_ptr->index(i_r, i_theta); + const int index = grid.index(i_r, i_theta); x[index] = temp[index]; }); Kokkos::fence(); @@ -80,8 +78,6 @@ void SmootherGive::solveWhiteCircleSection(Vector x, Vec #endif } - const PolarGrid* grid_ptr = &grid; - // Move updated values to x const int start_white_circles = is_inner_circle_white ? 0 : 1; const int num_white_circles = (grid.numberSmootherCircles() - start_white_circles + 1) / 2; @@ -90,7 +86,7 @@ void SmootherGive::solveWhiteCircleSection(Vector x, Vec Kokkos::MDRangePolicy>({0, 0}, {num_white_circles, grid.ntheta()}), KOKKOS_LAMBDA(const int circle_task, const int i_theta) { const int i_r = start_white_circles + circle_task * 2; - const int index = grid_ptr->index(i_r, i_theta); + const int index = grid.index(i_r, i_theta); x[index] = temp[index]; }); Kokkos::fence(); @@ -110,8 +106,6 @@ void SmootherGive::solveBlackRadialSection(Vector x, Vec int batch_stride = 2; radial_tridiagonal_solver_.solve(radial_section, batch_offset, batch_stride); - const PolarGrid* grid_ptr = &grid; - // Move updated values to x assert(grid.ntheta() % 2 == 0); const int start_black_radials = 0; @@ -121,7 +115,7 @@ void SmootherGive::solveBlackRadialSection(Vector x, Vec Kokkos::MDRangePolicy>({0, grid.numberSmootherCircles()}, {num_black_radial_lines, grid.nr()}), KOKKOS_LAMBDA(const int radial_task, const int i_r) { const int i_theta = start_black_radials + radial_task * 2; - const int index = grid_ptr->index(i_r, i_theta); + const int index = grid.index(i_r, i_theta); x[index] = temp[index]; }); Kokkos::fence(); @@ -141,8 +135,6 @@ void SmootherGive::solveWhiteRadialSection(Vector x, Vec int batch_stride = 2; radial_tridiagonal_solver_.solve(radial_section, batch_offset, batch_stride); - const PolarGrid* grid_ptr = &grid; - // Move updated values to x assert(grid.ntheta() % 2 == 0); const int start_white_radials = 1; @@ -152,7 +144,7 @@ void SmootherGive::solveWhiteRadialSection(Vector x, Vec Kokkos::MDRangePolicy>({0, grid.numberSmootherCircles()}, {num_white_radial_lines, grid.nr()}), KOKKOS_LAMBDA(const int radial_task, const int i_r) { const int i_theta = start_white_radials + radial_task * 2; - const int index = grid_ptr->index(i_r, i_theta); + const int index = grid.index(i_r, i_theta); x[index] = temp[index]; }); Kokkos::fence(); From 0c004cea16ef6ae0662f40160cf2f09c56518330 Mon Sep 17 00:00:00 2001 From: Julian Litz <91479202+julianlitz@users.noreply.github.com> Date: Wed, 6 May 2026 21:43:47 +0200 Subject: [PATCH 06/11] Update solveAscSystem.inl --- .../ExtrapolatedSmootherTake/solveAscSystem.inl | 16 ++++------------ 1 file changed, 4 insertions(+), 12 deletions(-) diff --git a/include/ExtrapolatedSmoother/ExtrapolatedSmootherTake/solveAscSystem.inl b/include/ExtrapolatedSmoother/ExtrapolatedSmootherTake/solveAscSystem.inl index bb1b43d1..28465b39 100644 --- a/include/ExtrapolatedSmoother/ExtrapolatedSmootherTake/solveAscSystem.inl +++ b/include/ExtrapolatedSmoother/ExtrapolatedSmootherTake/solveAscSystem.inl @@ -26,8 +26,6 @@ void ExtrapolatedSmootherTake::solveBlackCircleSection(Vector::solveBlackCircleSection(Vector>({0, 0}, {num_black_circles, grid.ntheta()}), KOKKOS_LAMBDA(const int circle_task, const int i_theta) { const int i_r = start_black_circles + circle_task * 2; - const int index = grid_ptr->index(i_r, i_theta); + const int index = grid.index(i_r, i_theta); x[index] = temp[index]; }); Kokkos::fence(); @@ -68,8 +66,6 @@ void ExtrapolatedSmootherTake::solveWhiteCircleSection(Vector::solveWhiteCircleSection(Vector>({0, 0}, {num_white_circles, grid.ntheta()}), KOKKOS_LAMBDA(const int circle_task, const int i_theta) { const int i_r = start_white_circles + circle_task * 2; - const int index = grid_ptr->index(i_r, i_theta); + const int index = grid.index(i_r, i_theta); x[index] = temp[index]; }); Kokkos::fence(); @@ -98,8 +94,6 @@ void ExtrapolatedSmootherTake::solveBlackRadialSection(Vector::solveBlackRadialSection(Vector>({0, grid.numberSmootherCircles()}, {num_black_radial_lines, grid.nr()}), KOKKOS_LAMBDA(const int radial_task, const int i_r) { const int i_theta = start_black_radials + radial_task * 2; - const int index = grid_ptr->index(i_r, i_theta); + const int index = grid.index(i_r, i_theta); x[index] = temp[index]; }); Kokkos::fence(); @@ -129,8 +123,6 @@ void ExtrapolatedSmootherTake::solveWhiteRadialSection(Vector::solveWhiteRadialSection(Vector>({0, grid.numberSmootherCircles()}, {num_white_radial_lines, grid.nr()}), KOKKOS_LAMBDA(const int radial_task, const int i_r) { const int i_theta = start_white_radials + radial_task * 2; - const int index = grid_ptr->index(i_r, i_theta); + const int index = grid.index(i_r, i_theta); x[index] = temp[index]; }); Kokkos::fence(); From 14c88983d2363781f14af47dd398dfc7cc2437a8 Mon Sep 17 00:00:00 2001 From: Julian Litz <91479202+julianlitz@users.noreply.github.com> Date: Wed, 6 May 2026 21:44:52 +0200 Subject: [PATCH 07/11] Update solveAscSystem.inl --- .../ExtrapolatedSmootherGive/solveAscSystem.inl | 16 ++++------------ 1 file changed, 4 insertions(+), 12 deletions(-) diff --git a/include/ExtrapolatedSmoother/ExtrapolatedSmootherGive/solveAscSystem.inl b/include/ExtrapolatedSmoother/ExtrapolatedSmootherGive/solveAscSystem.inl index 762c1144..2c9b0fc5 100644 --- a/include/ExtrapolatedSmoother/ExtrapolatedSmootherGive/solveAscSystem.inl +++ b/include/ExtrapolatedSmoother/ExtrapolatedSmootherGive/solveAscSystem.inl @@ -37,8 +37,6 @@ void ExtrapolatedSmootherGive::solveBlackCircleSection(Vector::solveBlackCircleSection(Vector>({0, 0}, {num_black_circles, grid.ntheta()}), KOKKOS_LAMBDA(const int circle_task, const int i_theta) { const int i_r = start_black_circles + circle_task * 2; - const int index = grid_ptr->index(i_r, i_theta); + const int index = grid.index(i_r, i_theta); x[index] = temp[index]; }); Kokkos::fence(); @@ -90,8 +88,6 @@ void ExtrapolatedSmootherGive::solveWhiteCircleSection(Vector::solveWhiteCircleSection(Vector>({0, 0}, {num_white_circles, grid.ntheta()}), KOKKOS_LAMBDA(const int circle_task, const int i_theta) { const int i_r = start_white_circles + circle_task * 2; - const int index = grid_ptr->index(i_r, i_theta); + const int index = grid.index(i_r, i_theta); x[index] = temp[index]; }); Kokkos::fence(); @@ -120,8 +116,6 @@ void ExtrapolatedSmootherGive::solveBlackRadialSection(Vector::solveBlackRadialSection(Vector>({0, grid.numberSmootherCircles()}, {num_black_radial_lines, grid.nr()}), KOKKOS_LAMBDA(const int radial_task, const int i_r) { const int i_theta = start_black_radials + radial_task * 2; - const int index = grid_ptr->index(i_r, i_theta); + const int index = grid.index(i_r, i_theta); x[index] = temp[index]; }); Kokkos::fence(); @@ -151,8 +145,6 @@ void ExtrapolatedSmootherGive::solveWhiteRadialSection(Vector::solveWhiteRadialSection(Vector>({0, grid.numberSmootherCircles()}, {num_white_radial_lines, grid.nr()}), KOKKOS_LAMBDA(const int radial_task, const int i_r) { const int i_theta = start_white_radials + radial_task * 2; - const int index = grid_ptr->index(i_r, i_theta); + const int index = grid.index(i_r, i_theta); x[index] = temp[index]; }); Kokkos::fence(); From f82f5812168d7620fc8ea15520c365550fa2c60c Mon Sep 17 00:00:00 2001 From: BOURNE Emily EPFL Date: Tue, 12 May 2026 13:18:28 +0200 Subject: [PATCH 08/11] Functions containing lambdas must be public --- .../ExtrapolatedSmootherGive/extrapolatedSmootherGive.h | 2 ++ .../ExtrapolatedSmootherTake/extrapolatedSmootherTake.h | 2 ++ include/Smoother/SmootherGive/smootherGive.h | 2 ++ include/Smoother/SmootherTake/smootherTake.h | 3 +-- 4 files changed, 7 insertions(+), 2 deletions(-) diff --git a/include/ExtrapolatedSmoother/ExtrapolatedSmootherGive/extrapolatedSmootherGive.h b/include/ExtrapolatedSmoother/ExtrapolatedSmootherGive/extrapolatedSmootherGive.h index 281a0242..d5ab54e2 100644 --- a/include/ExtrapolatedSmoother/ExtrapolatedSmootherGive/extrapolatedSmootherGive.h +++ b/include/ExtrapolatedSmoother/ExtrapolatedSmootherGive/extrapolatedSmootherGive.h @@ -187,6 +187,8 @@ class ExtrapolatedSmootherGive : public ExtrapolatedSmoother /* Line-wise solvers */ /* ----------------- */ + // Functions must be public due to cuda restriction +public: // Solve the linear system: // A_sc * u_sc = f_sc − A_sc^ortho * u_sc^ortho // Parameter mapping: diff --git a/include/ExtrapolatedSmoother/ExtrapolatedSmootherTake/extrapolatedSmootherTake.h b/include/ExtrapolatedSmoother/ExtrapolatedSmootherTake/extrapolatedSmootherTake.h index 8f5e78fe..7a386433 100644 --- a/include/ExtrapolatedSmoother/ExtrapolatedSmootherTake/extrapolatedSmootherTake.h +++ b/include/ExtrapolatedSmoother/ExtrapolatedSmootherTake/extrapolatedSmootherTake.h @@ -183,6 +183,8 @@ class ExtrapolatedSmootherTake : public ExtrapolatedSmoother /* Line-wise solvers */ /* ----------------- */ + // Functions must be public due to cuda restriction +public: // Solve the linear system: // A_sc * u_sc = f_sc − A_sc^ortho * u_sc^ortho // Parameter mapping: diff --git a/include/Smoother/SmootherGive/smootherGive.h b/include/Smoother/SmootherGive/smootherGive.h index d5e12ddd..1a7cc82a 100644 --- a/include/Smoother/SmootherGive/smootherGive.h +++ b/include/Smoother/SmootherGive/smootherGive.h @@ -180,6 +180,8 @@ class SmootherGive : public Smoother /* Line-wise solvers */ /* ----------------- */ + // Functions must be public due to cuda restriction +public: // Solve the linear system: // A_sc * u_sc = f_sc − A_sc^ortho * u_sc^ortho // Parameter mapping: diff --git a/include/Smoother/SmootherTake/smootherTake.h b/include/Smoother/SmootherTake/smootherTake.h index 63967288..e5f02849 100644 --- a/include/Smoother/SmootherTake/smootherTake.h +++ b/include/Smoother/SmootherTake/smootherTake.h @@ -165,7 +165,7 @@ class SmootherTake : public Smoother /* Orthogonal application */ /* ---------------------- */ - // Public is required as Cuda needs to be abe to get the address of functions enclosing lambda functions + // Public is required as Cuda needs to be able to get the address of functions enclosing lambda functions public: // Compute temp = f_sc − A_sc^ortho * u_sc^ortho (precomputed right-hand side) // where x = u_sc and rhs = f_sc @@ -174,7 +174,6 @@ class SmootherTake : public Smoother void applyAscOrthoBlackRadialSection(ConstVector x, ConstVector rhs, Vector temp); void applyAscOrthoWhiteRadialSection(ConstVector x, ConstVector rhs, Vector temp); -private: /* ----------------- */ /* Line-wise solvers */ /* ----------------- */ From 66774c22fed5896f21eb2b783136598115425010 Mon Sep 17 00:00:00 2001 From: Emily Bourne Date: Tue, 12 May 2026 13:26:23 +0200 Subject: [PATCH 09/11] Clang formatting --- include/DirectSolver/DirectSolverGive/buildSolverMatrix.inl | 4 ++-- include/DirectSolver/DirectSolverTake/buildSolverMatrix.inl | 4 ++-- .../ExtrapolatedSmootherGive/buildInnerBoundaryAsc.inl | 4 ++-- .../ExtrapolatedSmootherGive/extrapolatedSmootherGive.h | 2 +- .../ExtrapolatedSmootherTake/buildInnerBoundaryAsc.inl | 4 ++-- .../ExtrapolatedSmootherTake/extrapolatedSmootherTake.h | 2 +- include/Smoother/SmootherGive/buildInnerBoundaryAsc.inl | 4 ++-- include/Smoother/SmootherGive/smootherGive.h | 2 +- include/Smoother/SmootherGive/solveAscSystem.inl | 3 ++- include/Smoother/SmootherTake/buildInnerBoundaryAsc.inl | 6 +++--- 10 files changed, 18 insertions(+), 17 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/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/ExtrapolatedSmootherGive/extrapolatedSmootherGive.h b/include/ExtrapolatedSmoother/ExtrapolatedSmootherGive/extrapolatedSmootherGive.h index d5ab54e2..32f448d0 100644 --- a/include/ExtrapolatedSmoother/ExtrapolatedSmootherGive/extrapolatedSmootherGive.h +++ b/include/ExtrapolatedSmoother/ExtrapolatedSmootherGive/extrapolatedSmootherGive.h @@ -187,7 +187,7 @@ class ExtrapolatedSmootherGive : public ExtrapolatedSmoother /* Line-wise solvers */ /* ----------------- */ - // Functions must be public due to cuda restriction + // Functions must be public due to cuda restriction public: // Solve the linear system: // A_sc * u_sc = f_sc − A_sc^ortho * u_sc^ortho 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 7a386433..841d8203 100644 --- a/include/ExtrapolatedSmoother/ExtrapolatedSmootherTake/extrapolatedSmootherTake.h +++ b/include/ExtrapolatedSmoother/ExtrapolatedSmootherTake/extrapolatedSmootherTake.h @@ -183,7 +183,7 @@ class ExtrapolatedSmootherTake : public ExtrapolatedSmoother /* Line-wise solvers */ /* ----------------- */ - // Functions must be public due to cuda restriction + // Functions must be public due to cuda restriction public: // Solve the linear system: // A_sc * u_sc = f_sc − A_sc^ortho * u_sc^ortho 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/SmootherGive/smootherGive.h b/include/Smoother/SmootherGive/smootherGive.h index 1a7cc82a..c2988286 100644 --- a/include/Smoother/SmootherGive/smootherGive.h +++ b/include/Smoother/SmootherGive/smootherGive.h @@ -180,7 +180,7 @@ class SmootherGive : public Smoother /* Line-wise solvers */ /* ----------------- */ - // Functions must be public due to cuda restriction + // Functions must be public due to cuda restriction public: // Solve the linear system: // A_sc * u_sc = f_sc − A_sc^ortho * u_sc^ortho diff --git a/include/Smoother/SmootherGive/solveAscSystem.inl b/include/Smoother/SmootherGive/solveAscSystem.inl index 18f4eaf5..b3de3feb 100644 --- a/include/Smoother/SmootherGive/solveAscSystem.inl +++ b/include/Smoother/SmootherGive/solveAscSystem.inl @@ -119,7 +119,8 @@ void SmootherGive::solveWhiteRadialSection(Vector x, Vec const int num_white_radial_lines = grid.ntheta() / 2; Kokkos::parallel_for( "SmootherGive: moveUpdatedValues (White Radial)", - Kokkos::MDRangePolicy>({0, grid.numberSmootherCircles()}, {num_white_radial_lines, grid.nr()}), + Kokkos::MDRangePolicy>({0, grid.numberSmootherCircles()}, + {num_white_radial_lines, grid.nr()}), KOKKOS_LAMBDA(const int radial_task, const int i_r) { const int i_theta = start_white_radials + radial_task * 2; const int index = grid.index(i_r, i_theta); 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. From efb6da934bbbb11de47dcff9fe25f8f8aa96a5db Mon Sep 17 00:00:00 2001 From: BOURNE Emily EPFL Date: Tue, 12 May 2026 13:39:43 +0200 Subject: [PATCH 10/11] Run on Host for now --- .../ExtrapolatedSmootherGive/solveAscSystem.inl | 8 ++++---- .../ExtrapolatedSmootherTake/solveAscSystem.inl | 8 ++++---- include/Smoother/SmootherGive/solveAscSystem.inl | 6 +++--- include/Smoother/SmootherTake/solveAscSystem.inl | 8 ++++---- 4 files changed, 15 insertions(+), 15 deletions(-) diff --git a/include/ExtrapolatedSmoother/ExtrapolatedSmootherGive/solveAscSystem.inl b/include/ExtrapolatedSmoother/ExtrapolatedSmootherGive/solveAscSystem.inl index c6dad69c..624c11e1 100644 --- a/include/ExtrapolatedSmoother/ExtrapolatedSmootherGive/solveAscSystem.inl +++ b/include/ExtrapolatedSmoother/ExtrapolatedSmootherGive/solveAscSystem.inl @@ -31,7 +31,7 @@ void ExtrapolatedSmootherGive::solveBlackCircleSection(Vector>({0, 0}, {num_black_circles, grid.ntheta()}), + Kokkos::MDRangePolicy>({0, 0}, {num_black_circles, grid.ntheta()}), KOKKOS_LAMBDA(const int circle_task, const int i_theta) { const int i_r = start_black_circles + circle_task * 2; const int index = grid.index(i_r, i_theta); @@ -71,7 +71,7 @@ void ExtrapolatedSmootherGive::solveWhiteCircleSection(Vector>({0, 0}, {num_white_circles, grid.ntheta()}), + Kokkos::MDRangePolicy>({0, 0}, {num_white_circles, grid.ntheta()}), KOKKOS_LAMBDA(const int circle_task, const int i_theta) { const int i_r = start_white_circles + circle_task * 2; const int index = grid.index(i_r, i_theta); @@ -100,7 +100,7 @@ void ExtrapolatedSmootherGive::solveBlackRadialSection(Vector>({0, grid.numberSmootherCircles()}, {num_black_radial_lines, grid.nr()}), + Kokkos::MDRangePolicy>({0, grid.numberSmootherCircles()}, {num_black_radial_lines, grid.nr()}), KOKKOS_LAMBDA(const int radial_task, const int i_r) { const int i_theta = start_black_radials + radial_task * 2; const int index = grid.index(i_r, i_theta); @@ -129,7 +129,7 @@ void ExtrapolatedSmootherGive::solveWhiteRadialSection(Vector>({0, grid.numberSmootherCircles()}, {num_white_radial_lines, grid.nr()}), + Kokkos::MDRangePolicy>({0, grid.numberSmootherCircles()}, {num_white_radial_lines, grid.nr()}), KOKKOS_LAMBDA(const int radial_task, const int i_r) { const int i_theta = start_white_radials + radial_task * 2; const int index = grid.index(i_r, i_theta); diff --git a/include/ExtrapolatedSmoother/ExtrapolatedSmootherTake/solveAscSystem.inl b/include/ExtrapolatedSmoother/ExtrapolatedSmootherTake/solveAscSystem.inl index 28465b39..5d73a2da 100644 --- a/include/ExtrapolatedSmoother/ExtrapolatedSmootherTake/solveAscSystem.inl +++ b/include/ExtrapolatedSmoother/ExtrapolatedSmootherTake/solveAscSystem.inl @@ -31,7 +31,7 @@ void ExtrapolatedSmootherTake::solveBlackCircleSection(Vector>({0, 0}, {num_black_circles, grid.ntheta()}), + Kokkos::MDRangePolicy>({0, 0}, {num_black_circles, grid.ntheta()}), KOKKOS_LAMBDA(const int circle_task, const int i_theta) { const int i_r = start_black_circles + circle_task * 2; const int index = grid.index(i_r, i_theta); @@ -71,7 +71,7 @@ void ExtrapolatedSmootherTake::solveWhiteCircleSection(Vector>({0, 0}, {num_white_circles, grid.ntheta()}), + Kokkos::MDRangePolicy>({0, 0}, {num_white_circles, grid.ntheta()}), KOKKOS_LAMBDA(const int circle_task, const int i_theta) { const int i_r = start_white_circles + circle_task * 2; const int index = grid.index(i_r, i_theta); @@ -100,7 +100,7 @@ void ExtrapolatedSmootherTake::solveBlackRadialSection(Vector>({0, grid.numberSmootherCircles()}, {num_black_radial_lines, grid.nr()}), + Kokkos::MDRangePolicy>({0, grid.numberSmootherCircles()}, {num_black_radial_lines, grid.nr()}), KOKKOS_LAMBDA(const int radial_task, const int i_r) { const int i_theta = start_black_radials + radial_task * 2; const int index = grid.index(i_r, i_theta); @@ -129,7 +129,7 @@ void ExtrapolatedSmootherTake::solveWhiteRadialSection(Vector>({0, grid.numberSmootherCircles()}, {num_white_radial_lines, grid.nr()}), + Kokkos::MDRangePolicy>({0, grid.numberSmootherCircles()}, {num_white_radial_lines, grid.nr()}), KOKKOS_LAMBDA(const int radial_task, const int i_r) { const int i_theta = start_white_radials + radial_task * 2; const int index = grid.index(i_r, i_theta); diff --git a/include/Smoother/SmootherGive/solveAscSystem.inl b/include/Smoother/SmootherGive/solveAscSystem.inl index b3de3feb..b9fad306 100644 --- a/include/Smoother/SmootherGive/solveAscSystem.inl +++ b/include/Smoother/SmootherGive/solveAscSystem.inl @@ -26,7 +26,7 @@ void SmootherGive::solveBlackCircleSection(Vector x, Vec const int num_black_circles = (grid.numberSmootherCircles() - start_black_circles + 1) / 2; Kokkos::parallel_for( "SmootherGive: moveUpdatedValues (Black Circular)", - Kokkos::MDRangePolicy>({0, 0}, {num_black_circles, grid.ntheta()}), + Kokkos::MDRangePolicy>({0, 0}, {num_black_circles, grid.ntheta()}), KOKKOS_LAMBDA(const int circle_task, const int i_theta) { const int i_r = start_black_circles + circle_task * 2; const int index = grid.index(i_r, i_theta); @@ -61,7 +61,7 @@ void SmootherGive::solveWhiteCircleSection(Vector x, Vec const int num_white_circles = (grid.numberSmootherCircles() - start_white_circles + 1) / 2; Kokkos::parallel_for( "SmootherGive: moveUpdatedValues (White Circular)", - Kokkos::MDRangePolicy>({0, 0}, {num_white_circles, grid.ntheta()}), + Kokkos::MDRangePolicy>({0, 0}, {num_white_circles, grid.ntheta()}), KOKKOS_LAMBDA(const int circle_task, const int i_theta) { const int i_r = start_white_circles + circle_task * 2; const int index = grid.index(i_r, i_theta); @@ -90,7 +90,7 @@ void SmootherGive::solveBlackRadialSection(Vector x, Vec const int num_black_radial_lines = grid.ntheta() / 2; Kokkos::parallel_for( "SmootherGive: moveUpdatedValues (Black Radial)", - Kokkos::MDRangePolicy>({0, grid.numberSmootherCircles()}, {num_black_radial_lines, grid.nr()}), + Kokkos::MDRangePolicy>({0, grid.numberSmootherCircles()}, {num_black_radial_lines, grid.nr()}), KOKKOS_LAMBDA(const int radial_task, const int i_r) { const int i_theta = start_black_radials + radial_task * 2; const int index = grid.index(i_r, i_theta); diff --git a/include/Smoother/SmootherTake/solveAscSystem.inl b/include/Smoother/SmootherTake/solveAscSystem.inl index aae9b3c0..197b3a52 100644 --- a/include/Smoother/SmootherTake/solveAscSystem.inl +++ b/include/Smoother/SmootherTake/solveAscSystem.inl @@ -26,7 +26,7 @@ void SmootherTake::solveBlackCircleSection(Vector x, Vec const int num_black_circles = (grid.numberSmootherCircles() - start_black_circles + 1) / 2; Kokkos::parallel_for( "SmootherTake: moveUpdatedValues (Black Circular)", - Kokkos::MDRangePolicy>({0, 0}, {num_black_circles, grid.ntheta()}), + Kokkos::MDRangePolicy>({0, 0}, {num_black_circles, grid.ntheta()}), KOKKOS_LAMBDA(const int circle_task, const int i_theta) { const int i_r = start_black_circles + circle_task * 2; const int index = grid.index(i_r, i_theta); @@ -61,7 +61,7 @@ void SmootherTake::solveWhiteCircleSection(Vector x, Vec const int num_white_circles = (grid.numberSmootherCircles() - start_white_circles + 1) / 2; Kokkos::parallel_for( "SmootherTake: moveUpdatedValues (White Circular)", - Kokkos::MDRangePolicy>({0, 0}, {num_white_circles, grid.ntheta()}), + Kokkos::MDRangePolicy>({0, 0}, {num_white_circles, grid.ntheta()}), KOKKOS_LAMBDA(const int circle_task, const int i_theta) { const int i_r = start_white_circles + circle_task * 2; const int index = grid.index(i_r, i_theta); @@ -90,7 +90,7 @@ void SmootherTake::solveBlackRadialSection(Vector x, Vec const int num_black_radial_lines = grid.ntheta() / 2; Kokkos::parallel_for( "SmootherTake: moveUpdatedValues (Black Radial)", - Kokkos::MDRangePolicy>({0, grid.numberSmootherCircles()}, {num_black_radial_lines, grid.nr()}), + Kokkos::MDRangePolicy>({0, grid.numberSmootherCircles()}, {num_black_radial_lines, grid.nr()}), KOKKOS_LAMBDA(const int radial_task, const int i_r) { const int i_theta = start_black_radials + radial_task * 2; const int index = grid.index(i_r, i_theta); @@ -119,7 +119,7 @@ void SmootherTake::solveWhiteRadialSection(Vector x, Vec const int num_white_radial_lines = grid.ntheta() / 2; Kokkos::parallel_for( "SmootherTake: moveUpdatedValues (White Radial)", - Kokkos::MDRangePolicy>({0, grid.numberSmootherCircles()}, {num_white_radial_lines, grid.nr()}), + Kokkos::MDRangePolicy>({0, grid.numberSmootherCircles()}, {num_white_radial_lines, grid.nr()}), KOKKOS_LAMBDA(const int radial_task, const int i_r) { const int i_theta = start_white_radials + radial_task * 2; const int index = grid.index(i_r, i_theta); From 07371cc395532e57f3d52245e491bc366283c892 Mon Sep 17 00:00:00 2001 From: Emily Bourne Date: Tue, 12 May 2026 13:41:23 +0200 Subject: [PATCH 11/11] Clang formatting --- .../ExtrapolatedSmootherGive/solveAscSystem.inl | 12 ++++++++---- .../ExtrapolatedSmootherTake/solveAscSystem.inl | 12 ++++++++---- include/Smoother/SmootherGive/solveAscSystem.inl | 9 ++++++--- include/Smoother/SmootherTake/solveAscSystem.inl | 12 ++++++++---- 4 files changed, 30 insertions(+), 15 deletions(-) diff --git a/include/ExtrapolatedSmoother/ExtrapolatedSmootherGive/solveAscSystem.inl b/include/ExtrapolatedSmoother/ExtrapolatedSmootherGive/solveAscSystem.inl index 624c11e1..87ebfcb2 100644 --- a/include/ExtrapolatedSmoother/ExtrapolatedSmootherGive/solveAscSystem.inl +++ b/include/ExtrapolatedSmoother/ExtrapolatedSmootherGive/solveAscSystem.inl @@ -31,7 +31,8 @@ void ExtrapolatedSmootherGive::solveBlackCircleSection(Vector>({0, 0}, {num_black_circles, grid.ntheta()}), + Kokkos::MDRangePolicy>({0, 0}, + {num_black_circles, grid.ntheta()}), KOKKOS_LAMBDA(const int circle_task, const int i_theta) { const int i_r = start_black_circles + circle_task * 2; const int index = grid.index(i_r, i_theta); @@ -71,7 +72,8 @@ void ExtrapolatedSmootherGive::solveWhiteCircleSection(Vector>({0, 0}, {num_white_circles, grid.ntheta()}), + Kokkos::MDRangePolicy>({0, 0}, + {num_white_circles, grid.ntheta()}), KOKKOS_LAMBDA(const int circle_task, const int i_theta) { const int i_r = start_white_circles + circle_task * 2; const int index = grid.index(i_r, i_theta); @@ -100,7 +102,8 @@ void ExtrapolatedSmootherGive::solveBlackRadialSection(Vector>({0, grid.numberSmootherCircles()}, {num_black_radial_lines, grid.nr()}), + Kokkos::MDRangePolicy>({0, grid.numberSmootherCircles()}, + {num_black_radial_lines, grid.nr()}), KOKKOS_LAMBDA(const int radial_task, const int i_r) { const int i_theta = start_black_radials + radial_task * 2; const int index = grid.index(i_r, i_theta); @@ -129,7 +132,8 @@ void ExtrapolatedSmootherGive::solveWhiteRadialSection(Vector>({0, grid.numberSmootherCircles()}, {num_white_radial_lines, grid.nr()}), + Kokkos::MDRangePolicy>({0, grid.numberSmootherCircles()}, + {num_white_radial_lines, grid.nr()}), KOKKOS_LAMBDA(const int radial_task, const int i_r) { const int i_theta = start_white_radials + radial_task * 2; const int index = grid.index(i_r, i_theta); diff --git a/include/ExtrapolatedSmoother/ExtrapolatedSmootherTake/solveAscSystem.inl b/include/ExtrapolatedSmoother/ExtrapolatedSmootherTake/solveAscSystem.inl index 5d73a2da..6f5c7679 100644 --- a/include/ExtrapolatedSmoother/ExtrapolatedSmootherTake/solveAscSystem.inl +++ b/include/ExtrapolatedSmoother/ExtrapolatedSmootherTake/solveAscSystem.inl @@ -31,7 +31,8 @@ void ExtrapolatedSmootherTake::solveBlackCircleSection(Vector>({0, 0}, {num_black_circles, grid.ntheta()}), + Kokkos::MDRangePolicy>({0, 0}, + {num_black_circles, grid.ntheta()}), KOKKOS_LAMBDA(const int circle_task, const int i_theta) { const int i_r = start_black_circles + circle_task * 2; const int index = grid.index(i_r, i_theta); @@ -71,7 +72,8 @@ void ExtrapolatedSmootherTake::solveWhiteCircleSection(Vector>({0, 0}, {num_white_circles, grid.ntheta()}), + Kokkos::MDRangePolicy>({0, 0}, + {num_white_circles, grid.ntheta()}), KOKKOS_LAMBDA(const int circle_task, const int i_theta) { const int i_r = start_white_circles + circle_task * 2; const int index = grid.index(i_r, i_theta); @@ -100,7 +102,8 @@ void ExtrapolatedSmootherTake::solveBlackRadialSection(Vector>({0, grid.numberSmootherCircles()}, {num_black_radial_lines, grid.nr()}), + Kokkos::MDRangePolicy>({0, grid.numberSmootherCircles()}, + {num_black_radial_lines, grid.nr()}), KOKKOS_LAMBDA(const int radial_task, const int i_r) { const int i_theta = start_black_radials + radial_task * 2; const int index = grid.index(i_r, i_theta); @@ -129,7 +132,8 @@ void ExtrapolatedSmootherTake::solveWhiteRadialSection(Vector>({0, grid.numberSmootherCircles()}, {num_white_radial_lines, grid.nr()}), + Kokkos::MDRangePolicy>({0, grid.numberSmootherCircles()}, + {num_white_radial_lines, grid.nr()}), KOKKOS_LAMBDA(const int radial_task, const int i_r) { const int i_theta = start_white_radials + radial_task * 2; const int index = grid.index(i_r, i_theta); diff --git a/include/Smoother/SmootherGive/solveAscSystem.inl b/include/Smoother/SmootherGive/solveAscSystem.inl index b9fad306..2a010874 100644 --- a/include/Smoother/SmootherGive/solveAscSystem.inl +++ b/include/Smoother/SmootherGive/solveAscSystem.inl @@ -26,7 +26,8 @@ void SmootherGive::solveBlackCircleSection(Vector x, Vec const int num_black_circles = (grid.numberSmootherCircles() - start_black_circles + 1) / 2; Kokkos::parallel_for( "SmootherGive: moveUpdatedValues (Black Circular)", - Kokkos::MDRangePolicy>({0, 0}, {num_black_circles, grid.ntheta()}), + Kokkos::MDRangePolicy>({0, 0}, + {num_black_circles, grid.ntheta()}), KOKKOS_LAMBDA(const int circle_task, const int i_theta) { const int i_r = start_black_circles + circle_task * 2; const int index = grid.index(i_r, i_theta); @@ -61,7 +62,8 @@ void SmootherGive::solveWhiteCircleSection(Vector x, Vec const int num_white_circles = (grid.numberSmootherCircles() - start_white_circles + 1) / 2; Kokkos::parallel_for( "SmootherGive: moveUpdatedValues (White Circular)", - Kokkos::MDRangePolicy>({0, 0}, {num_white_circles, grid.ntheta()}), + Kokkos::MDRangePolicy>({0, 0}, + {num_white_circles, grid.ntheta()}), KOKKOS_LAMBDA(const int circle_task, const int i_theta) { const int i_r = start_white_circles + circle_task * 2; const int index = grid.index(i_r, i_theta); @@ -90,7 +92,8 @@ void SmootherGive::solveBlackRadialSection(Vector x, Vec const int num_black_radial_lines = grid.ntheta() / 2; Kokkos::parallel_for( "SmootherGive: moveUpdatedValues (Black Radial)", - Kokkos::MDRangePolicy>({0, grid.numberSmootherCircles()}, {num_black_radial_lines, grid.nr()}), + Kokkos::MDRangePolicy>({0, grid.numberSmootherCircles()}, + {num_black_radial_lines, grid.nr()}), KOKKOS_LAMBDA(const int radial_task, const int i_r) { const int i_theta = start_black_radials + radial_task * 2; const int index = grid.index(i_r, i_theta); diff --git a/include/Smoother/SmootherTake/solveAscSystem.inl b/include/Smoother/SmootherTake/solveAscSystem.inl index 197b3a52..c53abdbc 100644 --- a/include/Smoother/SmootherTake/solveAscSystem.inl +++ b/include/Smoother/SmootherTake/solveAscSystem.inl @@ -26,7 +26,8 @@ void SmootherTake::solveBlackCircleSection(Vector x, Vec const int num_black_circles = (grid.numberSmootherCircles() - start_black_circles + 1) / 2; Kokkos::parallel_for( "SmootherTake: moveUpdatedValues (Black Circular)", - Kokkos::MDRangePolicy>({0, 0}, {num_black_circles, grid.ntheta()}), + Kokkos::MDRangePolicy>({0, 0}, + {num_black_circles, grid.ntheta()}), KOKKOS_LAMBDA(const int circle_task, const int i_theta) { const int i_r = start_black_circles + circle_task * 2; const int index = grid.index(i_r, i_theta); @@ -61,7 +62,8 @@ void SmootherTake::solveWhiteCircleSection(Vector x, Vec const int num_white_circles = (grid.numberSmootherCircles() - start_white_circles + 1) / 2; Kokkos::parallel_for( "SmootherTake: moveUpdatedValues (White Circular)", - Kokkos::MDRangePolicy>({0, 0}, {num_white_circles, grid.ntheta()}), + Kokkos::MDRangePolicy>({0, 0}, + {num_white_circles, grid.ntheta()}), KOKKOS_LAMBDA(const int circle_task, const int i_theta) { const int i_r = start_white_circles + circle_task * 2; const int index = grid.index(i_r, i_theta); @@ -90,7 +92,8 @@ void SmootherTake::solveBlackRadialSection(Vector x, Vec const int num_black_radial_lines = grid.ntheta() / 2; Kokkos::parallel_for( "SmootherTake: moveUpdatedValues (Black Radial)", - Kokkos::MDRangePolicy>({0, grid.numberSmootherCircles()}, {num_black_radial_lines, grid.nr()}), + Kokkos::MDRangePolicy>({0, grid.numberSmootherCircles()}, + {num_black_radial_lines, grid.nr()}), KOKKOS_LAMBDA(const int radial_task, const int i_r) { const int i_theta = start_black_radials + radial_task * 2; const int index = grid.index(i_r, i_theta); @@ -119,7 +122,8 @@ void SmootherTake::solveWhiteRadialSection(Vector x, Vec const int num_white_radial_lines = grid.ntheta() / 2; Kokkos::parallel_for( "SmootherTake: moveUpdatedValues (White Radial)", - Kokkos::MDRangePolicy>({0, grid.numberSmootherCircles()}, {num_white_radial_lines, grid.nr()}), + Kokkos::MDRangePolicy>({0, grid.numberSmootherCircles()}, + {num_white_radial_lines, grid.nr()}), KOKKOS_LAMBDA(const int radial_task, const int i_r) { const int i_theta = start_white_radials + radial_task * 2; const int index = grid.index(i_r, i_theta);