NVHPC emits a faulting implicit reduction(*) for scalar self-multiplication inside offloaded loops

[sbryngelson]

Summary

Multiplying a scalar into itself inside an offloaded loop makes NVHPC infer an implicit
reduction(*:<scalar>) for that loop — even when the scalar is named in the directive's own
private clause. The reduction epilogue it then emits loads and CAS-updates a global accumulator
that was never allocated, so the kernel dies with CUDA_ERROR_ILLEGAL_ADDRESS.

The epilogue runs at kernel exit unconditionally, so the crash does not require the offending
statement to be reachable at runtime.

Where it bit

s_convert_conservative_to_primitive_variables in src/common/m_variables_conversion.fpp:

if (hypoelasticity) then
    if (cont_damage) G_K = G_K*max((1._wp - qK_cons_vf(eqn_idx%damage)%sf(j, k, l)), 0._wp)
    ...
        ... f_elastic_energy(real(qK_prim_vf(i)%sf(j, k, l), wp), G_K, any(i == shear_indices)) ...

G_K is listed in the GPU_PARALLEL_LOOP(... private='[... G_K ...]') of the same loop. NVHPC
disregards that:

s_convert_conservative_to_primitive_variables:
  441, !$omp target teams loop
  444, Loop parallelized across teams, threads(128) collapse(3)
  444, Generating implicit reduction(*:g_k)

hypoelasticity and cont_damage are runtime variables, not case-optimization constants, so the
block is compiled into the kernel for every case. Four unrelated case-optimized gpu-omp
benchmarks — 5eq_rk3_weno3_hllc, viscous_weno5_sgb_acoustic, hypo_hll, ibm — all aborted at
the first RHS evaluation, in code that is dead for three of them. igr was the only survivor.

Evidence

compute-sanitizer, 465 identical errors:

Invalid __global__ read of size 8 bytes
  at nvkernel_m_variables_conversion_s_convert_conservative_to_primitive_variables__F8L441_2_+0x11e00
  Access to 0x7e6176efa000 is out of bounds
  Host Frame: s_convert_conservative_to_primitive_variables in m_variables_conversion.fpp:441
  Host Frame: m_rhs_s_compute_rhs_ in m_rhs.fpp:639
  Host Frame: m_time_steppers_s_tvd_rk_ in m_time_steppers.fpp:460

The SASS at that offset is the reduction epilogue, not the loop body — warp/block product reduction
(SHFL.DOWN + DMUL, identity 1.0) followed by the atomic combine:

/*11de0*/  MOV R2, c[0x0][0x168]     ; accumulator pointer, kernel param
/*11e00*/  LDG.E.64 R8, [R2.64]      ; <-- the faulting 8-byte read
/*11e20*/  DMUL R10, R6, R8
/*11e40*/  ATOMG.E.CAS.64.STRONG.GPU PT, R10, [R2], R8, R10

Reduction instructions in that one kernel: 15 on the affected build, 0 on master.

Scope

• OpenMP offload only. The same source built with --gpu acc runs clean; !$acc parallel loop
  with an explicit private does not get the spurious reduction.
• Reproduced on two unrelated systems: Phoenix (H200, NVHPC 24.5) and a local box (A100,
  NVHPC 25.11). On 24.5 the fault is masked behind NVHPC's own assert
  (HX_CU_CALL_CHECK(p_cuStreamSynchronize(...))), which reports nothing useful; 25.11 reports
  CUDA_ERROR_ILLEGAL_ADDRESS directly.
• The same x = x*(...) shape appears in offloaded loops at src/simulation/m_hypoelastic.fpp:204,
  :405 and :565. Those did not produce an implicit reduction in the builds checked here —
  reduction(*: appears exactly twice in the whole simulation build, both from the site above — but
  they are the same hazard and are worth auditing. (m_variables_conversion.fpp:898 has the shape
  too but is host-only; that subroutine has no offloaded loop.)

Fix

Route the result into a separate variable so no scalar is self-assigned in the loop, which leaves
nothing for the compiler to pattern-match as a reduction. Fixed by #1762 (commit e85d162). After
the fix all four benchmarks complete and pass the NaN/Inf check, and the kernel's reduction
instruction count drops to 0, matching master.