ocean_vcoord_rho_target Subroutine

private pure subroutine ocean_vcoord_rho_target(nx, ny, nz, target_h, remap_h_old, total_h, eta, t_conc, s_conc, dsig, floor_dz, rho_target, eos, p_ref, h_min, h_floor_eff, hybrid, floor_mode, h_nominal)

Column kernel of the RHO / HYCOM regrid (algorithm: see ocean_vcoord_compute_target_h_rho_impl). Flat on purpose: every array is an explicit-shape dummy and every knob a scalar dummy — no derived-type component and no associate reaches the do concurrent.

This is load-bearing on nvfortran -stdpar=gpu -gpu=mem:separate. The kernel used to run inside associate (rho_ref_pressure => this%rho_ref_pressure, ...) and hand that name BY REFERENCE to the out-of-module !$acc routine seq eos_density_point; the device callee received the HOST address of the component and faulted (CUDA_ERROR_ILLEGAL_ADDRESS, compute-sanitizer: invalid __global__ read at rdb_eos.F90 p_plus_p0 = p + p_0) on the first regrid of every RHO/HYCOM × Wright (or Roquet) run. The linear branch never reads p, which is why it hid. Gated by tests/test_ocean_vcoord_wright_device.F90. It also retires the associate-over-do concurrent shape CLAUDE.md forbids for ifx.

Arguments

Type IntentOptional Attributes Name
integer, intent(in), value :: nx

i-extent of every horizontal array (total, incl. halos).

integer, intent(in), value :: ny

j-extent of every horizontal array (total, incl. halos).

integer, intent(in), value :: nz

Number of layers; k = 1 is the bed, k = nz the surface.

real(kind=wp), intent(inout) :: target_h(nx,ny,nz)

Target layer thickness (m), bottom-up.

real(kind=wp), intent(in) :: remap_h_old(nx,ny,nz)

Pre-remap layer thickness snapshot (m), bottom-up.

real(kind=wp), intent(in) :: total_h(nx,ny)

Column reference depth H (m).

real(kind=wp), intent(in) :: eta(nx,ny)

Free-surface anomaly η (m).

real(kind=wp), intent(in) :: t_conc(nx,ny,nz)

Layer-mean potential temperature (°C), bottom-up.

real(kind=wp), intent(in) :: s_conc(nx,ny,nz)

Layer-mean salinity (PSU), bottom-up.

real(kind=wp), intent(in) :: dsig(nz)

Nominal layer fractions, bottom-up (dsig(nz) = surface) — the HYCOM floor only under HYCOM_FLOOR_SIGMA.

real(kind=wp), intent(in) :: floor_dz(nz)

HYCOM z* nominal layer thicknesses (m), bottom-up (floor_dz(nz) = surface) — read only under HYCOM_FLOOR_PROFILE.

real(kind=wp), intent(in) :: rho_target(0:nz)

Target potential densities (kg/m³), 0 = lightest = surface.

type(eos_t), intent(in) :: eos

Shared EOS handle (flat POD).

real(kind=wp), intent(in), value :: p_ref

Coordinate reference pressure (Pa) — rho_ref_pressure.

real(kind=wp), intent(in), value :: h_min

Pre-compaction strip threshold (m) — zstar_h_min.

real(kind=wp), intent(in), value :: h_floor_eff

Min-thickness inflation floor (m), > H_VANISHED.

logical, intent(in), value :: hybrid

.true. = HYCOM deltas; .false. = pure RHO.

integer, intent(in), value :: floor_mode

HYCOM z* floor source: HYCOM_FLOOR_PROFILE (floor_dz), HYCOM_FLOOR_UNIFORM (h_nominal) or HYCOM_FLOOR_SIGMA (dsig·H, the unconfigured-slot fallback).

real(kind=wp), intent(in), value :: h_nominal

Uniform z* nominal thickness (m) for HYCOM_FLOOR_UNIFORM.


Calls

proc~~ocean_vcoord_rho_target~~CallsGraph proc~ocean_vcoord_rho_target ocean_vcoord_rho_target proc~ocean_vcoord_rho_target_column ocean_vcoord_rho_target_column proc~ocean_vcoord_rho_target->proc~ocean_vcoord_rho_target_column proc~eos_density_point eos_density_point proc~ocean_vcoord_rho_target_column->proc~eos_density_point proc~invert_density_targets invert_density_targets proc~ocean_vcoord_rho_target_column->proc~invert_density_targets proc~roquet_spv_value roquet_spv_value proc~eos_density_point->proc~roquet_spv_value rdb_roq_spv_p rdb_roq_spv_p proc~roquet_spv_value->rdb_roq_spv_p rdb_roq_ts_coeffs rdb_roq_ts_coeffs proc~roquet_spv_value->rdb_roq_ts_coeffs

Called by

proc~~ocean_vcoord_rho_target~~CalledByGraph proc~ocean_vcoord_rho_target ocean_vcoord_rho_target proc~ocean_vcoord_compute_target_h_rho_impl ocean_vcoord_compute_target_h_rho_impl proc~ocean_vcoord_compute_target_h_rho_impl->proc~ocean_vcoord_rho_target proc~ocean_vcoord_compute_target_h_rho ocean_vcoord_t%ocean_vcoord_compute_target_h_rho proc~ocean_vcoord_compute_target_h_rho->proc~ocean_vcoord_compute_target_h_rho_impl proc~ocean_apply_ale_remap_centres ocean_apply_ale_remap_centres proc~ocean_apply_ale_remap_centres->proc~ocean_vcoord_compute_target_h_rho proc~ocean_apply_ale_remap_step ocean_apply_ale_remap_step proc~ocean_apply_ale_remap_step->proc~ocean_vcoord_compute_target_h_rho proc~ocean_dyn_step_split ocean_dyn_step_split proc~ocean_dyn_step_split->proc~ocean_apply_ale_remap_step proc~engine_step engine_step proc~engine_step->proc~ocean_dyn_step_split

Variables

Type Visibility Attributes Name Initial
integer, private :: i
integer, private :: j

Source Code

   pure subroutine ocean_vcoord_rho_target(nx, ny, nz, target_h, remap_h_old, &
                                           total_h, eta, t_conc, s_conc, dsig, &
                                           floor_dz, rho_target, eos, p_ref, h_min, &
                                           h_floor_eff, hybrid, floor_mode, h_nominal)
      !! Column kernel of the RHO / HYCOM regrid (algorithm: see
      !! `ocean_vcoord_compute_target_h_rho_impl`).  Flat on purpose:
      !! every array is an explicit-shape dummy and every knob a scalar
      !! dummy — no derived-type component and no `associate` reaches the
      !! `do concurrent`.
      !!
      !! This is load-bearing on nvfortran `-stdpar=gpu -gpu=mem:separate`.
      !! The kernel used to run inside `associate (rho_ref_pressure =>
      !! this%rho_ref_pressure, ...)` and hand that name BY REFERENCE to the
      !! out-of-module `!$acc routine seq` `eos_density_point`; the device
      !! callee received the HOST address of the component and faulted
      !! (`CUDA_ERROR_ILLEGAL_ADDRESS`, compute-sanitizer: invalid
      !! `__global__` read at `rdb_eos.F90` `p_plus_p0 = p + p_0`) on the
      !! first regrid of every RHO/HYCOM × Wright (or Roquet) run.  The
      !! linear branch never reads `p`, which is why it hid.  Gated by
      !! `tests/test_ocean_vcoord_wright_device.F90`.  It also retires the
      !! `associate`-over-`do concurrent` shape CLAUDE.md forbids for ifx.
      integer, intent(in), value :: nx
         !! i-extent of every horizontal array (total, incl. halos).
      integer, intent(in), value :: ny
         !! j-extent of every horizontal array (total, incl. halos).
      integer, intent(in), value :: nz
         !! Number of layers; `k = 1` is the bed, `k = nz` the surface.
      real(wp), intent(inout) :: target_h(nx, ny, nz)
         !! Target layer thickness (m), bottom-up.
      real(wp), intent(in) :: remap_h_old(nx, ny, nz)
         !! Pre-remap layer thickness snapshot (m), bottom-up.
      real(wp), intent(in) :: total_h(nx, ny)
         !! Column reference depth H (m).
      real(wp), intent(in) :: eta(nx, ny)
         !! Free-surface anomaly η (m).
      real(wp), intent(in) :: t_conc(nx, ny, nz)
         !! Layer-mean potential temperature (°C), bottom-up.
      real(wp), intent(in) :: s_conc(nx, ny, nz)
         !! Layer-mean salinity (PSU), bottom-up.
      real(wp), intent(in) :: dsig(nz)
         !! Nominal layer fractions, bottom-up (`dsig(nz)` = surface) —
         !! the HYCOM floor only under `HYCOM_FLOOR_SIGMA`.
      real(wp), intent(in) :: floor_dz(nz)
         !! HYCOM z* nominal layer thicknesses (m), bottom-up
         !! (`floor_dz(nz)` = surface) — read only under `HYCOM_FLOOR_PROFILE`.
      real(wp), intent(in) :: rho_target(0:nz)
         !! Target potential densities (kg/m³), `0` = lightest = surface.
      type(eos_t), intent(in) :: eos
         !! Shared EOS handle (flat POD).
      real(wp), intent(in), value :: p_ref
         !! Coordinate reference pressure (Pa) — `rho_ref_pressure`.
      real(wp), intent(in), value :: h_min
         !! Pre-compaction strip threshold (m) — `zstar_h_min`.
      real(wp), intent(in), value :: h_floor_eff
         !! Min-thickness inflation floor (m), `> H_VANISHED`.
      logical, intent(in), value :: hybrid
         !! `.true.` = HYCOM deltas; `.false.` = pure RHO.
      integer, intent(in), value :: floor_mode
         !! HYCOM z* floor source: `HYCOM_FLOOR_PROFILE` (`floor_dz`),
         !! `HYCOM_FLOOR_UNIFORM` (`h_nominal`) or `HYCOM_FLOOR_SIGMA`
         !! (`dsig·H`, the unconfigured-slot fallback).
      real(wp), intent(in), value :: h_nominal
         !! Uniform z* nominal thickness (m) for `HYCOM_FLOOR_UNIFORM`.

      integer :: i, j

      ! One column per thread: the whole column walk is the same-module
      ! `!$acc routine seq` `ocean_vcoord_rho_target_column`, so its
      ! NZ_STACK_MAX work arrays are thread-private and no inner loop is
      ! spread across threads.  (Inlined, nvfortran 26.5 auto-collapsed the
      ! (j,i) nest, promoted the work arrays AND `nk` to shared memory and
      ! vectorised the inner loops with shared-memory reductions whose
      ! scratch fell outside the kernel's shared allocation —
      ! compute-sanitizer: invalid `__shared__` write in the `donate`
      ! reduction, every RHO/HYCOM run, both EOS.)
      do concurrent(j=1:ny, i=1:nx)
         call ocean_vcoord_rho_target_column(i, j, nx, ny, nz, target_h, remap_h_old, &
                                             total_h, eta, t_conc, s_conc, dsig, &
                                             floor_dz, rho_target, eos, p_ref, h_min, &
                                             h_floor_eff, hybrid, floor_mode, h_nominal)
      end do
   end subroutine ocean_vcoord_rho_target