Raising miscompiles mfem's shared-memory simplex diffusion apply (thread-0-guarded sections yield zeros)

[wsmoses]

The raised mfem::internal::SmemPADiffusionApplyTetrahedron<2,2> (mfem fem/integ/bilininteg_diffusion_pa_simplices.hpp) compiles strict but computes an all-zero output at runtime (beam-tet order 1, Q1D=2; the <2,1> single-point instantiation is exact, as are the 2D triangle kernels).

Post-pipeline module: https://gist.github.com/wsmoses/e438453754885ba06b6be3d81815477a

The suspicious shape in the raised IR:

%114 = stablehlo.reduce(%93 ...)         // tid.x + tid.y
%115 = stablehlo.compare LT, %114, %c_56 // thread-0 guard: tid sum < 1
...
%126 = stablehlo.select %115, %125#7, %cst_20  // carried 2x4xf64 accumulators ...
%127 = stablehlo.select %115, %125#6, %cst_20  // ... become ZERO on non-zero lanes

The kernel is a hard case: MFEM_FOREACH_THREAD_DIRECT sections, MFEM_SHARED scratch unioned under several shapes (sm0/sm1 reinterpreted as X/DDQ/DQQ/QQQ/QQD/QDD views across phases), and thread-0-only writes. The masked raise appears to zero the loop-carried accumulators for lanes outside the guard rather than preserving them.

Reproduction: mfem PA Simplices GPU unit test through the raising path (#2968); a minimal driver is

// beam-tet.mesh, H1 order 1 BasisType::Positive, DiffusionIntegrator with
// StroudIntRules order 2p+OrderW (nq=8): AddMultPA returns y == 0.

Worked around in the mfem harness by dispatching to the per-element fallback (PADiffusionApplyTetrahedron) under the XLA backend, which raises and runs bit-correct — so this is not blocking, but the shared-memory shape-punned scratch + direct-thread-section pattern will recur (mfem uses it widely).

🤖 Generated with Claude Code

https://claude.ai/code/session_016zErYp7upmqr4NHfhod9UD

[wsmoses]

Root cause found (very likely the same for this tet kernel): privatizeLaneScratch's 65k-element size budget silently skipped the per-lane register tensors of the 3D register-blocked kernels (24×24 derived lane bound × 126+ registers ≈ 72k elements). The un-privatized scratch then raised as a single shared tensor: every lane read the wrong lane's values — for VecDivPA's dynamic-fallback kernel that produced errors of ~3e-2; for this tet kernel the same collapse plausibly yields the all-zero output.

Fixed by raising the budget (#2996, and the harness verified VecDivPA 3D goes from 3e-2 to 1e-17 across trial/test order combinations). Two follow-ups remain:

1. A skipped privatization must fail the strict raise loudly — this bug compiled clean and was wrong only at runtime.
2. Retest the shared-memory tet apply without the per-element-fallback dispatch once the full re-sweep lands, and drop the mfem-side workaround if it is fixed.