refactor: Rename trans/transpose -> reorder.

Author: semi-h

src/CMakeLists.txt

@@ -17,7 +17,7 @@ set(CUDASRC
   cuda/allocator.f90
   cuda/exec_dist.f90
   cuda/kernels_dist.f90
-  cuda/kernels_trans.f90
+  cuda/kernels_reorder.f90
   cuda/sendrecv.f90
   cuda/tdsops.f90
 )

src/backend.f90

@@ -29,7 +29,7 @@ module m_base_backend
       procedure(transeq_ders), deferred :: transeq_y
       procedure(transeq_ders), deferred :: transeq_z
       procedure(tds_solve), deferred :: tds_solve
-      procedure(trans_d2d), deferred :: trans_d2d
+      procedure(reorder), deferred :: reorder
       procedure(sum9into3), deferred :: sum_yzintox
       procedure(vecadd), deferred :: vecadd
       procedure(get_fields), deferred :: get_fields
@@ -78,8 +78,8 @@ end subroutine tds_solve
    end interface
 
    abstract interface
-      subroutine trans_d2d(self, u_, u, direction)
-         !! transposer subroutines are straightforward, they rearrange
+      subroutine reorder(self, u_, u, direction)
+         !! reorder subroutines are straightforward, they rearrange
          !! data into our specialist data structure so that regardless
          !! of the direction tridiagonal systems are solved efficiently
          !! and fast.
@@ -91,7 +91,7 @@ subroutine trans_d2d(self, u_, u, direction)
          class(field_t), intent(inout) :: u_
          class(field_t), intent(in) :: u
          integer, intent(in) :: direction
-      end subroutine trans_d2d
+      end subroutine reorder
    end interface
 
    abstract interface

src/common.f90

@@ -4,8 +4,8 @@ module m_common
    integer, parameter :: dp=kind(0.0d0)
    real(dp), parameter :: pi = 4*atan(1.0_dp)
 
-   integer, parameter :: TRP_X2Y = 12, TRP_X2Z = 13, TRP_Y2X = 21, &
-                         TRP_Y2Z = 23, TRP_Z2Y = 32
+   integer, parameter :: RDR_X2Y = 12, RDR_X2Z = 13, RDR_Y2X = 21, &
+                         RDR_Y2Z = 23, RDR_Z2Y = 32
 
    type :: globs_t
       integer :: nx, ny, nz

src/cuda/backend.f90

@@ -3,7 +3,7 @@ module m_cuda_backend
 
    use m_allocator, only: allocator_t, field_t
    use m_base_backend, only: base_backend_t
-   use m_common, only: dp, globs_t, TRP_X2Y, TRP_X2Z, TRP_Y2X, TRP_Y2Z, TRP_Z2Y
+   use m_common, only: dp, globs_t, RDR_X2Y, RDR_X2Z, RDR_Y2X, RDR_Y2Z, RDR_Z2Y
    use m_tdsops, only: dirps_t, tdsops_t
 
    use m_cuda_allocator, only: cuda_allocator_t, cuda_field_t
@@ -12,8 +12,8 @@ module m_cuda_backend
    use m_cuda_sendrecv, only: sendrecv_fields, sendrecv_3fields
    use m_cuda_tdsops, only: cuda_tdsops_t
    use m_cuda_kernels_dist, only: transeq_3fused_dist, transeq_3fused_subs
-   use m_cuda_kernels_trans, only: trans_x2y_k, trans_x2z_k, trans_y2x_k, &
-                                   trans_y2z_k, trans_z2y_k
+   use m_cuda_kernels_reorder, only: reorder_x2y, reorder_x2z, reorder_y2x, &
+                                     reorder_y2z, reorder_z2y
 
    implicit none
 
@@ -34,7 +34,7 @@ module m_cuda_backend
       procedure :: transeq_y => transeq_y_cuda
       procedure :: transeq_z => transeq_z_cuda
       procedure :: tds_solve => tds_solve_cuda
-      procedure :: trans_d2d => trans_d2d_cuda
+      procedure :: reorder => reorder_cuda
       procedure :: sum_yzintox => sum_yzintox_cuda
       procedure :: vecadd => vecadd_cuda
       procedure :: set_fields => set_fields_cuda
@@ -417,7 +417,7 @@ subroutine tds_solve_dist(self, du, u, dirps, tdsops, blocks, threads)
 
    end subroutine tds_solve_dist
 
-   subroutine trans_d2d_cuda(self, u_o, u_i, direction)
+   subroutine reorder_cuda(self, u_o, u_i, direction)
       implicit none
 
       class(cuda_backend_t) :: self
@@ -432,35 +432,35 @@ subroutine trans_d2d_cuda(self, u_o, u_i, direction)
       select type(u_i); type is (cuda_field_t); u_i_d => u_i%data_d; end select
 
       select case (direction)
-      case (TRP_X2Y) ! x2y
+      case (RDR_X2Y) ! x2y
          blocks = dim3(self%nx_loc/SZ, self%nz_loc, self%ny_loc/SZ)
          threads = dim3(SZ, SZ, 1)
-         call trans_x2y_k<<<blocks, threads>>>(u_o_d, u_i_d, self%nz_loc)
-      case (TRP_X2Z) ! x2z
+         call reorder_x2y<<<blocks, threads>>>(u_o_d, u_i_d, self%nz_loc)
+      case (RDR_X2Z) ! x2z
          blocks = dim3(self%nx_loc, self%ny_loc/SZ, 1)
          threads = dim3(SZ, 1, 1)
-         call trans_x2z_k<<<blocks, threads>>>(u_o_d, u_i_d, self%nz_loc)
-      case (TRP_Y2X) ! y2x
+         call reorder_x2z<<<blocks, threads>>>(u_o_d, u_i_d, self%nz_loc)
+      case (RDR_Y2X) ! y2x
          blocks = dim3(self%nx_loc/SZ, self%ny_loc/SZ, self%nz_loc)
          threads = dim3(SZ, SZ, 1)
-         call trans_y2x_k<<<blocks, threads>>>(u_o_d, u_i_d, self%nz_loc)
-      case (TRP_Y2Z) ! y2z
+         call reorder_y2x<<<blocks, threads>>>(u_o_d, u_i_d, self%nz_loc)
+      case (RDR_Y2Z) ! y2z
          blocks = dim3(self%nx_loc/SZ, self%ny_loc/SZ, self%nz_loc)
          threads = dim3(SZ, SZ, 1)
-         call trans_y2z_k<<<blocks, threads>>>(u_o_d, u_i_d, &
+         call reorder_y2z<<<blocks, threads>>>(u_o_d, u_i_d, &
                                                self%nx_loc, self%nz_loc)
-      case (TRP_Z2Y) ! z2y
+      case (RDR_Z2Y) ! z2y
          blocks = dim3(self%nx_loc/SZ, self%ny_loc/SZ, self%nz_loc)
          threads = dim3(SZ, SZ, 1)
 
-         call trans_z2y_k<<<blocks, threads>>>(u_o_d, u_i_d, &
+         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
       end select
 
-   end subroutine trans_d2d_cuda
+   end subroutine reorder_cuda
 
    subroutine sum_yzintox_cuda(self, du, dv, dw, &
                                du_y, dv_y, dw_y, du_z, dv_z, dw_z)

src/cuda/kernels_trans.f90 src/cuda/kernels_reorder.f90src/cuda/kernels_trans.f90 renamed to src/cuda/kernels_reorder.f90

@@ -1,12 +1,12 @@
-module m_cuda_kernels_trans
+module m_cuda_kernels_reorder
    use cudafor
 
    use m_common, only: dp
    use m_cuda_common, only: SZ
 
 contains
 
-   attributes(global) subroutine trans_x2y_k(u_y, u_x, nz)
+   attributes(global) subroutine reorder_x2y(u_y, u_x, nz)
       implicit none
 
       real(dp), device, intent(out), dimension(:, :, :) :: u_y
@@ -27,9 +27,9 @@ attributes(global) subroutine trans_x2y_k(u_y, u_x, nz)
       ! copy into output array from shared
       u_y(i, j + (b_k - 1)*SZ, b_j + (b_i - 1)*nz) = tile(j, i)
 
-   end subroutine trans_x2y_k
+   end subroutine reorder_x2y
 
-   attributes(global) subroutine trans_x2z_k(u_z, u_x, nz)
+   attributes(global) subroutine reorder_x2z(u_z, u_x, nz)
       implicit none
 
       real(dp), device, intent(out), dimension(:, :, :) :: u_z
@@ -45,9 +45,9 @@ attributes(global) subroutine trans_x2z_k(u_z, u_x, nz)
          u_z(i, j, b_i + (b_j - 1)*nx) = u_x(i, b_i, j + (b_j - 1)*nz)
       end do
 
-   end subroutine trans_x2z_k
+   end subroutine reorder_x2z
 
-   attributes(global) subroutine trans_y2x_k(u_x, u_y, nz)
+   attributes(global) subroutine reorder_y2x(u_x, u_y, nz)
       implicit none
 
       real(dp), device, intent(out), dimension(:, :, :) :: u_x
@@ -68,9 +68,9 @@ attributes(global) subroutine trans_y2x_k(u_x, u_y, nz)
       ! copy into output array from shared
       u_x(i, (b_i - 1)*SZ + j, (b_j - 1)*nz + b_k) = tile(j, i)
 
-   end subroutine trans_y2x_k
+   end subroutine reorder_y2x
 
-   attributes(global) subroutine trans_y2z_k(u_z, u_y, nx, nz)
+   attributes(global) subroutine reorder_y2z(u_z, u_y, nx, nz)
       implicit none
 
       real(dp), device, intent(out), dimension(:, :, :) :: u_z
@@ -91,9 +91,9 @@ attributes(global) subroutine trans_y2z_k(u_z, u_y, nx, nz)
       ! copy into output array from shared
       u_z(i, b_k, (b_i - 1)*SZ + j + (b_j - 1)*nx) = tile(j, i)
 
-   end subroutine trans_y2z_k
+   end subroutine reorder_y2z
 
-   attributes(global) subroutine trans_z2y_k(u_y, u_z, nx, nz)
+   attributes(global) subroutine reorder_z2y(u_y, u_z, nx, nz)
       implicit none
 
       real(dp), device, intent(out), dimension(:, :, :) :: u_y
@@ -114,6 +114,6 @@ attributes(global) subroutine trans_z2y_k(u_y, u_z, nx, nz)
       ! copy into output array from shared
       u_y(i, (b_j - 1)*SZ + j, (b_i - 1)*nz + b_k) = tile(j, i)
 
-   end subroutine trans_z2y_k
+   end subroutine reorder_z2y
 
-end module m_cuda_kernels_trans
+end module m_cuda_kernels_reorder

src/omp/backend.f90

@@ -24,7 +24,7 @@ module m_omp_backend
       procedure :: transeq_y => transeq_y_omp
       procedure :: transeq_z => transeq_z_omp
       procedure :: tds_solve => tds_solve_omp
-      procedure :: trans_d2d => trans_d2d_omp
+      procedure :: reorder => reorder_omp
       procedure :: sum_yzintox => sum_yzintox_omp
       procedure :: vecadd => vecadd_omp
       procedure :: set_fields => set_fields_omp
@@ -160,15 +160,15 @@ subroutine tds_solve_omp(self, du, u, dirps, tdsops)
 
    end subroutine tds_solve_omp
 
-   subroutine trans_d2d_omp(self, u_, u, direction)
+   subroutine reorder_omp(self, u_, u, direction)
       implicit none
 
       class(omp_backend_t) :: self
       class(field_t), intent(inout) :: u_
       class(field_t), intent(in) :: u
       integer, intent(in) :: direction
 
-   end subroutine trans_d2d_omp
+   end subroutine reorder_omp
 
    subroutine sum_yzintox_omp(self, du, dv, dw, &
                                du_y, dv_y, dw_y, du_z, dv_z, dw_z)

src/solver.f90

@@ -1,7 +1,7 @@
 module m_solver
    use m_allocator, only: allocator_t, field_t
    use m_base_backend, only: base_backend_t
-   use m_common, only: dp, globs_t, TRP_X2Y, TRP_X2Z, TRP_Y2X, TRP_Y2Z, TRP_Z2Y
+   use m_common, only: dp, globs_t, RDR_X2Y, RDR_X2Z, RDR_Y2X, RDR_Y2Z, RDR_Z2Y
    use m_tdsops, only: tdsops_t, dirps_t
    use m_time_integrator, only: time_intg_t
 
@@ -185,9 +185,9 @@ subroutine transeq(self, du, dv, dw, u, v, w)
       dw_y => self%backend%allocator%get_block()
 
       ! reorder data from x orientation to y orientation
-      call self%backend%trans_d2d(u_y, u, TRP_X2Y)
-      call self%backend%trans_d2d(v_y, v, TRP_X2Y)
-      call self%backend%trans_d2d(w_y, w, TRP_X2Y)
+      call self%backend%reorder(u_y, u, RDR_X2Y)
+      call self%backend%reorder(v_y, v, RDR_X2Y)
+      call self%backend%reorder(w_y, w, RDR_X2Y)
 
       ! similar to the x direction, obtain derivatives in y.
       call self%backend%transeq_y(du_y, dv_y, dw_y, u_y, v_y, w_y, self%ydirps)
@@ -209,9 +209,9 @@ subroutine transeq(self, du, dv, dw, u, v, w)
       dw_z => self%backend%allocator%get_block()
 
       ! reorder from x to z
-      call self%backend%trans_d2d(u_z, u, TRP_X2Z)
-      call self%backend%trans_d2d(v_z, v, TRP_X2Z)
-      call self%backend%trans_d2d(w_z, w, TRP_X2Z)
+      call self%backend%reorder(u_z, u, RDR_X2Z)
+      call self%backend%reorder(v_z, v, RDR_X2Z)
+      call self%backend%reorder(w_z, w, RDR_X2Z)
 
       ! get the derivatives in z
       call self%backend%transeq_z(du_z, dv_z, dw_z, u_z, v_z, w_z, self%zdirps)
@@ -267,9 +267,9 @@ subroutine divergence(self, div_u, u, v, w)
       w_y => self%backend%allocator%get_block()
 
       ! reorder data from x orientation to y orientation
-      call self%backend%trans_d2d(u_y, du_x, TRP_X2Y)
-      call self%backend%trans_d2d(v_y, dv_x, TRP_X2Y)
-      call self%backend%trans_d2d(w_y, dw_x, TRP_X2Y)
+      call self%backend%reorder(u_y, du_x, RDR_X2Y)
+      call self%backend%reorder(v_y, dv_x, RDR_X2Y)
+      call self%backend%reorder(w_y, dw_x, RDR_X2Y)
 
       call self%backend%allocator%release_block(du_x)
       call self%backend%allocator%release_block(dv_x)
@@ -303,8 +303,8 @@ subroutine divergence(self, div_u, u, v, w)
       call self%backend%vecadd(1._dp, dw_y, 1._dp, dv_y)
 
       ! reorder from y to z
-      call self%backend%trans_d2d(u_z, du_y, TRP_Y2Z)
-      call self%backend%trans_d2d(w_z, dw_y, TRP_Y2Z)
+      call self%backend%reorder(u_z, du_y, RDR_Y2Z)
+      call self%backend%reorder(w_z, dw_y, RDR_Y2Z)
 
       ! release all the unnecessary blocks.
       call self%backend%allocator%release_block(du_y)
@@ -358,8 +358,8 @@ subroutine gradient(self, dpdx, dpdy, dpdz, pressure)
       dpdz_sxy_y => self%backend%allocator%get_block()
 
       ! reorder data from z orientation to y orientation
-      call self%backend%trans_d2d(p_sxy_y, p_sxy_z, TRP_Z2Y)
-      call self%backend%trans_d2d(dpdz_sxy_y, dpdz_sxy_z, TRP_Z2Y)
+      call self%backend%reorder(p_sxy_y, p_sxy_z, RDR_Z2Y)
+      call self%backend%reorder(dpdz_sxy_y, dpdz_sxy_z, RDR_Z2Y)
 
       call self%backend%allocator%release_block(p_sxy_z)
       call self%backend%allocator%release_block(dpdz_sxy_z)
@@ -386,9 +386,9 @@ subroutine gradient(self, dpdx, dpdy, dpdz, pressure)
       dpdz_sx_x => self%backend%allocator%get_block()
 
       ! reorder from y to x
-      call self%backend%trans_d2d(p_sx_x, p_sx_y, TRP_Y2X)
-      call self%backend%trans_d2d(dpdy_sx_x, dpdy_sx_y, TRP_Y2X)
-      call self%backend%trans_d2d(dpdz_sx_x, dpdz_sx_y, TRP_Y2X)
+      call self%backend%reorder(p_sx_x, p_sx_y, RDR_Y2X)
+      call self%backend%reorder(dpdy_sx_x, dpdy_sx_y, RDR_Y2X)
+      call self%backend%reorder(dpdz_sx_x, dpdz_sx_y, RDR_Y2X)
 
       ! release all the y directional fields.
       call self%backend%allocator%release_block(p_sx_y)