diff --git a/.github/workflows/health.yml b/.github/workflows/health.yml index f87e0784..3b0e5ed7 100644 --- a/.github/workflows/health.yml +++ b/.github/workflows/health.yml @@ -1,6 +1,6 @@ # health — the combinations of scripts/ww3-gpu.sh that CI keeps alive. For now one, by the owner's # word (2026-10-14): "h1 - --ww3-gpu + --cpu ... this combination is the only one that needs to be -# present in CI for now". h1 builds WW3 from the WW3/ submodule (h0ffmann/WW3, ww3-gpu/kokkos-wire-base) +# present in CI for now". h1 builds WW3 from the WW3/ submodule (h0ffmann/WW3, ww3-gpu/kokkos-pnt) # with the KOKKOS switch and the ww_kokkos package on OpenMP, runs regtest ww3_ts1 with the # Fortran and the Kokkos DIA, then builds --ww3 and compares pair to pair (nccmp-tol). The other # combinations and why they are not here: README.md, "Health". diff --git a/.gitmodules b/.gitmodules index fdf39ba7..103a3c4d 100644 --- a/.gitmodules +++ b/.gitmodules @@ -6,7 +6,7 @@ [submodule "WW3"] path = WW3 url = git@github.com:h0ffmann/WW3.git - branch = ww3-gpu/kokkos-wire-base + branch = ww3-gpu/kokkos-pnt [submodule "WW4"] path = WW4 url = git@github.com:h0ffmann/WW4.git diff --git a/AGENTS.md b/AGENTS.md index 3d67b069..c2759ada 100644 --- a/AGENTS.md +++ b/AGENTS.md @@ -308,7 +308,10 @@ those transfers. are not resident. - **How it is checked.** A change to the resident sync reports `WW_KOKKOS_VA_COUNT=1` (bytes host to device and device to host) on `full` before and after; the whole-spectrum copies per step - must not grow, or the change is refused. + must not grow, or the change is refused. Every Kokkos run of `scripts/ww3-gpu.sh` and of + `just ww3-gpu-matrix` has `WW_KOKKOS_VA_POISON=1`: the stale host `VA` is filled with NaN after + each kernel write, so a host reader that misses its sync changes an output file and the byte + comparison fails. A caller's `WW_KOKKOS_VA_POISON=0` turns it off; the Fortran runs never have it. ## Bit for bit with the Fortran (hard rule) diff --git a/CONTRIBUTING.md b/CONTRIBUTING.md index b08ca89d..14c3e8f6 100644 --- a/CONTRIBUTING.md +++ b/CONTRIBUTING.md @@ -146,7 +146,9 @@ A mistake in a figure is as welcome as one in the text: the gallery links an iss the host only before an output or a restart that needs it. A whole-spectrum copy per step, per grid or per `W3WAVE` call is refused, even when it is the simpler fix: on a GPU it is a bus transfer every step. A change to the sync reports `WW_KOKKOS_VA_COUNT=1` on `full` before and - after. + after. `scripts/ww3-gpu.sh` and `just ww3-gpu-matrix` set `WW_KOKKOS_VA_POISON=1` in every Kokkos + run, so a host reader that misses its sync reads NaN and fails the byte comparison; + `WW_KOKKOS_VA_POISON=0` in the caller's environment turns it off, and the Fortran runs never have it. - `codemeta.json` is generated from `CITATION.cff` and `.zenodo.json` (`just codemeta`); never edit it by hand. - Run `just kokkos-test serial-debug` after touching `kokkos/`, and `bash -n` plus diff --git a/WW3 b/WW3 index 50a10730..fac74069 160000 --- a/WW3 +++ b/WW3 @@ -1 +1 @@ -Subproject commit 50a107307aad05207f751f399f5f6481becfc2eb +Subproject commit fac74069ad3dad694e153d78ac08bec7a4b1a44b diff --git a/docs/GLOSSARY.md b/docs/GLOSSARY.md index b6e0bc1d..82968c6d 100644 --- a/docs/GLOSSARY.md +++ b/docs/GLOSSARY.md @@ -703,10 +703,10 @@ and [kokkos/src/fortran_iface/PATCH.md](../kokkos/src/fortran_iface/PATCH.md). | `WW_KOKKOS_SRCE_STAGED` / input shadow | a diagnostic of `ww_srce_batch_va`, read once: `=1` makes a host backend stage the per-step inputs (`BJSEA`, `BWN1`, `BCG1`, `BPIN`, `BPSC`, `BWC`, `BTB`, `BTI`) into the batch buffers as a device does, skipping an input whose bits equal its *shadow* (the host copy of what the buffer holds, compared with `memcmp`), instead of reading them in place; a wrong skip would show as a stale input in the output. With `WW_KOKKOS_SRCE_COUNT=1` the shim prints the input copies and bytes H2D, the unchanged inputs, and the copies back D2H, projected for a device on a host backend `(v)` [kokkos/src/fortran_iface/srce_shim.cpp](../kokkos/src/fortran_iface/srce_shim.cpp) | Default unset: on Serial and OpenMP the kernel reads and writes the caller's packed arrays in place, with no copy (WFIP-0005 unit 4). | [kokkos/src/fortran_iface/PATCH.md](../kokkos/src/fortran_iface/PATCH.md) | | resident VA / `ww_va_host_read` / `ww_va_host_write` / `ww_va_device` / `ww_va_report` / `ww_srce_batch_va` | the device copy of the spectrum `VA(NSPEC,0:NSEALM)` that lives for the run (WFIP-0005), with a state per grid saying which copy holds the newer bits (`HOST`, `DEVICE`, `BOTH`); the C calls W3WAVE makes before a host read or write of `VA`, the one a kernel shim makes before reading the device copy, the traffic report, and the batched `W3SRCE` that works in place on the device copy `(v)` [kokkos/src/fortran_iface/resident_c.hpp](../kokkos/src/fortran_iface/resident_c.hpp) | The protocol of `Kokkos::DualView` (modify, sync) over WW3's own array; with no kernel on the device copy no call copies anything. Wrapped by `W3KOKKOS_VA_HOST_READ` / `_WRITE` in `w3kokkos_residentmd.F90` on the fork. | [kokkos/src/fortran_iface/PATCH.md](../kokkos/src/fortran_iface/PATCH.md) | | `WW_KOKKOS_VA_RESIDENT` / `WW_KOKKOS_VA_COUNT` | diagnostics of the resident VA, ignoring `WW_KOKKOS_ALL`: `=1` makes every host sync call push a host-newer `VA` to the device and copy it back, so a host write without its sync call shows in the output; the count prints the traffic line and the copy counters at exit `(v)` [kokkos/src/fortran_iface/resident_shim.cpp](../kokkos/src/fortran_iface/resident_shim.cpp) | Default unset. On the CPU both copies are `memcpy`, so the bits cannot change. | [kokkos/src/fortran_iface/PATCH.md](../kokkos/src/fortran_iface/PATCH.md) | -| `WW_KOKKOS_VA_CHECK` / `WW_KOKKOS_VA_POISON` | diagnostics of the resident VA, read once, ignoring `WW_KOKKOS_ALL`, which compare or overwrite only the stale copy: `VA_CHECK=1` makes every sync call check what the state claims before it copies (unless `HOST`, the host `VA` still holds the bits of the last sync, so a host write without its `W3KOKKOS_VA_HOST_WRITE` stops the model; with `BOTH`, `d_va` holds them too, so a kernel that wrote `d_va` without marking it does); `VA_POISON=1` fills the host `VA` with a quiet NaN after every kernel write of `d_va`, so a host read without its sync, or a column the partial read misses, puts NaN in the output `(v)` [kokkos/src/fortran_iface/resident_shim.cpp](../kokkos/src/fortran_iface/resident_shim.cpp) | Default unset; a correct run writes the same bits with either. Added by the mutation audit of 2026-10-17. | [docs/log/2026-10-17-mutation-audit.md](log/2026-10-17-mutation-audit.md) | +| `WW_KOKKOS_VA_CHECK` / `WW_KOKKOS_VA_POISON` | diagnostics of the resident VA, read once, ignoring `WW_KOKKOS_ALL`, which compare or overwrite only the stale copy: `VA_CHECK=1` makes every sync call check what the state claims before it copies (unless `HOST`, the host `VA` still holds the bits of the last sync, so a host write without its `W3KOKKOS_VA_HOST_WRITE` stops the model; with `BOTH`, `d_va` holds them too, so a kernel that wrote `d_va` without marking it does); `VA_POISON=1` fills the host `VA` with a quiet NaN after every kernel write of `d_va`, so a host read without its sync, or a column the partial read misses, puts NaN in the output `(v)` [kokkos/src/fortran_iface/resident_shim.cpp](../kokkos/src/fortran_iface/resident_shim.cpp) | Default unset in the shim; a correct run writes the same bits with either. Since 2026-10-18 `scripts/ww3-gpu.sh` and the regtest matrix set `WW_KOKKOS_VA_POISON=1` in every Kokkos run (a caller's `=0` turns it off) and unset it in the Fortran runs. Added by the mutation audit of 2026-10-17. | [docs/log/2026-10-17-mutation-audit.md](log/2026-10-17-mutation-audit.md) | | `WW_KOKKOS_SRCE_LANES` | the most sea points per team of the batched `W3SRCE` launch on a host backend, 1 to 8, read once; each point runs the statements it runs alone, the points of a team take turns `(v)` [kokkos/src/fortran_iface/srce_shim.cpp](../kokkos/src/fortran_iface/srce_shim.cpp) | Default 8 (the owner, 2026-10-18), 1 for one point per team; a launch with fewer points than lanes times threads uses 1, and a device always uses 1. The lane layer of plan 005, which `W3SIN4`'s stress scan vectorises across (L2). | [docs/compound/plans/2026-10-18-005-srce-points-in-lanes-plan.md](compound/plans/2026-10-18-005-srce-points-in-lanes-plan.md) | | `WW_KOKKOS_OUTG_STAGED` | a diagnostic of the outg unit, read once: `=1` makes a host backend take the device's path of `ww_outg_va`, the inputs copied into the shim's own buffers, the fields computed into buffers, and only the fields a host reader needs copied back, so a field the copy-back selection misses keeps its old bits in WW3's array and shows in `out_grd.ww3` `(v)` [kokkos/src/fortran_iface/outg_shim.cpp](../kokkos/src/fortran_iface/outg_shim.cpp) | Default unset: on Serial and OpenMP the kernel writes WW3's arrays in place. | [docs/log/2026-10-17-mutation-audit.md](log/2026-10-17-mutation-audit.md) | -| `WW_KOKKOS_STRICT` | environment variable read once by `ww_kokkos_strict()`: `=1` makes every place where a part asked on Kokkos would run the Fortran stop the model with `EXTCDE` and a `NDSE` line naming the part and the reason: `KOKKOS_PART` under `WW_KOKKOS_ALL=1`, the two `W3SRCE` set-up fallbacks and the per-step `W3SRCE` gate (`W3KOKKOS_SRCE_STEP`). Unset, those places run the Fortran as before and the shim prints one line at exit, `ww_kokkos: calls left on the Fortran (...): SRCE= ...` `(v)` [kokkos/src/fortran_iface/part_flags.cpp](../kokkos/src/fortran_iface/part_flags.cpp) | Default unset. The regtest matrix's Kokkos run sets it, its Fortran runs unset it, so a pair cannot compare the Fortran with itself unseen. | [kokkos/src/fortran_iface/PATCH.md](../kokkos/src/fortran_iface/PATCH.md), section 24 | +| `WW_KOKKOS_STRICT` | environment variable read once by `ww_kokkos_strict()`: `=1` makes every place where a part asked on Kokkos would run the Fortran stop the model with `EXTCDE` and a `NDSE` line naming the part and the reason: `KOKKOS_PART` under `WW_KOKKOS_ALL=1`, the two `W3SRCE` set-up fallbacks, the per-step `W3SRCE` gate (`W3KOKKOS_SRCE_STEP`), `W3KOKKOS_CPRT`'s two fallbacks and `W3KOKKOS_OUTG`'s call of the Fortran `W3OUTG` for a field its translation leaves to the Fortran (U3 Stokes, skewness, WBT, an unallocated 3-D field). Unset, those places run the Fortran as before and the shim prints one line at exit, `ww_kokkos: calls left on the Fortran (...): SRCE= ...` `(v)` [kokkos/src/fortran_iface/part_flags.cpp](../kokkos/src/fortran_iface/part_flags.cpp) | Default unset. The regtest matrix's Kokkos run sets it when it names its parts (a preset or a part list), not under `WW3GPU_KOKKOS=all`, where a fallback is counted; its Fortran runs unset it, so a pair cannot compare the Fortran with itself unseen. | [kokkos/src/fortran_iface/PATCH.md](../kokkos/src/fortran_iface/PATCH.md), section 24 | | `WW_KOKKOS_PRESET` / `presets.txt` | environment variable naming a line `PRESET ` of `kokkos/src/fortran_iface/presets.txt`: `ww_kokkos_flag` turns its parts on as if each `WW_KOKKOS_` were `1`; a set `WW_KOKKOS_` still wins, and `W3INIT` stops on a name the file lacks. `resident` is the only preset (`SRCE,XYP3,KTP3,OUTG,IOPE,UPDV,CPRT`) `(v)` [kokkos/src/fortran_iface/presets.txt](../kokkos/src/fortran_iface/presets.txt) | The file is the one list: CMake compiles it into the shim, `scripts/ww3-gpu.sh --kokkos resident` and the matrix's `WW3GPU_KOKKOS=resident` read it. | [kokkos/src/fortran_iface/PATCH.md](../kokkos/src/fortran_iface/PATCH.md), section 25 | | mutation audit / mutant (equivalent, surviving) | a deliberate, minimal breakage of a piece of the port, applied in a scratch tree and never committed, to ask whether the checks that call the piece correct would fail without it: a mutant is *detected* when a check fails, *surviving* when every check passes, and *equivalent* when it cannot change any bit the case computes (it is not a defect the case can show) `(v)` [docs/log/2026-10-17-mutation-audit.md](log/2026-10-17-mutation-audit.md) | First run 2026-10-17, on the CPU. | [docs/log/2026-10-17-mutation-audit.md](log/2026-10-17-mutation-audit.md) | | `WW_KOKKOS_ALL` / `ww_kokkos_flag` | the master run-time flag of every Kokkos part, and the one C function that resolves each `WW_KOKKOS_`: a set part flag wins (`1` on, anything else off), an unset one follows `WW_KOKKOS_ALL=1` (on) or stays off `(v)` [kokkos/src/fortran_iface/part_flags.cpp](../kokkos/src/fortran_iface/part_flags.cpp) | "100 %" means every wired part, not all of WW3; under the master a part whose translation does not cover the switch file stays Fortran. `W3INIT` logs one "Kokkos parts" line; `scripts/ww3-gpu.sh --kokkos all\|none\|` sets it. | [kokkos/src/fortran_iface/PATCH.md](../kokkos/src/fortran_iface/PATCH.md) | diff --git a/docs/NOAA_PROGRESS_RECORD_202610.md b/docs/NOAA_PROGRESS_RECORD_202610.md index 5fa21e28..d270ccbd 100644 --- a/docs/NOAA_PROGRESS_RECORD_202610.md +++ b/docs/NOAA_PROGRESS_RECORD_202610.md @@ -181,7 +181,7 @@ tables in `docs/data/ww3gpu_matrix_*_pin-50a10730_2026-10-18.tsv`. `ww3_grid` namelist read against the test switch). None is marked expected-to-fail. - **Checked (v).** The DIA (`W3SNL1`) alone passes the `NL1` entries of the regtest matrix. - **Not shown.** Parity with the Intel compiler; parity on a GPU (judged by tolerance); restart - reproduction (`scripts/restart_pair.sh` exists and has not run); `ww3_multi` under MPI. + reproduction beyond one `SHRD` case (`scripts/restart_pair.sh` ran on `full` at `WW3/` `50a10730` and `fac74069`, OpenMP backend, 12 hours, every comparison identical; [log](log/2026-10-18-resident-restart-pair.md)); `ww3_multi` under MPI. ## 5. Where the CPU time goes @@ -380,7 +380,7 @@ Sections 3 and 4 of `W3SNL1` read the spectrum at the quadruplet bins through th | Median of three, alternated with NOAA WW3 | add `--repeat 3` | | Stop on any fallback; catch a missed sync | `WW_KOKKOS_STRICT=1 WW_KOKKOS_VA_POISON=1` in the environment | | The regtest entries of section 4 | the commands of `docs/log/2026-10-18-u3-regtests-v2.md`, "What was done" | -| Restart reproduction, `SHRD` | `scripts/restart_pair.sh` (not run yet) | +| Restart reproduction, `SHRD` | `scripts/restart_pair.sh` (ran at `WW3/` `50a10730` and `fac74069`, [log](log/2026-10-18-resident-restart-pair.md)) | | The same on the point case | `bash scripts/ww3-gpu.sh --ww3-gpu smoke --cpu --kokkos resident` | | Per-part timers | add `WW_KOKKOS_TIMERS=1` to the Kokkos run, or `--timers` | | DIA kernel alone | `OMP_PROC_BIND=false OMP_NUM_THREADS= kokkos/build/openmp-release/tests/ww_bench_snl1` | diff --git a/docs/PRODUCTION_READINESS_202610.md b/docs/PRODUCTION_READINESS_202610.md index ca0a684a..5fed01d6 100644 --- a/docs/PRODUCTION_READINESS_202610.md +++ b/docs/PRODUCTION_READINESS_202610.md @@ -47,11 +47,11 @@ regtest matrix yet"). 10 of the 310 rows of `kokkos/ports.tsv` are regtested and | 3 | MPI: four of the six resident parts are `SHRD` only, the preset stops on `DIST MPI`, the communicator is ignored, the device id is not per rank | Blocker | Major | A resident preset without `XYP3`, `OUTG`, `IOPE`, `UPDV`; one run on `switch_lab_mpi` | | 4 | Bit parity measured with gfortran and g++ only; the operational compiler and the GPU are not measured | Blocker | Major | Static argument for the Intel build's math calls; one smoke with it | | 5 | Build and packaging: a separate static library, C++20 and CMake 3.25; the install-only configure and the exact Kokkos check are done (U6) | Major | Major | A shared-library build and a recipe for the module systems of an operational centre | -| 6 | `WW_KOKKOS_ALL=1` and `W3KOKKOS_SRCE_ON` fall back to the Fortran without a word | Major | Minor | Done, at the pin `50a10730` since #191: `WW_KOKKOS_STRICT=1` stops, otherwise one fallback-count line at exit; the matrix's Kokkos runs set it | -| 7 | Configuration is 61 environment variables, no namelist; the preset lives in a lab script | Major | Major | Smallest form done, at the pin `50a10730` since #191: `WW_KOKKOS_PRESET=resident` from one table, `presets.txt`; next, a namelist group read by `W3INIT` | +| 6 | `WW_KOKKOS_ALL=1` and `W3KOKKOS_SRCE_ON` fall back to the Fortran without a word | Major | Minor | Done: `WW_KOKKOS_STRICT=1` stops, otherwise one fallback-count line at exit, `W3OUTG`'s fallback included since `ww3-gpu/strict-outg`; the matrix's Kokkos runs set it when they name their parts (section 6) | +| 7 | Configuration is 61 environment variables, no namelist; the preset lives in a lab script | Major | Major | Smallest form done on `ww3-gpu/u2-strict-preset`: `WW_KOKKOS_PRESET=resident` from one table, `presets.txt`; next, a namelist group read by `W3INIT` | | 8 | Observability: log lines on every rank, counters at exit to `stderr`, no provenance in the output files | Major | Minor | Gate the lines by `NAPLOG`; one NetCDF global attribute | | 9 | OpenMP inside MPI: the per-point parts are serialised by `!$OMP CRITICAL`; the thread budget per rank is not set by the shim | Major | Minor | Measure the hybrid entries of `ww3_ufs1.1` before claiming a speed for them | -| 10 | Restart with the resident spectrum: covered by the sync calls, never regtested; every `matrix.base` restart pair is `MPI` | Major | Minor | Run `scripts/restart_pair.sh` (a `SHRD` restart pair on `full`, written 2026-10-18, not run) | +| 10 | Restart with the resident spectrum: covered by the sync calls, never regtested; every `matrix.base` restart pair is `MPI` | Major | Minor | Run `scripts/restart_pair.sh` (a `SHRD` restart pair on `full`); it passed at `WW3/` `50a10730` and again at `fac74069` (strict mode there stops a run that asks for `WBT` or `SKW`; this case asks for neither), [log](log/2026-10-18-resident-restart-pair.md) | | 11 | `PDLIB` and unstructured grids: translated, not wired, gated out | Minor ⚠ | Major | None for GFS-Wave v16; a wiring unit for `PDLIB` when the owner asks | | 12 | The fork patch: 4 082 lines of new Fortran, 819 changed lines in three WW3 files, 58 218 lines of C++ in another repository, most of it LGPL-3.0-or-later, some MIT, 76 files with no licence line | Minor | Blocker | ADR-0018 (Proposed) names the three places `ww_kokkos` could live; the owner chooses, after a legal reading ⚠ | | 13 | A stale contract remark: "one spectral grid per process"; edited on the fork branch `ww3-gpu/u1-multi-va` (section 13) | Minor | Minor | Closed: at the pin `50a10730` since #191 | @@ -307,6 +307,17 @@ asked for Kokkos and got the Fortran learns it only from one log line or a `*_CO `W3KOKKOS_SRCE_ON` is unchanged. The regtest matrix's Kokkos run sets `WW_KOKKOS_STRICT=1` (`regtests/cmake/ww3gpu_common.cmake`, `ww3gpu_kokkos_env`). `PATCH.md` section 24 has the runs that show it. +- **Update, 2026-10-18 (the follow-up of #191, its code review's findings 1 to 3).** Three changes, + [log](log/2026-10-18-strict-outg-poison.md): + - `W3KOKKOS_OUTG`'s call of the Fortran `W3OUTG` for a field its translation leaves to the + Fortran (U3 Stokes, skewness, `WBT`, an unallocated 3-D field) now goes through + `KOKKOS_FALLBACK`: it stops under `WW_KOKKOS_STRICT=1` and is counted otherwise (fork branch + `ww3-gpu/strict-outg`, `PATCH.md` section 31). Before, strict mode did not see it. + - The matrix sets `WW_KOKKOS_STRICT=1` only for a Kokkos run that names its parts + (`WW3GPU_KOKKOS=resident` or a part list); under `WW3GPU_KOKKOS=all` a part the entry's + switch does not cover is counted, not stopped, so an entry without `PR3` runs. + - `WW_KOKKOS_VA_POISON=1`, the missed-sync detector, is on in every Kokkos run of + `scripts/ww3-gpu.sh` and of the matrix, off in the Fortran runs; a caller's `=0` turns it off. ## 7. Configuration surface @@ -393,7 +404,16 @@ the UFS case (`work_c`, then `work_d` from its restart, compared by `test.comp`; restarts a `SHRD` case. [`scripts/restart_pair.sh`](../scripts/restart_pair.sh) is the smallest equivalent on `full` with `switch_lab_shrd`: a 00:00 to 12:00 run and a 06:00 to 12:00 run from its own 06:00 restart, for the Fortran and the resident Kokkos of one executable, compared with -`cmp`. It was not run ⚠ ([log](log/2026-10-18-u3-regtests-v2.md)); the gap stays Major for (a). +`cmp`. It was not run when written ⚠ ([log](log/2026-10-18-u3-regtests-v2.md)). + +**Update, 2026-10-18 (restart pair run).** The pair ran once at `b828041` with `WW3/` at `50a10730`, +4 threads, `WW_KOKKOS_VA_POISON=1`, `WW_KOKKOS_STRICT=1`: the 06:00-to-12:00 run from the 06:00 +restart gives the bytes of the run cut from the cold start (`out_grd.ww3` and the 12:00 restart), +for the Fortran and for the resident Kokkos, and every file of the Kokkos runs equals the +Fortran's (v, `cmp`, 9 comparisons, exit 0; +[log](log/2026-10-18-resident-restart-pair.md)). It covers one `SHRD` case on the OpenMP backend, +one restart and 12 hours; `MPI` restart, the CUDA backend and `ww3_multi` are not covered ⚠. The +gap stays open for those, and for (a) until the owner accepts one case as enough. ## 11. `PDLIB` and unstructured grids diff --git a/docs/compound/plans/2026-10-18-005-srce-points-in-lanes-plan.md b/docs/compound/plans/2026-10-18-005-srce-points-in-lanes-plan.md index 38681659..32158da8 100644 --- a/docs/compound/plans/2026-10-18-005-srce-points-in-lanes-plan.md +++ b/docs/compound/plans/2026-10-18-005-srce-points-in-lanes-plan.md @@ -88,7 +88,7 @@ change a bit. | S0 | One-thread profile of today's batch, libm separated from the rest | the shares below are measured; the go/no-go of S0 is applied | | L1 | The lane layer: teams of `W` points (`WW_KOKKOS_SRCE_LANES`, 1 to 8, 1 reproduces today, default 8 since L2), each section run for the team's points in turn with today's scalar routines | **Done** 2026-10-18: `full` and `smoke` byte-identical with 1 and 8 (`kokkos/port_report/W3SRCE_KOKKOS.md`, "Points per team"). The point-minor scratch moves to L2, the first unit that reads it | | L2 | `W3SIN4` on lanes (its stress loop, `st4_sin4.hpp:198-249`) | **Done** 2026-10-18: byte-identical on `smoke` and `full`; one-thread `W3SRCE` kernel 105.17 s to 99.77 s, 5.1 % ([log](../../log/2026-10-18-srce-lanes-sin4.md)) | -| L3 | `W3SDS4` on lanes, the libmvec sites left as lane loops | the same | +| L3 | `W3SDS4` on lanes, the libmvec sites left as lane loops | **Built** 2026-10-18: byte-identical on `smoke` and `full` with 1 and 8 points per team; timing and regtests in the [log](../../log/2026-10-18-srce-lanes-sds4.md) | | L4 | The integration and the initialisation (`srce_kernel.hpp:177-400`) | the same | | L5 | `W3SPR4`, `W3SLN1`, `W3SDB1`, `W3SBT1` | the same | | L6 | The DIA reads its input from the point-minor layout, or the layer transposes around it, whichever the timing favours | the same | diff --git a/docs/log/2026-10-18-resident-restart-pair.md b/docs/log/2026-10-18-resident-restart-pair.md new file mode 100644 index 00000000..7627e080 --- /dev/null +++ b/docs/log/2026-10-18-resident-restart-pair.md @@ -0,0 +1,79 @@ +# Restart with the resident spectrum: the `SHRD` restart pair, run for the first time + +2026-10-18, an agent for M. Hoffmann (goal `5afd9379-8152`); gap 10 of +[`PRODUCTION_READINESS_202610.md`](../PRODUCTION_READINESS_202610.md). The pair was written in +[the `abbb0456` handover](2026-10-18-goal-abbb0456-handover.md) ("Waiting on the owner") and listed +again in [the `c3e7eecf` handover](2026-10-18-goal-c3e7eecf-handover.md); neither had run it. + +## Question + +Does a run restarted from its own restart file give the bytes of the run it was cut from, with the +resident Kokkos preset, and do the Kokkos runs give the Fortran's bytes of the same executable? +No `matrix.base` pair answers it: every restart-reproduction entry at the pin is `MPI` (`ww3_tp2.17`, +`ww3_ufs1.1`), where the resident parts do not run (ADR-0017). + +## What was done + +On the owner's machine (Intel Core i9-14900, shared with two other sessions), in the pinned `ww3` +shell (gfortran 16.2.0), Kokkos OpenMP backend, never `--gpu`, from a clean worktree at `b828041` +(`main`) with `WW3/` at `50a10730`. The wrapper `test-191-restart.sh` of the earlier handovers +points at the checkout of PR #191, so `scripts/restart_pair.sh` ran from this worktree instead, with +the wrapper's arguments: + +```bash +WW3GPU_SWITCH=switch_lab_shrd nix develop ./picasso/labs/pratico#ww3 --command \ + bash scripts/ww3-gpu.sh --ww3-gpu full --cpu --kokkos resident --build-only +WW_KOKKOS_VA_POISON=1 OMP_NUM_THREADS=4 nix develop ./picasso/labs/pratico#ww3 --command \ + bash scripts/restart_pair.sh +``` + +The build's switch was `SHRD ... KOKKOS` (`build/ww3-gpu-openmp/switch`). The script runs +`bench_small` (120 x 80, 9 520 sea points, 12 hours) four times with that one executable: `a` from +00:00 with restarts at 06:00 and 12:00, and `b` from 06:00 on `a`'s restart, each for the Fortran +(`WW_KOKKOS_ALL=0`) and for the resident preset (`WW_KOKKOS_STRICT=1`, every `*_COUNT=1`). It +compares with `cmp`. + +The first attempt exited 1 after `fortran/b` with no message. The cause is the script's plumbing: +it read the "Kokkos parts" line from `ww3_shel.out`, where WW3 does not print it (it is in +`log.ww3`, as `scripts/ww3-gpu.sh:658` reads it), and `grep` under `pipefail` ended the script +silently. The fix reads `log.ww3` and tolerates a missing line, which the next test then reports. +What is compared did not change. + +## Result + +Exit 0 (v, the command above, 4 threads; the four runs took about 50 s each): + +| Comparison | Files | Result | +|---|---|---| +| restart reproduction, Fortran: `b` against `a` | `out_grd.ww3` (2 594 586 bytes), the 12:00 restart (21 618 432 bytes) | identical | +| restart reproduction, resident Kokkos: `b` against `a` | the same two | identical | +| Kokkos against Fortran, `a` | `out_grd.ww3`, `restart001.ww3`, `restart002.ww3` | identical | +| Kokkos against Fortran, `b` | `out_grd.ww3`, `restart001.ww3` | identical | + +The Kokkos runs had `SRCE`, `CPRT`, `OUTG`, `IOPE`, `XYP3`, `KTP3`, `UPDV` on Kokkos (7 of 20, the +"Kokkos parts" line of `log.ww3`), no call left on the Fortran (`SRCE=0 CPRT=0`), one host-to-device +copy of `VA` per run, two device-to-host copies in `a` (the two restart writes) and one in `b`; 81 +host copies poisoned with NaN in `a` and none read back. So a missed sync after the restart read +would have shown as NaN or a different byte. + +## Open questions + +- One `SHRD` case, one restart, OpenMP backend. The CUDA backend, `ww3_multi`, `MPI` and a restart + that moves the grid are not covered ⚠; ADR-0017 stops the resident preset under `MPI`. +- The pair compares bytes of a 12-hour case at 4 threads; other thread counts were not run ⚠. +- The matrix's `rstrt_b4b` pairs remain `MPI` and out of reach of the preset. +- `README.md` and `docs/NOAA_PROGRESS_RECORD_202610.md` may still say the pair was never run; not + edited here. + +## Rerun at the combined pin `fac74069`, 2026-10-18 + +Branch `goal/abbb0456-8241-strict-outg-poison` at `debb1b0`, `WW3/` at `fac74069` (strict mode +stops a run when `W3OUTG` falls back to the Fortran). The same two commands as above, 4 threads, +`WW_KOKKOS_VA_POISON=1`; the script was not edited and keeps `WW_KOKKOS_STRICT=1` for the Kokkos +runs. + +Exit 0 (v): the same nine comparisons are identical (`out_grd.ww3` 2 594 586 bytes, the restarts +21 618 432 bytes). Kokkos runs: one host-to-device and two (`a`) or one (`b`) device-to-host copies +of `VA`, calls left on the Fortran `SRCE=0 CPRT=0 OUTG=0`. The strict stop did not fire because +this case requests neither `WBT` nor `SKW`. A case that does request them (`smoke`) stops at +`fac74069`: see the [lanes log](2026-10-18-srce-lanes-sds4.md#rerun-at-the-combined-pin-fac74069-2026-10-18). diff --git a/docs/log/2026-10-18-srce-lanes-sds4.md b/docs/log/2026-10-18-srce-lanes-sds4.md new file mode 100644 index 00000000..092a5979 --- /dev/null +++ b/docs/log/2026-10-18-srce-lanes-sds4.md @@ -0,0 +1,130 @@ +# `W3SDS4` with one sea point per vector lane: bit-identical, 19 % off the one-thread `W3SRCE` kernel after L2 + +2026-10-18, for M. Hoffmann (goal `5afd9379-4417`), unit L3 of +[the points-in-lanes plan](../compound/plans/2026-10-18-005-srce-points-in-lanes-plan.md), +after [L2](2026-10-18-srce-lanes-sin4.md) (PR #199). + +## Question + +Does `W3SDS4` (`w3src4md.F90:2043-2659` at `WW3@50a10730`, 24.8 % of the one-thread batch in +[S0](2026-10-18-srce-lanes-s0-profile.md)) give the Fortran's bits and take less time when eight +sea points run its loops at once, one per vector lane? + +## What was done + +- `w3sds4_point` (`kokkos/src/ww_kokkos/st4_sds4.hpp`) split into `sds4_prep` (sections 0 and 1, + `:2203-2224`), `sds4_mss` (`:2228-2254`), `sds4_sat` (case 1, `:2262-2290`), `sds4_case1` (the + rest of case 1, `:2291-2331`), `sds4_case2`, `sds4_case3`, `sds4_turb` (`:2533-2575`) and + `sds4_whitecap` (`:2581-2650`). (v) Every line of the old body is in the new file, case for + case, with the whitespace removed (a line-multiset comparison: only the `switch` scaffolding is + not in a part, and it is in `w3sds4_point`); the case bodies are dedented. The arguments and the + work arrays travel in `Sds4Args` and `Sds4Work`, the locals that cross sections in `Sds4Mid`. +- `w3sds4_lanes<8>` runs the parts for up to eight points. Three loops have a lane version: + - `mss_lanes`: the three mean-square-slope sums, one accumulator per lane, ITH order; `powf` and + `atan2f` per lane. + - `sat_lanes`: the saturation integral, one accumulator per ITH and lane, added in J order from + zero over all `2*SDSNTH+1` terms (the zero weights are multiplied, as the scalar code does); + the sum of `A` and the `MAXVAL` per lane. Without the transposed tables each point goes + through `sds4_sat`. + - `turb_lanes`: the cheap isotropic sum and the default O(NK^2 NTH^2) branch (`:2554-2563`) + with the branch on `BTH0` as a select; `COS/SIN(USDIR)` computed once per lane instead of once + per ITH (the same call, same arguments, so the same bits). + - The two libmvec sites of `sds4_case1` and `sds4_case3` and every other loop stay per point: + each point makes today's grouped call. +- The vector types moved to `st4_vec8.hpp`, shared with `W3SIN4`. +- `w3srce_lanes` calls `w3sds4_lanes` once per step for the points still in the loop + (`srce_kernel.hpp`, `sds4_lanes`); `step_c_main` is `sds4` then `step_c_after_sds4`, so the + scalar path runs the same statements in the same order. + +## Result + +- (v) Bits, `WW_KOKKOS_STRICT=1 WW_KOKKOS_VA_POISON=1`, `WW_KOKKOS_SRCE_LANES=1` and `=8`: + `bash scripts/ww3-gpu.sh --ww3-gpu smoke --cpu --kokkos resident` and the same on `full` give + every model file byte-identical to the Fortran of the same executable (`smoke`: `out_grd.ww3`, + `out_pnt.ww3`, `restart.ww3`, `ww3.196806.nc`; `full`: `out_grd.ww3`, `ww3.field.nc`), every + file but the NetCDF byte-identical to `--ww3`'s, `nccmp-tol` PASS. +- (v) What `full` exercises, read from a temporary print (removed): case 1, `SSDSC(3)` = -0.40344 + (the cumulative branch of `:2554-2563`), `SSDSC(18)` = 400 (so section 1.b runs), `SDSDTH` = 80 + (the weighted branch), the transposed tables present. Not exercised by it: cases 2 and 3 (their + parts are the old code called per point), `SSDSC(3) > 0`, `SDSDTH >= 180`, no transposed tables. +- (v) Those branches checked against the scalar path with a temporary hack (removed, not + committed) that overwrote a parameter from an environment variable in both the scalar and the + lane entry: `SSDSC(3)` = 0.4, `SDSDTH` = 200, case 2, case 3 and the sums of those. For each, + `full`, `WW_KOKKOS_SRCE_LANES=1` and `=8`, `out_grd.ww3` byte-identical between the two (and + different from the normal run, so the hack took effect). This compares the lane path with the + scalar path of the same code, not with the Fortran, and these parameter values are not WW3's. +- (v) `just kokkos-test serial-debug` and `openmp-release`: 11/11 each, no warning. +- (v) `bash scripts/ww3-gpu.sh --ww3-gpu smoke --gpu --kokkos resident --build-only`: compiles and + links, no warning; nothing ran on a device. The lane code is host-only. +- (v) `-fopt-info-vec-all` on `srce_shim.cpp` (openmp-release flags, the pinned `g++` 16.2.0): the + vector-type statements are vector instructions by construction; the plain loops inside the lane + functions (gathers of `A` through eight pointers, the per-lane stores) are scalar, as written. +- (v) WW3's regtests, `just regtest-selected` at `958b7bc` plus this log's uncommitted edits + (`WW3@50a10730`, resident, lab switch, strict, 6 threads per run): `ww3_tp2.15` default, + `-g 5km`, `input_const`, `input_rho` and `ww3_ts2` `ST4_PR3_UQ` give `test.comp` 0/23, 0/23, + 0/23, 0/25, 0/22, `nccmp-tol` PASS, `smoke` byte-identical, 8/8 ctest tests passed. +- (v) Time: `srce.batch.kernel` (`WW_KOKKOS_TIMERS=1`), `full`, one thread pinned to a performance + core (`taskset -c 2`), runs alternated, Intel Core i9-14900, 2026-10-18 22:52 to 23:02 UTC + (the machine's clock), load read before each run: + +| `WW_KOKKOS_SRCE_LANES` | Runs (s) | Median | Load before each run | +|---|---|---|---| +| 1 | 108.71, 104.99, 104.43 | 104.99 s | 7.85, 2.51, 1.44 | +| 8 | 83.39, 80.51, 80.55 | 80.55 s | 4.17, 2.06, 1.45 | + +- (v) Every run's `out_grd.ww3` is byte-identical to the Fortran-matched output of the build. +- ⚠ 24.4 s of 105.0 s, 23.3 %, for L2 and L3 together; L2 alone was 5.1 % in its log (105.17 s to + 99.77 s), so L3's part is about 19 s (19 %) of the kernel, against 24.8 % of the batch that + `W3SDS4` was in S0 ⚠ (the two logs ran on different builds and days; L2's lanes=8 figure was not + repeated here). The first lanes=1 run began at load 7.85, and its 108.71 s is the slowest. +- ⚠ No 32-thread figure: not measured in this unit. + +## Open questions for the owner + +1. The host lane loops call the same glibc `libm` and `libmvec` as the scalar code, per lane. The + transcribed device `libm` (`libm_device.hpp`) is not used on the host; whether the `powf`, + `atan2f` and `sincosf` calls (libm and libmvec together were 8.9 % of the batch in S0) may use a + vector `libm` is decision 2 of `docs/GPU_LIBM_BITS_202610.md`. +2. Zero weights of the saturation integral are multiplied, not skipped: 12 of the 23 products per + direction in the default configuration multiply zero. Skipping them is not bit-neutral (0 x Inf + is NaN) and would be its own PR with its own evidence. +3. The loop nesting: IK and ITH outside, the lane loop inside, one accumulator per lane and ITH, + and the vector types (GCC vector extensions, not `std::simd`). Review asked. +4. `COS/SIN(USDIR)` is computed once per lane in `turb_lanes`, not per ITH as the Fortran text + reads. It is a pure call on the same argument, so the bits are the same; say if "translate, + don't improve" should keep the per-ITH call. + +## How to repeat + +`lanes_1t.sh` of [the L2 entry](2026-10-18-srce-lanes-sin4.md#how-to-repeat), with `repo` set to +the checkout, after `bash scripts/ww3-gpu.sh --ww3-gpu full --cpu --kokkos resident`. + +## Rerun at the combined pin `fac74069`, 2026-10-18 + +Branch `goal/abbb0456-8241-strict-outg-poison` at `debb1b0`, `WW3/` at `fac74069` (strict mode +stops a run when `W3OUTG` falls back to the Fortran), Kokkos OpenMP backend, never `--gpu`. For +`WW_KOKKOS_SRCE_LANES` 1 and 8, `WW_KOKKOS_VA_POISON=1`, with and without `WW_KOKKOS_STRICT=1`: + +```bash +WW_KOKKOS_SRCE_LANES=<1|8> WW_KOKKOS_VA_POISON=1 [WW_KOKKOS_STRICT=1] \ + bash scripts/ww3-gpu.sh --ww3-gpu --cpu --kokkos resident +``` + +| Case | Strict | Lanes 1 and 8 | +|---|---|---| +| `smoke` | off | exit 0; every model file byte-identical to the Fortran of the same executable (`out_grd.ww3`, `out_pnt.ww3`, `restart.ww3`, `ww3.196806.nc`), every file but the NetCDF byte-identical to `--ww3`'s, `nccmp-tol` PASS | +| `smoke` | on | exit 1, the run stops in `W3OUTG` (below); no comparison was made | +| `full` | off | exit 0; `out_grd.ww3`, `ww3.field.nc` byte-identical to the Fortran's, every file but the NetCDF byte-identical to `--ww3`'s, `nccmp-tol` PASS | +| `full` | on | exit 0; the same as off | + +(v) The stop, `smoke`, both lane counts: +`*** WAVEWATCH III ERROR IN W3KOKKOSMD : WW_KOKKOS_STRICT=1 AND PART OUTG WOULD RUN ON THE FORTRAN: +A SKEWNESS FIELD IS REQUESTED (FLOLOC(8,7:9), SKEWNESS), WHICH THE TRANSLATION LEAVES TO THE +FORTRAN W3OUTG`. This is the follow-up of `fac74069` doing what it was written to do: `smoke` +requests a skewness field, which `W3OUTG` leaves to the Fortran. The earlier byte-identical +`smoke` result under strict (above) was at `50a10730`, where strict did not stop on this. The bit +verdict for `smoke` at `fac74069` is the non-strict run; its fallback count printed at exit shows +no call left on the Fortran for `SRCE` and `CPRT` (`W3CPRT` fallbacks 0). `full` is unchanged by +the new pin. Also at the combined HEAD: `just kokkos-test serial-debug` and `openmp-release` exit +0; `just regtest-selected` passed; `bash scripts/ww3-gpu.sh --ww3-gpu smoke --gpu --kokkos +resident --build-only` exit 0, nothing run on a device. diff --git a/docs/log/2026-10-18-strict-outg-poison.md b/docs/log/2026-10-18-strict-outg-poison.md new file mode 100644 index 00000000..3441f2fa --- /dev/null +++ b/docs/log/2026-10-18-strict-outg-poison.md @@ -0,0 +1,88 @@ +# Strict mode sees `W3OUTG`'s Fortran fallback, stops only where a run names its parts, and the missed-sync poison is on by default + +2026-10-18, an agent for M. Hoffmann (goal `abbb0456-8241`), the follow-up of PR #191: findings +1, 2 and 3 of its code review (`ce-code-review` run `20261018-154456-d3a4d3cf`), approved by the +owner the same day. [`PATCH.md`](../../kokkos/src/fortran_iface/PATCH.md) section 31, +[`PRODUCTION_READINESS_202610.md`](../PRODUCTION_READINESS_202610.md) gap 6. + +## Question + +Three gaps in what a pass of `scripts/ww3-gpu.sh` or `just ww3-gpu-matrix` means after #191: + +1. `W3KOKKOS_OUTG` runs the Fortran `W3OUTG` whole for a field its translation leaves to the + Fortran, without telling `WW_KOKKOS_STRICT` (finding 3). Does routing it through + `KOKKOS_FALLBACK` stop a strict run and count a non-strict one, with nothing else changed? +2. `WW_KOKKOS_VA_POISON=1`, the detector of a host reader that misses its sync, was set by no + script (finding 2). Can every Kokkos run have it without changing a correct run's bytes? +3. Under `WW3GPU_KOKKOS=all` the matrix set strict mode too, so every entry whose switch a part + does not cover stopped in `W3INIT` (finding 1). Does an `all` entry without `PR3` run again + when strict mode is set only for a run that names its parts? + +## What was done + +The owner's Intel Core i9-14900 (32 hardware threads, shared with other sessions: load average +about 12 at 17:45), the pinned `ww3` shell (gfortran and g++ 16.2.0), Kokkos' OpenMP backend, at +most 8 threads or jobs, never `--gpu`. + +- **Fork, `ww3-gpu/strict-outg`** (h0ffmann/WW3), one commit `fac74069` on `50a10730`: only + `model/src/w3kokkos_outgmd.F90` changes (PATCH.md section 31). `WW3/` is pinned to it in this + change. +- **`scripts/ww3-gpu.sh`**: `va_poison` (`${WW_KOKKOS_VA_POISON:-1}`) goes into the Kokkos run's + flags in `run_model`; the Fortran run unsets `WW_KOKKOS_VA_POISON`; `--self-test` checks both. +- **`regtests/cmake/ww3gpu_common.cmake`, `ww3gpu_kokkos_env`**: the K run gets + `WW_KOKKOS_VA_POISON=1` (or `0` when `run_entry.cmake` read a caller's `0` before the runs) and + `WW_KOKKOS_STRICT=1` only when its flags do not hold `ALL`; F and G unset both. + `selftest.cmake` has the new expectations. +- **Branch names**: `.gitmodules` and the comments and messages that named older fork branches now + name `ww3-gpu/kokkos-pnt`, which holds the pin. + +Commands, in the order run (the outputs are in the session's scratch directory, not committed): + +```bash +gfortran -E -P > model/src/*.F90 # at 50a10730 and fac74069 +bash scripts/ww3-gpu.sh --ww3-gpu smoke --cpu --kokkos resident # JOBS=8 OMP_NUM_THREADS=8 +bash scripts/ww3-gpu.sh --ww3-gpu full --cpu --kokkos resident +# a scratch copy of build/runs/ww3-gpu-full-openmp/kokkos with TYPE%FIELD%LIST = 'HS SKW', +# ww3_shel with the resident parts, WW_KOKKOS_VA_POISON=1, WW_KOKKOS_OUTG_COUNT=1, with and +# without WW_KOKKOS_STRICT=1 +WW3GPU_KOKKOS=all just ww3-gpu-matrix-list serial +WW3GPU_KOKKOS=all WW3GPU_SWITCH_WORDS="PR0;SHRD" WW3GPU_OMP_THREADS=8 just ww3-gpu-matrix 1 serial -R ww3_tnc1 +WW3GPU_KOKKOS=resident WW3GPU_LAB_SWITCH=ON WW3GPU_OMP_THREADS=2 just ww3-gpu-matrix 4 serial +``` + +## Result + +| Check | Result | +|---|---| +| `gfortran -E -P`, no `KOKKOS`, 143 files | 142 identical apart from blank lines; the other is `w3kokkos_outgmd.F90`, compiled only under `WW_KOKKOS` (v) | +| `smoke`, resident | every model file byte-identical to the Fortran of the same executable, every file but the NetCDF to `--ww3`'s; POISON fills 2161; `calls left on the Fortran: SRCE=0 CPRT=0 OUTG=37` (v) | +| `full`, resident | byte-identical to the Fortran and (but the NetCDF) to `--ww3`; POISON fills 97; `SRCE=0 CPRT=0 OUTG=0`, 13 `W3OUTG` Kokkos calls; ww3_shel 34.46 s Kokkos, 60.41 s Fortran, one run each, regtest pending ⚠ (v) | +| `full` copy with `SKW`, `WW_KOKKOS_STRICT=1` | exit 1 at the first output time: `... PART OUTG WOULD RUN ON THE FORTRAN: A SKEWNESS FIELD IS REQUESTED (FLOLOC(8,7:9), SKEWNESS), ...` (v) | +| `full` copy with `SKW`, strict unset, domain cut to 2 h | exit 0, `End of program`, `SRCE=0 CPRT=0 OUTG=3` (v) | +| `WW3GPU_KOKKOS=all just ww3-gpu-matrix-list serial` | 363 entries; 38 with `PR0;SHRD` (v) | +| `all`, `ww3_tnc1` (entry 38, `PR0 SHRD`) | runs and passes: test.comp 0/30, 10 of 16 parts on Kokkos, `calls left on the Fortran: MAP3=1 MAPT=1 SRCE=1`; before this change strict mode would have stopped it in `W3INIT` (v, run; the stop itself not re-run) | +| resident matrix, 4 serial jobs, merged with `origin/main` at `0cfb595`, pin `fac74069`: 5 of 7 entries PASS (`ww3_tp2.15` x4, `test.comp` 0 of 23, 23, 23, 25 files different, `nccmp-tol` PASS; `ww3_ts2:ST4_PR3_UQ` 0 of 22, no `nccmp-tol` reference), 2 FAIL by design (`ww3_tp2.6` `ST4` and `ST4_ASCII`: unstructured mesh, `W3INIT` stops on `XYP3`, `ran=fortran-only`; the owner keeps them FAIL); every passing Kokkos run ends `calls left on the Fortran: SRCE=0 CPRT=0 OUTG=0`; 729.7 s wall, load average 12 to 25 on the shared machine; the run asked for no `WBT` or `SKW`, so strict mode's new `W3OUTG` stop did not fire in it (v, `build/ww3-gpu-matrix-serial-kokkos-resident-lab/summary.tsv`) | + +Two things the runs show beyond the question: + +- The `smoke` case asks for `WBT` and `SKW`, so its `OUTG` has always taken the Fortran branch; + 37 output times fell back, now counted. A strict run of `smoke` with the resident preset would + stop (not run ⚠). +- The Fortran `SKEWNESS` is slow on `full`: with `SKW` requested, about four output times took 24 + minutes on the shared machine, where the Kokkos run of `HS` alone takes 34 s for all thirteen. + That twelve-hour run was stopped and the counted run repeated over two hours. Not profiled ⚠. + +## What it changes + +- A pass of the matrix under the resident preset now means `W3OUTG` ran on Kokkos at every output + time, and every Kokkos pass of either tool ran with the stale host `VA` poisoned. +- `all` is a coverage run again: what it cannot put on Kokkos is counted in the exit line, not a + stop. +- Units 2 and 3 change a gate's invocation, and unit 1 adds a stop: a person reviews them before + the merge (AGENTS.md, "Never cheat a regtest tool"). +- The matrix command, run in the worktree of this branch: `WW3GPU_KOKKOS=resident + WW3GPU_LAB_SWITCH=ON WW3GPU_OMP_THREADS=2 just ww3-gpu-matrix 4 serial` (exit 8 from the two + `ww3_tp2.6` entries only). Regtest verdict of this commit: the entries above; the full ADR-0016 + validation phase is still the owner's. +- Next: a supervised GPU run reads the POISON count with the others; `smoke`'s `WBT` and `SKW` + stay on the Fortran until those fields are translated. diff --git a/docs/log/README.md b/docs/log/README.md index 5c7c7b79..f787655d 100644 --- a/docs/log/README.md +++ b/docs/log/README.md @@ -72,7 +72,9 @@ Date, author, issue/PR/WFIP. | 2026-10-18 | [`ww3_multi` and the resident `VA`: the defect reproduced, column syncs at the exchanges, no read at `W3WAVE`'s return](2026-10-18-ww3-multi-resident-va.md) | plan 2026-10-18-002 U1 + U7, WFIP-0005 | | 2026-10-18 | [`W3CPRT` on the resident `VA`: partitioning without copying the spectrum back](2026-10-18-cprt-resident.md) | plan 2026-10-18-002 U9, WFIP-0005 | | 2026-10-18 | [U3 on the integrated pin: the resident serial, OMPG, `ww3_multi` and `ww3_tnc1` entries; the restart pair designed, not run](2026-10-18-u3-regtests-v2.md) | plan 2026-10-18-002 U3, #191 | +| 2026-10-18 | [Strict mode sees `W3OUTG`'s Fortran fallback, stops only where a run names its parts, and the missed-sync poison is on by default](2026-10-18-strict-outg-poison.md) | #191 review findings 1 to 3 | | 2026-10-18 | [One thread of the resident `W3SRCE` batch: 91 % in the port's own code, 9 % in libm](2026-10-18-srce-lanes-s0-profile.md) | plan 2026-10-18-005 S0, #197 | | 2026-10-18 | [Handover of goal `abbb0456`: what #191 delivered, the open branches (the review follow-up, plan 006, the Fortran A/B harness), what waits on the owner, the traps met](2026-10-18-goal-abbb0456-handover.md) | PR #191, PR #198 | | 2026-10-18 | [`W3SIN4`'s stress scan with one sea point per vector lane: bit-identical, 5.1 % off the one-thread `W3SRCE` kernel](2026-10-18-srce-lanes-sin4.md) | plan 2026-10-18-005 L2 | | 2026-10-18 | [Handover of goal `c3e7eecf`: the NOAA record and its ten-run median, plan 005 up to L2 (#199), what waits on the owner, the traps met](2026-10-18-goal-c3e7eecf-handover.md) | PR #197, PR #199 | +| 2026-10-18 | [`W3SDS4` with one sea point per vector lane: bit-identical, 19 % off the one-thread `W3SRCE` kernel after L2](2026-10-18-srce-lanes-sds4.md) | plan 2026-10-18-005 L3 | diff --git a/justfile b/justfile index e1e91638..0799a161 100644 --- a/justfile +++ b/justfile @@ -62,7 +62,7 @@ toolchain-smoke: # --------------------------------------------------------------------- # WW3 two ways (scripts/ww3-gpu.sh): --ww3, pure NOAA WW3 from WW3-noaa/; --ww3-gpu, WW3 built -# with the KOKKOS switch from the h0ffmann/WW3 branch ww3-gpu/kokkos-wire-base and linked to kokkos/. +# with the KOKKOS switch from the h0ffmann/WW3 branch ww3-gpu/kokkos-pnt and linked to kokkos/. # --ww3-gpu runs on the CPU (Kokkos OpenMP) unless asked: `just smoke-ww3-gpu gpu` runs CUDA, and only # with a human supervising (AGENTS.md hard rule; it asks at the terminal). # --------------------------------------------------------------------- diff --git a/kokkos/PORT_STATUS.md b/kokkos/PORT_STATUS.md index acf904e5..9d216458 100644 --- a/kokkos/PORT_STATUS.md +++ b/kokkos/PORT_STATUS.md @@ -49,7 +49,7 @@ Profile: full (bench_small, 120 x 80, 9 520 sea points, 12 h), Release -pg, swit | ported | **99.3 %** | 17270 of 263375 (6.56 %) | 63 of 1023 (6.16 %) | 25 of 147 | WW3 at the WW3/ pin calls the Kokkos code, behind its run-time flag | | regtested | **48.6 %** | 2689 of 263375 (1.02 %) | 12 of 1023 (1.17 %) | 6 of 147 | WW3's regtests ran with it on Kokkos and agreed with the Fortran | -The port itself: 60037 lines in 350 files, the C++ under kokkos/src (every .cpp and .hpp line). +The port itself: 60563 lines in 351 files, the C++ under kokkos/src (every .cpp and .hpp line). On a GPU: none yet. The CUDA build links; the RTX 4090 run, `just smoke-ww3-gpu gpu`, is the owner's, supervised. diff --git a/kokkos/port_report/W3SRCE_KOKKOS.md b/kokkos/port_report/W3SRCE_KOKKOS.md index ad155043..12852eda 100644 --- a/kokkos/port_report/W3SRCE_KOKKOS.md +++ b/kokkos/port_report/W3SRCE_KOKKOS.md @@ -237,6 +237,23 @@ lane's operations those of the scalar statement in its order, each lane's stress | `ww3_shel`, `full`, 32 threads, 3 alternated runs each | 15.88 s with 1, 15.47 s with the default 8: 2.6 % less (v, the same log) | | `bash scripts/ww3-gpu.sh --ww3-gpu smoke --gpu --kokkos resident --build-only` | compiles and links, no warning; the lane path is host-only and a device keeps one point per team; nothing run on a device (v) | +### `W3SDS4` with one point per vector lane (plan 005, unit L3) + +`w3srce_lanes` calls `st4::w3sds4_lanes<8>` once per step for the points still in the loop +(`:1303`). `w3sds4_point` is now the parts `sds4_prep`, `sds4_mss`, `sds4_sat`, `sds4_case1`, +`sds4_case2`, `sds4_case3`, `sds4_turb` and `sds4_whitecap` (`kokkos/src/ww_kokkos/st4_sds4.hpp`), +its lines moved without a change (v, a line comparison). The mean-square-slope sums +(`w3src4md.F90:2228-2254`), the saturation integral (`:2262-2290`) and the wave-turbulence and +cumulative loops (`:2533-2575`) run on vectors of eight `REAL`, one point per lane, each lane's +sums its own and in the routine's order; the libmvec sites stay one grouped call per point. + +| Check, 2026-10-18 | Result | +|---|---| +| `smoke`, `full`, `WW_KOKKOS_SRCE_LANES` 1 and 8, strict, poisoned | every model file byte-identical to the Fortran of the same executable; every file but the NetCDF to `--ww3`'s (v) | +| `just regtest-selected` | `ww3_tp2.15` x4 and `ww3_ts2` `ST4_PR3_UQ`: `test.comp` 0 files different, `nccmp-tol` PASS (v) | +| `srce.batch.kernel`, `full`, one thread on a performance core, 3 alternated runs each | 104.99 s with 1 point per team, 80.55 s with 8: 23.3 % less, L2 and L3 together (v, [log](../../docs/log/2026-10-18-srce-lanes-sds4.md)) | +| `--gpu --build-only` | compiles and links, no warning; nothing run on a device (v) | + ## Benchmarks One CPU measurement of the batched path, after the replay above passed on the same tree diff --git a/kokkos/src/fortran_iface/CMakeLists.txt b/kokkos/src/fortran_iface/CMakeLists.txt index 9b2a5549..0fbeba82 100644 --- a/kokkos/src/fortran_iface/CMakeLists.txt +++ b/kokkos/src/fortran_iface/CMakeLists.txt @@ -2,7 +2,7 @@ # The boundary between WAVEWATCH III and the Kokkos kernels: the C ABI # (ww_kokkos_c.hpp + snl1_shim.cpp) here, and the Fortran module that binds it # (w3kokkosmd.F90), which lives in WW3 itself: the h0ffmann/WW3 branch -# ww3-gpu/kokkos-wire-base, at the commit the WW3/ submodule pins. This repository keeps +# ww3-gpu/kokkos-pnt, at the commit the WW3/ submodule pins. This repository keeps # no WW3 Fortran; this build compiles the module from WW3/model/src. # # The C++ half goes into ww_kokkos itself, so it is always built; the benchmark @@ -44,14 +44,14 @@ file(CONFIGURE OUTPUT "${CMAKE_CURRENT_BINARY_DIR}/ww_kokkos_presets.inc" target_include_directories(ww_kokkos PRIVATE "${CMAKE_CURRENT_BINARY_DIR}") set(WW_W3KOKKOSMD "${CMAKE_CURRENT_SOURCE_DIR}/../../../WW3/model/src/w3kokkosmd.F90" - CACHE FILEPATH "WW3's w3kokkosmd.F90 (the WW3/ submodule, branch ww3-gpu/kokkos-wire-base)") + CACHE FILEPATH "WW3's w3kokkosmd.F90 (the WW3/ submodule, branch ww3-gpu/kokkos-pnt)") if(WW_ENABLE_FORTRAN) # Loud, not skipped: a missing module means the WW3/ pin and this C ABI can drift # apart unseen. if(NOT EXISTS "${WW_W3KOKKOSMD}") message(FATAL_ERROR "No ${WW_W3KOKKOSMD}: run `git submodule update --init WW3` " - "(the ww3-gpu/kokkos-wire-base branch of h0ffmann/WW3), or set WW_W3KOKKOSMD, " + "(the ww3-gpu/kokkos-pnt branch of h0ffmann/WW3), or set WW_W3KOKKOSMD, " "or configure with -DWW_ENABLE_FORTRAN=OFF") endif() # A separate target, not sources of ww_kokkos: WW3 compiles the module in its own diff --git a/kokkos/src/fortran_iface/PATCH.md b/kokkos/src/fortran_iface/PATCH.md index 7d095976..9dd8a4ad 100644 --- a/kokkos/src/fortran_iface/PATCH.md +++ b/kokkos/src/fortran_iface/PATCH.md @@ -7,6 +7,7 @@ Read with `git -C WW3 fetch origin && git -C WW3 branch -r` and `git rev-list -- | h0ffmann/WW3 branch | Head | Against the `WW3/` pin | What it wires | |---|---|---|---| +| `ww3-gpu/strict-outg` | `fac74069` | one commit on `50a10730` (the pin before it, also the head of `ww3-gpu/kokkos-pnt`), pushed 2026-10-18 (v, `ls-remote`) | `W3KOKKOS_OUTG`'s call of the Fortran `W3OUTG` reports to `WW_KOKKOS_STRICT` ([section 31](#31-w3outgs-fortran-fallback-reports-to-strict-mode-ww3-gpustrict-outg)) | | `ww3-gpu/integ-fixes` | `711afb80` | one commit on `ww3-gpu/integ-abbb0456` (`b733066b`, the merge of `u1-multi-va`, `u2-strict-preset`, `u8-out-readers` and `u9-cprt` onto `80644ba9`), pushed 2026-10-18 (v, `ls-remote`) | `W3KOKKOS_CPRT`'s two fallbacks to the Fortran `W3CPRT` report to `WW_KOKKOS_STRICT` ([section 28](#28-the-integration-fixes-ww3-gpuinteg-fixes)) | | `ww3-gpu/u1-multi-va` | `f330dc53` | five commits on `80644ba9`, pushed 2026-10-18 (v, `ls-remote`) | `ww3_multi`'s exchanges sync only the columns of `VA` they read or write (`W3KOKKOS_VA_READ_COLS`, `W3KOKKOS_VA_WRITE_COLS`, `ww_va_host_wrote_cols`), no read at `W3WAVE`'s return, the ABI `0.10` ([section 26](#26-ww3_multi-and-the-resident-va-column-syncs-no-read-at-w3waves-return-ww3-gpuu1-multi-va)) | | `ww3-gpu/kokkos-outg` | `e0cc8ee9` | the `WW3/` pin of h0ffmann/ww3-gpu `goal/42-gpu-step-units45`; one commit on `ww3-gpu/kokkos-prop3dev`, pushed 2026-10-17 (v) | `W3OUTG`'s integrals on the device copy of `VA` (`ww_outg_va`, `W3KOKKOS_OUTGMD`), the host read of `VA` at an output time split per reader, the ABI `0.7` ([section 15](#15-w3outg-on-the-resident-spectrum-ww3-gpukokkos-outg)) | @@ -2249,3 +2250,52 @@ binds it on the first sync, `w3kokkos_residentmd.F90:68-70`), so it needs no syn `wav_restart_mod.F90`), which is not in the fork at `a304018c`; `production/GFS.v17` has it, with 16 and 13 uses of `VA` in the last two (v, `git show`). The pin move to that branch ([plan](../../../docs/compound/plans/2026-10-18-004-pin-gfs-v17-plan.md)) lists and syncs them. + +## 31. `W3OUTG`'s Fortran fallback reports to strict mode, `ww3-gpu/strict-outg` + +One fork commit, `fac74069`, on `50a10730`, 2026-10-18; it touches only `w3kokkos_outgmd.F90`. No +ABI change and no new entry point: `KOKKOS_FALLBACK` and `ww_kokkos_fallback_slot` are section 24's. +Section 30 is held by the unstructured-propagation port (h0ffmann/ww3-gpu#196). + +`W3KOKKOS_OUTG` runs the Fortran `W3OUTG` whole, after a host read of `VA`, when a requested field +is one its translation leaves to the Fortran: U3 Stokes (`FLOLOC(6,8)`, `(6,12)`, +`CALC_U3STOKES`), skewness (`FLOLOC(8,7:9)`, `SKEWNESS`), `WBT` (`FLOLOC(2,17)`, `CALC_WBT`), or a +requested 3-D or `P2SMS` field that is not allocated (`E3DF(1,K)` or `P2MSF(1)` `<= 0`). Section 24 +left that branch out: under `WW_KOKKOS_STRICT=1` it neither stopped nor counted, and the parts line +showed `OUTG=K` while the Fortran wrote the fields (the code review of #191, finding 3). + +- **The branch calls `KOKKOS_FALLBACK ( 'OUTG', REASON )`** before `W3KOKKOS_VA_HOST_READ` + (`w3kokkos_outgmd.F90:138-157` at `fac74069`), with one reason per kind of field, in the order + above: `A U3 STOKES FIELD IS REQUESTED ...`, `A SKEWNESS FIELD IS REQUESTED ...`, `THE WBT FIELD + IS REQUESTED ...`, `A REQUESTED 3-D OR P2SMS FIELD IS NOT ALLOCATED ...`. With + `WW_KOKKOS_STRICT=1` the model stops on `... PART OUTG WOULD RUN ON THE FORTRAN: `; + without it each output time that falls back is counted. +- **`W3KOKKOS_OUTG_SETUP` makes the `OUTG` counter** when the part is on (`:89-91`), as + `W3KOKKOS_CPRT_SETUP` does, so the line at exit shows `OUTG=0` when nothing fell back. +- **Nothing changes on the kernel path or per point.** The condition `FORT` is the one before; + only its `THEN` branch gained the call. +- **A build without `KOKKOS` is unchanged.** `gfortran -E -P` (16.2.0, the `#ww3` shell's + compiler) of all 143 `model/src/*.F90` at `50a10730` and at `fac74069`, with `switch_lab_shrd`'s + words as `-DW3_`, blank lines dropped: 142 identical, and the one that differs is + `w3kokkos_outgmd.F90`, which `model/src/CMakeLists.txt` compiles only under `WW_KOKKOS` (its + glob of `w3kokkos_*md.F90`, `:25-85`) (v). With `-DW3_KOKKOS` that file gains 30 lines and loses + only the four it rewrites: two lines of the header comment and the two `USE` lines that now + import `C_NULL_CHAR`, `KOKKOS_FALLBACK` and `WW_KOKKOS_FALLBACK_SLOT` (v). + +The stop, shown on a scratch copy of the `full` case's Kokkos run directory whose `ww3_shel.nml` +asks for `TYPE%FIELD%LIST = 'HS SKW'` (the fixture `kokkos/tests/fixtures/bench_small` is +unchanged), with the parts of the `resident` preset, `WW_KOKKOS_VA_POISON=1` and +`WW_KOKKOS_OUTG_COUNT=1`, 2026-10-18 (OpenMP backend, 8 threads, gfortran 16.2.0; never `--gpu`): + +| Run | Result | +|---|---| +| `WW_KOKKOS_STRICT=1` | `ww3_shel` exit 1 at the first output time, `*** WAVEWATCH III ERROR IN W3KOKKOSMD : WW_KOKKOS_STRICT=1 AND PART OUTG WOULD RUN ON THE FORTRAN: A SKEWNESS FIELD IS REQUESTED (FLOLOC(8,7:9), SKEWNESS), WHICH THE TRANSLATION LEAVES TO THE FORTRAN W3OUTG` (v) | +| strict unset, the domain cut to `20200101 020000` (three output times; the Fortran `SKEWNESS` at every point made the twelve-hour run take over an hour on the shared machine, so it was stopped) | `ww3_shel` exit 0, `End of program`, `calls left on the Fortran (WW_KOKKOS_STRICT=1 stops instead): SRCE=0 CPRT=0 OUTG=3`, `W3OUTG Kokkos calls: 0` (v) | + +The `smoke` case (`ww3_ts1`) asks for `WBT` and `SKW` among its fields, so its `OUTG` has always +taken this branch: `scripts/ww3-gpu.sh --ww3-gpu smoke --cpu --kokkos resident` at `fac74069` ends +with `calls left on the Fortran (...): SRCE=0 CPRT=0 OUTG=37`, one per output time, with every +model file byte-identical to the Fortran of the same executable and every file but the NetCDF to +`--ww3`'s (v). At `50a10730` no `OUTG` counter exists, so the line names `SRCE` and `CPRT` alone (v, by reading the set-up; that run was not repeated). The `full` case asks for `HS` +alone and ends with `OUTG=0` and 13 `W3OUTG` Kokkos calls (v). + diff --git a/kokkos/src/fortran_iface/resident_c.hpp b/kokkos/src/fortran_iface/resident_c.hpp index a3a4eb6b..de6872ee 100644 --- a/kokkos/src/fortran_iface/resident_c.hpp +++ b/kokkos/src/fortran_iface/resident_c.hpp @@ -4,7 +4,7 @@ // write of ww3_multi's exchanges, section 26): declarations only. // resident_shim.cpp implements the ww_va_* calls, srce_shim.cpp ww_srce_batch_va; // model/src/w3kokkos_residentmd.F90 and w3kokkosmd.F90 on the fork -// (h0ffmann/WW3 ww3-gpu/kokkos-resident) mirror them. The rules of ww_kokkos_c.hpp hold: +// (h0ffmann/WW3 ww3-gpu/kokkos-pnt) mirror them. The rules of ww_kokkos_c.hpp hold: // Fortran order, caller-owned arrays, float is WW3's REAL, nothing throws. // // One context per grid `igrid` (W3GDATMD's IGRID): VA(NSPEC,0:NSEALM) of that grid, diff --git a/kokkos/src/fortran_iface/srce_shim.cpp b/kokkos/src/fortran_iface/srce_shim.cpp index b6bcb744..2aca2e4e 100644 --- a/kokkos/src/fortran_iface/srce_shim.cpp +++ b/kokkos/src/fortran_iface/srce_shim.cpp @@ -902,7 +902,7 @@ void ww_srce_batch(int igrid, int it, int inflags24, int flogrd57, int flogrd58, } // ww_srce_batch_va (resident_c.hpp), called by W3KOKKOS_SRCE_BATCH on the fork branch -// ww3-gpu/kokkos-resident: ww_srce_batch with the spectra read and written in the resident +// ww3-gpu/kokkos-pnt: ww_srce_batch with the spectra read and written in the resident // VA (WFIP-0005 unit 2). The arguments are ww_srce_batch's, spec replaced by: // | C | Fortran | Type | Intent | Shape | // | va | VA(1,0) (W3WDATMD) | float* | inout | nspec x nsealm+1 | diff --git a/kokkos/src/ww_kokkos/srce_kernel.hpp b/kokkos/src/ww_kokkos/srce_kernel.hpp index 0681fdea..2097a735 100644 --- a/kokkos/src/ww_kokkos/srce_kernel.hpp +++ b/kokkos/src/ww_kokkos/srce_kernel.hpp @@ -311,14 +311,27 @@ struct SrcePoint { sin4(); // :2007 step_c_end(); } + KOKKOS_INLINE_FUNCTION void sds4() const { + st4::w3sds4_point(q, f, spec, wn1, cg1, o.ustar, o.ustdir, st.depth, in.dair, + st.dlwmean, in.flogrd57, in.flogrd58, vsds, vdds, brlambda, whitecap, + st4w, st4iw); // :1303 + } + /// The same arguments, for st4::w3sds4_lanes. + KOKKOS_INLINE_FUNCTION st4::Sds4Args sds4_args() const { + return st4::Sds4Args{spec, wn1, cg1, o.ustar, o.ustdir, st.depth, in.dair, st.dlwmean, + in.flogrd57, in.flogrd58, vsds, vdds, brlambda, whitecap}; + } + KOKKOS_INLINE_FUNCTION st4::Sds4Work sds4_work() const { return st4::sds4_layout(q, st4w, st4iw); } + /// The step after the DIA, W3SDS4 (:1303) and then step_c_after_sds4. KOKKOS_INLINE_FUNCTION void step_c_main() const { + sds4(); + step_c_after_sds4(); + } + KOKKOS_INLINE_FUNCTION void step_c_after_sds4() const { using srce_detail::mn; using srce_detail::mni; using srce_detail::mx; using srce_detail::mxi; - st4::w3sds4_point(q, f, spec, wn1, cg1, o.ustar, o.ustdir, st.depth, in.dair, - st.dlwmean, in.flogrd57, in.flogrd58, vsds, vdds, brlambda, whitecap, - st4w, st4iw); // :1303 bool lbreak = false; w3sdb1_point(p, g.sig, g.dden, spec, st.depth, st.emean, st.fmean, o.wnmean, cg1, lbreak, vsdb, vddb, eb); // :1321; LBREAK is not read by W3SRCE @@ -584,6 +597,17 @@ KOKKOS_INLINE_FUNCTION void w3srce_lanes(const TeamMember& team, const snl1::Con st4::w3sin4_lanes(x[on[0]].q, x[on[0]].f, n, a, cg, k, u, ustar, drat, as, usdir, z0, tauwx, tauwy, tauwnx, tauwny, s, d, llws, work); }; + // W3SDS4 for the points listed in `on`, one point per vector lane (st4::w3sds4_lanes). + const auto sds4_lanes = [&](const int* on, const int n) { + if (n == 0) return; + st4::Sds4Args pt[kSrceMaxLanes]; + st4::Sds4Work wk[kSrceMaxLanes]; + for (int j = 0; j < n; ++j) { + pt[j] = x[on[j]].sds4_args(); + wk[j] = x[on[j]].sds4_work(); + } + st4::w3sds4_lanes(x[on[0]].q, x[on[0]].f, n, pt, wk); + }; while (true) { // st.done changes only in step_c_end, so the points in the loop are the same for the whole // step. @@ -601,7 +625,8 @@ KOKKOS_INLINE_FUNCTION void w3srce_lanes(const TeamMember& team, const snl1::Con team.team_barrier(); } Kokkos::single(Kokkos::PerTeam(team), [&]() { - for (int j = 0; j < non; ++j) x[on[j]].step_c_main(); + sds4_lanes(on, non); // :1303 + for (int j = 0; j < non; ++j) x[on[j]].step_c_after_sds4(); sin4_lanes(on, non); // :2007 for (int j = 0; j < non; ++j) x[on[j]].step_c_end(); }); diff --git a/kokkos/src/ww_kokkos/st4_sds4.hpp b/kokkos/src/ww_kokkos/st4_sds4.hpp index 567e42a7..ed8c811c 100644 --- a/kokkos/src/ww_kokkos/st4_sds4.hpp +++ b/kokkos/src/ww_kokkos/st4_sds4.hpp @@ -25,6 +25,7 @@ #include "st4_arrays.hpp" #include "st4_libmvec.hpp" #include "st4_params.hpp" +#include "st4_vec8.hpp" namespace ww::st4 { @@ -64,48 +65,88 @@ inline int sds4_work_reals(const Params& p) { return 18 * p.nk + 3 * p.nspec + p /// Ints: NTIMES and IKSUP, NK each. NSMOOTH (:2163) is zeroed and never read: not kept. inline int sds4_work_ints(const Params& p) { return 2 * p.nk; } -/// W3SDS4 for one point. a = A(1:NSPEC), k = K(1:NK), cg = CG(1:NK), srhs, -/// ddiag, brlambda = (1:NSPEC), whitecap = WHITECAP(1:4), all 0-based. -/// flogrd57, flogrd58 are FLOGRD(5,7), FLOGRD(5,8) as 0/1. IX, IY are not -/// read and not passed. The OUT arrays arrive holding the caller's values: -/// the early return at w3src4md.F90:2581 leaves WHITECAP as it was. -KOKKOS_INLINE_FUNCTION void w3sds4_point(const Params& p, const Fields& f, const Real* a, - const Real* k, const Real* cg, Real ustar, Real usdir, - Real depth, Real dair, Real dlwmean, int flogrd57, - int flogrd58, Real* srhs, Real* ddiag, Real* brlambda, - Real* whitecap, Real* work, int* iwork) { - using ww::st4::max_r; - using ww::st4::min_r; +/// The arguments of w3sds4_point that its parts read and write (IX, IY are not read and not +/// passed), so that the same parts serve one point (w3sds4_point) and the points of a team +/// (w3sds4_lanes). +struct Sds4Args { + const Real* a; + const Real* k; + const Real* cg; + Real ustar, usdir, depth, dair, dlwmean; + int flogrd57, flogrd58; + Real *srhs, *ddiag, *brlambda, *whitecap; +}; + +/// The per-point work arrays of w3sds4_point, by name (the layout of sds4_work_reals). +struct Sds4Work { + Real *coef4, *coef5, *s1, *e1, *dk, *hs, *kbar, *dck, *efdf, *bth0, *qb, *s2, *brm12, *msssum, + *bth, *pb, *pb2, *lmodulation; + int *ntimes, *iksup; +}; + +/// The locals that cross from section 1 to the later sections. +struct Sds4Mid { + Real facturb; + int dikcumul, ik1; +}; + +KOKKOS_INLINE_FUNCTION Sds4Work sds4_layout(const Params& p, Real* work, int* iwork) { + const int nk = p.nk, nspec = p.nspec; + Sds4Work w; + w.coef4 = work; + w.coef5 = w.coef4 + nk; + w.s1 = w.coef5 + nk; + w.e1 = w.s1 + nk; + w.dk = w.e1 + nk; + w.hs = w.dk + nk; + w.kbar = w.hs + nk; + w.dck = w.kbar + nk; + w.efdf = w.dck + nk; + w.bth0 = w.efdf + nk; + w.qb = w.bth0 + nk; + w.s2 = w.qb + nk; + w.brm12 = w.s2 + nk; + w.msssum = w.brm12 + nk; // MSSSUM(IK,J) = msssum[(J-1)*nk + IK-1] + w.bth = w.msssum + 5 * nk; + w.pb = w.bth + nspec; + w.pb2 = w.pb + nspec; + w.lmodulation = w.pb2 + nspec; + w.ntimes = iwork; + w.iksup = w.ntimes + nk; + return w; +} + +/// Sections 0 and 1 of W3SDS4 (:2203-2224): the pre-initialisation, FACTURB, DIKCUMUL, IK1. +KOKKOS_INLINE_FUNCTION void sds4_prep(const Params& p, const Fields& f, const Sds4Args& pt, + const Sds4Work& wk, Sds4Mid& mid) { + (void)f; using ww::st4::nint; const Real zero = static_cast(0.); const Real one = static_cast(1.); - const int nk = p.nk, nth = p.nth, nspec = p.nspec; - const Real dth = p.dth; + const int nk = p.nk; + const int nspec = p.nspec; const Real* ssdsc = p.ssdsc; // SSDSC(i) is ssdsc[i-1] - - // Work storage, 1-based helpers below take the Fortran index. - Real* coef4 = work; - Real* coef5 = coef4 + nk; - Real* s1 = coef5 + nk; - Real* e1 = s1 + nk; - Real* dk = e1 + nk; - Real* hs = dk + nk; - Real* kbar = hs + nk; - Real* dck = kbar + nk; - Real* efdf = dck + nk; - Real* bth0 = efdf + nk; - Real* qb = bth0 + nk; - Real* s2 = qb + nk; - Real* brm12 = s2 + nk; - Real* msssum = brm12 + nk; // MSSSUM(IK,J) = msssum[(J-1)*nk + IK-1] - Real* bth = msssum + 5 * nk; - Real* pb = bth + nspec; - Real* pb2 = pb + nspec; - Real* lmodulation = pb2 + nspec; - int* ntimes = iwork; - int* iksup = ntimes + nk; - auto MSS = [&](int ik, int j) -> Real& { return msssum[(j - 1) * nk + (ik - 1)]; }; - + const Real ustar = pt.ustar; + const Real dair = pt.dair; + Real* const srhs = pt.srhs; + Real* const ddiag = pt.ddiag; + Real* const s1 = wk.s1; + Real* const e1 = wk.e1; + Real* const dk = wk.dk; + Real* const hs = wk.hs; + Real* const kbar = wk.kbar; + Real* const dck = wk.dck; + Real* const efdf = wk.efdf; + Real* const bth0 = wk.bth0; + Real* const qb = wk.qb; + Real* const s2 = wk.s2; + Real* const brm12 = wk.brm12; + Real* const msssum = wk.msssum; + Real* const bth = wk.bth; + Real* const pb = wk.pb; + Real* const pb2 = wk.pb2; + int* const ntimes = wk.ntimes; + int* const iksup = wk.iksup; // 0. Pre-initialization (:2203-2213). for (int i = 0; i < nk; ++i) { s1[i] = zero; e1[i] = zero; ntimes[i] = 0; iksup[i] = 0; @@ -121,10 +162,26 @@ KOKKOS_INLINE_FUNCTION void w3sds4_point(const Params& p, const Fields& f, const const Real facturb = (((ssdsc[4] * (ustar * ustar)) / kGrav) * dair) / kDwat; // DIKCUMUL is a W3SRC4MD module variable written here (:2218); kept local, not written back. const int dikcumul = nint(p.ssdsbrf1 / (p.xfr - one)); - Real renewalfreq = zero; int ik1 = 1; if (p.ig1) ik1 = nint(p.igpars5) + 1; // W3_IG1, :2223 + mid.facturb = facturb; + mid.dikcumul = dikcumul; + mid.ik1 = ik1; +} +/// Section 1.b (:2228-2254): the mean-square slope for the modulation of lambda. +KOKKOS_INLINE_FUNCTION void sds4_mss(const Params& p, const Fields& f, const Sds4Args& pt, + const Sds4Work& wk, const Sds4Mid& mid) { + (void)mid; + const Real zero = static_cast(0.); + const int nk = p.nk; + const int nth = p.nth; + const Real* ssdsc = p.ssdsc; // SSDSC(i) is ssdsc[i-1] + const Real* const a = pt.a; + const Real* const k = pt.k; + const Real* const cg = pt.cg; + Real* const msssum = wk.msssum; + auto MSS = [&](int ik, int j) -> Real& { return msssum[(j - 1) * nk + (ik - 1)]; }; // 1.b MSS parameters for the modulation of lambda, Romero (:2228-2254). if (ssdsc[7] > zero || ssdsc[10] > zero || ssdsc[17] > zero) { for (int ik = 1; ik <= nk; ++ik) { @@ -147,286 +204,390 @@ KOKKOS_INLINE_FUNCTION void w3sds4_point(const Params& p, const Fields& f, const MSS(ik, 2) = mssd; } } +} - // 2. Spontaneous breaking from local saturation (:2259). - switch (nint(ssdsc[0])) { - case 1: { - // 2.a Direction-dependent breaking, Ardhuin et al. 2010 (:2260-2331). - const Real epsr = Kokkos::sqrt(p.ssdsbr); - for (int i = 0; i < nspec; ++i) bth[i] = zero; - // 2.a.1 Saturation. - const int nsat = static_cast(f.satweights.extent(0)); - for (int ik = ik1; ik <= nk; ++ik) { - const Real kk = k[ik - 1]; - const Real facsat = (f.sig(ik) * ((kk * kk) * kk)) * dth; - const int is0 = (ik - 1) * nth; - bth[is0] = zero; - Real asum = zero; - for (int ith = 1; ith <= nth; ++ith) asum = asum + a[is0 + ith - 1]; - bth0[ik - 1] = asum * facsat; - if (p.ssdsdth >= static_cast(180)) { - for (int i = is0; i < is0 + nth; ++i) bth[i] = bth0[ik - 1]; - } else { - // Full length 2*(MTH/2-1)+1 for any SDSNTH (:2283): the extra terms have weight 0 - // and index 1, and 0*x is not 0 for an infinite x; the length sets the order. - // J outside, ITH inside: each BTH(IS) still adds its products in the order - // J = 1..NSAT, starting from 0, so its bits are DOT_PRODUCT's; only the NTH sums are - // interleaved, which removes the add-latency chain of one sum. bth and a are distinct - // work arrays, so the ITH loop carries no dependence. - // The transposed tables make the ITH loop's loads contiguous; without them (a Fields not - // built by st4_shim) the same products are read from the originals. - for (int ith = 0; ith < nth; ++ith) bth[is0 + ith] = zero; - const bool tr = f.satweights_t.extent(1) == static_cast(nsat); - for (int j = 0; j < nsat; ++j) { - if (tr) { - const Real* const w = &f.satweights_t(0, j); - const int* const x = &f.satindices_t(0, j); +/// Case 1, 2.a.1 (:2262-2290): the saturation BTH and BTH0. +KOKKOS_INLINE_FUNCTION void sds4_sat(const Params& p, const Fields& f, const Sds4Args& pt, + const Sds4Work& wk, const Sds4Mid& mid) { + const Real zero = static_cast(0.); + const int nk = p.nk; + const int nth = p.nth; + const int nspec = p.nspec; + const Real dth = p.dth; + const Real* const a = pt.a; + const Real* const k = pt.k; + Real* const bth0 = wk.bth0; + Real* const bth = wk.bth; + const int ik1 = mid.ik1; + for (int i = 0; i < nspec; ++i) bth[i] = zero; + // 2.a.1 Saturation. + const int nsat = static_cast(f.satweights.extent(0)); + for (int ik = ik1; ik <= nk; ++ik) { + const Real kk = k[ik - 1]; + const Real facsat = (f.sig(ik) * ((kk * kk) * kk)) * dth; + const int is0 = (ik - 1) * nth; + bth[is0] = zero; + Real asum = zero; + for (int ith = 1; ith <= nth; ++ith) asum = asum + a[is0 + ith - 1]; + bth0[ik - 1] = asum * facsat; + if (p.ssdsdth >= static_cast(180)) { + for (int i = is0; i < is0 + nth; ++i) bth[i] = bth0[ik - 1]; + } else { + // Full length 2*(MTH/2-1)+1 for any SDSNTH (:2283): the extra terms have weight 0 + // and index 1, and 0*x is not 0 for an infinite x; the length sets the order. + // J outside, ITH inside: each BTH(IS) still adds its products in the order + // J = 1..NSAT, starting from 0, so its bits are DOT_PRODUCT's; only the NTH sums are + // interleaved, which removes the add-latency chain of one sum. bth and a are distinct + // work arrays, so the ITH loop carries no dependence. + // The transposed tables make the ITH loop's loads contiguous; without them (a Fields not + // built by st4_shim) the same products are read from the originals. + for (int ith = 0; ith < nth; ++ith) bth[is0 + ith] = zero; + const bool tr = f.satweights_t.extent(1) == static_cast(nsat); + for (int j = 0; j < nsat; ++j) { + if (tr) { + const Real* const w = &f.satweights_t(0, j); + const int* const x = &f.satindices_t(0, j); #if defined(_OPENMP) && !defined(__CUDACC__) && !defined(__HIPCC__) #pragma omp simd #endif - for (int ith = 0; ith < nth; ++ith) - bth[is0 + ith] = bth[is0 + ith] + w[ith] * a[is0 + x[ith] - 1]; - } else { - for (int ith = 0; ith < nth; ++ith) - bth[is0 + ith] = bth[is0 + ith] + f.satweights(j, ith) * a[is0 + f.satindices(j, ith) - 1]; - } - } - for (int ith = 0; ith < nth; ++ith) bth[is0 + ith] = bth[is0 + ith] * facsat; - // MAXVAL (:2287), as gfortran lowers it: first element >= -Inf, then strictly greater. - Real lim = -Kokkos::Experimental::infinity_v; - int i = is0; - bool found = false; - for (; i < is0 + nth; ++i) { - if (bth[i] >= lim) { lim = bth[i]; ++i; found = true; break; } - } - if (!found) lim = Kokkos::Experimental::quiet_NaN_v; - for (; i < is0 + nth; ++i) - if (bth[i] > lim) lim = bth[i]; - bth0[ik - 1] = lim; - } - } - // 2.a.2 Spontaneous breaking dissipation rate. - for (int ik = ik1; ik <= nk; ++ik) { - Real miche; - if (p.ssdsbm[0] == one) { - miche = one; + for (int ith = 0; ith < nth; ++ith) + bth[is0 + ith] = bth[is0 + ith] + w[ith] * a[is0 + x[ith] - 1]; } else { - const Real x = dm::tanhf(min_r(k[ik - 1] * depth, static_cast(10.))); - const Real t = - x * (p.ssdsbm[1] + x * (p.ssdsbm[2] + x * (p.ssdsbm[3] + x * p.ssdsbm[4]))); - miche = t * t; + for (int ith = 0; ith < nth; ++ith) + bth[is0 + ith] = bth[is0 + ith] + f.satweights(j, ith) * a[is0 + f.satindices(j, ith) - 1]; } - const Real coef1 = p.ssdsbr * miche; - const Real m = max_r(bth0[ik - 1] / coef1 - one, zero); - const Real sdiagiso = ((ssdsc[1] * f.sig(ik)) * ssdsc[5]) * (m * m); - const Real coef2 = ((ssdsc[1] * f.sig(ik)) * (one - ssdsc[5])) / (coef1 * coef1); - // LIBMVEC SITE: -O3 uses _ZGVbN4vv_powf here (INVENTORY.md, :2315), lane i = ITH-1. - const int is0 = (ik - 1) * nth; - const Real ssdsp = p.ssdsp; - mvec_powf( - nth, [&](int i) { return max_r(bth[is0 + i] - coef1, zero); }, - [&](int) { return ssdsp; }, - [&](int i, Real v) { ddiag[is0 + i] = sdiagiso + coef2 * v; }); } - // Breaking probability (:2320-2327), lambda (:2330), source (:2331). - for (int i = 0; i < nspec; ++i) { - const Real m = max_r(Kokkos::sqrt(bth[i]) - epsr, zero); - pb[i] = m * m; + for (int ith = 0; ith < nth; ++ith) bth[is0 + ith] = bth[is0 + ith] * facsat; + // MAXVAL (:2287), as gfortran lowers it: first element >= -Inf, then strictly greater. + Real lim = -Kokkos::Experimental::infinity_v; + int i = is0; + bool found = false; + for (; i < is0 + nth; ++i) { + if (bth[i] >= lim) { lim = bth[i]; ++i; found = true; break; } } - for (int i = 0; i < nspec; ++i) pb[i] = pb[i] * static_cast(28.16); - for (int i = 0; i < nspec; ++i) brlambda[i] = pb[i] / sds4_detail::kTwoPiSq; - for (int i = 0; i < nspec; ++i) srhs[i] = ddiag[i] * a[i]; - break; + if (!found) lim = Kokkos::Experimental::quiet_NaN_v; + for (; i < is0 + nth; ++i) + if (bth[i] > lim) lim = bth[i]; + bth0[ik - 1] = lim; } - case 2: { - // 2.b Spontaneous breaking for T500, Filipot et al. JGR 2010 (:2338-2477). - for (int i = 0; i < nk; ++i) { e1[i] = zero; hs[i] = zero; } - for (int i = 0; i < nspec; ++i) { srhs[i] = zero; ddiag[i] = zero; pb2[i] = zero; } - for (int ik = ik1; ik <= nk; ++ik) { - e1[ik - 1] = zero; - for (int ith = 1; ith <= nth; ++ith) { - const int is = ith + (ik - 1) * nth; - e1[ik - 1] = e1[ik - 1] + (a[is - 1] * f.sig(ik)) * dth; - } - dk[ik - 1] = f.dden(ik - 1) / ((dth * f.sig(ik)) * cg[ik - 1]); - } - int id = nint(depth); - if (id > kNdtab) id = kNdtab; - if (id < 1) { - id = 1; - } else if (id > kNdtab) { - id = kNdtab; - } - for (int i = 0; i < nk; ++i) { hs[i] = zero; efdf[i] = zero; kbar[i] = zero; } - int nkl = 0; - for (int ikl = 1; ikl <= nk; ++ikl) { - iksup[ikl - 1] = f.iktab(ikl - 1, id - 1); - if (iksup[ikl - 1] <= nk) { - Real v = zero; - for (int i = ikl; i <= iksup[ikl - 1] - 1; ++i) v = v + e1[i - 1] * dk[i - 1]; - efdf[ikl - 1] = v; - if (efdf[ikl - 1] != zero) { - Real w = zero; - for (int i = ikl; i <= iksup[ikl - 1] - 1; ++i) - w = w + (k[i - 1] * e1[i - 1]) * dk[i - 1]; - kbar[ikl - 1] = w / efdf[ikl - 1]; - } else { - kbar[ikl - 1] = zero; - } - hs[ikl - 1] = static_cast(4) * Kokkos::sqrt(efdf[ikl - 1]); - nkl = nkl + 1; - } - } - // Dissipation and breaking probability per scale. The loop runs to the count NKL, - // not over the scales that passed IKSUP <= NK (:2392). - for (int i = 0; i < nk; ++i) { dck[i] = zero; qb[i] = zero; } - for (int ikl = 1; ikl <= nkl; ++ikl) { - const Real kb = kbar[ikl - 1]; - if (hs[ikl - 1] != zero && kb != zero) { - // FAC_KD2 + ANINT(...) is REAL, truncated on assignment to INTEGER. - int ikd = fint(Kokkos::round(dm::logf(kb * depth) / sds4_detail::kLogFacKd1) + - static_cast(kFacKd2)); - int ikhs = fint(one + Kokkos::round((kb * hs[ikl - 1]) / sds4_detail::kDkhs)); - if (ikd > kNkd) { - ikd = kNkd; - } else if (ikd < 1) { - ikd = 1; - } - if (ikhs > kNkhs) { - ikhs = kNkhs; - } else if (ikhs < 1) { - ikhs = 1; - } - const Real xt = dm::tanhf(kb * depth); - // GAM (:2411) is computed and never used. - const Real gam = ((((xt * xt) * xt) * static_cast(1.0314) - - (xt * xt) * static_cast(1.9958)) + - xt * static_cast(1.5522)) + - static_cast(0.1885); - (void)gam; - dck[ikl - 1] = (dm::powf(kb, static_cast(-2.5)) * (kb / sds4_detail::kTwoPi)) * - f.dcki(ikhs - 1, ikd - 1); - qb[ikl - 1] = f.qbi(ikhs - 1, ikd - 1); - } else { - dck[ikl - 1] = zero; - qb[ikl - 1] = zero; - } - } - // Distributes scale dissipation over the frequency spectrum (:2429-2452). - for (int i = 0; i < nk; ++i) { s1[i] = zero; s2[i] = zero; ntimes[i] = 0; } - for (int ikl = 1; ikl <= nkl; ++ikl) { - if (efdf[ikl - 1] > zero) { - for (int i = ikl; i <= iksup[ikl - 1]; ++i) - s1[i - 1] = s1[i - 1] + (dck[ikl - 1] * e1[i - 1]) / efdf[ikl - 1]; - for (int i = ikl; i <= iksup[ikl - 1]; ++i) - s2[i - 1] = s2[i - 1] + (qb[ikl - 1] * e1[i - 1]) / efdf[ikl - 1]; - for (int i = ikl; i <= iksup[ikl - 1]; ++i) ntimes[i - 1] = ntimes[i - 1] + 1; - } + } +} + +/// Case 1, 2.a.2 and the rest (:2291-2331): the dissipation rate, PB, BRLAMBDA, SRHS. +KOKKOS_INLINE_FUNCTION void sds4_case1(const Params& p, const Fields& f, const Sds4Args& pt, + const Sds4Work& wk, const Sds4Mid& mid) { + using ww::st4::max_r; + using ww::st4::min_r; + const Real zero = static_cast(0.); + const Real one = static_cast(1.); + const int nk = p.nk; + const int nth = p.nth; + const int nspec = p.nspec; + const Real* ssdsc = p.ssdsc; // SSDSC(i) is ssdsc[i-1] + const Real* const a = pt.a; + const Real* const k = pt.k; + const Real depth = pt.depth; + Real* const srhs = pt.srhs; + Real* const ddiag = pt.ddiag; + Real* const brlambda = pt.brlambda; + Real* const bth0 = wk.bth0; + Real* const bth = wk.bth; + Real* const pb = wk.pb; + const int ik1 = mid.ik1; + const Real epsr = Kokkos::sqrt(p.ssdsbr); + // 2.a.2 Spontaneous breaking dissipation rate. + for (int ik = ik1; ik <= nk; ++ik) { + Real miche; + if (p.ssdsbm[0] == one) { + miche = one; + } else { + const Real x = dm::tanhf(min_r(k[ik - 1] * depth, static_cast(10.))); + const Real t = + x * (p.ssdsbm[1] + x * (p.ssdsbm[2] + x * (p.ssdsbm[3] + x * p.ssdsbm[4]))); + miche = t * t; + } + const Real coef1 = p.ssdsbr * miche; + const Real m = max_r(bth0[ik - 1] / coef1 - one, zero); + const Real sdiagiso = ((ssdsc[1] * f.sig(ik)) * ssdsc[5]) * (m * m); + const Real coef2 = ((ssdsc[1] * f.sig(ik)) * (one - ssdsc[5])) / (coef1 * coef1); + // LIBMVEC SITE: -O3 uses _ZGVbN4vv_powf here (INVENTORY.md, :2315), lane i = ITH-1. + const int is0 = (ik - 1) * nth; + const Real ssdsp = p.ssdsp; + mvec_powf( + nth, [&](int i) { return max_r(bth[is0 + i] - coef1, zero); }, + [&](int) { return ssdsp; }, + [&](int i, Real v) { ddiag[is0 + i] = sdiagiso + coef2 * v; }); + } + // Breaking probability (:2320-2327), lambda (:2330), source (:2331). + for (int i = 0; i < nspec; ++i) { + const Real m = max_r(Kokkos::sqrt(bth[i]) - epsr, zero); + pb[i] = m * m; + } + for (int i = 0; i < nspec; ++i) pb[i] = pb[i] * static_cast(28.16); + for (int i = 0; i < nspec; ++i) brlambda[i] = pb[i] / sds4_detail::kTwoPiSq; + for (int i = 0; i < nspec; ++i) srhs[i] = ddiag[i] * a[i]; +} + +/// Case 2, 2.b (:2338-2477), unchanged. +KOKKOS_INLINE_FUNCTION void sds4_case2(const Params& p, const Fields& f, const Sds4Args& pt, + const Sds4Work& wk, const Sds4Mid& mid) { + using ww::st4::nint; + const Real zero = static_cast(0.); + const Real one = static_cast(1.); + const int nk = p.nk; + const int nth = p.nth; + const int nspec = p.nspec; + const Real dth = p.dth; + const Real* ssdsc = p.ssdsc; // SSDSC(i) is ssdsc[i-1] + const Real* const a = pt.a; + const Real* const k = pt.k; + const Real* const cg = pt.cg; + const Real depth = pt.depth; + Real* const srhs = pt.srhs; + Real* const ddiag = pt.ddiag; + Real* const brlambda = pt.brlambda; + Real* const s1 = wk.s1; + Real* const e1 = wk.e1; + Real* const dk = wk.dk; + Real* const hs = wk.hs; + Real* const kbar = wk.kbar; + Real* const dck = wk.dck; + Real* const efdf = wk.efdf; + Real* const bth0 = wk.bth0; + Real* const qb = wk.qb; + Real* const s2 = wk.s2; + Real* const pb = wk.pb; + Real* const pb2 = wk.pb2; + int* const ntimes = wk.ntimes; + int* const iksup = wk.iksup; + const int ik1 = mid.ik1; + // 2.b Spontaneous breaking for T500, Filipot et al. JGR 2010 (:2338-2477). + for (int i = 0; i < nk; ++i) { e1[i] = zero; hs[i] = zero; } + for (int i = 0; i < nspec; ++i) { srhs[i] = zero; ddiag[i] = zero; pb2[i] = zero; } + for (int ik = ik1; ik <= nk; ++ik) { + e1[ik - 1] = zero; + for (int ith = 1; ith <= nth; ++ith) { + const int is = ith + (ik - 1) * nth; + e1[ik - 1] = e1[ik - 1] + (a[is - 1] * f.sig(ik)) * dth; + } + dk[ik - 1] = f.dden(ik - 1) / ((dth * f.sig(ik)) * cg[ik - 1]); + } + int id = nint(depth); + if (id > kNdtab) id = kNdtab; + if (id < 1) { + id = 1; + } else if (id > kNdtab) { + id = kNdtab; + } + for (int i = 0; i < nk; ++i) { hs[i] = zero; efdf[i] = zero; kbar[i] = zero; } + int nkl = 0; + for (int ikl = 1; ikl <= nk; ++ikl) { + iksup[ikl - 1] = f.iktab(ikl - 1, id - 1); + if (iksup[ikl - 1] <= nk) { + Real v = zero; + for (int i = ikl; i <= iksup[ikl - 1] - 1; ++i) v = v + e1[i - 1] * dk[i - 1]; + efdf[ikl - 1] = v; + if (efdf[ikl - 1] != zero) { + Real w = zero; + for (int i = ikl; i <= iksup[ikl - 1] - 1; ++i) + w = w + (k[i - 1] * e1[i - 1]) * dk[i - 1]; + kbar[ikl - 1] = w / efdf[ikl - 1]; + } else { + kbar[ikl - 1] = zero; } - for (int i = 0; i < nk; ++i) { - if (ntimes[i] > 0) { - s1[i] = s1[i] / static_cast(ntimes[i]); - s2[i] = s2[i] / static_cast(ntimes[i]); - } else { - s1[i] = zero; - s2[i] = zero; - } + hs[ikl - 1] = static_cast(4) * Kokkos::sqrt(efdf[ikl - 1]); + nkl = nkl + 1; + } + } + // Dissipation and breaking probability per scale. The loop runs to the count NKL, + // not over the scales that passed IKSUP <= NK (:2392). + for (int i = 0; i < nk; ++i) { dck[i] = zero; qb[i] = zero; } + for (int ikl = 1; ikl <= nkl; ++ikl) { + const Real kb = kbar[ikl - 1]; + if (hs[ikl - 1] != zero && kb != zero) { + // FAC_KD2 + ANINT(...) is REAL, truncated on assignment to INTEGER. + int ikd = fint(Kokkos::round(dm::logf(kb * depth) / sds4_detail::kLogFacKd1) + + static_cast(kFacKd2)); + int ikhs = fint(one + Kokkos::round((kb * hs[ikl - 1]) / sds4_detail::kDkhs)); + if (ikd > kNkd) { + ikd = kNkd; + } else if (ikd < 1) { + ikd = 1; } - for (int ik = 1; ik <= nk; ++ik) s1[ik - 1] = s1[ik - 1] / f.sig(ik); - // Isotropic distribution (:2456-2471). - for (int ik = 1; ik <= nk; ++ik) { - Real v = zero; - for (int is = (ik - 1) * nth + 1; is <= ik * nth; ++is) v = v + a[is - 1]; - const Real asum = v * dth; - if (asum > static_cast(1.E-8)) { - for (int is = 1 + (ik - 1) * nth; is <= ik * nth; ++is) ddiag[is - 1] = s1[ik - 1] / asum; - for (int is = 1 + (ik - 1) * nth; is <= ik * nth; ++is) pb2[is - 1] = s2[ik - 1] / asum; - } else { - for (int is = 1 + (ik - 1) * nth; is <= ik * nth; ++is) ddiag[is - 1] = zero; - for (int is = 1 + (ik - 1) * nth; is <= ik * nth; ++is) pb2[is - 1] = zero; - } - if (pb2[(ik - 1) * nth] > static_cast(0.001)) { - bth0[ik - 1] = static_cast(2.) * p.ssdsbr; - } else { - bth0[ik - 1] = zero; - } + if (ikhs > kNkhs) { + ikhs = kNkhs; + } else if (ikhs < 1) { + ikhs = 1; } - // Kept as written (:2473): PB is 0 here, so this is (1-SSDSC(1))*PB2*A + SSDSC(1)*0. - for (int i = 0; i < nspec; ++i) - pb[i] = ((one - ssdsc[0]) * pb2[i]) * a[i] + ssdsc[0] * pb[i]; - for (int i = 0; i < nspec; ++i) brlambda[i] = pb[i] / sds4_detail::kTwoPiSq; - for (int i = 0; i < nspec; ++i) srhs[i] = ddiag[i] * a[i]; - break; + const Real xt = dm::tanhf(kb * depth); + // GAM (:2411) is computed and never used. + const Real gam = ((((xt * xt) * xt) * static_cast(1.0314) - + (xt * xt) * static_cast(1.9958)) + + xt * static_cast(1.5522)) + + static_cast(0.1885); + (void)gam; + dck[ikl - 1] = (dm::powf(kb, static_cast(-2.5)) * (kb / sds4_detail::kTwoPi)) * + f.dcki(ikhs - 1, ikd - 1); + qb[ikl - 1] = f.qbi(ikhs - 1, ikd - 1); + } else { + dck[ikl - 1] = zero; + qb[ikl - 1] = zero; } - case 3: { - // 2.c Romero (GRL 2019) (:2481-2520). Division by USTAR**2 kept: Inf at USTAR = 0 (:2485). - const Real t28 = static_cast(28.) / ssdsc[15]; - const Real ko = (kGrav / (ustar * ustar)) / (t28 * t28); - for (int ik = 1; ik <= nk; ++ik) { - const int is0 = (ik - 1) * nth; - const Real kk = k[ik - 1]; - const Real kloc = dm::powf(kk, static_cast(2.) - ssdsc[19]); - const Real k3 = (kk * kk) * kk; - // BTH(1:NTH) is the scratch of this IK. - for (int ith = 1; ith <= nth; ++ith) - bth[ith - 1] = - max_r((a[is0 + ith - 1] * f.sig(ik)) * k3, static_cast(.00000000000001)); - Real dirforcum = dlwmean; - if (ssdsc[10] > zero) dirforcum = MSS(ik, 2); - const Real c = f.sig(ik) / kk; - Real sb = zero; - for (int ith = 1; ith <= nth; ++ith) sb = sb + bth[ith - 1] * dth; - bth0[ik - 1] = sb; - if (ssdsc[17] > zero) { - Real cd, sd; - dm::sincosf(dirforcum, sd, cd); - const Real pm = dm::powf(MSS(ik, 1) * kloc, ssdsc[13]); - for (int ith = 1; ith <= nth; ++ith) { - const Real q = f.ecos(ith - 1) * cd + f.esin(ith - 1) * sd; - lmodulation[ith - 1] = one + ssdsc[17] * (pm * (q * q)); // FACSTRAINL - } - // LIBMVEC SITE: -O3 uses _ZGVbN4vv_powf here (INVENTORY.md, :2500), lane i = ITH-1. - const Real e19 = ssdsc[18]; - mvec_powf( - nth, [&](int i) { return lmodulation[i]; }, [&](int) { return e19; }, - [&](int i, Real v) { lmodulation[i] = v; }); - } else { - for (int ith = 1; ith <= nth; ++ith) lmodulation[ith - 1] = one; - } - // Lambda (:2507-2509). - const Real fk = one + ssdsc[12] * dm::powf(max_r(kk / ko, one), ssdsc[14]); - // LIBMVEC SITE: -O3 uses _ZGVbN4v_expf here (INVENTORY.md, :2509), lane i = ITH-1; - // the exponential first, then (((e*SSDSC(9))*F)/(SSDSC(13)+1))*LMODULATION per element. - const Real ssdsbr = p.ssdsbr, c9 = ssdsc[8], c13p1 = ssdsc[12] + one; - mvec_expf( - nth, [&](int i) { return -(ssdsbr / bth[i]); }, - [&](int i, Real v) { brlambda[is0 + i] = (((c9 * v) * fk) / c13p1) * lmodulation[i]; }); - // Breaking strength, a generalisation of Duncan's b (:2511-2512). - const Real btover = Kokkos::sqrt(bth0[ik - 1]) - Kokkos::sqrt(p.ssdsbt); - // MAX(0.,BTOVER) computes 0 > b ? 0 : b (INVENTORY.md, :2512 assembly). - brm12[ik - 1] = - ssdsc[1] * dm::powf(zero > btover ? zero : btover, static_cast(2.5)); - // MAX(0.,SIGN(x,BTOVER)) computes x > 0 ? x : +0. (INVENTORY.md, :2514 assembly). - for (int ith = 1; ith <= nth; ++ith) { - const Real x = Kokkos::copysign(brlambda[is0 + ith - 1], btover); - brlambda[is0 + ith - 1] = x > zero ? x : zero; - } - const Real c2 = c * c; - const Real c5 = (c * c2) * c2; - for (int ith = 1; ith <= nth; ++ith) - srhs[is0 + ith - 1] = - (((brm12[ik - 1] / sds4_detail::kGrav2) * brlambda[is0 + ith - 1]) * c5) / f.sig(ik); - for (int ith = 1; ith <= nth; ++ith) - ddiag[is0 + ith - 1] = - ((srhs[is0 + ith - 1] * p.ssdsbr) / max_r(bth[ith - 1], static_cast(1.e-20))) / - max_r(a[is0 + ith - 1], static_cast(1e-20)); + } + // Distributes scale dissipation over the frequency spectrum (:2429-2452). + for (int i = 0; i < nk; ++i) { s1[i] = zero; s2[i] = zero; ntimes[i] = 0; } + for (int ikl = 1; ikl <= nkl; ++ikl) { + if (efdf[ikl - 1] > zero) { + for (int i = ikl; i <= iksup[ikl - 1]; ++i) + s1[i - 1] = s1[i - 1] + (dck[ikl - 1] * e1[i - 1]) / efdf[ikl - 1]; + for (int i = ikl; i <= iksup[ikl - 1]; ++i) + s2[i - 1] = s2[i - 1] + (qb[ikl - 1] * e1[i - 1]) / efdf[ikl - 1]; + for (int i = ikl; i <= iksup[ikl - 1]; ++i) ntimes[i - 1] = ntimes[i - 1] + 1; + } + } + for (int i = 0; i < nk; ++i) { + if (ntimes[i] > 0) { + s1[i] = s1[i] / static_cast(ntimes[i]); + s2[i] = s2[i] / static_cast(ntimes[i]); + } else { + s1[i] = zero; + s2[i] = zero; + } + } + for (int ik = 1; ik <= nk; ++ik) s1[ik - 1] = s1[ik - 1] / f.sig(ik); + // Isotropic distribution (:2456-2471). + for (int ik = 1; ik <= nk; ++ik) { + Real v = zero; + for (int is = (ik - 1) * nth + 1; is <= ik * nth; ++is) v = v + a[is - 1]; + const Real asum = v * dth; + if (asum > static_cast(1.E-8)) { + for (int is = 1 + (ik - 1) * nth; is <= ik * nth; ++is) ddiag[is - 1] = s1[ik - 1] / asum; + for (int is = 1 + (ik - 1) * nth; is <= ik * nth; ++is) pb2[is - 1] = s2[ik - 1] / asum; + } else { + for (int is = 1 + (ik - 1) * nth; is <= ik * nth; ++is) ddiag[is - 1] = zero; + for (int is = 1 + (ik - 1) * nth; is <= ik * nth; ++is) pb2[is - 1] = zero; + } + if (pb2[(ik - 1) * nth] > static_cast(0.001)) { + bth0[ik - 1] = static_cast(2.) * p.ssdsbr; + } else { + bth0[ik - 1] = zero; + } + } + // Kept as written (:2473): PB is 0 here, so this is (1-SSDSC(1))*PB2*A + SSDSC(1)*0. + for (int i = 0; i < nspec; ++i) + pb[i] = ((one - ssdsc[0]) * pb2[i]) * a[i] + ssdsc[0] * pb[i]; + for (int i = 0; i < nspec; ++i) brlambda[i] = pb[i] / sds4_detail::kTwoPiSq; + for (int i = 0; i < nspec; ++i) srhs[i] = ddiag[i] * a[i]; +} + +/// Case 3, 2.c (:2481-2520), unchanged. +KOKKOS_INLINE_FUNCTION void sds4_case3(const Params& p, const Fields& f, const Sds4Args& pt, + const Sds4Work& wk, const Sds4Mid& mid) { + (void)mid; + using ww::st4::max_r; + const Real zero = static_cast(0.); + const Real one = static_cast(1.); + const int nk = p.nk; + const int nth = p.nth; + const Real dth = p.dth; + const Real* ssdsc = p.ssdsc; // SSDSC(i) is ssdsc[i-1] + const Real* const a = pt.a; + const Real* const k = pt.k; + const Real ustar = pt.ustar; + const Real dlwmean = pt.dlwmean; + Real* const srhs = pt.srhs; + Real* const ddiag = pt.ddiag; + Real* const brlambda = pt.brlambda; + Real* const bth0 = wk.bth0; + Real* const brm12 = wk.brm12; + Real* const msssum = wk.msssum; + Real* const bth = wk.bth; + Real* const lmodulation = wk.lmodulation; + auto MSS = [&](int ik, int j) -> Real& { return msssum[(j - 1) * nk + (ik - 1)]; }; + // 2.c Romero (GRL 2019) (:2481-2520). Division by USTAR**2 kept: Inf at USTAR = 0 (:2485). + const Real t28 = static_cast(28.) / ssdsc[15]; + const Real ko = (kGrav / (ustar * ustar)) / (t28 * t28); + for (int ik = 1; ik <= nk; ++ik) { + const int is0 = (ik - 1) * nth; + const Real kk = k[ik - 1]; + const Real kloc = dm::powf(kk, static_cast(2.) - ssdsc[19]); + const Real k3 = (kk * kk) * kk; + // BTH(1:NTH) is the scratch of this IK. + for (int ith = 1; ith <= nth; ++ith) + bth[ith - 1] = + max_r((a[is0 + ith - 1] * f.sig(ik)) * k3, static_cast(.00000000000001)); + Real dirforcum = dlwmean; + if (ssdsc[10] > zero) dirforcum = MSS(ik, 2); + const Real c = f.sig(ik) / kk; + Real sb = zero; + for (int ith = 1; ith <= nth; ++ith) sb = sb + bth[ith - 1] * dth; + bth0[ik - 1] = sb; + if (ssdsc[17] > zero) { + Real cd, sd; + dm::sincosf(dirforcum, sd, cd); + const Real pm = dm::powf(MSS(ik, 1) * kloc, ssdsc[13]); + for (int ith = 1; ith <= nth; ++ith) { + const Real q = f.ecos(ith - 1) * cd + f.esin(ith - 1) * sd; + lmodulation[ith - 1] = one + ssdsc[17] * (pm * (q * q)); // FACSTRAINL } - break; + // LIBMVEC SITE: -O3 uses _ZGVbN4vv_powf here (INVENTORY.md, :2500), lane i = ITH-1. + const Real e19 = ssdsc[18]; + mvec_powf( + nth, [&](int i) { return lmodulation[i]; }, [&](int) { return e19; }, + [&](int i, Real v) { lmodulation[i] = v; }); + } else { + for (int ith = 1; ith <= nth; ++ith) lmodulation[ith - 1] = one; } - default: - break; + // Lambda (:2507-2509). + const Real fk = one + ssdsc[12] * dm::powf(max_r(kk / ko, one), ssdsc[14]); + // LIBMVEC SITE: -O3 uses _ZGVbN4v_expf here (INVENTORY.md, :2509), lane i = ITH-1; + // the exponential first, then (((e*SSDSC(9))*F)/(SSDSC(13)+1))*LMODULATION per element. + const Real ssdsbr = p.ssdsbr, c9 = ssdsc[8], c13p1 = ssdsc[12] + one; + mvec_expf( + nth, [&](int i) { return -(ssdsbr / bth[i]); }, + [&](int i, Real v) { brlambda[is0 + i] = (((c9 * v) * fk) / c13p1) * lmodulation[i]; }); + // Breaking strength, a generalisation of Duncan's b (:2511-2512). + const Real btover = Kokkos::sqrt(bth0[ik - 1]) - Kokkos::sqrt(p.ssdsbt); + // MAX(0.,BTOVER) computes 0 > b ? 0 : b (INVENTORY.md, :2512 assembly). + brm12[ik - 1] = + ssdsc[1] * dm::powf(zero > btover ? zero : btover, static_cast(2.5)); + // MAX(0.,SIGN(x,BTOVER)) computes x > 0 ? x : +0. (INVENTORY.md, :2514 assembly). + for (int ith = 1; ith <= nth; ++ith) { + const Real x = Kokkos::copysign(brlambda[is0 + ith - 1], btover); + brlambda[is0 + ith - 1] = x > zero ? x : zero; + } + const Real c2 = c * c; + const Real c5 = (c * c2) * c2; + for (int ith = 1; ith <= nth; ++ith) + srhs[is0 + ith - 1] = + (((brm12[ik - 1] / sds4_detail::kGrav2) * brlambda[is0 + ith - 1]) * c5) / f.sig(ik); + for (int ith = 1; ith <= nth; ++ith) + ddiag[is0 + ith - 1] = + ((srhs[is0 + ith - 1] * p.ssdsbr) / max_r(bth[ith - 1], static_cast(1.e-20))) / + max_r(a[is0 + ith - 1], static_cast(1e-20)); } +} +/// Wave-turbulence interaction and cumulative effect (:2533-2575). +KOKKOS_INLINE_FUNCTION void sds4_turb(const Params& p, const Fields& f, const Sds4Args& pt, + const Sds4Work& wk, const Sds4Mid& mid) { + using ww::st4::max_r; + const Real zero = static_cast(0.); + const int nk = p.nk; + const int nth = p.nth; + const Real dth = p.dth; + const Real* ssdsc = p.ssdsc; // SSDSC(i) is ssdsc[i-1] + const Real* const a = pt.a; + const Real* const k = pt.k; + const Real usdir = pt.usdir; + Real* const srhs = pt.srhs; + Real* const ddiag = pt.ddiag; + Real* const brlambda = pt.brlambda; + Real* const bth0 = wk.bth0; + const Real facturb = mid.facturb; + const int dikcumul = mid.dikcumul; + const int ik1 = mid.ik1; + Real renewalfreq = zero; // Wave-turbulence interaction and cumulative effect (:2533-2575). if ((ssdsc[2] != zero) || (ssdsc[4] != zero) || (ssdsc[20] != zero)) { for (int ik = ik1; ik <= nk; ++ik) { @@ -472,7 +633,26 @@ KOKKOS_INLINE_FUNCTION void w3sds4_point(const Params& p, const Fields& f, const } } } +} +/// Whitecap parameters (:2581-2650). Early return (:2581): WHITECAP keeps the caller's values. +KOKKOS_INLINE_FUNCTION void sds4_whitecap(const Params& p, const Fields& f, const Sds4Args& pt, + const Sds4Work& wk, const Sds4Mid& mid) { + const Real zero = static_cast(0.); + const Real one = static_cast(1.); + const int nk = p.nk; + const int nth = p.nth; + const Real* ssdsc = p.ssdsc; // SSDSC(i) is ssdsc[i-1] + const Real* const k = pt.k; + const Real* const cg = pt.cg; + const int flogrd57 = pt.flogrd57; + const int flogrd58 = pt.flogrd58; + Real* const brlambda = pt.brlambda; + Real* const whitecap = pt.whitecap; + Real* const coef4 = wk.coef4; + Real* const coef5 = wk.coef5; + Real* const brm12 = wk.brm12; + const int ik1 = mid.ik1; // Whitecap parameters. Early return (:2581): WHITECAP keeps the caller's values; // WHITECAP(3) is never set. if (!(flogrd57 || flogrd58)) return; @@ -528,4 +708,308 @@ KOKKOS_INLINE_FUNCTION void w3sds4_point(const Params& p, const Fields& f, const } } +/// W3SDS4 for one point. a = A(1:NSPEC), k = K(1:NK), cg = CG(1:NK), srhs, +/// ddiag, brlambda = (1:NSPEC), whitecap = WHITECAP(1:4), all 0-based. +/// flogrd57, flogrd58 are FLOGRD(5,7), FLOGRD(5,8) as 0/1. IX, IY are not +/// read and not passed. The OUT arrays arrive holding the caller's values: +/// the early return at w3src4md.F90:2581 leaves WHITECAP as it was. +/// The routine's sections are the sds4_* parts above, called in the routine's order. +KOKKOS_INLINE_FUNCTION void w3sds4_point(const Params& p, const Fields& f, const Real* a, + const Real* k, const Real* cg, Real ustar, Real usdir, + Real depth, Real dair, Real dlwmean, int flogrd57, + int flogrd58, Real* srhs, Real* ddiag, Real* brlambda, + Real* whitecap, Real* work, int* iwork) { + const Sds4Args pt{a, k, cg, ustar, usdir, depth, dair, dlwmean, flogrd57, flogrd58, + srhs, ddiag, brlambda, whitecap}; + const Sds4Work w = sds4_layout(p, work, iwork); + Sds4Mid mid; + sds4_prep(p, f, pt, w, mid); + sds4_mss(p, f, pt, w, mid); + switch (ww::st4::nint(p.ssdsc[0])) { // :2259 + case 1: + sds4_sat(p, f, pt, w, mid); + sds4_case1(p, f, pt, w, mid); + break; + case 2: + sds4_case2(p, f, pt, w, mid); + break; + case 3: + sds4_case3(p, f, pt, w, mid); + break; + default: + break; + } + sds4_turb(p, f, pt, w, mid); + sds4_whitecap(p, f, pt, w, mid); +} + +namespace sds4_detail { + +using lanes::M8; +using lanes::V8; +using lanes::vload; + +/// Section 1.b (:2228-2254) for the `n` points of a team, one point per vector lane. The loops +/// over IK and ITH stay outside, each lane keeps its own three sums in ITH order; powf and +/// atan2f are called per lane. Lanes n..W-1 repeat point 0's inputs and are not stored. +template +inline void mss_lanes(const Params& p, const Fields& f, const int n, const Sds4Args* pt, + const Sds4Work* w) { + static_assert(W == 8, "the ITH loops are written for V8"); + const Real zero = static_cast(0.); + const Real* ssdsc = p.ssdsc; + if (!(ssdsc[7] > zero || ssdsc[10] > zero || ssdsc[17] > zero)) return; + const int nk = p.nk, nth = p.nth; + const Real *a[W], *k[W], *cg[W]; + for (int l = 0; l < W; ++l) { + const int q = l < n ? l : 0; + a[l] = pt[q].a; + k[l] = pt[q].k; + cg[l] = pt[q].cg; + } + for (int ik = 1; ik <= nk; ++ik) { + Real kp[W], cgk[W]; + for (int l = 0; l < W; ++l) { + kp[l] = dm::powf(k[l][ik - 1], ssdsc[19]); + cgk[l] = cg[l][ik - 1]; + } + V8 kp_v, cg_v; + vload(kp_v, kp); + vload(cg_v, cgk); + const Real dden = f.dden(ik - 1); + V8 mssp = {}, msspc2 = {}, mssps2 = {}; + for (int ith = 1; ith <= nth; ++ith) { + const int is = ith + (ik - 1) * nth; + const Real ecos = f.ecos(ith - 1), esin = f.esin(ith - 1); + V8 av; + for (int l = 0; l < W; ++l) av[l] = a[l][is - 1]; + const V8 msslong = ((kp_v * av) * dden) / cg_v; + msspc2 = msspc2 + msslong * ecos; + mssps2 = mssps2 + msslong * esin; + mssp = mssp + msslong; + } + for (int l = 0; l < n; ++l) { + Real* const msssum = w[l].msssum; + auto MSS = [&](int i, int j) -> Real& { return msssum[(j - 1) * nk + (i - 1)]; }; + for (int i = ik; i <= nk; ++i) { + MSS(i, 1) = MSS(i, 1) + mssp[l]; + MSS(i, 3) = MSS(i, 3) + msspc2[l]; + MSS(i, 4) = MSS(i, 4) + mssps2[l]; + } + Real mssd = dm::atan2f(MSS(ik, 4), MSS(ik, 3)); + if (mssd < zero) mssd = mssd + kPi; + MSS(ik, 2) = mssd; + } + } +} + +/// Case 1, 2.a.1 (:2262-2290) for the `n` points of a team. The sum over the SATWEIGHTS terms +/// is DOT_PRODUCT's, from zero in J order, one accumulator per ITH and lane, so each BTH(IS) +/// has the bits of sds4_sat's; the zero weights are multiplied like the others. Without the +/// transposed tables the points go one after the other through sds4_sat. +template +inline void sat_lanes(const Params& p, const Fields& f, const int n, const Sds4Args* pt, + const Sds4Work* w, const Sds4Mid* mid) { + const int nsat = static_cast(f.satweights.extent(0)); + const bool tr = f.satweights_t.extent(1) == static_cast(nsat); + if (!tr) { + for (int l = 0; l < n; ++l) sds4_sat(p, f, pt[l], w[l], mid[l]); + return; + } + const Real zero = static_cast(0.); + const int nk = p.nk, nth = p.nth, nspec = p.nspec; + const Real dth = p.dth; + const Real *a[W], *k[W]; + Real *bth[W], *bth0[W]; + for (int l = 0; l < W; ++l) { + const int q = l < n ? l : 0; + a[l] = pt[q].a; + k[l] = pt[q].k; + bth[l] = w[q].bth; + bth0[l] = w[q].bth0; + } + for (int l = 0; l < n; ++l) + for (int i = 0; i < nspec; ++i) bth[l][i] = zero; + const int ik1 = mid[0].ik1; + for (int ik = ik1; ik <= nk; ++ik) { + Real facsat[W]; + for (int l = 0; l < W; ++l) { + const Real kk = k[l][ik - 1]; + facsat[l] = (f.sig(ik) * ((kk * kk) * kk)) * dth; + } + V8 facsat_v; + vload(facsat_v, facsat); + const int is0 = (ik - 1) * nth; + for (int l = 0; l < n; ++l) bth[l][is0] = zero; + V8 asum = {}; + for (int ith = 1; ith <= nth; ++ith) { + V8 av; + for (int l = 0; l < W; ++l) av[l] = a[l][is0 + ith - 1]; + asum = asum + av; + } + const V8 b0 = asum * facsat_v; + for (int l = 0; l < n; ++l) bth0[l][ik - 1] = b0[l]; + if (p.ssdsdth >= static_cast(180)) { + for (int l = 0; l < n; ++l) + for (int i = is0; i < is0 + nth; ++i) bth[l][i] = bth0[l][ik - 1]; + } else { + for (int ith = 0; ith < nth; ++ith) { + V8 acc = {}; + for (int j = 0; j < nsat; ++j) { + const Real wj = f.satweights_t(ith, j); + const int xj = f.satindices_t(ith, j); + V8 av; + for (int l = 0; l < W; ++l) av[l] = a[l][is0 + xj - 1]; + acc = acc + wj * av; + } + acc = acc * facsat_v; + for (int l = 0; l < n; ++l) bth[l][is0 + ith] = acc[l]; + } + for (int l = 0; l < n; ++l) { + Real* const bl = bth[l]; + // MAXVAL (:2287), as gfortran lowers it: first element >= -Inf, then strictly greater. + Real lim = -Kokkos::Experimental::infinity_v; + int i = is0; + bool found = false; + for (; i < is0 + nth; ++i) { + if (bl[i] >= lim) { lim = bl[i]; ++i; found = true; break; } + } + if (!found) lim = Kokkos::Experimental::quiet_NaN_v; + for (; i < is0 + nth; ++i) + if (bl[i] > lim) lim = bl[i]; + bth0[l][ik - 1] = lim; + } + } + } +} + +/// The wave-turbulence interaction and cumulative effect (:2533-2575) for the `n` points of a +/// team. Each lane keeps its own RENEWALFREQ and its own sums in the routine's order; the +/// branch on BTH0 is a select of the value the taken side computes; COS/SIN(USDIR) is computed +/// once per lane, the same call the routine repeats for each ITH. +template +inline void turb_lanes(const Params& p, const Fields& f, const int n, const Sds4Args* pt, + const Sds4Work* w, const Sds4Mid* mid) { + const Real zero = static_cast(0.); + const Real* ssdsc = p.ssdsc; + if (!((ssdsc[2] != zero) || (ssdsc[4] != zero) || (ssdsc[20] != zero))) return; + const int nk = p.nk, nth = p.nth; + const Real dth = p.dth; + const int dikcumul = mid[0].dikcumul, ik1 = mid[0].ik1; + const Real *a[W], *k[W]; + Real *ddiag[W], *srhs[W], *brlambda[W], *bth0[W]; + Real su[W], cu[W], facturb[W]; + for (int l = 0; l < W; ++l) { + const int q = l < n ? l : 0; + a[l] = pt[q].a; + k[l] = pt[q].k; + ddiag[l] = pt[q].ddiag; + srhs[l] = pt[q].srhs; + brlambda[l] = pt[q].brlambda; + bth0[l] = w[q].bth0; + facturb[l] = mid[q].facturb; + dm::sincosf(pt[q].usdir, su[l], cu[l]); + } + V8 su_v, cu_v; + vload(su_v, su); + vload(cu_v, cu); + for (int ik = ik1; ik <= nk; ++ik) { + V8 renewalfreq = {}; + Real facturb2[W], dvisc[W], c[W]; + for (int l = 0; l < W; ++l) { + const Real kk = k[l][ik - 1]; + facturb2[l] = -(((f.sig(ik) * static_cast(2.)) * kk) * facturb[l]); + dvisc[l] = -(((ssdsc[20] * static_cast(4.)) * kk) * kk); + c[l] = f.sig(ik) / kk; + } + V8 facturb2_v, dvisc_v; + vload(facturb2_v, facturb2); + vload(dvisc_v, dvisc); + if (ssdsc[2] > zero && ik > dikcumul) { + // The cheap isotropic version. + for (int ik2 = ik1; ik2 <= ik - dikcumul; ++ik2) { + Real cw[W]; + for (int l = 0; l < W; ++l) { + const Real c2 = f.sig(ik2) / k[l][ik2 - 1]; + cw[l] = ((Kokkos::fabs(c2 - c[l]) * f.dsip(ik2)) / (c2 * static_cast(0.5))) * dth; + } + const int is2 = (ik2 - 1) * nth; + V8 v = {}; + for (int i = is2; i < is2 + nth; ++i) { + V8 bl; + for (int l = 0; l < W; ++l) bl[l] = brlambda[l][i]; + v = v + bl; + } + V8 cw_v; + vload(cw_v, cw); + renewalfreq = renewalfreq - cw_v * v; + } + } + for (int ith = 1; ith <= nth; ++ith) { + const int is = ith + (ik - 1) * nth; + if (ssdsc[2] < zero && ik > dikcumul) { + // The default (SSDSC(3) < 0), O(NK^2 NTH^2) branch (:2554-2563). + renewalfreq = V8{}; + for (int ik2 = ik1; ik2 <= ik - dikcumul; ++ik2) { + M8 on; + for (int l = 0; l < W; ++l) on[l] = bth0[l][ik2 - 1] > p.ssdsbr ? -1 : 0; + // No lane takes the branch: the sum below would be discarded. + bool any = false; + for (int l = 0; l < W; ++l) any = any || on[l] != 0; + if (!any) continue; + const int is2 = (ik2 - 1) * nth; + V8 v = {}; + for (int j = 1; j <= nth; ++j) { + const Real cw = f.cumulw(is2 + j - 1, is - 1); + V8 bl; + for (int l = 0; l < W; ++l) bl[l] = brlambda[l][is2 + j - 1]; + v = v + cw * bl; + } + renewalfreq = on ? renewalfreq + v : renewalfreq; + } + } + const Real ecos = f.ecos(ith - 1), esin = f.esin(ith - 1); + const V8 coswind = ecos * cu_v + esin * su_v; + // MAX(COSWIND, 0.) is coswind > 0 ? coswind : 0 (max_r). + const V8 dturb = facturb2_v * (coswind > zero ? coswind : V8{}); + const V8 diag2 = (ssdsc[2] * renewalfreq + dturb) + dvisc_v; + for (int l = 0; l < n; ++l) { + ddiag[l][is - 1] = ddiag[l][is - 1] + diag2[l]; + srhs[l][is - 1] = srhs[l][is - 1] + a[l][is - 1] * diag2[l]; + } + } + } +} + +} // namespace sds4_detail + +/// W3SDS4 for the `n` points of a team (1 <= n <= W), `pt[l]` and `w[l]` being the arguments and +/// work arrays of point l: the sections of w3sds4_point in its order, the loops that gain from it +/// with one point per vector lane (plan 2026-10-18-005, unit L3), the rest point after point. The +/// libmvec sites (sds4_case1, sds4_case3) stay one grouped call per point. Every point gets the +/// bits w3sds4_point gives it. +template +inline void w3sds4_lanes(const Params& p, const Fields& f, const int n, const Sds4Args* pt, + const Sds4Work* w) { + Sds4Mid mid[W] = {}; + for (int l = 0; l < n; ++l) sds4_prep(p, f, pt[l], w[l], mid[l]); + sds4_detail::mss_lanes(p, f, n, pt, w); + switch (ww::st4::nint(p.ssdsc[0])) { // :2259 + case 1: + sds4_detail::sat_lanes(p, f, n, pt, w, mid); + for (int l = 0; l < n; ++l) sds4_case1(p, f, pt[l], w[l], mid[l]); + break; + case 2: + for (int l = 0; l < n; ++l) sds4_case2(p, f, pt[l], w[l], mid[l]); + break; + case 3: + for (int l = 0; l < n; ++l) sds4_case3(p, f, pt[l], w[l], mid[l]); + break; + default: + break; + } + sds4_detail::turb_lanes(p, f, n, pt, w, mid); + for (int l = 0; l < n; ++l) sds4_whitecap(p, f, pt[l], w[l], mid[l]); +} + } // namespace ww::st4 diff --git a/kokkos/src/ww_kokkos/st4_sin4.hpp b/kokkos/src/ww_kokkos/st4_sin4.hpp index 87a80b9b..003cd60d 100644 --- a/kokkos/src/ww_kokkos/st4_sin4.hpp +++ b/kokkos/src/ww_kokkos/st4_sin4.hpp @@ -17,6 +17,7 @@ #include "real.hpp" #include "st4_arrays.hpp" #include "st4_params.hpp" +#include "st4_vec8.hpp" namespace ww::st4 { @@ -50,15 +51,10 @@ struct Mid { }; -// 8 lanes of REAL and their masks, GCC's vector extensions (host code only: the lane path runs -// on host backends). Arithmetic on them is lane-wise IEEE; with -ffp-contract=off nothing fuses. -using V8 = Real __attribute__((vector_size(8 * sizeof(Real)))); -using M8 = int __attribute__((vector_size(8 * sizeof(int)))); -static_assert(sizeof(Real) == sizeof(int), "a mask lane per REAL lane"); -// By reference, not by value: a function returning a 32-byte vector warns (-Wpsabi) in a build -// without AVX (serial-debug has no -march). -inline void vload(V8& v, const Real* x) { __builtin_memcpy(&v, x, sizeof v); } -inline void vstore(Real* x, const V8& v) { __builtin_memcpy(x, &v, sizeof v); } +using lanes::M8; +using lanes::V8; +using lanes::vload; +using lanes::vstore; } // namespace sin4_detail diff --git a/kokkos/src/ww_kokkos/st4_vec8.hpp b/kokkos/src/ww_kokkos/st4_vec8.hpp new file mode 100644 index 00000000..fce71b3d --- /dev/null +++ b/kokkos/src/ww_kokkos/st4_vec8.hpp @@ -0,0 +1,21 @@ +// kokkos/src/ww_kokkos/st4_vec8.hpp +// Eight lanes of REAL and their masks, GCC's vector extensions, for the routines that run one +// sea point per lane (plan 2026-10-18-005). Host code only: the lane path runs on host backends. +// Arithmetic on these types is lane-wise IEEE; with -ffp-contract=off nothing fuses. +#pragma once + +#include + +#include "real.hpp" + +namespace ww::st4::lanes { + +using V8 = Real __attribute__((vector_size(8 * sizeof(Real)))); +using M8 = int __attribute__((vector_size(8 * sizeof(int)))); +static_assert(sizeof(Real) == sizeof(int), "a mask lane per REAL lane"); +// By reference, not by value: a function returning a 32-byte vector warns (-Wpsabi) in a build +// without AVX (serial-debug has no -march). +inline void vload(V8& v, const Real* x) { __builtin_memcpy(&v, x, sizeof v); } +inline void vstore(Real* x, const V8& v) { __builtin_memcpy(x, &v, sizeof v); } + +} // namespace ww::st4::lanes diff --git a/regtests/CMakeLists.txt b/regtests/CMakeLists.txt index 53c7faec..5cb351f4 100644 --- a/regtests/CMakeLists.txt +++ b/regtests/CMakeLists.txt @@ -38,7 +38,7 @@ # DB1 MLIM ...), the only set the Kokkos W3SRCE translates; no matrix.base # switch has it. The entry keeps its inputs, grid and propagation; its outputs # are no longer WW3's reference outputs, and only the pair is compared -# WW3GPU_WW3 the h0ffmann/WW3 ww3-gpu/kokkos-st4 tree; never written (the inputs are staged +# WW3GPU_WW3 the h0ffmann/WW3 ww3-gpu/kokkos-pnt tree; never written (the inputs are staged # in /regtests/, and record:summary fails when `git status` shows a change) # WW3GPU_DATA WW3's ftp data, /data_regtests: the store path of # `nix build ./picasso/labs/pratico#ww3-regtest-data` (`just ww3-gpu-matrix` passes @@ -73,7 +73,7 @@ include("${CMAKE_CURRENT_SOURCE_DIR}/cmake/ww3gpu_common.cmake") set(WW3GPU_SELECTION serial CACHE STRING "Which matrix.base entries: serial, omp, mpi or all") set_property(CACHE WW3GPU_SELECTION PROPERTY STRINGS serial omp mpi all) set(WW3GPU_ROUTINES SNL1 CACHE STRING "The routines the Kokkos run turns on: a list of WW3GPU_ALL_ROUTINES") -set(WW3GPU_WW3 "${WW3GPU_ROOT}/WW3" CACHE PATH "The WW3 tree (h0ffmann/WW3, ww3-gpu/kokkos-st4), read-only") +set(WW3GPU_WW3 "${WW3GPU_ROOT}/WW3" CACHE PATH "The WW3 tree (h0ffmann/WW3, ww3-gpu/kokkos-pnt), read-only") set(WW3GPU_BACKEND openmp CACHE STRING "The Kokkos backend: openmp or cuda") set_property(CACHE WW3GPU_BACKEND PROPERTY STRINGS openmp cuda) if(NOT WW3GPU_BACKEND MATCHES "^(openmp|cuda)$") @@ -144,7 +144,7 @@ if(NOT EXISTS "${_base}") message(FATAL_ERROR "${WW3GPU_WW3} has no regtests/bin/matrix.base: git submodule update --init WW3") endif() if(NOT EXISTS "${WW3GPU_WW3}/model/src/w3kokkosmd.F90") - message(FATAL_ERROR "${WW3GPU_WW3} is not the ww3-gpu/kokkos-st4 branch of h0ffmann/WW3 (no model/src/w3kokkosmd.F90)") + message(FATAL_ERROR "${WW3GPU_WW3} is not the ww3-gpu/kokkos-pnt branch of h0ffmann/WW3 (no model/src/w3kokkosmd.F90)") endif() ww3gpu_ftp_version(_ftpver _ftpurl "${WW3GPU_WW3}/model/bin/ww3_from_ftp.sh") set(WW3GPU_DATA "" CACHE PATH "WW3's ftp data, /data_regtests (pratico's ww3-regtest-data)") diff --git a/regtests/cmake/run_entry.cmake b/regtests/cmake/run_entry.cmake index 36c0b856..1fcd87a1 100644 --- a/regtests/cmake/run_entry.cmake +++ b/regtests/cmake/run_entry.cmake @@ -14,13 +14,19 @@ # WW3GPU_KOKKOS the Kokkos run takes that preset's flags instead (ww3gpu_kokkos_env), kokkos_log # is W3INIT's parts line showing each part on Kokkos, and kokkos_spectra lists every # "W3 spectra|points computed" counter. Every run sets WW_KOKKOS_ALL and unsets the parts it -# does not turn on, so no flag is inherited from the shell. The Kokkos run also sets -# WW_KOKKOS_STRICT=1: a part it asks for that would run the Fortran stops it, so a pair never -# compares the Fortran with itself unseen (kokkos/src/fortran_iface/PATCH.md section 24). +# does not turn on, so no flag is inherited from the shell. A Kokkos run that names its parts (not +# ALL) also sets WW_KOKKOS_STRICT=1: a part it asks for that would run the Fortran stops it, so a +# pair never compares the Fortran with itself unseen (kokkos/src/fortran_iface/PATCH.md section +# 24). Every Kokkos run sets WW_KOKKOS_VA_POISON=1 (section 21) unless the caller's environment +# has WW_KOKKOS_VA_POISON=0; the Fortran runs never have it. cmake_minimum_required(VERSION 3.21) include("${CONFIG}") include("${CMAKE_CURRENT_LIST_DIR}/ww3gpu_common.cmake") ww3gpu_require_gpu_confirmation() +# Read once, before run() rewrites the environment of the runs. +if("$ENV{WW_KOKKOS_VA_POISON}" STREQUAL "0") + set(WW3GPU_VA_POISON 0) +endif() # Not string(TIMESTAMP): it honours SOURCE_DATE_EPOCH, which the Nix shell sets to 1980. execute_process(COMMAND date +%s OUTPUT_VARIABLE t0 OUTPUT_STRIP_TRAILING_WHITESPACE) diff --git a/regtests/cmake/selftest.cmake b/regtests/cmake/selftest.cmake index 10e681e8..75af7a84 100644 --- a/regtests/cmake/selftest.cmake +++ b/regtests/cmake/selftest.cmake @@ -193,13 +193,17 @@ if(NOT f MATCHES "^ERROR:'bogus' is not a preset") endif() set(WW3GPU_PRESETS_TXT "${_real_presets}") ww3gpu_kokkos_env(e K "SRCE" "SRCE;XYP3" "WW_KOKKOS_SRCE_COUNT") -expect("${e}" "WW_KOKKOS_ALL=0;WW_KOKKOS_STRICT=1;WW_KOKKOS_PRESET=;WW_KOKKOS_SRCE=1;WW_KOKKOS_XYP3=;WW_KOKKOS_SRCE_COUNT=1" "env, Kokkos run of a part list") +expect("${e}" "WW_KOKKOS_ALL=0;WW_KOKKOS_STRICT=1;WW_KOKKOS_VA_POISON=1;WW_KOKKOS_PRESET=;WW_KOKKOS_SRCE=1;WW_KOKKOS_XYP3=;WW_KOKKOS_SRCE_COUNT=1" "env, Kokkos run of a part list: strict, poison on") +set(WW3GPU_VA_POISON 0) +ww3gpu_kokkos_env(e K "SRCE" "SRCE;XYP3" "") +expect("${e}" "WW_KOKKOS_ALL=0;WW_KOKKOS_STRICT=1;WW_KOKKOS_VA_POISON=0;WW_KOKKOS_PRESET=;WW_KOKKOS_SRCE=1;WW_KOKKOS_XYP3=" "env, a caller's WW_KOKKOS_VA_POISON=0 turns the poison off") +unset(WW3GPU_VA_POISON) ww3gpu_kokkos_env(e F "SRCE" "SRCE;XYP3" "WW_KOKKOS_SRCE_COUNT") -expect("${e}" "WW_KOKKOS_ALL=0;WW_KOKKOS_STRICT=;WW_KOKKOS_PRESET=;WW_KOKKOS_SRCE=;WW_KOKKOS_XYP3=;WW_KOKKOS_SRCE_COUNT=" "env, Fortran run: ALL=0, nothing inherited") +expect("${e}" "WW_KOKKOS_ALL=0;WW_KOKKOS_STRICT=;WW_KOKKOS_VA_POISON=;WW_KOKKOS_PRESET=;WW_KOKKOS_SRCE=;WW_KOKKOS_XYP3=;WW_KOKKOS_SRCE_COUNT=" "env, Fortran run: ALL=0, nothing inherited") ww3gpu_kokkos_env(e K "ALL" "SRCE" "") -expect("${e}" "WW_KOKKOS_ALL=1;WW_KOKKOS_STRICT=1;WW_KOKKOS_PRESET=;WW_KOKKOS_SRCE=" "env, Kokkos run of all") +expect("${e}" "WW_KOKKOS_ALL=1;WW_KOKKOS_STRICT=;WW_KOKKOS_VA_POISON=1;WW_KOKKOS_PRESET=;WW_KOKKOS_SRCE=" "env, Kokkos run of all: fallbacks counted, not stopped; poison on") ww3gpu_kokkos_env(e G "ALL" "SRCE" "") -expect("${e}" "WW_KOKKOS_ALL=0;WW_KOKKOS_STRICT=;WW_KOKKOS_PRESET=;WW_KOKKOS_SRCE=" "env, the repeated Fortran run of all") +expect("${e}" "WW_KOKKOS_ALL=0;WW_KOKKOS_STRICT=;WW_KOKKOS_VA_POISON=;WW_KOKKOS_PRESET=;WW_KOKKOS_SRCE=" "env, the repeated Fortran run of all") set(pl " Kokkos parts (K Kokkos, F Fortran): SNL1=F SRCE=K GATH=F XYP3=K KTP3=K (3 of 5 on Kokkos)") ww3gpu_parts_line_check(m "${pl}" "SRCE;XYP3;KTP3") expect("${m}" "" "parts line, every part on Kokkos") diff --git a/regtests/cmake/ww3gpu_common.cmake b/regtests/cmake/ww3gpu_common.cmake index 7f6e0bbd..bd3a829c 100644 --- a/regtests/cmake/ww3gpu_common.cmake +++ b/regtests/cmake/ww3gpu_common.cmake @@ -141,10 +141,15 @@ endfunction() # ww3gpu_kokkos_env( ): the environment of one run, a list of # NAME=VALUE, NAME= meaning unset. Every WW_KOKKOS_ of is unset unless the Kokkos run # turns it on, WW_KOKKOS_ALL is 0 except in a Kokkos run whose has ALL, and the Kokkos run -# sets every counter of . The Kokkos run sets WW_KOKKOS_STRICT=1, so a part it asks for -# that would run the Fortran stops the run instead of comparing the Fortran with itself (PATCH.md -# section 24); the Fortran runs unset it. No run takes WW_KOKKOS_PRESET from the shell. The -# Fortran runs never inherit a flag from the shell. +# sets every counter of . A Kokkos run that names its parts (a preset or a part list, not +# ALL) sets WW_KOKKOS_STRICT=1, so a part it asks for that would run the Fortran stops the run +# instead of comparing the Fortran with itself (PATCH.md section 24). Under ALL a part the entry's +# switch does not cover is counted at exit, not stopped: ALL asks for every part the build covers, +# and ww3gpu_parts_line_check accepts partial coverage. The Fortran runs unset it. Every Kokkos +# run sets WW_KOKKOS_VA_POISON=1, the missed-sync detector (PATCH.md section 21), unless +# WW3GPU_VA_POISON is 0 (run_entry.cmake sets it from a caller's WW_KOKKOS_VA_POISON=0); the +# Fortran runs unset it. No run takes WW_KOKKOS_PRESET from the shell. The Fortran runs never +# inherit a flag from the shell. function(ww3gpu_kokkos_env out k flags parts counts) set(e "") if(k STREQUAL "K" AND "ALL" IN_LIST flags) @@ -152,11 +157,20 @@ function(ww3gpu_kokkos_env out k flags parts counts) else() list(APPEND e "WW_KOKKOS_ALL=0") endif() - if(k STREQUAL "K") + if(k STREQUAL "K" AND NOT "ALL" IN_LIST flags) list(APPEND e "WW_KOKKOS_STRICT=1") else() list(APPEND e "WW_KOKKOS_STRICT=") endif() + if(k STREQUAL "K") + if(DEFINED WW3GPU_VA_POISON AND WW3GPU_VA_POISON STREQUAL "0") + list(APPEND e "WW_KOKKOS_VA_POISON=0") + else() + list(APPEND e "WW_KOKKOS_VA_POISON=1") + endif() + else() + list(APPEND e "WW_KOKKOS_VA_POISON=") + endif() list(APPEND e "WW_KOKKOS_PRESET=") foreach(p IN LISTS parts) if(k STREQUAL "K" AND p IN_LIST flags) diff --git a/scripts/restart_pair.sh b/scripts/restart_pair.sh index c04018e1..4cc58a70 100755 --- a/scripts/restart_pair.sh +++ b/scripts/restart_pair.sh @@ -90,7 +90,7 @@ run() { # grep -q "End of program" "$d/ww3_shel.out" || { echo "restart_pair: ww3_shel did not finish in $d" >&2; return 1; } if [ "$1" = kokkos ]; then local p line - line="$(grep -h "Kokkos parts" "$d/ww3_shel.out" | head -1)" + line="$(grep -h "Kokkos parts" "$d/log.ww3" | head -1 || true)" for p in ${preset//,/ }; do [[ "$line" == *"$p=K"* ]] || { echo "restart_pair: $d: $p is not K in '$line'" >&2; return 1; } done diff --git a/scripts/ww3-gpu.sh b/scripts/ww3-gpu.sh index c8400cf8..b051fde5 100755 --- a/scripts/ww3-gpu.sh +++ b/scripts/ww3-gpu.sh @@ -1,7 +1,7 @@ #!/usr/bin/env bash # ww3-gpu — run WAVEWATCH three ways. WW3's Fortran comes from two submodules: WW3-noaa/ is pure # NOAA-EMC/WW3 at the commit the fork branched from, and WW3/ is the fork (the h0ffmann/WW3 branch -# ww3-gpu/kokkos-wire-base: WW3-noaa plus the KOKKOS switch and the shim calls). WW4's C++ comes only from +# ww3-gpu/kokkos-pnt, which holds the pin: WW3-noaa plus the KOKKOS switch and the shim calls). WW4's C++ comes only from # the WW4/ submodule (h0ffmann/WW4, develop); this repository holds the Kokkos port and nothing of # either model's source: # @@ -67,6 +67,14 @@ # ww3_shel.out says how many calls each part made. It fails when the log's "Kokkos parts" line # (W3INIT) does not show a listed part as K. # +# Every Kokkos run, with --kokkos or WW3GPU_ROUTINES, has WW_KOKKOS_VA_POISON=1: after each kernel +# write of the resident VA the stale host copy is filled with NaN, so a host reader that misses +# its sync changes an output file and the byte comparison fails (PATCH.md section 21). It fills +# only a copy no correct reader reads, so it cannot change the bits of a correct run; the line +# "VA host copies poisoned" of ww3_shel.out counts the fills. WW_KOKKOS_VA_POISON=0 in the +# caller's environment turns it off (a --repeat timing without the fills, for instance). The +# Fortran runs never have it. +# # --kokkos gpu is the part list for a run on Kokkos' CUDA backend, SRCE alone (ADR-0013): phase 1 # copies one sea point per call, so W3SRCE routed whole makes one round trip per point and time # step with the DIA and ST4 inside it, where SNL1, SDS4, SIN4 and SPR4 routed one by one make one @@ -168,6 +176,8 @@ kokkos_env() { for p in "${parts[@]}"; do printf ' %s' "WW_KOKKOS_$p=1"; done echo } +# va_poison: the WW_KOKKOS_VA_POISON of the Kokkos run, 1 unless the caller set it. +va_poison() { echo "${WW_KOKKOS_VA_POISON:-1}"; } # model_files : the files ww3_shel and ww3_ounf write in a run directory, the gridded and # point output, the restarts and the NetCDF fields, one name per line. Each is compared byte for # byte: a defect can reach the point output or a restart and leave the field file alone @@ -262,6 +272,15 @@ if [ "$mode" = self-test ]; then || { echo "ww3-gpu: --kokkos gpu: $(kokkos_env 1 "$(kokkos_preset gpu)")" >&2; exit 1; } [ "$(kokkos_env 1 "$(kokkos_preset resident)")" = "WW_KOKKOS_ALL=0 WW_KOKKOS_SRCE=1 WW_KOKKOS_XYP3=1 WW_KOKKOS_KTP3=1 WW_KOKKOS_OUTG=1 WW_KOKKOS_IOPE=1 WW_KOKKOS_UPDV=1 WW_KOKKOS_CPRT=1" ] \ || { echo "ww3-gpu: --kokkos resident: $(kokkos_env 1 "$(kokkos_preset resident)")" >&2; exit 1; } + [ "$(unset WW_KOKKOS_VA_POISON; va_poison)" = 1 ] || { echo "ww3-gpu: WW_KOKKOS_VA_POISON is not on by default" >&2; exit 1; } + [ "$(WW_KOKKOS_VA_POISON=0 va_poison)" = 0 ] || { echo "ww3-gpu: WW_KOKKOS_VA_POISON=0 did not turn it off" >&2; exit 1; } + # The two lines of run_model, read as text: run_model itself runs the model. + # shellcheck disable=SC2016 + grep -qxF ' [ "$k" = 0 ] || flags+=("WW_KOKKOS_VA_POISON=$(va_poison)")' "${BASH_SOURCE[0]}" \ + || { echo "ww3-gpu: run_model lost the Kokkos run's WW_KOKKOS_VA_POISON" >&2; exit 1; } + # shellcheck disable=SC2016 + grep -qxF ' [ "$k" = 1 ] || unset+=(-u WW_KOKKOS_STRICT -u WW_KOKKOS_VA_POISON)' "${BASH_SOURCE[0]}" \ + || { echo "ww3-gpu: run_model no longer unsets WW_KOKKOS_VA_POISON in the Fortran run" >&2; exit 1; } [ "$(kokkos_preset SNL1,GATH)" = "SNL1,GATH" ] || { echo "ww3-gpu: kokkos_preset changed a part list" >&2; exit 1; } # cmp_files: a point output or restart that differs, or exists on one side only, fails; the # NetCDF is compared only when asked; logs are never compared. @@ -484,8 +503,8 @@ else fi fi if [ "$mode" = ww3-gpu ] && [ ! -f "$ww3/model/src/w3kokkosmd.F90" ]; then - echo "ww3-gpu: WW3/ is not the ww3-gpu/kokkos-wire-base branch of h0ffmann/WW3 (no model/src/w3kokkosmd.F90):" >&2 - echo " git -C WW3 fetch origin ww3-gpu/kokkos-wire-base && git -C WW3 switch ww3-gpu/kokkos-wire-base" >&2 + echo "ww3-gpu: WW3/ is not the ww3-gpu/kokkos-pnt branch of h0ffmann/WW3 (no model/src/w3kokkosmd.F90):" >&2 + echo " git submodule update --init WW3 (the pin; h0ffmann/WW3 ww3-gpu/kokkos-pnt follows it)" >&2 exit 3 fi jobs="${JOBS:-$(nproc)}" @@ -584,9 +603,9 @@ run_model() { # <0|1>: grid, strt, shel, ounf; prints ww3_shel's wall time local -a flags=() unset=() for r in "${all_parts[@]}"; do unset+=(-u "WW_KOKKOS_$r"); done # No preset from the shell in either run (the Fortran run would turn its parts on); strict mode - # reaches the Kokkos run only (PATCH.md sections 24 and 25). + # reaches the Kokkos run only (PATCH.md sections 24 and 25), and so does the missed-sync poison. unset+=(-u WW_KOKKOS_PRESET) - [ "$k" = 1 ] || unset+=(-u WW_KOKKOS_STRICT) + [ "$k" = 1 ] || unset+=(-u WW_KOKKOS_STRICT -u WW_KOKKOS_VA_POISON) if [ -n "$kokkos" ]; then read -r -a flags <<<"$(kokkos_env "$k" "$kokkos")" [ "$k" = 0 ] || flags+=("${count_flags[@]}") @@ -596,6 +615,7 @@ run_model() { # <0|1>: grid, strt, shel, ounf; prints ww3_shel's wall time if [ "$k" = 1 ] && [[ " ${routines[*]} " == *" $r "* ]]; then flags+=("WW_KOKKOS_$r=1"); else flags+=("WW_KOKKOS_$r=0"); fi done fi + [ "$k" = 0 ] || flags+=("WW_KOKKOS_VA_POISON=$(va_poison)") [ "$k" = 0 ] || [ -z "$timers" ] || flags+=(WW_KOKKOS_TIMERS=1) (cd "$d" && "$bin/ww3_grid" > ww3_grid.out 2>&1) || { tail -20 "$d/ww3_grid.out" >&2; return 1; } # The forcing and the boundary spectra (the coast case), made by the Fortran in every run, so @@ -650,7 +670,7 @@ if [ -n "$timers" ]; then fi # W3WAVE no longer reads VA back at its return (PATCH.md section 26): each host reader outside it # syncs what it reads, and ww3_shel reads no VA after its last W3WAVE, so a run may end with the -# device copy newer. A host reader that misses its sync shows under WW_KOKKOS_VA_POISON=1. +# device copy newer. A host reader that misses its sync shows under WW_KOKKOS_VA_POISON=1, on here. if grep -q "still newer on the device at finalize" "$d1/ww3_shel.out"; then echo ">> $(grep -h "still newer on the device at finalize" "$d1/ww3_shel.out" | head -1) (expected: no read at W3WAVE's return)" fi diff --git a/site/data/port_score.tsv b/site/data/port_score.tsv index 4911a9e1..59425277 100644 --- a/site/data/port_score.tsv +++ b/site/data/port_score.tsv @@ -7,4 +7,4 @@ of the run time, regtested 48.6 % 12 of 1 023 routines WW3's regtests ran with i of the lines of code, translated 25.80 % 67 947 of 263 375 Fortran lines of the translated routines, of every line in WW3's model/src, in 65 of 147 files of the lines of code, ported 6.56 % 17 270 of 263 375 Fortran lines of the ported routines, of every line in WW3's model/src, in 25 of 147 files of the lines of code, regtested 1.02 % 2 689 of 263 375 Fortran lines of the regtested routines, of every line in WW3's model/src, in 6 of 147 files -lines of Kokkos C++ 60 037 in 350 files under kokkos/src: the port itself +lines of Kokkos C++ 60 563 in 351 files under kokkos/src: the port itself