Restore HIP feature parity with CUDA; make hipify.sh maintainable

[susilehtola]

Problem

The HIP backend does not build. scheme1_base.cxx is compiled for both backends — it is listed in src/xc_integrator/local_work_driver/device/CMakeLists.txt:16, before the conditional cuda/ and hip/ subdirectories — and it calls zmat_*_fxc, increment_exc_grad_*, exx_ek_screening_bfn_stats and the shell-to-task collocation. None of those had HIP definitions, so any -DGAUXC_ENABLE_HIP=ON build fails to link.

The HIP sources that did exist were stale as well:

• kernels/zmat_vxc.hip still had the pre-UKS/GKS kernel names (zmat_lda_vxc_kernel rather than zmat_lda_vxc_{rks,uks,gks}_kernel);
• kernels/uvvars.hip had no mGGA path at all (zero tau references against 19 on the CUDA side);
• collocation_shell_to_task_kernels.hpp and the split uvvars_{lda,gga,mgga}.hpp headers were absent.

In effect the HIP backend has been frozen at roughly its 2022 LDA/GGA-VXC feature level while mGGA, FXC contraction, EXC gradients and sn-LinK screening were added on the CUDA side.

Root cause

hip/hipify.sh covered only 8 of the 24 current files under cuda/kernels/ and was never updated as kernels were added.

What this PR does

Rewrites hipify.sh to cover all current CUDA kernel sources, regenerates the HIP tree from it, and adds zmat_fxc, increment_exc_grad and exx_ek_screening_bfn_stats to hip/CMakeLists.txt.

The hipify-perl invocations are replaced with explicit sed rules. Two reasons: regeneration then works on machines without a ROCm installation, and the handful of decisions that are not mechanical end up documented in one place rather than in a maintainer's head:

• Wavefront size is never hardcoded; kernels take it from GauXC::cuda::warp_size → GauXC::hip::warp_size (32 vs 64), so launch geometry adapts. No literal 32s were introduced.
• __syncwarp() has no HIP equivalent (wavefronts are lockstep on AMD) and is commented out — the convention the existing 2022-era files in this directory already used.
• __shfl_*_sync(mask, …) → __shfl_*(…), since HIP shuffles take no mask argument.
• cub → hipcub (cub::DeviceScan in the EXX screening kernel).
• CUDA's two-tier shared-memory limit has no AMD counterpart — LDS per workgroup is a single fixed limit — so both cudaDevAttrMaxSharedMemoryPerBlock and …Optin map to hipDeviceAttributeMaxSharedMemoryPerBlock. The overflow guard in exx_ek_screening is preserved; the opt-in branch becomes a no-op.
• __stcs() and its inline-PTX fallback in pack_submat are lowered to a plain store. CUDART_VERSION is undefined under HIP, so the preprocessor would otherwise select the PTX branch and fail to compile. The cache hint is an optimization, not semantics.
• CUTLASS is deliberately not translated; it has no HIP counterpart and GAUXC_ENABLE_CUTLASS is already a CUDA-only dependent option.

Verification

Static checks only (see caveat):

• every kernel symbol referenced by scheme1_base.cxx now resolves against the HIP tree;
• no compile-breaking CUDA token remains in the generated sources;
• per-file function counts match the CUDA originals;
• the hand-maintained hip_aos_scheme1{,_data}.cxx were checked to have the same method sets as their CUDA counterparts (the line-count difference is CUTLASS-only code), so they did not need regeneration.

Caveat — please read before merging

This has not been compiled. I have neither a ROCm installation nor AMD hardware, so the changes have never been through hipcc. Expect a round of compile fixes.

Beyond compilation, the one correctness risk worth a reviewer's eyes is wavefront width: nothing hardcodes 32 and launch geometry derives from hip::warp_size, but any intra-wavefront reduction whose trip count is written as a literal rather than derived from warp_size should be checked before the numbers are trusted on AMD. hipify.sh prints a CHECK_WAVEFRONT reminder to that effect.

I am happy to iterate on this if someone with an AMD GPU can run the test suite, or to split it into "script + build wiring" and "regenerated kernels" commits if that is easier to review.