From 77560c2609eecf434c41a305b298caf718db03d3 Mon Sep 17 00:00:00 2001 From: "M.Hoffmann" Date: Sun, 18 Oct 2026 17:32:36 -0300 Subject: [PATCH 1/9] ww3-gpu.sh, regtest matrix: WW_KOKKOS_VA_POISON=1 in every Kokkos run Goal abbb0456-8241, the follow-up of #191 (code review finding 2). Since #191 the "still newer on the device at finalize" check only prints, and its replacement, WW_KOKKOS_VA_POISON=1, was set by no script: a host reader that missed its column sync passed unless the stale bytes changed an output file. Every Kokkos run of scripts/ww3-gpu.sh (run_model, with --kokkos or WW3GPU_ROUTINES) and of the regtest matrix (ww3gpu_kokkos_env, the K case) now has WW_KOKKOS_VA_POISON=1, so the stale host VA is NaN after each kernel write and a missed sync fails the byte comparison. The Fortran runs (F, G, and ww3-gpu.sh's k=0) unset it. A caller's WW_KOKKOS_VA_POISON=0 still turns it off: ww3-gpu.sh reads it with ${WW_KOKKOS_VA_POISON:-1}, run_entry.cmake reads it once before the runs rewrite the environment. The poison fills only a copy a correct reader never reads, so a correct run's bits do not change. This changes a gate's invocation (AGENTS.md, "Never cheat a regtest tool"): a person reviews it. Tested: bash -n and shellcheck scripts/ww3-gpu.sh (clean); bash scripts/ww3-gpu.sh --self-test (ok, with the new va_poison and run_model checks); just ww3-gpu-matrix-selftest (ok, with the new POISON expectations); ww3-gpu.sh --ww3-gpu smoke and full --cpu --kokkos resident on the pinned fork commit: byte-identical to the Fortran and to --ww3 with POISON on (2161 and 97 fills) Co-Authored-By: Claude Opus 5.5 --- regtests/cmake/run_entry.cmake | 10 ++++++++-- regtests/cmake/selftest.cmake | 12 ++++++++---- regtests/cmake/ww3gpu_common.cmake | 16 ++++++++++++++-- scripts/ww3-gpu.sh | 26 +++++++++++++++++++++++--- 4 files changed, 53 insertions(+), 11 deletions(-) diff --git a/regtests/cmake/run_entry.cmake b/regtests/cmake/run_entry.cmake index 36c0b856..1d984dc5 100644 --- a/regtests/cmake/run_entry.cmake +++ b/regtests/cmake/run_entry.cmake @@ -15,12 +15,18 @@ # 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). +# 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..74ae2c23 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=1;WW_KOKKOS_VA_POISON=1;WW_KOKKOS_PRESET=;WW_KOKKOS_SRCE=" "env, Kokkos run of all: 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..497c06cd 100644 --- a/regtests/cmake/ww3gpu_common.cmake +++ b/regtests/cmake/ww3gpu_common.cmake @@ -143,8 +143,11 @@ endfunction() # 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. +# section 24); 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) @@ -157,6 +160,15 @@ function(ww3gpu_kokkos_env out k flags parts counts) 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/ww3-gpu.sh b/scripts/ww3-gpu.sh index c8400cf8..6dc50a2a 100755 --- a/scripts/ww3-gpu.sh +++ b/scripts/ww3-gpu.sh @@ -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. @@ -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 From 156ba8ccd18f41a56757be7f17e0578a0eec7580 Mon Sep 17 00:00:00 2001 From: "M.Hoffmann" Date: Sun, 18 Oct 2026 18:09:29 -0300 Subject: [PATCH 2/9] regtest matrix: WW_KOKKOS_STRICT=1 only for a Kokkos run that names its parts Goal abbb0456-8241, the follow-up of #191 (code review finding 1). With WW3GPU_KOKKOS=all the Kokkos run had WW_KOKKOS_ALL=1 and WW_KOKKOS_STRICT=1 together, so KOKKOS_PART's fallback for a part the entry's switch does not cover (MAP3 and MAPT with PR3, SRCE without LN1, BT1 or DB1) stopped every such entry in W3INIT, while ww3gpu_parts_line_check(ALL) accepts partial coverage and ww3-gpu.sh --kokkos all sets no strict mode. ww3gpu_kokkos_env now sets WW_KOKKOS_STRICT=1 for the K run only when its flags do not hold ALL (WW3GPU_KOKKOS=resident or a part list); under ALL the fallback is counted at exit. The selftest expectation "env, Kokkos run of all" follows. This changes a gate's invocation (AGENTS.md, "Never cheat a regtest tool"): a person reviews it. Tested: just ww3-gpu-matrix-selftest (ok); WW3GPU_KOKKOS=all just ww3-gpu-matrix-list serial (363 entries; 38 with WW3GPU_SWITCH_WORDS=PR0;SHRD); WW3GPU_KOKKOS=all WW3GPU_SWITCH_WORDS="PR0;SHRD" WW3GPU_OMP_THREADS=8 just ww3-gpu-matrix 1 serial -R ww3_tnc1: entry 38 runs and passes, test.comp 0/30, 10 of 16 parts on Kokkos, "calls left on the Fortran: MAP3=1 MAPT=1 SRCE=1" Co-Authored-By: Claude Opus 5.5 --- regtests/cmake/run_entry.cmake | 4 ++-- regtests/cmake/selftest.cmake | 2 +- regtests/cmake/ww3gpu_common.cmake | 10 ++++++---- 3 files changed, 9 insertions(+), 7 deletions(-) diff --git a/regtests/cmake/run_entry.cmake b/regtests/cmake/run_entry.cmake index 1d984dc5..1fcd87a1 100644 --- a/regtests/cmake/run_entry.cmake +++ b/regtests/cmake/run_entry.cmake @@ -14,8 +14,8 @@ # 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 +# 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. diff --git a/regtests/cmake/selftest.cmake b/regtests/cmake/selftest.cmake index 74ae2c23..75af7a84 100644 --- a/regtests/cmake/selftest.cmake +++ b/regtests/cmake/selftest.cmake @@ -201,7 +201,7 @@ 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_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_VA_POISON=1;WW_KOKKOS_PRESET=;WW_KOKKOS_SRCE=" "env, Kokkos run of all: poison on") +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_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)") diff --git a/regtests/cmake/ww3gpu_common.cmake b/regtests/cmake/ww3gpu_common.cmake index 497c06cd..bd3a829c 100644 --- a/regtests/cmake/ww3gpu_common.cmake +++ b/regtests/cmake/ww3gpu_common.cmake @@ -141,9 +141,11 @@ 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. Every Kokkos +# 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 @@ -155,7 +157,7 @@ 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=") From 7e9d1094c73a99924de2d35d440047c923609745 Mon Sep 17 00:00:00 2001 From: "M.Hoffmann" Date: Sun, 18 Oct 2026 18:09:53 -0300 Subject: [PATCH 3/9] Name ww3-gpu/kokkos-pnt as the fork branch that holds the WW3 pin Goal abbb0456-8241. .gitmodules still had branch = ww3-gpu/kokkos-wire-base for WW3, so git submodule update --remote would have moved the pin back to 4a94a8f6; it now names ww3-gpu/kokkos-pnt, which the owner fast-forwarded to the pin 50a10730. The comments and messages that named older fork branches (ww3-gpu/kokkos-wire-base, kokkos-st4, kokkos-resident) as "the fork" now name kokkos-pnt: ww3-gpu.sh's header and its "not the fork" message, whose hint is now git submodule update --init WW3 (the pin, not a branch head), the justfile, health.yml, regtests/CMakeLists.txt, kokkos/src/fortran_iface/CMakeLists.txt, resident_c.hpp and srce_shim.cpp. Text only; no logic changes. PATCH.md's dated history keeps the names it recorded. Tested: bash -n and shellcheck scripts/ww3-gpu.sh (clean); bash scripts/ww3-gpu.sh --self-test (ok); just ww3-gpu-matrix-selftest (ok); git grep -nE 'ww3-gpu/kokkos-(wire-base|st4|resident)' -- .gitmodules scripts justfile regtests .github kokkos/src finds only PATCH.md's history Co-Authored-By: Claude Opus 5.5 --- .github/workflows/health.yml | 2 +- .gitmodules | 2 +- justfile | 2 +- kokkos/src/fortran_iface/CMakeLists.txt | 6 +++--- kokkos/src/fortran_iface/resident_c.hpp | 2 +- kokkos/src/fortran_iface/srce_shim.cpp | 2 +- regtests/CMakeLists.txt | 6 +++--- scripts/ww3-gpu.sh | 6 +++--- 8 files changed, 14 insertions(+), 14 deletions(-) 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/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/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/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 105d3336..f406685d 100644 --- a/kokkos/src/fortran_iface/srce_shim.cpp +++ b/kokkos/src/fortran_iface/srce_shim.cpp @@ -817,7 +817,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/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/scripts/ww3-gpu.sh b/scripts/ww3-gpu.sh index 6dc50a2a..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: # @@ -503,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)}" From 04634cf58256d3f711220518fb4727aed0ff974c Mon Sep 17 00:00:00 2001 From: "M.Hoffmann" Date: Sun, 18 Oct 2026 18:20:20 -0300 Subject: [PATCH 4/9] wip: docs, PATCH.md 31 and the WW3 pin fac74069 for the follow-up of #191 Stopped by the owner before the last step; another session takes it over. Done and recorded in docs/log/2026-10-18-strict-outg-poison.md: the fork commit fac74069 (W3KOKKOS_OUTG's fallback through KOKKOS_FALLBACK), POISON by default in Kokkos runs, strict only where a run names its parts, .gitmodules on ww3-gpu/kokkos-pnt. Not done: the resident serial matrix result (MATRIXRESULT placeholder in the log), port_score --write after the pin bump, the security review, the PR. Tested: per the log entry: gfortran -E -P of 143 files without KOKKOS; smoke and full resident byte-identical to the Fortran and to --ww3 with POISON on; full with SKW stops under WW_KOKKOS_STRICT=1 and counts OUTG=3 without; all-mode ww3_tnc1 test.comp 0/30. Not run here: the resident matrix Co-Authored-By: Claude Opus 5.5 --- AGENTS.md | 5 +- CONTRIBUTING.md | 4 +- WW3 | 2 +- docs/GLOSSARY.md | 4 +- docs/PRODUCTION_READINESS_202610.md | 13 +++- docs/log/2026-10-18-strict-outg-poison.md | 84 +++++++++++++++++++++++ docs/log/README.md | 1 + kokkos/src/fortran_iface/PATCH.md | 50 ++++++++++++++ 8 files changed, 157 insertions(+), 6 deletions(-) create mode 100644 docs/log/2026-10-18-strict-outg-poison.md 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 096ca35c..9651799e 100644 --- a/docs/GLOSSARY.md +++ b/docs/GLOSSARY.md @@ -702,9 +702,9 @@ 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_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/PRODUCTION_READINESS_202610.md b/docs/PRODUCTION_READINESS_202610.md index 52b3d126..d1de1ab3 100644 --- a/docs/PRODUCTION_READINESS_202610.md +++ b/docs/PRODUCTION_READINESS_202610.md @@ -42,7 +42,7 @@ 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 on fork branch `ww3-gpu/u2-strict-preset`, not yet at the pin: `WW_KOKKOS_STRICT=1` stops, otherwise one fallback-count line at exit; the matrix's Kokkos runs set it | +| 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 | @@ -302,6 +302,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 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..fc66a4ba --- /dev/null +++ b/docs/log/2026-10-18-strict-outg-poison.md @@ -0,0 +1,84 @@ +# 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 | MATRIXRESULT | + +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"). +- 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 62dc56a1..7669d485 100644 --- a/docs/log/README.md +++ b/docs/log/README.md @@ -72,3 +72,4 @@ 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 | 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). + From 1d0e184b325ae910fe2b89f211152eb3d231d495 Mon Sep 17 00:00:00 2001 From: "M.Hoffmann" Date: Sun, 18 Oct 2026 19:40:24 -0300 Subject: [PATCH 5/9] docs/log: resident serial matrix of the strict-outg-poison branch Fills the matrix row of the 2026-10-18 strict-outg-poison entry with the real verdicts at the pin fac74069 merged with origin/main: five of seven entries pass, the two ww3_tp2.6 entries fail by design (unstructured mesh, kept FAIL by the owner). The branch's earlier wip: commit holds the docs, PATCH.md section 31 and the pin; this commit completes it. Tested: WW3GPU_KOKKOS=resident WW3GPU_LAB_SWITCH=ON WW3GPU_OMP_THREADS=2 just ww3-gpu-matrix 4 serial (5 PASS, 2 FAIL by design, 729.7 s); bash -n and bash scripts/ww3-gpu.sh --self-test ok; just ww3-gpu-matrix-selftest ok; just ww3-gpu-matrix-list 7 entries; port_score.py --check current; wfip.py check 0 problems; python3 -m unittest discover tests ok; just kokkos-test serial-debug and openmp-release 100% of 11; shellcheck and vale not installed here (not run) Co-Authored-By: Claude Sonnet 5.5 --- docs/log/2026-10-18-strict-outg-poison.md | 6 +++++- 1 file changed, 5 insertions(+), 1 deletion(-) diff --git a/docs/log/2026-10-18-strict-outg-poison.md b/docs/log/2026-10-18-strict-outg-poison.md index fc66a4ba..3441f2fa 100644 --- a/docs/log/2026-10-18-strict-outg-poison.md +++ b/docs/log/2026-10-18-strict-outg-poison.md @@ -61,7 +61,7 @@ WW3GPU_KOKKOS=resident WW3GPU_LAB_SWITCH=ON WW3GPU_OMP_THREADS=2 just ww3-gpu-ma | `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 | MATRIXRESULT | +| 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: @@ -80,5 +80,9 @@ Two things the runs show beyond the question: 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. From d3dcb7e7e3f3d0104eeb2f1a9ee6de6c75c5c930 Mon Sep 17 00:00:00 2001 From: "M.Hoffmann" Date: Sun, 18 Oct 2026 19:51:30 -0300 Subject: [PATCH 6/9] restart pair: run it for the first time, fix the parts-line plumbing, record gap 10 Goal 5afd9379-8152. scripts/restart_pair.sh read the "Kokkos parts" line from ww3_shel.out, where WW3 does not print it (log.ww3 has it), and grep under pipefail ended the script silently with exit 1. It now reads log.ww3. What the script compares is unchanged. The pair then ran on main b828041 with WW3 at 50a10730 (SHRD KOKKOS build, OpenMP, 4 threads, WW_KOKKOS_VA_POISON=1, WW_KOKKOS_STRICT=1): b against a identical for the Fortran and for the resident Kokkos (out_grd.ww3 and the 12:00 restart), and every file of Kokkos against Fortran identical. Gap 10 of the readiness list and a dated log say what that covers and what it does not. Tested: scripts/restart_pair.sh exit 0 (all comparisons identical); restart_pair.sh --self-test passed; bash -n and shellcheck clean; just vale 0 errors. Co-Authored-By: Claude Sonnet 5.5 --- docs/PRODUCTION_READINESS_202610.md | 11 +++- docs/log/2026-10-18-resident-restart-pair.md | 66 ++++++++++++++++++++ scripts/restart_pair.sh | 2 +- 3 files changed, 77 insertions(+), 2 deletions(-) create mode 100644 docs/log/2026-10-18-resident-restart-pair.md diff --git a/docs/PRODUCTION_READINESS_202610.md b/docs/PRODUCTION_READINESS_202610.md index a4efc1a1..139829ae 100644 --- a/docs/PRODUCTION_READINESS_202610.md +++ b/docs/PRODUCTION_READINESS_202610.md @@ -404,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/log/2026-10-18-resident-restart-pair.md b/docs/log/2026-10-18-resident-restart-pair.md new file mode 100644 index 00000000..6190a488 --- /dev/null +++ b/docs/log/2026-10-18-resident-restart-pair.md @@ -0,0 +1,66 @@ +# 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. 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 From 9cef132f3182731efe73906f9587a278865de77e Mon Sep 17 00:00:00 2001 From: "M.Hoffmann" Date: Sun, 18 Oct 2026 19:51:39 -0300 Subject: [PATCH 7/9] kokkos: W3SDS4 split into parts and run with one sea point per vector lane (plan 005 L3) Unit L3 of plan 005, goal 5afd9379-4417. w3sds4_point is split into sds4_prep, sds4_mss, sds4_sat, sds4_case1, sds4_case2, sds4_case3, sds4_turb and sds4_whitecap, the routine's lines moved without a change (only the case bodies are dedented), and w3sds4_point calls them in the routine's order. w3sds4_lanes<8> runs them for up to eight points: the mean-square-slope sums (w3src4md.F90:2228-2254), the saturation integral (:2262-2290) and the wave-turbulence and cumulative loops (:2533-2575) with one point per lane of GCC vector types, each lane keeping its own sums in the routine's order and every zero weight multiplied; branches on BTH0 are selects; powf, atan2f and sincosf are called per lane and the libmvec sites stay one grouped call per point. COS/SIN(USDIR) is computed once per lane instead of per ITH (the same pure call). The vector types move to st4_vec8.hpp for W3SIN4 and W3SDS4. w3srce_lanes calls W3SDS4 once per step for the points still in the loop; the scalar path is unchanged. Tested: just kokkos-test serial-debug and openmp-release, 11/11 each, no warning; bash scripts/ww3-gpu.sh --ww3-gpu smoke and full --cpu --kokkos resident with WW_KOKKOS_STRICT=1 WW_KOKKOS_VA_POISON=1 and WW_KOKKOS_SRCE_LANES=1 and 8: every model file byte-identical to the Fortran of the same executable, nccmp-tol PASS; bash scripts/ww3-gpu.sh --ww3-gpu smoke --gpu --kokkos resident --build-only compiles, no warning Co-Authored-By: Claude Sonnet 5.5 --- kokkos/src/ww_kokkos/srce_kernel.hpp | 33 +- kokkos/src/ww_kokkos/st4_sds4.hpp | 1088 +++++++++++++++++++------- kokkos/src/ww_kokkos/st4_sin4.hpp | 14 +- kokkos/src/ww_kokkos/st4_vec8.hpp | 21 + 4 files changed, 841 insertions(+), 315 deletions(-) create mode 100644 kokkos/src/ww_kokkos/st4_vec8.hpp 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 From debb1b04b0fae9f0ad6d50735ef8bc2785695e21 Mon Sep 17 00:00:00 2001 From: "M.Hoffmann" Date: Sun, 18 Oct 2026 20:13:00 -0300 Subject: [PATCH 8/9] docs/log: W3SDS4 with one sea point per vector lane (plan 005 L3), port report and plan updated The log records the question, the parts and the three lane loops, the bit verdicts on smoke and full with 1 and 8 points per team, the branches the full case does not reach (checked lane path against scalar path with a temporary parameter override, removed), WW3's regtests of the selected entries, the one-thread kernel timing (104.99 s to 80.55 s, L2 and L3 together) and the four open questions for the owner. The port report gains the L3 section and the plan marks L3. Tested: just regtest-selected (5 entries PASS, 8/8 ctest, smoke byte-identical); python3 scripts/port_score.py --write then --check, current; vale 3.17.1 on the four changed Markdown files, 0 errors, 0 warnings Co-Authored-By: Claude Sonnet 5.5 --- ...026-10-18-005-srce-points-in-lanes-plan.md | 2 +- docs/log/2026-10-18-srce-lanes-sds4.md | 100 ++++++++++++++++++ docs/log/README.md | 1 + kokkos/PORT_STATUS.md | 2 +- kokkos/port_report/W3SRCE_KOKKOS.md | 17 +++ site/data/port_score.tsv | 2 +- 6 files changed, 121 insertions(+), 3 deletions(-) create mode 100644 docs/log/2026-10-18-srce-lanes-sds4.md 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-srce-lanes-sds4.md b/docs/log/2026-10-18-srce-lanes-sds4.md new file mode 100644 index 00000000..adb0cdc7 --- /dev/null +++ b/docs/log/2026-10-18-srce-lanes-sds4.md @@ -0,0 +1,100 @@ +# `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`. diff --git a/docs/log/README.md b/docs/log/README.md index a52fd5f3..f787655d 100644 --- a/docs/log/README.md +++ b/docs/log/README.md @@ -77,3 +77,4 @@ Date, author, issue/PR/WFIP. | 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/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/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 From 02e14dd0160886bbb06865aafd010fdb4507874b Mon Sep 17 00:00:00 2001 From: "M.Hoffmann" Date: Sun, 18 Oct 2026 20:49:17 -0300 Subject: [PATCH 9/9] docs/log: restart pair and SDS4 lanes rerun at WW3 fac74069, stale "not run" lines fixed The three pieces of the branch were verified at WW3 50a10730. At fac74069 strict mode stops a run when W3OUTG falls back to the Fortran. The restart pair (full) and full pass under strict; smoke stops in W3OUTG because it requests a skewness field, so its bit verdict at fac74069 is the non-strict run. Both logs get a dated section; gap 10 and the two stale lines of the NOAA progress record now say the pair ran. Tested: just kokkos-test serial-debug and openmp-release, exit 0; WW_KOKKOS_SRCE_LANES=1 and 8 with WW_KOKKOS_VA_POISON=1, smoke and full --cpu --kokkos resident: non-strict byte-identical to the Fortran and nccmp-tol PASS, strict full the same, strict smoke stops in W3OUTG (exit 1); scripts/restart_pair.sh at fac74069 exit 0, every comparison identical; just regtest-selected passed; smoke --gpu --build-only exit 0; bash -n, shellcheck, both --self-test, just ww3-gpu-matrix-selftest, port_score.py --check, unittest discover tests, just wfip check, figures.py check, just vale (3.17.1) all exit 0 Co-Authored-By: Claude Sonnet 5.5 --- docs/NOAA_PROGRESS_RECORD_202610.md | 4 +-- docs/PRODUCTION_READINESS_202610.md | 2 +- docs/log/2026-10-18-resident-restart-pair.md | 13 +++++++++ docs/log/2026-10-18-srce-lanes-sds4.md | 30 ++++++++++++++++++++ 4 files changed, 46 insertions(+), 3 deletions(-) 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 139829ae..5fed01d6 100644 --- a/docs/PRODUCTION_READINESS_202610.md +++ b/docs/PRODUCTION_READINESS_202610.md @@ -51,7 +51,7 @@ regtest matrix yet"). 10 of the 310 rows of `kokkos/ports.tsv` are regtested and | 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 | diff --git a/docs/log/2026-10-18-resident-restart-pair.md b/docs/log/2026-10-18-resident-restart-pair.md index 6190a488..7627e080 100644 --- a/docs/log/2026-10-18-resident-restart-pair.md +++ b/docs/log/2026-10-18-resident-restart-pair.md @@ -64,3 +64,16 @@ would have shown as NaN or a different byte. - 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 index adb0cdc7..092a5979 100644 --- a/docs/log/2026-10-18-srce-lanes-sds4.md +++ b/docs/log/2026-10-18-srce-lanes-sds4.md @@ -98,3 +98,33 @@ sea points run its loops at once, one per vector lane? `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.