[14091] | 1 | MODULE diahsb |
---|
| 2 | !!====================================================================== |
---|
| 3 | !! *** MODULE diahsb *** |
---|
| 4 | !! Ocean diagnostics: Heat, salt and volume budgets |
---|
| 5 | !!====================================================================== |
---|
| 6 | !! History : 3.3 ! 2010-09 (M. Leclair) Original code |
---|
| 7 | !! ! 2012-10 (C. Rousset) add iom_put |
---|
| 8 | !!---------------------------------------------------------------------- |
---|
| 9 | |
---|
| 10 | !!---------------------------------------------------------------------- |
---|
| 11 | !! dia_hsb : Diagnose the conservation of ocean heat and salt contents, and volume |
---|
| 12 | !! dia_hsb_rst : Read or write DIA file in restart file |
---|
| 13 | !! dia_hsb_init : Initialization of the conservation diagnostic |
---|
| 14 | !!---------------------------------------------------------------------- |
---|
| 15 | USE oce ! ocean dynamics and tracers |
---|
| 16 | USE dom_oce ! ocean space and time domain |
---|
| 17 | USE phycst ! physical constants |
---|
| 18 | USE sbc_oce ! surface thermohaline fluxes |
---|
| 19 | USE isf_oce ! ice shelf fluxes |
---|
| 20 | USE sbcrnf ! river runoff |
---|
| 21 | USE domvvl ! vertical scale factors |
---|
| 22 | USE traqsr ! penetrative solar radiation |
---|
| 23 | USE trabbc ! bottom boundary condition |
---|
| 24 | USE trabbc ! bottom boundary condition |
---|
| 25 | USE restart ! ocean restart |
---|
| 26 | USE bdy_oce , ONLY : ln_bdy |
---|
| 27 | ! |
---|
| 28 | USE iom ! I/O manager |
---|
| 29 | USE in_out_manager ! I/O manager |
---|
| 30 | |
---|
| 31 | USE gpu_manager ! GPU manager |
---|
| 32 | USE cudafor ! CUDA toolkit libs |
---|
| 33 | USE cuda_fortran ! CUDA routines |
---|
| 34 | !USE nvtx ! CUDA profiling/DEGUG tools |
---|
| 35 | |
---|
| 36 | USE lib_fortran ! glob_sum |
---|
| 37 | USE lib_mpp ! distributed memory computing library |
---|
| 38 | USE timing ! preformance summary |
---|
| 39 | |
---|
| 40 | IMPLICIT NONE |
---|
| 41 | PRIVATE |
---|
| 42 | |
---|
| 43 | PUBLIC dia_hsb ! routine called by step.F90 |
---|
| 44 | PUBLIC dia_hsb_init ! routine called by nemogcm.F90 |
---|
| 45 | |
---|
| 46 | LOGICAL, PUBLIC :: ln_diahsb !: check the heat and salt budgets |
---|
| 47 | |
---|
| 48 | REAL(wp) :: surf_tot ! ocean surface |
---|
| 49 | |
---|
| 50 | |
---|
| 51 | |
---|
| 52 | |
---|
| 53 | REAL(wp) , DIMENSION(2), SAVE :: frc_t, frc_s, frc_v ! global forcing trends |
---|
| 54 | REAL(wp) , DIMENSION(2), SAVE :: frc_wn_t, frc_wn_s ! global forcing trends |
---|
| 55 | |
---|
| 56 | REAL(wp), DIMENSION(:,:) , ALLOCATABLE :: surf |
---|
| 57 | REAL(wp), DIMENSION(:,:) , ALLOCATABLE :: surf_ini , ssh_ini ! |
---|
| 58 | REAL(wp), DIMENSION(:,:) , ALLOCATABLE :: ssh_hc_loc_ini, ssh_sc_loc_ini ! |
---|
| 59 | |
---|
| 60 | |
---|
| 61 | |
---|
| 62 | REAL(wp), DIMENSION(:,:,:), ALLOCATABLE, PINNED :: hc_loc_ini, sc_loc_ini ! |
---|
| 63 | REAL(wp), DIMENSION(:,:,:), ALLOCATABLE :: e3t_ini ! |
---|
| 64 | REAL(wp), DIMENSION(:) , ALLOCATABLE, PINNED, SAVE :: h_ztmpv, h_ztmph, h_ztmps, h_ztmp ! |
---|
| 65 | |
---|
| 66 | REAL(wp), DIMENSION(:,:,:), ALLOCATABLE :: tmask_ini |
---|
| 67 | |
---|
| 68 | |
---|
| 69 | !Device data associate to PUBLIC arrays |
---|
| 70 | REAL(8), DIMENSION(:,:,:,:) , ALLOCATABLE, DEVICE :: d_e3t ! |
---|
| 71 | REAL(8), DIMENSION(:,:,:) , ALLOCATABLE, DEVICE :: d_tmask ! |
---|
| 72 | REAL(8), DIMENSION(:,:) , ALLOCATABLE, DEVICE :: d_tmask_h ! |
---|
| 73 | REAL(8), DIMENSION(:,:,:) , ALLOCATABLE, DEVICE :: d_tmask_ini ! |
---|
| 74 | REAL(8), DIMENSION(:,:,:,:,:), ALLOCATABLE, DEVICE :: d_ts ! |
---|
| 75 | !Device data associate to LOCAL/DEVICE arrays |
---|
| 76 | REAL(8), DEVICE , DIMENSION(:,:) , ALLOCATABLE :: d_surf ! |
---|
| 77 | REAL(8), DEVICE , DIMENSION(:,:) , ALLOCATABLE :: d_surf_ini ! |
---|
| 78 | REAL(8), DEVICE , DIMENSION(:,:,:) , ALLOCATABLE :: d_hc_loc_ini ! |
---|
| 79 | REAL(8), DEVICE , DIMENSION(:,:,:) , ALLOCATABLE :: d_sc_loc_ini ! |
---|
| 80 | REAL(8), DEVICE , DIMENSION(:,:,:) , ALLOCATABLE :: d_e3t_ini ! |
---|
| 81 | REAL(8), DEVICE , DIMENSION(:,:,:) , ALLOCATABLE :: d_zwrkv, d_zwrkh, d_zwrks, d_zwrk ! 3D GPU workspace |
---|
| 82 | REAL(8), DEVICE :: ztmpv, ztmph, ztmps, ztmp ! Device Reduction |
---|
| 83 | ! |
---|
| 84 | INTEGER :: globsize ! 3D workspace size |
---|
| 85 | type(dim3) :: dimGrid, dimBlock ! cuda parameters |
---|
| 86 | INTEGER, parameter :: nstreams = 3 ! Streams Number |
---|
| 87 | INTEGER(kind=cuda_stream_kind) :: stream(nstreams), str ! Stream ID |
---|
| 88 | !DEBUG |
---|
| 89 | !REAL(8) , save , DIMENSION(:,:,:) , ALLOCATABLE :: prev_3d |
---|
| 90 | !REAL(8) :: accum |
---|
| 91 | |
---|
| 92 | |
---|
| 93 | |
---|
| 94 | |
---|
| 95 | !! * Substitutions |
---|
| 96 | !# include "domzgr_substitute.h90" |
---|
| 97 | !!---------------------------------------------------------------------- |
---|
| 98 | !! NEMO/OCE 4.0 , NEMO Consortium (2018) |
---|
| 99 | !! $Id$ |
---|
| 100 | !! Software governed by the CeCILL license (see ./LICENSE) |
---|
| 101 | !!---------------------------------------------------------------------- |
---|
| 102 | CONTAINS |
---|
| 103 | |
---|
| 104 | SUBROUTINE dia_hsb( kt, Kbb, Kmm ) |
---|
| 105 | !!--------------------------------------------------------------------------- |
---|
| 106 | !! *** ROUTINE dia_hsb *** |
---|
| 107 | !! |
---|
| 108 | !! ** Purpose: Compute the ocean global heat content, salt content and volume conservation |
---|
| 109 | !! |
---|
| 110 | !! ** Method : - Compute the deviation of heat content, salt content and volume |
---|
| 111 | !! at the current time step from their values at nit000 |
---|
| 112 | !! - Compute the contribution of forcing and remove it from these deviations |
---|
| 113 | !! |
---|
| 114 | !!--------------------------------------------------------------------------- |
---|
| 115 | INTEGER, INTENT(in) :: kt ! ocean time-step index |
---|
| 116 | INTEGER, INTENT(in) :: Kbb, Kmm ! ocean time level indices |
---|
| 117 | ! |
---|
| 118 | |
---|
| 119 | INTEGER, VALUE :: ji, jj, jk, kts ! dummy loop indice |
---|
| 120 | INTEGER, VALUE :: localsize ! jpi * jpj * jpk |
---|
| 121 | INTEGER :: istat ! CUDA error check |
---|
| 122 | COMPLEX :: ctmp ! dummy complex number |
---|
| 123 | INTEGER(kind=cuda_stream_kind) :: str ! dummy kernel stream |
---|
| 124 | INTEGER :: tile_n, tile_b ! tile indexe. _n now, _b before |
---|
| 125 | REAL(wp) , DIMENSION(2), SAVE :: zdiff_hc1, zdiff_sc1 ! heat and salt content variations |
---|
| 126 | REAL(wp) , DIMENSION(2), SAVE :: zdiff_hc, zdiff_sc ! - - - - |
---|
| 127 | REAL(wp) , DIMENSION(2), SAVE :: zdiff_v2 ! volume variation |
---|
| 128 | REAL(wp) , DIMENSION(2), SAVE :: zdiff_v1 ! volume variation |
---|
| 129 | REAL(wp) , DIMENSION(2), SAVE :: zerr_hc1, zerr_sc1 ! heat and salt content misfit |
---|
| 130 | REAL(wp) , DIMENSION(2), SAVE :: zvol_tot ! volume |
---|
| 131 | REAL(wp) , DIMENSION(2), SAVE :: z_frc_trd_t, z_frc_trd_s ! - - |
---|
| 132 | REAL(wp) , DIMENSION(2), SAVE :: z_frc_trd_v ! - - |
---|
| 133 | REAL(wp) , DIMENSION(2), SAVE :: z_wn_trd_t, z_wn_trd_s ! - - |
---|
| 134 | REAL(wp) , DIMENSION(2), SAVE :: z_ssh_hc, z_ssh_sc ! - - |
---|
| 135 | # 147 "diahsb_new.F90" |
---|
| 136 | REAL(wp), DIMENSION(jpi,jpj) :: z2d0, z2d1 ! 2D workspace |
---|
| 137 | REAL(wp), DIMENSION(jpi,jpj,jpkm1) :: zwrk ! 3D workspace |
---|
| 138 | !!--------------------------------------------------------------------------- |
---|
| 139 | IF( ln_timing ) CALL timing_start('dia_hsb') |
---|
| 140 | |
---|
| 141 | localsize = jpi * jpj * jpk |
---|
| 142 | kts = kt |
---|
| 143 | IF (kts == 1) THEN |
---|
| 144 | tile_n = 1 |
---|
| 145 | tile_b = 1 |
---|
| 146 | ELSE |
---|
| 147 | IF( MOD(kts,2) == 0) THEN |
---|
| 148 | tile_n = 2 |
---|
| 149 | tile_b = 1 |
---|
| 150 | ELSE IF( MOD(kts,2) == 1 ) THEN |
---|
| 151 | tile_n = 1 |
---|
| 152 | tile_b = 2 |
---|
| 153 | END IF |
---|
| 154 | END IF |
---|
| 155 | |
---|
| 156 | ! |
---|
| 157 | ts(:,:,:,1,Kmm) = ts(:,:,:,1,Kmm) * tmask(:,:,:) ; ts(:,:,:,1,Kbb) = ts(:,:,:,1,Kbb) * tmask(:,:,:) ; |
---|
| 158 | ts(:,:,:,2,Kmm) = ts(:,:,:,2,Kmm) * tmask(:,:,:) ; ts(:,:,:,2,Kbb) = ts(:,:,:,2,Kbb) * tmask(:,:,:) ; |
---|
| 159 | ! ------------------------- ! |
---|
| 160 | ! 1 - Trends due to forcing ! |
---|
| 161 | ! ------------------------- ! |
---|
| 162 | |
---|
| 163 | z_frc_trd_v(tile_n) = r1_rho0 * glob_sum( 'diahsb', - ( emp(:,:) - rnf(:,:) + fwfisf_cav(:,:) + fwfisf_par(:,:) ) * surf(:,:) )! volume fluxes |
---|
| 164 | z_frc_trd_t(tile_n) = glob_sum( 'diahsb', sbc_tsc(:,:,jp_tem) * surf(:,:) ) ! heat fluxes |
---|
| 165 | z_frc_trd_s(tile_n) = glob_sum( 'diahsb', sbc_tsc(:,:,jp_sal) * surf(:,:) ) ! salt fluxes |
---|
| 166 | ! ! Add runoff heat & salt input |
---|
| 167 | IF( ln_rnf ) z_frc_trd_t(tile_n) = z_frc_trd_t(tile_n) + glob_sum( 'diahsb', rnf_tsc(:,:,jp_tem) * surf(:,:) ) |
---|
| 168 | IF( ln_rnf_sal) z_frc_trd_s(tile_n) = z_frc_trd_s(tile_n) + glob_sum( 'diahsb', rnf_tsc(:,:,jp_sal) * surf(:,:) ) ! Add ice shelf heat & salt input |
---|
| 169 | IF( ln_isf ) z_frc_trd_t(tile_n) = z_frc_trd_t(tile_n) & |
---|
| 170 | & + glob_sum( 'diahsb', ( risf_cav_tsc(:,:,jp_tem) + risf_par_tsc(:,:,jp_tem) ) * surf(:,:) ) ! Add penetrative solar radiation |
---|
| 171 | IF( ln_traqsr ) z_frc_trd_t(tile_n) = z_frc_trd_t(tile_n) + r1_rho0_rcp * glob_sum( 'diahsb', qsr (:,:) * surf(:,:) ) ! Add geothermal heat flux |
---|
| 172 | IF( ln_trabbc ) z_frc_trd_t(tile_n) = z_frc_trd_t(tile_n) + glob_sum( 'diahsb', qgh_trd0(:,:) * surf(:,:) ) |
---|
| 173 | ! |
---|
| 174 | IF( ln_linssh ) THEN |
---|
| 175 | IF( ln_isfcav ) THEN |
---|
| 176 | DO ji=1,jpi |
---|
| 177 | DO jj=1,jpj |
---|
| 178 | z2d0(ji,jj) = surf(ji,jj) * ww(ji,jj,mikt(ji,jj)) * ts(ji,jj,mikt(ji,jj),jp_tem,Kbb) |
---|
| 179 | z2d1(ji,jj) = surf(ji,jj) * ww(ji,jj,mikt(ji,jj)) * ts(ji,jj,mikt(ji,jj),jp_sal,Kbb) |
---|
| 180 | END DO |
---|
| 181 | END DO |
---|
| 182 | ELSE |
---|
| 183 | z2d0(:,:) = surf(:,:) * ww(:,:,1) * ts(:,:,1,jp_tem,Kbb) |
---|
| 184 | z2d1(:,:) = surf(:,:) * ww(:,:,1) * ts(:,:,1,jp_sal,Kbb) |
---|
| 185 | END IF |
---|
| 186 | |
---|
| 187 | z_wn_trd_t(tile_n) = - glob_sum( 'diahsb', z2d0 ) |
---|
| 188 | z_wn_trd_s(tile_n) = - glob_sum( 'diahsb', z2d1 ) |
---|
| 189 | |
---|
| 190 | |
---|
| 191 | |
---|
| 192 | |
---|
| 193 | ENDIF |
---|
| 194 | |
---|
| 195 | IF (kts>1) THEN |
---|
| 196 | frc_v(tile_n) = frc_v(tile_b) |
---|
| 197 | frc_t(tile_n) = frc_t(tile_b) |
---|
| 198 | frc_s(tile_n) = frc_s(tile_b) |
---|
| 199 | frc_wn_t(tile_n) = frc_wn_t(tile_b) |
---|
| 200 | frc_wn_s(tile_n) = frc_wn_s(tile_b) |
---|
| 201 | END IF |
---|
| 202 | frc_v(tile_n) = frc_v(tile_n) + z_frc_trd_v(tile_n) * rn_Dt |
---|
| 203 | frc_t(tile_n) = frc_t(tile_n) + z_frc_trd_t(tile_n) * rn_Dt |
---|
| 204 | frc_s(tile_n) = frc_s(tile_n) + z_frc_trd_s(tile_n) * rn_Dt |
---|
| 205 | ! ! Advection flux through fixed surface (z=0) |
---|
| 206 | IF( ln_linssh ) THEN |
---|
| 207 | frc_wn_t(tile_n) = frc_wn_t(tile_n) + z_wn_trd_t(tile_n) * rn_Dt |
---|
| 208 | frc_wn_s(tile_n) = frc_wn_s(tile_n) + z_wn_trd_s(tile_n) * rn_Dt |
---|
| 209 | ENDIF |
---|
| 210 | ! ------------------------ ! |
---|
| 211 | ! 2 - Content variations ! |
---|
| 212 | ! ------------------------ ! |
---|
| 213 | ! glob_sum_full is needed because you keep the full interior domain to compute the sum (iscpl) |
---|
| 214 | |
---|
| 215 | ! ! volume variation (calculated with ssh) |
---|
| 216 | |
---|
| 217 | zdiff_v1(tile_n) = glob_sum_full( 'diahsb', surf(:,:)*ssh(:,:,Kmm) - surf_ini(:,:)*ssh_ini(:,:) ) |
---|
| 218 | |
---|
| 219 | |
---|
| 220 | |
---|
| 221 | ! ! heat & salt content variation (associated with ssh) |
---|
| 222 | IF( ln_linssh ) THEN ! linear free surface case |
---|
| 223 | IF( ln_isfcav ) THEN ! ISF case |
---|
| 224 | DO ji = 1, jpi |
---|
| 225 | DO jj = 1, jpj |
---|
| 226 | z2d0(ji,jj) = surf(ji,jj) * ( ts(ji,jj,mikt(ji,jj),jp_tem,Kmm) * ssh(ji,jj,Kmm) - ssh_hc_loc_ini(ji,jj) ) |
---|
| 227 | z2d1(ji,jj) = surf(ji,jj) * ( ts(ji,jj,mikt(ji,jj),jp_sal,Kmm) * ssh(ji,jj,Kmm) - ssh_sc_loc_ini(ji,jj) ) |
---|
| 228 | END DO |
---|
| 229 | END DO |
---|
| 230 | ELSE ! no under ice-shelf seas |
---|
| 231 | z2d0(:,:) = surf(:,:) * ( ts(:,:,1,jp_tem,Kmm) * ssh(:,:,Kmm) - ssh_hc_loc_ini(:,:) ) |
---|
| 232 | z2d1(:,:) = surf(:,:) * ( ts(:,:,1,jp_sal,Kmm) * ssh(:,:,Kmm) - ssh_sc_loc_ini(:,:) ) |
---|
| 233 | END IF |
---|
| 234 | |
---|
| 235 | z_ssh_hc(tile_n) = glob_sum_full( 'diahsb', z2d0 ) |
---|
| 236 | z_ssh_sc(tile_n) = glob_sum_full( 'diahsb', z2d1 ) |
---|
| 237 | |
---|
| 238 | |
---|
| 239 | |
---|
| 240 | |
---|
| 241 | ENDIF |
---|
| 242 | |
---|
| 243 | str = stream(tile_n) |
---|
| 244 | istat = 0 |
---|
| 245 | istat = cudaMemcpyAsync( d_e3t, e3t , jpi*jpj*jpk*jpt , str ) + istat |
---|
| 246 | istat = cudaMemcpyAsync(d_hc_loc_ini, hc_loc_ini, jpi*jpj*jpk , str ) + istat |
---|
| 247 | istat = cudaMemcpyAsync(d_sc_loc_ini, sc_loc_ini, jpi*jpj*jpk , str ) + istat |
---|
| 248 | istat = cudaMemcpyAsync( d_ts, ts , jpi*jpj*jpk*2*jpt, str ) + istat |
---|
| 249 | IF( istat /= 0 ) THEN |
---|
| 250 | CALL ctl_stop( 'dia_hsb: unable to async GPU copy H2D' ) ; RETURN |
---|
| 251 | ENDIF |
---|
| 252 | dimBlock = dim3(4,4,4) |
---|
| 253 | dimGrid = dim3( ceiling( real( jpi ) / dimBlock%x ) , ceiling( real( jpj ) / dimBlock%y ) , & |
---|
| 254 | ceiling( real( jpkm1 ) / dimBlock%z ) ) |
---|
| 255 | ! |
---|
| 256 | CALL dia_hsb_kernel<<<dimGrid, dimBlock, 0, str>>> (d_surf , d_e3t, d_surf_ini, d_e3t_ini, & |
---|
| 257 | & d_ts, d_hc_loc_ini, d_sc_loc_ini, d_tmask, d_tmask_ini, d_zwrkv, d_zwrkh, d_zwrks, d_zwrk, jpi, jpj, jpk, jpt, Kmm) |
---|
| 258 | |
---|
| 259 | |
---|
| 260 | CALL filter_cuda<<<dimGrid, dimBlock, 0, str>>>(d_zwrkv , d_tmask_h , jpi, jpj, jpk) |
---|
| 261 | CALL filter_cuda<<<dimGrid, dimBlock, 0, str>>>(d_zwrkh , d_tmask_h , jpi, jpj, jpk) |
---|
| 262 | CALL filter_cuda<<<dimGrid, dimBlock, 0, str>>>(d_zwrks , d_tmask_h , jpi, jpj, jpk) |
---|
| 263 | CALL filter_cuda<<<dimGrid, dimBlock, 0, str>>>(d_zwrk , d_tmask_h , jpi, jpj, jpk) |
---|
| 264 | |
---|
| 265 | ztmpv = 0.e0 |
---|
| 266 | ztmph = 0.e0 |
---|
| 267 | ztmps = 0.e0 |
---|
| 268 | ztmp = 0.e0 |
---|
| 269 | |
---|
| 270 | istat = 0 |
---|
| 271 | !$cuf kernel do <<< *, *, stream=str >>> |
---|
| 272 | do ji = 1, localsize |
---|
| 273 | ztmpv = ztmpv + d_zwrkv(ji) |
---|
| 274 | end do |
---|
| 275 | istat = cudaMemcpyAsync( h_ztmpv(tile_n) , ztmpv , 1 , str ) + istat |
---|
| 276 | !$cuf kernel do <<< *, *, stream=str >>> |
---|
| 277 | do ji = 1, localsize |
---|
| 278 | ztmph = ztmph + d_zwrkh(ji) |
---|
| 279 | end do |
---|
| 280 | istat = cudaMemcpyAsync( h_ztmph(tile_n) , ztmph , 1 , str ) + istat |
---|
| 281 | !$cuf kernel do <<< *, *, stream=str >>> |
---|
| 282 | do ji = 1, localsize |
---|
| 283 | ztmps = ztmps + d_zwrks(ji) |
---|
| 284 | end do |
---|
| 285 | istat = cudaMemcpyAsync( h_ztmps(tile_n) , ztmps , 1 , str ) + istat |
---|
| 286 | !$cuf kernel do <<< *, *, stream=str >>> |
---|
| 287 | do ji = 1, localsize |
---|
| 288 | ztmp = ztmp + d_zwrk(ji) |
---|
| 289 | end do |
---|
| 290 | istat = cudaMemcpyAsync( h_ztmp (tile_n) , ztmp , 1 , str ) + istat |
---|
| 291 | ! |
---|
| 292 | IF( istat /= 0 ) THEN |
---|
| 293 | CALL ctl_stop( 'dia_hsb: unable to async GPU copy D2H' ) ; RETURN |
---|
| 294 | ENDIF |
---|
| 295 | ! |
---|
| 296 | istat = cudaStreamSynchronize(stream(tile_b)) |
---|
| 297 | ! |
---|
| 298 | IF( istat /= 0 ) THEN |
---|
| 299 | CALL ctl_stop( 'dia_hsb: unable to stream synchronize' ) ; RETURN |
---|
| 300 | ENDIF |
---|
| 301 | ! |
---|
| 302 | ctmp = CMPLX( h_ztmpv(tile_b) , 0.e0, 8 ) |
---|
| 303 | CALL mpp_sum('diahsb', ctmp ) |
---|
| 304 | zdiff_v2(tile_b) = REAL( ctmp, 8 ) |
---|
| 305 | |
---|
| 306 | ctmp = CMPLX( h_ztmph(tile_b) , 0.e0, 8 ) |
---|
| 307 | CALL mpp_sum('diahsb', ctmp ) |
---|
| 308 | zdiff_hc(tile_b) = REAL( ctmp, 8 ) |
---|
| 309 | |
---|
| 310 | ctmp = CMPLX( h_ztmps(tile_b) , 0.e0, 8 ) |
---|
| 311 | CALL mpp_sum('diahsb', ctmp ) |
---|
| 312 | zdiff_sc(tile_b) = REAL( ctmp, 8 ) |
---|
| 313 | |
---|
| 314 | ctmp = CMPLX( h_ztmp(tile_b) , 0.e0, 8 ) |
---|
| 315 | CALL mpp_sum('diahsb', ctmp ) |
---|
| 316 | zvol_tot(tile_b) = REAL( ctmp, 8 ) |
---|
| 317 | |
---|
| 318 | IF ( kt == nitend ) THEN |
---|
| 319 | ! |
---|
| 320 | istat = cudaStreamSynchronize(stream(tile_n)) |
---|
| 321 | IF( istat /= 0 ) THEN |
---|
| 322 | CALL ctl_stop( 'dia_hsb: unable to stream synchronize' ) ; RETURN |
---|
| 323 | ENDIF |
---|
| 324 | ! |
---|
| 325 | ctmp = CMPLX( h_ztmpv(tile_n) , 0.e0, 8 ) |
---|
| 326 | CALL mpp_sum('diahsb', ctmp ) |
---|
| 327 | zdiff_v2(tile_n) = REAL( ctmp, 8 ) |
---|
| 328 | |
---|
| 329 | ctmp = CMPLX( h_ztmph(tile_n) , 0.e0, 8 ) |
---|
| 330 | CALL mpp_sum('diahsb', ctmp ) |
---|
| 331 | zdiff_hc(tile_n) = REAL( ctmp, 8 ) |
---|
| 332 | |
---|
| 333 | ctmp = CMPLX( h_ztmps(tile_n) , 0.e0, 8 ) |
---|
| 334 | CALL mpp_sum('diahsb', ctmp ) |
---|
| 335 | zdiff_sc(tile_n) = REAL( ctmp, 8 ) |
---|
| 336 | |
---|
| 337 | ctmp = CMPLX( h_ztmp(tile_n) , 0.e0, 8 ) |
---|
| 338 | CALL mpp_sum('diahsb', ctmp ) |
---|
| 339 | zvol_tot(tile_n) = REAL( ctmp, 8 ) |
---|
| 340 | ENDIF |
---|
| 341 | ! ------------------------ ! |
---|
| 342 | ! 3 - Drifts ! |
---|
| 343 | ! ------------------------ ! |
---|
| 344 | |
---|
| 345 | IF ( kt > 1 ) THEN |
---|
| 346 | kts = kts - 1 |
---|
| 347 | zdiff_v1(tile_b) = zdiff_v1(tile_b) - frc_v(tile_b) |
---|
| 348 | IF( .NOT.ln_linssh ) zdiff_v2(tile_b) = zdiff_v2(tile_b) - frc_v(tile_b) |
---|
| 349 | zdiff_hc(tile_b) = zdiff_hc(tile_b) - frc_t(tile_b) |
---|
| 350 | zdiff_sc(tile_b) = zdiff_sc(tile_b) - frc_s(tile_b) |
---|
| 351 | IF( ln_linssh ) THEN |
---|
| 352 | zdiff_hc1(tile_b) = zdiff_hc (tile_b) + z_ssh_hc(tile_b) |
---|
| 353 | zdiff_sc1(tile_b) = zdiff_sc (tile_b) + z_ssh_sc(tile_b) |
---|
| 354 | zerr_hc1 (tile_b) = z_ssh_hc(tile_b) - frc_wn_t(tile_b) |
---|
| 355 | zerr_sc1 (tile_b) = z_ssh_sc(tile_b) - frc_wn_s(tile_b) |
---|
| 356 | ENDIF |
---|
| 357 | !!gm to be added ? |
---|
| 358 | ! IF( ln_linssh ) THEN ! fixed volume, add the ssh contribution |
---|
| 359 | ! zvol_tot = zvol_tot + glob_sum( 'diahsb', surf(:,:) * sshn(:,:) ) |
---|
| 360 | ! ENDIF |
---|
| 361 | !!gm end |
---|
| 362 | |
---|
| 363 | CALL iom_put( 'bgfrcvol' , frc_v(tile_b) * 1.e-9 ) ! vol - surface forcing (km3) |
---|
| 364 | CALL iom_put( 'bgfrctem' , frc_t(tile_b) * rho0 * rcp * 1.e-20 ) ! hc - surface forcing (1.e20 J) |
---|
| 365 | CALL iom_put( 'bgfrchfx' , frc_t(tile_b) * rho0 * rcp / & ! hc - surface forcing (W/m2) |
---|
| 366 | & ( surf_tot * kts * rn_Dt ) ) |
---|
| 367 | CALL iom_put( 'bgfrcsal' , frc_s(tile_b) * 1.e-9 ) ! sc - surface forcing (psu*km3) |
---|
| 368 | IF( .NOT. ln_linssh ) THEN |
---|
| 369 | CALL iom_put( 'bgtemper' , zdiff_hc(tile_b) / zvol_tot(tile_b) ) ! Temperature drift (C) |
---|
| 370 | CALL iom_put( 'bgsaline' , zdiff_sc(tile_b) / zvol_tot(tile_b) ) ! Salinity drift (PSU) |
---|
| 371 | CALL iom_put( 'bgheatco' , zdiff_hc(tile_b) * 1.e-20 * rho0 * rcp ) ! Heat content drift (1.e20 J) |
---|
| 372 | CALL iom_put( 'bgheatfx' , zdiff_hc(tile_b) * rho0 * rcp / & ! Heat flux drift (W/m2) |
---|
| 373 | & ( surf_tot * kts * rn_Dt ) ) |
---|
| 374 | CALL iom_put( 'bgsaltco' , zdiff_sc(tile_b) * 1.e-9 ) ! Salt content drift (psu*km3) |
---|
| 375 | CALL iom_put( 'bgvolssh' , zdiff_v1(tile_b) * 1.e-9 ) ! volume ssh drift (km3) |
---|
| 376 | CALL iom_put( 'bgvole3t' , zdiff_v2(tile_b) * 1.e-9 ) ! volume e3t drift (km3) |
---|
| 377 | ! |
---|
| 378 | ! IF( lwp ) THEN |
---|
| 379 | ! WRITE(numout,*) |
---|
| 380 | ! WRITE(numout,*) 'dia_hsb : last time step hsb diagnostics: at it= ', kt,' date= ', ndastp |
---|
| 381 | ! WRITE(numout,*) '~~~~~~~' |
---|
| 382 | ! WRITE(numout,*) ' Temperature drift = ', zdiff_hc(tile_b) / zvol_tot(tile_b), ' C' |
---|
| 383 | ! WRITE(numout,*) ' Salinity12 drift = ', zdiff_sc(tile_b) / zvol_tot(tile_b), ' PSU' |
---|
| 384 | ! WRITE(numout,*) ' volume ssh drift = ', zdiff_v1(tile_b) * 1.e-9 , ' km^3' |
---|
| 385 | ! WRITE(numout,*) ' volume e3t drift = ', zdiff_v2(tile_b) * 1.e-9 , ' km^3' |
---|
| 386 | ! ENDIF |
---|
| 387 | ELSE |
---|
| 388 | CALL iom_put( 'bgtemper' , zdiff_hc1(tile_b) / zvol_tot(tile_b)) ! Heat content drift (C) |
---|
| 389 | CALL iom_put( 'bgsaline' , zdiff_sc1(tile_b) / zvol_tot(tile_b)) ! Salt content drift (PSU) |
---|
| 390 | CALL iom_put( 'bgheatco' , zdiff_hc1(tile_b) * 1.e-20 * rho0 * rcp ) ! Heat content drift (1.e20 J) |
---|
| 391 | CALL iom_put( 'bgheatfx' , zdiff_hc1(tile_b) * rho0 * rcp / & ! Heat flux drift (W/m2) |
---|
| 392 | & ( surf_tot * kts * rn_Dt ) ) |
---|
| 393 | CALL iom_put( 'bgsaltco' , zdiff_sc1(tile_b) * 1.e-9 ) ! Salt content drift (psu*km3) |
---|
| 394 | CALL iom_put( 'bgvolssh' , zdiff_v1(tile_b) * 1.e-9 ) ! volume ssh drift (km3) |
---|
| 395 | CALL iom_put( 'bgmistem' , zerr_hc1(tile_b) / zvol_tot(tile_b) ) ! hc - error due to free surface (C) |
---|
| 396 | CALL iom_put( 'bgmissal' , zerr_sc1(tile_b) / zvol_tot(tile_b) ) ! sc - error due to free surface (psu) |
---|
| 397 | ENDIF |
---|
| 398 | ! |
---|
| 399 | IF( lrst_oce ) CALL dia_hsb_rst( kts, Kmm, tile_n, 'WRITE' ) |
---|
| 400 | ! |
---|
| 401 | END IF |
---|
| 402 | IF ( kt == nitend ) THEN |
---|
| 403 | |
---|
| 404 | zdiff_v1(tile_n) = zdiff_v1(tile_n) - frc_v(tile_n) |
---|
| 405 | IF( .NOT.ln_linssh ) zdiff_v2(tile_n) = zdiff_v2(tile_n) - frc_v(tile_n) |
---|
| 406 | zdiff_hc(tile_n) = zdiff_hc(tile_n) - frc_t(tile_n) |
---|
| 407 | zdiff_sc(tile_n) = zdiff_sc(tile_n) - frc_s(tile_n) |
---|
| 408 | IF( ln_linssh ) THEN |
---|
| 409 | zdiff_hc1(tile_n) = zdiff_hc (tile_n) + z_ssh_hc(tile_n) |
---|
| 410 | zdiff_sc1(tile_n) = zdiff_sc (tile_n) + z_ssh_sc(tile_n) |
---|
| 411 | zerr_hc1 (tile_n) = z_ssh_hc(tile_n) - frc_wn_t(tile_n) |
---|
| 412 | zerr_sc1 (tile_n) = z_ssh_sc(tile_n) - frc_wn_s(tile_n) |
---|
| 413 | ENDIF |
---|
| 414 | !!gm to be added ? |
---|
| 415 | ! IF( ln_linssh ) THEN ! fixed volume, add the ssh contribution |
---|
| 416 | ! zvol_tot = zvol_tot + glob_sum( 'diahsb', surf(:,:) * sshn(:,:) ) |
---|
| 417 | ! ENDIF |
---|
| 418 | !!gm end |
---|
| 419 | |
---|
| 420 | CALL iom_put( 'bgfrcvol' , frc_v(tile_n) * 1.e-9 ) ! vol - surface forcing (km3) |
---|
| 421 | CALL iom_put( 'bgfrctem' , frc_t(tile_n) * rho0 * rcp * 1.e-20 ) ! hc - surface forcing (1.e20 J) |
---|
| 422 | CALL iom_put( 'bgfrchfx' , frc_t(tile_n) * rho0 * rcp / & ! hc - surface forcing (W/m2) |
---|
| 423 | & ( surf_tot * kt * rn_Dt ) ) |
---|
| 424 | CALL iom_put( 'bgfrcsal' , frc_s(tile_n) * 1.e-9 ) ! sc - surface forcing (psu*km3) |
---|
| 425 | IF( .NOT. ln_linssh ) THEN |
---|
| 426 | CALL iom_put( 'bgtemper' , zdiff_hc(tile_n) / zvol_tot(tile_n) ) ! Temperature drift (C) |
---|
| 427 | CALL iom_put( 'bgsaline' , zdiff_sc(tile_n) / zvol_tot(tile_n) ) ! Salinity drift (PSU) |
---|
| 428 | CALL iom_put( 'bgheatco' , zdiff_hc(tile_n) * 1.e-20 * rho0 * rcp ) ! Heat content drift (1.e20 J) |
---|
| 429 | CALL iom_put( 'bgheatfx' , zdiff_hc(tile_n) * rho0 * rcp / & ! Heat flux drift (W/m2) |
---|
| 430 | & ( surf_tot * kt * rn_Dt ) ) |
---|
| 431 | CALL iom_put( 'bgsaltco' , zdiff_sc(tile_n) * 1.e-9 ) ! Salt content drift (psu*km3) |
---|
| 432 | CALL iom_put( 'bgvolssh' , zdiff_v1(tile_n) * 1.e-9 ) ! volume ssh drift (km3) |
---|
| 433 | CALL iom_put( 'bgvole3t' , zdiff_v2(tile_n) * 1.e-9 ) ! volume e3t drift (km3) |
---|
| 434 | ! |
---|
| 435 | IF( kt == nitend .AND. lwp ) THEN |
---|
| 436 | WRITE(numout,*) |
---|
| 437 | WRITE(numout,*) 'dia_hsb : last time step hsb diagnostics: at it= ', kt,' date= ', ndastp |
---|
| 438 | WRITE(numout,*) '~~~~~~~' |
---|
| 439 | WRITE(numout,*) ' Temperature drift = ', zdiff_hc(tile_n) / zvol_tot(tile_n), ' C' |
---|
| 440 | WRITE(numout,*) ' Salinity drift = ', zdiff_sc(tile_n) / zvol_tot(tile_n), ' PSU' |
---|
| 441 | WRITE(numout,*) ' volume ssh drift = ', zdiff_v1(tile_n) * 1.e-9 , ' km^3' |
---|
| 442 | WRITE(numout,*) ' volume e3t drift = ', zdiff_v2(tile_n) * 1.e-9 , ' km^3' |
---|
| 443 | ! |
---|
| 444 | ENDIF |
---|
| 445 | ! |
---|
| 446 | ELSE |
---|
| 447 | CALL iom_put( 'bgtemper' , zdiff_hc1(tile_n) / zvol_tot(tile_n)) ! Heat content drift (C) |
---|
| 448 | CALL iom_put( 'bgsaline' , zdiff_sc1(tile_n) / zvol_tot(tile_n)) ! Salt content drift (PSU) |
---|
| 449 | CALL iom_put( 'bgheatco' , zdiff_hc1(tile_n) * 1.e-20 * rho0 * rcp ) ! Heat content drift (1.e20 J) |
---|
| 450 | CALL iom_put( 'bgheatfx' , zdiff_hc1(tile_n) * rho0 * rcp / & ! Heat flux drift (W/m2) |
---|
| 451 | & ( surf_tot * kt * rn_Dt ) ) |
---|
| 452 | CALL iom_put( 'bgsaltco' , zdiff_sc1(tile_n) * 1.e-9 ) ! Salt content drift (psu*km3) |
---|
| 453 | CALL iom_put( 'bgvolssh' , zdiff_v1(tile_n) * 1.e-9 ) ! volume ssh drift (km3) |
---|
| 454 | CALL iom_put( 'bgmistem' , zerr_hc1(tile_n) / zvol_tot(tile_n) ) ! hc - error due to free surface (C) |
---|
| 455 | CALL iom_put( 'bgmissal' , zerr_sc1(tile_n) / zvol_tot(tile_n) ) ! sc - error due to free surface (psu) |
---|
| 456 | ENDIF |
---|
| 457 | ! |
---|
| 458 | !Last step, don't need restart |
---|
| 459 | !IF( lrst_oce ) CALL dia_hsb_rst( kts, Kmm, tile_n, 'WRITE' ) |
---|
| 460 | ! |
---|
| 461 | END IF |
---|
| 462 | IF( ln_timing ) CALL timing_stop('dia_hsb') |
---|
| 463 | ! |
---|
| 464 | END SUBROUTINE dia_hsb |
---|
| 465 | |
---|
| 466 | |
---|
| 467 | SUBROUTINE dia_hsb_rst( kt, Kmm, tile, cdrw ) |
---|
| 468 | |
---|
| 469 | |
---|
| 470 | |
---|
| 471 | !!--------------------------------------------------------------------- |
---|
| 472 | !! *** ROUTINE dia_hsb_rst *** |
---|
| 473 | !! |
---|
| 474 | !! ** Purpose : Read or write DIA file in restart file |
---|
| 475 | !! |
---|
| 476 | !! ** Method : use of IOM library |
---|
| 477 | !!---------------------------------------------------------------------- |
---|
| 478 | INTEGER , INTENT(in) :: kt ! ocean time-step |
---|
| 479 | INTEGER , INTENT(in) :: Kmm ! ocean time level index |
---|
| 480 | |
---|
| 481 | INTEGER , INTENT(in) :: tile ! host tile |
---|
| 482 | |
---|
| 483 | CHARACTER(len=*), INTENT(in) :: cdrw ! "READ"/"WRITE" flag |
---|
| 484 | ! |
---|
| 485 | INTEGER :: ji, jj, jk ! dummy loop indices |
---|
| 486 | !!---------------------------------------------------------------------- |
---|
| 487 | ! |
---|
| 488 | IF( TRIM(cdrw) == 'READ' ) THEN ! Read/initialise |
---|
| 489 | IF( ln_rstart ) THEN !* Read the restart file |
---|
| 490 | ! |
---|
| 491 | IF(lwp) WRITE(numout,*) |
---|
| 492 | IF(lwp) WRITE(numout,*) ' dia_hsb_rst : read hsb restart at it= ', kt,' date= ', ndastp |
---|
| 493 | IF(lwp) WRITE(numout,*) |
---|
| 494 | |
---|
| 495 | CALL iom_get( numror, 'frc_v', frc_v(tile), ldxios = lrxios ) |
---|
| 496 | CALL iom_get( numror, 'frc_t', frc_t(tile), ldxios = lrxios ) |
---|
| 497 | CALL iom_get( numror, 'frc_s', frc_s(tile), ldxios = lrxios ) |
---|
| 498 | IF( ln_linssh ) THEN |
---|
| 499 | CALL iom_get( numror, 'frc_wn_t', frc_wn_t(tile), ldxios = lrxios ) |
---|
| 500 | CALL iom_get( numror, 'frc_wn_s', frc_wn_s(tile), ldxios = lrxios ) |
---|
| 501 | ENDIF |
---|
| 502 | CALL iom_get( numror, jpdom_auto, 'surf_ini' , surf_ini , ldxios = lrxios ) ! ice sheet coupling |
---|
| 503 | CALL iom_get( numror, jpdom_auto, 'ssh_ini' , ssh_ini , ldxios = lrxios ) |
---|
| 504 | CALL iom_get( numror, jpdom_auto, 'e3t_ini' , e3t_ini , ldxios = lrxios ) |
---|
| 505 | CALL iom_get( numror, jpdom_auto, 'tmask_ini' , tmask_ini , ldxios = lrxios ) |
---|
| 506 | CALL iom_get( numror, jpdom_auto, 'hc_loc_ini', hc_loc_ini, ldxios = lrxios ) |
---|
| 507 | CALL iom_get( numror, jpdom_auto, 'sc_loc_ini', sc_loc_ini, ldxios = lrxios ) |
---|
| 508 | IF( ln_linssh ) THEN |
---|
| 509 | CALL iom_get( numror, jpdom_auto, 'ssh_hc_loc_ini', ssh_hc_loc_ini, ldxios = lrxios ) |
---|
| 510 | CALL iom_get( numror, jpdom_auto, 'ssh_sc_loc_ini', ssh_sc_loc_ini, ldxios = lrxios ) |
---|
| 511 | ENDIF |
---|
| 512 | ELSE |
---|
| 513 | IF(lwp) WRITE(numout,*) |
---|
| 514 | IF(lwp) WRITE(numout,*) ' dia_hsb_rst : initialise hsb at initial state ' |
---|
| 515 | IF(lwp) WRITE(numout,*) |
---|
| 516 | surf_ini(:,:) = e1e2t(:,:) * tmask_i(:,:) ! initial ocean surface |
---|
| 517 | ssh_ini(:,:) = ssh(:,:,Kmm) ! initial ssh |
---|
| 518 | DO jk = 1, jpk |
---|
| 519 | ! if ice sheet/oceqn coupling, need to mask ini variables here (mask could change at the next NEMO instance). |
---|
| 520 | e3t_ini (:,:,jk) = e3t(:,:,jk,Kmm) * tmask(:,:,jk) ! initial vertical scale factors |
---|
| 521 | tmask_ini (:,:,jk) = tmask(:,:,jk) ! initial mask |
---|
| 522 | hc_loc_ini(:,:,jk) = ts(:,:,jk,jp_tem,Kmm) * e3t(:,:,jk,Kmm) * tmask(:,:,jk) ! initial heat content |
---|
| 523 | sc_loc_ini(:,:,jk) = ts(:,:,jk,jp_sal,Kmm) * e3t(:,:,jk,Kmm) * tmask(:,:,jk) ! initial salt content |
---|
| 524 | END DO |
---|
| 525 | |
---|
| 526 | d_surf_ini = surf_ini |
---|
| 527 | d_e3t_ini = e3t_ini |
---|
| 528 | d_tmask_ini = tmask_ini |
---|
| 529 | d_hc_loc_ini = hc_loc_ini |
---|
| 530 | d_sc_loc_ini = sc_loc_ini |
---|
| 531 | frc_v(tile) = 0._wp ! volume trend due to forcing |
---|
| 532 | frc_t(tile) = 0._wp ! heat content - - - - |
---|
| 533 | frc_s(tile) = 0._wp ! salt content - - - - |
---|
| 534 | |
---|
| 535 | |
---|
| 536 | |
---|
| 537 | |
---|
| 538 | |
---|
| 539 | IF( ln_linssh ) THEN |
---|
| 540 | IF( ln_isfcav ) THEN |
---|
| 541 | DO ji = 1, jpi |
---|
| 542 | DO jj = 1, jpj |
---|
| 543 | ssh_hc_loc_ini(ji,jj) = ts(ji,jj,mikt(ji,jj),jp_tem,Kmm) * ssh(ji,jj,Kmm) ! initial heat content in ssh |
---|
| 544 | ssh_sc_loc_ini(ji,jj) = ts(ji,jj,mikt(ji,jj),jp_sal,Kmm) * ssh(ji,jj,Kmm) ! initial salt content in ssh |
---|
| 545 | END DO |
---|
| 546 | END DO |
---|
| 547 | ELSE |
---|
| 548 | ssh_hc_loc_ini(:,:) = ts(:,:,1,jp_tem,Kmm) * ssh(:,:,Kmm) ! initial heat content in ssh |
---|
| 549 | ssh_sc_loc_ini(:,:) = ts(:,:,1,jp_sal,Kmm) * ssh(:,:,Kmm) ! initial salt content in ssh |
---|
| 550 | END IF |
---|
| 551 | |
---|
| 552 | frc_wn_t(tile) = 0._wp ! initial heat content misfit due to free surface |
---|
| 553 | frc_wn_s(tile) = 0._wp ! initial salt content misfit due to free surface |
---|
| 554 | |
---|
| 555 | |
---|
| 556 | |
---|
| 557 | |
---|
| 558 | ENDIF |
---|
| 559 | ENDIF |
---|
| 560 | ! |
---|
| 561 | ELSEIF( TRIM(cdrw) == 'WRITE' ) THEN ! Create restart file |
---|
| 562 | ! ! ------------------- |
---|
| 563 | IF(lwp) WRITE(numout,*) |
---|
| 564 | IF(lwp) WRITE(numout,*) ' dia_hsb_rst : write restart at it= ', kt,' date= ', ndastp |
---|
| 565 | IF(lwp) WRITE(numout,*) |
---|
| 566 | ! |
---|
| 567 | |
---|
| 568 | IF( lwxios ) CALL iom_swap( cwxios_context ) |
---|
| 569 | CALL iom_rstput( kt, nitrst, numrow, 'frc_v', frc_v(tile), ldxios = lwxios ) |
---|
| 570 | CALL iom_rstput( kt, nitrst, numrow, 'frc_t', frc_t(tile), ldxios = lwxios ) |
---|
| 571 | CALL iom_rstput( kt, nitrst, numrow, 'frc_s', frc_s(tile), ldxios = lwxios ) |
---|
| 572 | IF( ln_linssh ) THEN |
---|
| 573 | CALL iom_rstput( kt, nitrst, numrow, 'frc_wn_t', frc_wn_t(tile), ldxios = lwxios ) |
---|
| 574 | CALL iom_rstput( kt, nitrst, numrow, 'frc_wn_s', frc_wn_s(tile), ldxios = lwxios ) |
---|
| 575 | ENDIF |
---|
| 576 | CALL iom_rstput( kt, nitrst, numrow, 'surf_ini' , surf_ini , ldxios = lwxios ) ! ice sheet coupling |
---|
| 577 | CALL iom_rstput( kt, nitrst, numrow, 'ssh_ini' , ssh_ini , ldxios = lwxios ) |
---|
| 578 | CALL iom_rstput( kt, nitrst, numrow, 'e3t_ini' , e3t_ini , ldxios = lwxios ) |
---|
| 579 | CALL iom_rstput( kt, nitrst, numrow, 'tmask_ini' , tmask_ini , ldxios = lwxios ) |
---|
| 580 | CALL iom_rstput( kt, nitrst, numrow, 'hc_loc_ini', hc_loc_ini, ldxios = lwxios ) |
---|
| 581 | CALL iom_rstput( kt, nitrst, numrow, 'sc_loc_ini', sc_loc_ini, ldxios = lwxios ) |
---|
| 582 | IF( ln_linssh ) THEN |
---|
| 583 | CALL iom_rstput( kt, nitrst, numrow, 'ssh_hc_loc_ini', ssh_hc_loc_ini, ldxios = lwxios ) |
---|
| 584 | CALL iom_rstput( kt, nitrst, numrow, 'ssh_sc_loc_ini', ssh_sc_loc_ini, ldxios = lwxios ) |
---|
| 585 | ENDIF |
---|
| 586 | IF( lwxios ) CALL iom_swap( cxios_context ) |
---|
| 587 | ! |
---|
| 588 | ENDIF |
---|
| 589 | ! |
---|
| 590 | END SUBROUTINE dia_hsb_rst |
---|
| 591 | |
---|
| 592 | |
---|
| 593 | SUBROUTINE dia_hsb_init( Kmm ) |
---|
| 594 | !!--------------------------------------------------------------------------- |
---|
| 595 | !! *** ROUTINE dia_hsb *** |
---|
| 596 | !! |
---|
| 597 | !! ** Purpose: Initialization for the heat salt volume budgets |
---|
| 598 | !! |
---|
| 599 | !! ** Method : Compute initial heat content, salt content and volume |
---|
| 600 | !! |
---|
| 601 | !! ** Action : - Compute initial heat content, salt content and volume |
---|
| 602 | !! - Initialize forcing trends |
---|
| 603 | !! - Compute coefficients for conversion |
---|
| 604 | !!--------------------------------------------------------------------------- |
---|
| 605 | INTEGER, INTENT(in) :: Kmm ! time level index |
---|
| 606 | ! |
---|
| 607 | INTEGER :: ierror, ios ! local integer |
---|
| 608 | |
---|
| 609 | INTEGER :: i, istat ! local integer |
---|
| 610 | |
---|
| 611 | !! |
---|
| 612 | NAMELIST/namhsb/ ln_diahsb |
---|
| 613 | !!---------------------------------------------------------------------- |
---|
| 614 | ! |
---|
| 615 | IF(lwp) THEN |
---|
| 616 | WRITE(numout,*) |
---|
| 617 | WRITE(numout,*) 'dia_hsb_init : heat and salt budgets diagnostics' |
---|
| 618 | WRITE(numout,*) '~~~~~~~~~~~~ ' |
---|
| 619 | ENDIF |
---|
| 620 | READ ( numnam_ref, namhsb, IOSTAT = ios, ERR = 901) |
---|
| 621 | 901 IF( ios /= 0 ) CALL ctl_nam ( ios , 'namhsb in reference namelist' ) |
---|
| 622 | READ ( numnam_cfg, namhsb, IOSTAT = ios, ERR = 902 ) |
---|
| 623 | 902 IF( ios > 0 ) CALL ctl_nam ( ios , 'namhsb in configuration namelist' ) |
---|
| 624 | |
---|
| 625 | IF(lwm) WRITE( numond, namhsb ) |
---|
| 626 | |
---|
| 627 | IF(lwp) THEN |
---|
| 628 | WRITE(numout,*) ' Namelist namhsb :' |
---|
| 629 | WRITE(numout,*) ' check the heat and salt budgets (T) or not (F) ln_diahsb = ', ln_diahsb |
---|
| 630 | ENDIF |
---|
| 631 | ! |
---|
| 632 | IF( .NOT. ln_diahsb ) RETURN |
---|
| 633 | |
---|
| 634 | IF(lwxios) THEN |
---|
| 635 | ! define variables in restart file when writing with XIOS |
---|
| 636 | CALL iom_set_rstw_var_active('frc_v') |
---|
| 637 | CALL iom_set_rstw_var_active('frc_t') |
---|
| 638 | CALL iom_set_rstw_var_active('frc_s') |
---|
| 639 | CALL iom_set_rstw_var_active('surf_ini') |
---|
| 640 | CALL iom_set_rstw_var_active('ssh_ini') |
---|
| 641 | CALL iom_set_rstw_var_active('e3t_ini') |
---|
| 642 | CALL iom_set_rstw_var_active('hc_loc_ini') |
---|
| 643 | CALL iom_set_rstw_var_active('sc_loc_ini') |
---|
| 644 | IF( ln_linssh ) THEN |
---|
| 645 | CALL iom_set_rstw_var_active('ssh_hc_loc_ini') |
---|
| 646 | CALL iom_set_rstw_var_active('ssh_sc_loc_ini') |
---|
| 647 | CALL iom_set_rstw_var_active('frc_wn_t') |
---|
| 648 | CALL iom_set_rstw_var_active('frc_wn_s') |
---|
| 649 | ENDIF |
---|
| 650 | ENDIF |
---|
| 651 | ! ------------------- ! |
---|
| 652 | ! 1 - Allocate memory ! |
---|
| 653 | ! ------------------- ! |
---|
| 654 | |
---|
| 655 | CALL setdevice() |
---|
| 656 | !Device data associate to PUBLIC arrays |
---|
| 657 | ALLOCATE(d_e3t (jpi,jpj,jpk,jpt) ) ! |
---|
| 658 | ALLOCATE(d_tmask (jpi,jpj,jpk) ) ! |
---|
| 659 | ALLOCATE(d_tmask_ini (jpi,jpj,jpk) ) ! |
---|
| 660 | ALLOCATE(d_tmask_h (jpi,jpj) ) ! |
---|
| 661 | ALLOCATE(d_ts (jpi,jpj,jpk,2,jpj) ) ! |
---|
| 662 | !Device data associate to LOCAL/DEVICE arrays ! |
---|
| 663 | ALLOCATE(d_surf (jpi,jpj) ) ! |
---|
| 664 | ALLOCATE(d_surf_ini (jpi,jpj) ) ! |
---|
| 665 | ALLOCATE(d_hc_loc_ini (jpi,jpj,jpk) ) ! |
---|
| 666 | ALLOCATE(d_sc_loc_ini (jpi,jpj,jpk) ) ! |
---|
| 667 | ALLOCATE(d_e3t_ini (jpi,jpj,jpk) ) ! |
---|
| 668 | ALLOCATE(d_zwrkv (jpi,jpj,jpkm1) ) ! |
---|
| 669 | ALLOCATE(d_zwrkh (jpi,jpj,jpkm1) ) ! |
---|
| 670 | ALLOCATE(d_zwrks (jpi,jpj,jpkm1) ) ! |
---|
| 671 | ALLOCATE(d_zwrk (jpi,jpj,jpkm1) ) ! |
---|
| 672 | ALLOCATE(h_ztmpv(2),h_ztmph(2),h_ztmps(2),h_ztmp(2)) ! |
---|
| 673 | |
---|
| 674 | DO i = 1, nstreams !Create Streams |
---|
| 675 | istat = cudaStreamCreate(stream(i)) |
---|
| 676 | IF( istat /= 0 ) THEN |
---|
| 677 | CALL ctl_stop( 'dia_hsb_init: error in Stream creation' ) ; RETURN |
---|
| 678 | ENDIF |
---|
| 679 | END DO |
---|
| 680 | ! |
---|
| 681 | !Pinned reallocation step non constant |
---|
| 682 | istat = cudaHostRegister(C_LOC(ts ), sizeof(ts ), cudaHostRegisterMapped) |
---|
| 683 | istat = cudaHostRegister(C_LOC(e3t), sizeof(e3t), cudaHostRegisterMapped) |
---|
| 684 | IF( istat /= 0 ) THEN |
---|
| 685 | CALL ctl_stop( 'dia_hsb_init: unable to pin host memory to GPU' ) ; RETURN |
---|
| 686 | ENDIF |
---|
| 687 | |
---|
| 688 | ALLOCATE( hc_loc_ini(jpi,jpj,jpk), sc_loc_ini(jpi,jpj,jpk), surf_ini(jpi,jpj), & |
---|
| 689 | & e3t_ini(jpi,jpj,jpk), surf(jpi,jpj), ssh_ini(jpi,jpj), tmask_ini(jpi,jpj,jpk),STAT=ierror ) |
---|
| 690 | |
---|
| 691 | IF( ierror > 0 ) THEN |
---|
| 692 | CALL ctl_stop( 'dia_hsb_init: unable to allocate hc_loc_ini' ) ; RETURN |
---|
| 693 | ENDIF |
---|
| 694 | |
---|
| 695 | IF( ln_linssh ) ALLOCATE( ssh_hc_loc_ini(jpi,jpj), ssh_sc_loc_ini(jpi,jpj),STAT=ierror ) |
---|
| 696 | IF( ierror > 0 ) THEN |
---|
| 697 | CALL ctl_stop( 'dia_hsb: unable to allocate ssh_hc_loc_ini' ) ; RETURN |
---|
| 698 | ENDIF |
---|
| 699 | |
---|
| 700 | ! ----------------------------------------------- ! |
---|
| 701 | ! 2 - Time independant variables and file opening ! |
---|
| 702 | ! ----------------------------------------------- ! |
---|
| 703 | surf(:,:) = e1e2t(:,:) * tmask_i(:,:) ! masked surface grid cell area |
---|
| 704 | surf_tot = glob_sum( 'diahsb', surf(:,:) ) ! total ocean surface area |
---|
| 705 | |
---|
| 706 | d_surf = surf |
---|
| 707 | d_surf_ini = surf_ini |
---|
| 708 | d_e3t_ini = e3t_ini |
---|
| 709 | d_tmask = tmask |
---|
| 710 | d_tmask_ini = tmask_ini |
---|
| 711 | d_tmask_h = tmask_h |
---|
| 712 | h_ztmp = 0.0 |
---|
| 713 | |
---|
| 714 | IF( ln_bdy ) CALL ctl_warn( 'dia_hsb_init: heat/salt budget does not consider open boundary fluxes' ) |
---|
| 715 | ! |
---|
| 716 | ! ---------------------------------- ! |
---|
| 717 | ! 4 - initial conservation variables ! |
---|
| 718 | ! ---------------------------------- ! |
---|
| 719 | |
---|
| 720 | CALL dia_hsb_rst( nit000, Kmm, 1, 'READ' ) !* read or initialize all required files |
---|
| 721 | |
---|
| 722 | |
---|
| 723 | |
---|
| 724 | ! |
---|
| 725 | END SUBROUTINE dia_hsb_init |
---|
| 726 | |
---|
| 727 | !!====================================================================== |
---|
| 728 | END MODULE diahsb |
---|