fix: Minor changes.

Author: semi-h

src/cuda/backend.f90

@@ -1,4 +1,5 @@
 module m_cuda_backend
+   use iso_fortran_env, only: stderr => error_unit
    use cudafor
 
    use m_allocator, only: allocator_t, field_t
@@ -452,12 +453,10 @@ subroutine reorder_cuda(self, u_o, u_i, direction)
       case (RDR_Z2Y) ! z2y
          blocks = dim3(self%nx_loc/SZ, self%ny_loc/SZ, self%nz_loc)
          threads = dim3(SZ, SZ, 1)
-
          call reorder_z2y<<<blocks, threads>>>(u_o_d, u_i_d, &
                                                self%nx_loc, self%nz_loc)
       case default
-         print *, 'Transpose direction is undefined.'
-         stop
+         error stop 'Reorder direction is undefined.'
       end select
 
    end subroutine reorder_cuda

src/cuda/kernels_reorder.f90

@@ -36,11 +36,13 @@ attributes(global) subroutine reorder_x2z(u_z, u_x, nz)
       real(dp), device, intent(in), dimension(:, :, :) :: u_x
       integer, value, intent(in) :: nz
 
-      integer :: i, j, b_i, b_j, nx!, nz
+      integer :: i, j, b_i, b_j, nx
 
       i = threadIdx%x; b_i = blockIdx%x; b_j = blockIdx%y
       nx = gridDim%x
 
+      ! Data access pattern for reordering between x and z is quite nice
+      ! thus we don't need to use shared memory for this operation.
       do j = 1, nz
          u_z(i, j, b_i + (b_j - 1)*nx) = u_x(i, b_i, j + (b_j - 1)*nz)
       end do