Skip to content

Commit

Permalink
CUDA C SCC-HOIST: Loki generated (slightly manually adapted and fixed…
Browse files Browse the repository at this point in the history 
…) and copied to ecwam/src
  • Loading branch information
@MichaelSt98
MichaelSt98 committed Mar 21, 2024
1 parent 0f522dd commit ffdf4ca
Show file tree
Hide file tree
Showing 192 changed files with 21,220 additions and 0 deletions.
193 changes: 193 additions & 0 deletions src/phys-scc-cuda/airsea.c_hoist.F90
View file
Original file line number Diff line number Diff line change
@@ -0,0 +1,193 @@
! (C) Copyright 1989- ECMWF.
!
! This software is licensed under the terms of the Apache Licence Version 2.0
! which can be obtained at http://www.apache.org/licenses/LICENSE-2.0.
! In applying this licence, ECMWF does not waive the privileges and immunities
! granted to it by virtue of its status as an intergovernmental organisation
! nor does it submit to any jurisdiction.
!
ATTRIBUTES(DEVICE) SUBROUTINE AIRSEA_FC (KIJS, KIJL, HALP, U10, U10DIR, TAUW, TAUWDIR, RNFAC, US, Z0, Z0B, CHRNCK, ICODE_WND, &
& IUSFG, ACD, ALPHA, ALPHAMAX, ALPHAMIN, ANG_GC_A, ANG_GC_B, ANG_GC_C, BCD, BETAMAXOXKAPPA2, BMAXOKAP, C2OSQRTVG_GC, CDMAX, &
& CHNKMIN_U, CM_GC, DELKCC_GC_NS, DELKCC_OMXKM3_GC, EPS1, EPSMIN, EPSUS, G, GM1, LLCAPCHNK, LLGCBZ0, LLNORMAGAM, NWAV_GC, &
& OM3GMKM_GC, OMXKM3_GC, RN1_RN, RNU, RNUM, SQRTGOSURFT, WSPMIN, XKAPPA, XKMSQRTVGOC2_GC, XKM_GC, XK_GC, XLOGKRATIOM1_GC, XNLEV, &
& ZALP, ICHNK, NCHNK, IJ)

! ----------------------------------------------------------------------

!**** *AIRSEA* - DETERMINE TOTAL STRESS IN SURFACE LAYER.

! P.A.E.M. JANSSEN KNMI AUGUST 1990
! JEAN BIDLOT ECMWF FEBRUARY 1999 : TAUT is already
! SQRT(TAUT)
! JEAN BIDLOT ECMWF OCTOBER 2004: QUADRATIC STEP FOR
! TAUW

!* PURPOSE.
! --------

! COMPUTE TOTAL STRESS.

!** INTERFACE.
! ----------

! *CALL* *AIRSEA (KIJS, KIJL, FL1, WAVNUM,
! HALP, U10, U10DIR, TAUW, TAUWDIR, RNFAC,
! US, Z0, Z0B, CHRNCK, ICODE_WND, IUSFG)*

! *KIJS* - INDEX OF FIRST GRIDPOINT.
! *KIJL* - INDEX OF LAST GRIDPOINT.
! *FL1* - SPECTRA
! *WAVNUM* - WAVE NUMBER
! *HALP* - 1/2 PHILLIPS PARAMETER
! *U10* - WINDSPEED U10.
! *U10DIR* - WINDSPEED DIRECTION.
! *TAUW* - WAVE STRESS.
! *TAUWDIR* - WAVE STRESS DIRECTION.
! *RNFAC* - WIND DEPENDENT FACTOR USED IN THE GROWTH RENORMALISATION.
! *US* - OUTPUT OR OUTPUT BLOCK OF FRICTION VELOCITY.
! *Z0* - OUTPUT BLOCK OF ROUGHNESS LENGTH.
! *Z0B* - BACKGROUND ROUGHNESS LENGTH.
! *CHRNCK* - CHARNOCK COEFFICIENT
! *ICODE_WND* SPECIFIES WHICH OF U10 OR US HAS BEEN FILED UPDATED:
! U10: ICODE_WND=3 --> US will be updated
! US: ICODE_WND=1 OR 2 --> U10 will be updated
! *IUSFG* - IF = 1 THEN USE THE FRICTION VELOCITY (US) AS FIRST GUESS in TAUT_Z0
! 0 DO NOT USE THE FIELD US


! ----------------------------------------------------------------------

USE PARKIND_WAVE, ONLY: JWRU, JWIM, JWRB

USE YOWPARAM, ONLY: NFRE, NANG
USE YOWTEST, ONLY: IU06


! ----------------------------------------------------------------------
IMPLICIT NONE
INTERFACE
SUBROUTINE TAUT_Z0_FC (KIJS, KIJL, IUSFG, HALP, UTOP, UDIR, TAUW, TAUWDIR, RNFAC, USTAR, Z0, Z0B, CHRNCK)
USE parkind_wave, ONLY: jwim, jwrb
INTEGER(KIND=JWIM), INTENT(IN) :: KIJS, KIJL, IUSFG
REAL(KIND=JWRB), INTENT(IN), DIMENSION(KIJL) :: HALP, UTOP, UDIR, TAUW, TAUWDIR, RNFAC
REAL(KIND=JWRB), INTENT(INOUT), DIMENSION(KIJL) :: USTAR
REAL(KIND=JWRB), INTENT(OUT), DIMENSION(KIJL) :: Z0, Z0B, CHRNCK
END SUBROUTINE TAUT_Z0_FC
END INTERFACE
INTERFACE
SUBROUTINE Z0WAVE_FC (KIJS, KIJL, US, TAUW, UTOP, Z0, Z0B, CHRNCK)
USE parkind_wave, ONLY: jwim, jwrb
INTEGER(KIND=JWIM), INTENT(IN) :: KIJS, KIJL
REAL(KIND=JWRB), INTENT(IN), DIMENSION(KIJL) :: US, TAUW, UTOP
REAL(KIND=JWRB), INTENT(OUT), DIMENSION(KIJL) :: Z0, Z0B, CHRNCK
END SUBROUTINE Z0WAVE_FC
END INTERFACE
INTEGER(KIND=JWIM), VALUE, INTENT(IN) :: KIJS
INTEGER(KIND=JWIM), VALUE, INTENT(IN) :: KIJL
INTEGER(KIND=JWIM), VALUE, INTENT(IN) :: ICODE_WND
INTEGER(KIND=JWIM), VALUE, INTENT(IN) :: IUSFG
REAL(KIND=JWRB), TARGET, INTENT(IN) :: HALP(:)
REAL(KIND=JWRB), TARGET, INTENT(IN) :: RNFAC(:)
REAL(KIND=JWRB), TARGET, INTENT(IN) :: U10DIR(:, :)
REAL(KIND=JWRB), TARGET, INTENT(IN) :: TAUW(:, :)
REAL(KIND=JWRB), TARGET, INTENT(IN) :: TAUWDIR(:, :)
REAL(KIND=JWRB), TARGET, INTENT(INOUT) :: U10(:, :)
REAL(KIND=JWRB), TARGET, INTENT(INOUT) :: US(:, :)
REAL(KIND=JWRB), TARGET, INTENT(OUT) :: Z0(:, :)
REAL(KIND=JWRB), TARGET, INTENT(OUT) :: Z0B(:, :)
REAL(KIND=JWRB), TARGET, INTENT(OUT) :: CHRNCK(:, :)

INTEGER(KIND=JWIM), VALUE, INTENT(IN) :: IJ
INTEGER(KIND=JWIM) :: I
INTEGER(KIND=JWIM) :: J

REAL(KIND=JWRB) :: XI
REAL(KIND=JWRB) :: XJ
REAL(KIND=JWRB) :: DELI1
REAL(KIND=JWRB) :: DELI2
REAL(KIND=JWRB) :: DELJ1
REAL(KIND=JWRB) :: DELJ2
REAL(KIND=JWRB) :: UST2
REAL(KIND=JWRB) :: ARG
REAL(KIND=JWRB) :: SQRTCDM1
REAL(KIND=JWRB) :: XKAPPAD
REAL(KIND=JWRB) :: XLOGLEV
REAL(KIND=JWRB) :: XLEV
REAL(KIND=JWRB), VALUE, INTENT(IN) :: ACD
REAL(KIND=JWRB), VALUE, INTENT(IN) :: ALPHA
REAL(KIND=JWRB), VALUE, INTENT(IN) :: ALPHAMAX
REAL(KIND=JWRB), VALUE, INTENT(IN) :: ALPHAMIN
REAL(KIND=JWRB), VALUE, INTENT(IN) :: ANG_GC_A
REAL(KIND=JWRB), VALUE, INTENT(IN) :: ANG_GC_B
REAL(KIND=JWRB), VALUE, INTENT(IN) :: ANG_GC_C
REAL(KIND=JWRB), VALUE, INTENT(IN) :: BCD
REAL(KIND=JWRB), VALUE, INTENT(IN) :: BETAMAXOXKAPPA2
REAL(KIND=JWRB), VALUE, INTENT(IN) :: BMAXOKAP
REAL(KIND=JWRB), TARGET, INTENT(IN) :: C2OSQRTVG_GC(:)
REAL(KIND=JWRB), VALUE, INTENT(IN) :: CDMAX
REAL(KIND=JWRB), VALUE, INTENT(IN) :: CHNKMIN_U
REAL(KIND=JWRB), TARGET, INTENT(IN) :: CM_GC(:)
REAL(KIND=JWRB), TARGET, INTENT(IN) :: DELKCC_GC_NS(:)
REAL(KIND=JWRB), TARGET, INTENT(IN) :: DELKCC_OMXKM3_GC(:)
REAL(KIND=JWRB), VALUE, INTENT(IN) :: EPS1
REAL(KIND=JWRB), VALUE, INTENT(IN) :: EPSMIN
REAL(KIND=JWRB), VALUE, INTENT(IN) :: EPSUS
REAL(KIND=JWRB), VALUE, INTENT(IN) :: G
REAL(KIND=JWRB), VALUE, INTENT(IN) :: GM1
LOGICAL, VALUE, INTENT(IN) :: LLCAPCHNK
LOGICAL, VALUE, INTENT(IN) :: LLGCBZ0
LOGICAL, VALUE, INTENT(IN) :: LLNORMAGAM
INTEGER(KIND=JWIM), VALUE, INTENT(IN) :: NWAV_GC
REAL(KIND=JWRB), TARGET, INTENT(IN) :: OM3GMKM_GC(:)
REAL(KIND=JWRB), TARGET, INTENT(IN) :: OMXKM3_GC(:)
REAL(KIND=JWRB), VALUE, INTENT(IN) :: RN1_RN
REAL(KIND=JWRB), VALUE, INTENT(IN) :: RNU
REAL(KIND=JWRB), VALUE, INTENT(IN) :: RNUM
REAL(KIND=JWRB), VALUE, INTENT(IN) :: SQRTGOSURFT
REAL(KIND=JWRB), VALUE, INTENT(IN) :: WSPMIN
REAL(KIND=JWRB), VALUE, INTENT(IN) :: XKAPPA
REAL(KIND=JWRB), TARGET, INTENT(IN) :: XKMSQRTVGOC2_GC(:)
REAL(KIND=JWRB), TARGET, INTENT(IN) :: XKM_GC(:)
REAL(KIND=JWRB), TARGET, INTENT(IN) :: XK_GC(:)
REAL(KIND=JWRB), VALUE, INTENT(IN) :: XLOGKRATIOM1_GC
REAL(KIND=JWRB), VALUE, INTENT(IN) :: XNLEV
REAL(KIND=JWRB), VALUE, INTENT(IN) :: ZALP
INTEGER(KIND=JWIM), VALUE, INTENT(IN) :: ICHNK
INTEGER, VALUE, INTENT(IN) :: NCHNK

! ----------------------------------------------------------------------

!* 2. DETERMINE TOTAL STRESS (if needed)
! ----------------------------------

IF (ICODE_WND == 3) THEN

CALL TAUT_Z0_FC(KIJS, KIJL, IUSFG, HALP(:), U10(:, :), U10DIR(:, :), TAUW(:, :), TAUWDIR(:, :), RNFAC(:), US(:, :), &
& Z0(:, :), Z0B(:, :), CHRNCK(:, :), ACD, ALPHA, ALPHAMAX, ALPHAMIN, ANG_GC_A, ANG_GC_B, ANG_GC_C, BCD, BETAMAXOXKAPPA2, &
& BMAXOKAP, C2OSQRTVG_GC(:), CDMAX, CHNKMIN_U, CM_GC(:), DELKCC_GC_NS(:), DELKCC_OMXKM3_GC(:), EPS1, EPSMIN, EPSUS, G, GM1, &
& LLCAPCHNK, LLGCBZ0, LLNORMAGAM, NWAV_GC, OM3GMKM_GC(:), OMXKM3_GC(:), RN1_RN, RNU, RNUM, SQRTGOSURFT, XKAPPA, &
& XKMSQRTVGOC2_GC(:), XKM_GC(:), XK_GC(:), XLOGKRATIOM1_GC, XNLEV, ZALP, ICHNK, NCHNK, IJ)

ELSE IF (ICODE_WND == 1 .or. ICODE_WND == 2) THEN

!* 3. DETERMINE ROUGHNESS LENGTH (if needed).
! ---------------------------

CALL Z0WAVE_FC(KIJS, KIJL, US(:, :), TAUW(:, :), U10(:, :), Z0(:, :), Z0B(:, :), CHRNCK(:, :), ALPHA, ALPHAMIN, CHNKMIN_U, &
& EPS1, G, GM1, LLCAPCHNK, ICHNK, NCHNK, IJ)

!* 3. DETERMINE U10 (if needed).
! ---------------------------

XKAPPAD = 1.0_JWRB / XKAPPA
XLOGLEV = LOG(XNLEV)


U10(IJ, ICHNK) = XKAPPAD*US(IJ, ICHNK)*(XLOGLEV - LOG(Z0(IJ, ICHNK)))
U10(IJ, ICHNK) = MAX(U10(IJ, ICHNK), WSPMIN)


END IF


END SUBROUTINE AIRSEA_FC
69 changes: 69 additions & 0 deletions src/phys-scc-cuda/airsea_c.c
View file
Original file line number Diff line number Diff line change
@@ -0,0 +1,69 @@
#include <stdio.h>
#include <stdbool.h>
#include <float.h>
#include <math.h>
#include <cuda.h>
#include <cuda_runtime.h>
#include "airsea_c.h"
#include "z0wave_c.h"
#include "taut_z0_c.h"

__device__ void airsea_c(int kijs, int kijl, const double * halp, double * u10,
const double * u10dir, const double * tauw, const double * tauwdir,
const double * rnfac, double * us, double * z0, double * z0b, double * chrnck,
int icode_wnd, int iusfg, double acd, double alpha, double alphamax, double alphamin,
double ang_gc_a, double ang_gc_b, double ang_gc_c, double bcd, double betamaxoxkappa2,
double bmaxokap, const double * c2osqrtvg_gc, double cdmax, double chnkmin_u,
const double * cm_gc, const double * delkcc_gc_ns, const double * delkcc_omxkm3_gc,
double eps1, double epsmin, double epsus, double g, double gm1, int llcapchnk,
int llgcbz0, int llnormagam, int nwav_gc, const double * om3gmkm_gc,
const double * omxkm3_gc, double rn1_rn, double rnu, double rnum, double sqrtgosurft,
double wspmin, double xkappa, const double * xkmsqrtvgoc2_gc, const double * xkm_gc,
const double * xk_gc, double xlogkratiom1_gc, double xnlev, double zalp, int ichnk,
int nchnk, int ij) {





int i;
int j;

double xi;
double xj;
double deli1;
double deli2;
double delj1;
double delj2;
double ust2;
double arg;
double sqrtcdm1;
double xkappad;
double xloglev;
double xlev;
if (icode_wnd == 3) {

taut_z0_c(kijs, kijl, iusfg, halp, u10, u10dir, tauw, tauwdir, rnfac, us, z0, z0b,
chrnck, acd, alpha, alphamax, alphamin, ang_gc_a, ang_gc_b, ang_gc_c, bcd,
betamaxoxkappa2, bmaxokap, c2osqrtvg_gc, cdmax, chnkmin_u, cm_gc, delkcc_gc_ns,
delkcc_omxkm3_gc, eps1, epsmin, epsus, g, gm1, llcapchnk, llgcbz0, llnormagam,
nwav_gc, om3gmkm_gc, omxkm3_gc, rn1_rn, rnu, rnum, sqrtgosurft, xkappa,
xkmsqrtvgoc2_gc, xkm_gc, xk_gc, xlogkratiom1_gc, xnlev, zalp, ichnk, nchnk, ij);

} else if (icode_wnd == 1 || icode_wnd == 2) {
z0wave_c(kijs, kijl, us, tauw, u10, z0, z0b, chrnck, alpha, alphamin, chnkmin_u,
eps1, g, gm1, llcapchnk, ichnk, nchnk, ij);
xkappad = (double) 1.0 / xkappa;
xloglev = log(xnlev);


u10[ij - 1 + kijl*(ichnk - 1)] = xkappad*us[ij - 1 + kijl*(ichnk - 1)]*(xloglev -
log(z0[ij - 1 + kijl*(ichnk - 1)]));
u10[ij - 1 + kijl*(ichnk - 1)] =
max((double) (u10[ij - 1 + kijl*(ichnk - 1)]), (double) (wspmin));


}


}
22 changes: 22 additions & 0 deletions src/phys-scc-cuda/airsea_c.h
View file
Original file line number Diff line number Diff line change
@@ -0,0 +1,22 @@
#include <stdio.h>
#include <stdbool.h>
#include <float.h>
#include <math.h>
#include <cuda.h>
#include <cuda_runtime.h>
#include "z0wave_c.h"
#include "taut_z0_c.h"

__device__ void airsea_c(int kijs, int kijl, const double * halp, double * u10,
const double * u10dir, const double * tauw, const double * tauwdir,
const double * rnfac, double * us, double * z0, double * z0b, double * chrnck,
int icode_wnd, int iusfg, double acd, double alpha, double alphamax, double alphamin,
double ang_gc_a, double ang_gc_b, double ang_gc_c, double bcd, double betamaxoxkappa2,
double bmaxokap, const double * c2osqrtvg_gc, double cdmax, double chnkmin_u,
const double * cm_gc, const double * delkcc_gc_ns, const double * delkcc_omxkm3_gc,
double eps1, double epsmin, double epsus, double g, double gm1, int llcapchnk,
int llgcbz0, int llnormagam, int nwav_gc, const double * om3gmkm_gc,
const double * omxkm3_gc, double rn1_rn, double rnu, double rnum, double sqrtgosurft,
double wspmin, double xkappa, const double * xkmsqrtvgoc2_gc, const double * xkm_gc,
const double * xk_gc, double xlogkratiom1_gc, double xnlev, double zalp, int ichnk,
int nchnk, int ij);
Loading

0 comments on commit ffdf4ca

Please sign in to comment.