New URL for NEMO forge!   http://forge.nemo-ocean.eu

Since March 2022 along with NEMO 4.2 release, the code development moved to a self-hosted GitLab.
This present forge is now archived and remained online for history.
diahsb_gpu.h90 in NEMO/branches/2020/dev_r13747_HPC-10_HPDAonline_DiagGPU/tests/DIA_GPU/MY_SRC – NEMO

source: NEMO/branches/2020/dev_r13747_HPC-10_HPDAonline_DiagGPU/tests/DIA_GPU/MY_SRC/diahsb_gpu.h90 @ 14091

Last change on this file since 14091 was 14091, checked in by mcastril, 4 years ago

Add DIA_GPU test case
File size: 32.8 KB
Line 
1MODULE 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   !!----------------------------------------------------------------------
102CONTAINS
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)
621901 IF( ios /= 0 ) CALL ctl_nam ( ios , 'namhsb in reference namelist' )
622      READ ( numnam_cfg, namhsb, IOSTAT = ios, ERR = 902 )
623902 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   !!======================================================================
728END MODULE diahsb
Note: See TracBrowser for help on using the repository browser.