[ggorman]

Fixes #3093.

A device parallel loop over a SubDimension or MultiSubDimension has bounds that depend on runtime thicknesses, so its iteration space can be empty, sometimes with a negative size. This happens, for example, in the overlap2 remainder regions outside a SubDomainSet. _make_guard only checked the size of non-affine regions, so these affine loops were launched unguarded. With nvc, an OpenACC tile/collapse nest with one negative and one positive trip count can launch 2^31-1 empty gangs. Each such launch takes about 1.3 s on an H200, which is why overlap2 with a SubDomainSet ran more than 100x slower than basic. #3093 has the details and a standalone C reproducer.

Fix: emit the existing symbolic_size > 0 guard, a C if around the region, also for affine regions that iterate over a SubDimension. Device loops over root Dimensions, sparse loops, CPU code, and the CUDA/HIP backends are unchanged.

For a SubDomainSet loop, the generated code becomes:

if (x_M - x_m - x_ltkn0 - x_rtkn0 + 1 > 0 && y_M - y_m - y_ltkn0 - y_rtkn0 + 1 > 0)
{
  #pragma acc parallel loop collapse(2) present(f)
  for (int x = x_m + x_ltkn0; x <= x_M - x_rtkn0; x += 1)
  ...

The guard is evaluated on the host, once per launch.

Tests:

• New TestMultiSubDomain::test_device_guard_empty_ispace[openacc|openmp]. It checks the guard in the generated code and runs without a GPU. It fails on main and passes here.
• test_composite_buffering_tasking and test_composite_buffering_tasking_multi_output count iteration trees. op0 now has one more, because the guard Conditional sits in the time loop.

Evidence:

• On an H200 with NVHPC 25.9, through DevitoPRO's OpenACC overlap2 path, 1 rank, ~9 steps, Operator time:
  • SubDomainSet: overlap2 went from 23.22 s to 0.01 s; basic is 0.01 s.
  • DevitoPRO ABox: overlap2 went from 58.07 s to 0.01 s; basic is 0.01 s.
  • Receiver norms are unchanged.
• With 4 ranks and the full-size test (~480 steps):
  • ABox overlap2 takes 3.78 s against 2.06 s for basic; it was killed at 300 s before.
  • Per-rank receiver norms match basic and the whole-domain run.
• Lint and tests:
  • ruff (preview), isort, flake8 and pre-commit are clean.
  • Under C, CXX and openmp: test_gpu_common.py, test_gpu_openacc.py, test_gpu_openmp.py, test_mpi.py, test_subdomains.py, test_operator.py and test_iet.py pass (CPU; the GPU modules skip).
  • The MPI tests of test_mpi.py, test_subdomains.py and test_operator.py pass serially under openmp.
  • test_adjoint.py, test_gpu_openacc.py, test_gpu_common.py, test_subdomains.py::TestMultiSubDomain and test_operator.py::TestEstimateMemory under OpenACC on the H200: main gives 289 passed and 2 failed, this branch 291 passed and 2 failed. The 2 extra passes are the new test, and the 2 failures are the same on both (TestMultiSubDomain::test_3D, test_subdomainset_mpi[(4, overlap)]).

🤖 Generated with Claude Code

[codecov]

Codecov Report

✅ All modified and coverable lines are covered by tests.
✅ Project coverage is 84.40%. Comparing base (61a2745) to head (01adabf).

Additional details and impacted files

@@            Coverage Diff             @@
##             main    #3096      +/-   ##
==========================================
- Coverage   84.40%   84.40%   -0.01%     
==========================================
  Files         260      260              
  Lines       58153    58164      +11     
  Branches     4948     4948              
==========================================
+ Hits        49086    49093       +7     
- Misses       8245     8248       +3     
- Partials      822      823       +1     

Flag	Coverage Δ
pytest-gpu-aomp-amdgpuX	68.45% <100.00%> (+<0.01%)	⬆️
pytest-gpu-gcc-	79.12% <88.23%> (+0.03%)	⬆️
pytest-gpu-icx-	79.04% <88.23%> (+0.02%)	⬆️
pytest-gpu-nvc-nvidiaX	69.14% <100.00%> (+<0.01%)	⬆️

Flags with carried forward coverage won't be shown. Click here to find out more.

☔ View full report in Codecov by Harness.
📢 Have feedback on the report? Share it here.

🚀 New features to boost your workflow:

• ❄️ Test Analytics: Detect flaky tests, report on failures, and find test suite problems.