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<L> && is_expr_constant_v<R>) {
         return false;
       } else if constexpr (is_expr_constant_v<L>) {
@@ -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..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>
         <<<grid, block, 0, stream>>>(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>
+        <<<grid, block, 0, stream>>>(in, out, phs, nbatches, nblocks);
   } else if (Nz == 64) {
     constexpr int FFTS_PER_BLOCK = 4;
     constexpr int THREADS_PER_FFT = 64;
@@ -489,6 +501,26 @@ void ShiftedMetric::calcParallelSlices(Field3D& f) {
   f.splitParallelSlices();
 
 #if BOUT_HAS_CUDA
+  const bool cuda_fft_supported =
+      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) {
+      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()