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.
| Type | Intent | Optional | 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; |
|
| 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 ( |
||
| real(kind=wp), | intent(in) | :: | floor_dz(nz) |
HYCOM z* nominal layer thicknesses (m), bottom-up
( |
||
| real(kind=wp), | intent(in) | :: | rho_target(0:nz) |
Target potential densities (kg/m³), |
||
| type(eos_t), | intent(in) | :: | eos |
Shared EOS handle (flat POD). |
||
| real(kind=wp), | intent(in), | value | :: | p_ref |
Coordinate reference pressure (Pa) — |
|
| real(kind=wp), | intent(in), | value | :: | h_min |
Pre-compaction strip threshold (m) — |
|
| real(kind=wp), | intent(in), | value | :: | h_floor_eff |
Min-thickness inflation floor (m), |
|
| logical, | intent(in), | value | :: | hybrid |
|
|
| integer, | intent(in), | value | :: | floor_mode |
HYCOM z* floor source: |
|
| real(kind=wp), | intent(in), | value | :: | h_nominal |
Uniform z* nominal thickness (m) for |
| Type | Visibility | Attributes | Name | Initial | |||
|---|---|---|---|---|---|---|---|
| integer, | private | :: | i | ||||
| integer, | private | :: | j |
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