From da1aa0431eb11f6042ab0cf0e5c40ed766b67c81 Mon Sep 17 00:00:00 2001 From: Jony Castagna Date: Mon, 7 Sep 2026 15:50:01 +0100 Subject: [PATCH 1/2] Fix CI pipeline failures when UMPIRE is enabled Resolved compilation errors and test failures that occurred specifically when building with the UMPIRE configuration turned ON. --- include/bout/fieldops.hxx | 4 ++-- src/mesh/parallel/shiftedmetric.cxx | 20 ++++++++++++++++++++ tests/MMS/diffusion/diffusion.cxx | 7 +++++-- tests/unit/CMakeLists.txt | 2 +- 4 files changed, 28 insertions(+), 5 deletions(-) diff --git a/include/bout/fieldops.hxx b/include/bout/fieldops.hxx index b12c1d1046..dcdb4bc5c1 100644 --- a/include/bout/fieldops.hxx +++ b/include/bout/fieldops.hxx @@ -453,7 +453,7 @@ struct BinaryExpr { this->div = div; return *this; } - BOUT_HOST_DEVICE BOUT_FORCEINLINE bool hasParallelSlices() const { + BOUT_FORCEINLINE bool hasParallelSlices() const { if constexpr (is_expr_constant_v && is_expr_constant_v) { return false; } else if constexpr (is_expr_constant_v) { @@ -464,7 +464,7 @@ struct BinaryExpr { return lhs.hasParallelSlices() && rhs.hasParallelSlices(); } } - BOUT_HOST_DEVICE BOUT_FORCEINLINE int numberParallelSlices() const { + BOUT_FORCEINLINE int numberParallelSlices() const { if (!hasParallelSlices()) { return 0; } diff --git a/src/mesh/parallel/shiftedmetric.cxx b/src/mesh/parallel/shiftedmetric.cxx index 972cb58330..3d21723203 100644 --- a/src/mesh/parallel/shiftedmetric.cxx +++ b/src/mesh/parallel/shiftedmetric.cxx @@ -489,6 +489,26 @@ void ShiftedMetric::calcParallelSlices(Field3D& f) { f.splitParallelSlices(); #if BOUT_HAS_CUDA + const bool cuda_fft_supported = + mesh.LocalNz == 16 || mesh.LocalNz == 64 || mesh.LocalNz == 128 + || mesh.LocalNz == 256 || mesh.LocalNz == 512; + + if (!cuda_fft_supported) { + for (const auto& phase : parallel_slice_phases) { + auto& f_slice = f.ynext(phase.y_offset); + f_slice.allocate(); + + BOUT_FOR(i, mesh.getRegion2D("RGN_NOY")) { + const int ix = i.x(); + const int iy = i.y(); + const int iy_offset = iy + phase.y_offset; + shiftZ(&(f(ix, iy_offset, 0)), &(phase.phase_shift(ix, iy, 0)), + &(f_slice(ix, iy_offset, 0))); + } + } + return; + } + auto& region = mesh.getRegion2D("RGN_NOY"); static size_t nblocks = region.getBlocks().size(); if (nblocks != region.getBlocks().size()) { diff --git a/tests/MMS/diffusion/diffusion.cxx b/tests/MMS/diffusion/diffusion.cxx index 3353767e35..571e86b953 100644 --- a/tests/MMS/diffusion/diffusion.cxx +++ b/tests/MMS/diffusion/diffusion.cxx @@ -10,12 +10,15 @@ class Diffusion : public PhysicsModel { protected: int init(bool UNUSED(restarting)) override; int rhs(BoutReal t) override; + +private: + // Keep Umpire-backed field storage owned by the model so that it is + // released before BoutFinalise() tears down BOUT++'s array storage. + Field3D N; }; using bout::globals::mesh; -Field3D N; - BoutReal mu_N; // Parallel collisional diffusion coefficient BoutReal Lx, Ly, Lz; diff --git a/tests/unit/CMakeLists.txt b/tests/unit/CMakeLists.txt index 4384440c74..3be15363fd 100644 --- a/tests/unit/CMakeLists.txt +++ b/tests/unit/CMakeLists.txt @@ -151,5 +151,5 @@ add_dependencies(build-check-unit-tests serial_tests) if(BOUT_HAS_CUDA) set_source_files_properties(${serial_tests_source} PROPERTIES LANGUAGE CUDA) - set_target_properties(serial_tests PROPERTIES CUDA_STANDARD 14) + set_target_properties(serial_tests PROPERTIES CUDA_STANDARD 20) endif() From 33f6acd5f49831719dbf10892ceb8fb6edc494ae Mon Sep 17 00:00:00 2001 From: Jony Castagna Date: Tue, 8 Sep 2026 15:58:51 +0100 Subject: [PATCH 2/2] Support Nz=32 in CUDA shifted-metric FFT --- src/mesh/parallel/shiftedmetric.cxx | 20 ++++++++++++++++---- 1 file changed, 16 insertions(+), 4 deletions(-) diff --git a/src/mesh/parallel/shiftedmetric.cxx b/src/mesh/parallel/shiftedmetric.cxx index 3d21723203..c9c9e956d8 100644 --- a/src/mesh/parallel/shiftedmetric.cxx +++ b/src/mesh/parallel/shiftedmetric.cxx @@ -257,6 +257,8 @@ __global__ void fft_block_cooperative(const BoutReal** __restrict__ in, const double2* twiddles; if constexpr (NZ == 16) { twiddles = c_twiddle_16; + } else if constexpr (NZ == 32) { + twiddles = c_twiddle_32; } else if constexpr (NZ == 64) { twiddles = c_twiddle_64; } else if constexpr (NZ == 128) { @@ -266,8 +268,9 @@ __global__ void fft_block_cooperative(const BoutReal** __restrict__ in, } else if constexpr (NZ == 512) { twiddles = c_twiddle_512; } else { - static_assert(NZ == 16 || NZ == 64 || NZ == 128 || NZ == 256 || NZ == 512, - "Unsupported NZ"); + static_assert( + NZ == 16 || NZ == 32 || NZ == 64 || NZ == 128 || NZ == 256 || NZ == 512, + "Unsupported NZ"); } // Each block processes FFTS_PER_BLOCK FFTs @@ -429,6 +432,15 @@ static void shiftZ_block_fft(const int Nz, const BoutReal** in, BoutReal** out, fft_block_cooperative<16, FFTS_PER_BLOCK> <<>>(in, out, phs, nbatches, nblocks); + } else if (Nz == 32) { + constexpr int FFTS_PER_BLOCK = 8; + constexpr int THREADS_PER_FFT = 32; + + dim3 block(THREADS_PER_FFT, FFTS_PER_BLOCK); + dim3 grid((total_ffts + FFTS_PER_BLOCK - 1) / FFTS_PER_BLOCK); + + fft_block_cooperative<32, FFTS_PER_BLOCK> + <<>>(in, out, phs, nbatches, nblocks); } else if (Nz == 64) { constexpr int FFTS_PER_BLOCK = 4; constexpr int THREADS_PER_FFT = 64; @@ -490,8 +502,8 @@ void ShiftedMetric::calcParallelSlices(Field3D& f) { #if BOUT_HAS_CUDA const bool cuda_fft_supported = - mesh.LocalNz == 16 || mesh.LocalNz == 64 || mesh.LocalNz == 128 - || mesh.LocalNz == 256 || mesh.LocalNz == 512; + mesh.LocalNz == 16 || mesh.LocalNz == 32 || mesh.LocalNz == 64 || + mesh.LocalNz == 128 || mesh.LocalNz == 256 || mesh.LocalNz == 512; if (!cuda_fft_supported) { for (const auto& phase : parallel_slice_phases) {