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