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.

[wavefunction91]

Hey @susilehtola, just wanted to let you know that I am taking a look at this. Currently seeing if I can get access to some appropriate AMD GPUs to test this myself. I'll also look into some options for CI/CD if I can get my hands on some quota from one of the big cloud providers.

[Copilot]

🟡 Changes recommended

Missing generated headers, untranslated CUDA dependencies, and multiple 32-to-64-lane indexing errors prevent a correct HIP build.

Once you've addressed the issues Copilot identified, you can request another Copilot review.

Pull request overview

Restores HIP backend parity with CUDA by regenerating missing kernels and expanding HIP build coverage.

Changes:

• Reworks hipify.sh to translate current CUDA kernels.
• Adds mGGA, FXC, gradient, screening, and shell-to-task support.
• Updates HIP build wiring and wavefront-dependent kernels.

File summaries

File	Description
hip/kernels/zmat_vxc.hip	Adds UKS/GKS and mGGA VXC kernels.
hip/kernels/zmat_fxc.hip	Adds FXC contraction kernels.
hip/kernels/uvvars.hip	Dispatches split LDA/GGA/mGGA kernels.
hip/kernels/uvvars_mgga.hpp	Adds mGGA variable kernels.
hip/kernels/uvvars_lda.hpp	Adds LDA scheme variants.
hip/kernels/uvvars_gga.hpp	Adds GGA scheme variants.
hip/kernels/symmetrize_mat.hip	Updates matrix symmetrization.
hip/kernels/pack_submat.hip	Adds symmetric/asymmetric packing.
hip/kernels/hipblas_extensions.hip	Expands HIP BLAS helpers.
hip/kernels/hip_ssf_1d.hpp	Declares SSF gradient support.
hip/kernels/hip_ssf_1d.hip	Adds SSF weight derivatives.
hip/kernels/hip_inc_potential.hip	Adds symmetric/asymmetric increments.
hip/kernels/hip_extensions.hpp	Adds wavefront reductions.
hip/kernels/grid_to_center.hpp	Updates HIP declarations.
hip/kernels/grid_to_center.hip	Reworks distance computation.
hip/kernels/exx_ek_screening_bfn_stats.hip	Adds sn-LinK screening.
hip/kernels/collocation/collocation_device_constants.hpp	Expands angular constants.
hip/kernels/collocation/collocation_angular_spherical_unnorm.hpp	Adds angular momentum-4 support.
hip/kernels/collocation/collocation_angular_cartesian.hpp	Adds Cartesian momentum-4 support.
hip/kernels/collocation_shell_to_task_kernels.hpp	Aggregates shell-to-task kernels.
hip/kernels/collocation_masked_combined_kernels.hpp	Updates masked collocation.
hip/kernels/collocation_device.hip	Adds shell-to-task dispatch.
hip/hipify.sh	Rewrites HIP generation workflow.
hip/CMakeLists.txt	Registers new HIP sources.

Review details

Suppressed comments (3)

src/xc_integrator/local_work_driver/device/hip/kernels/pack_submat.hip:226

• As in the symmetric path, this 32×32 launch is incompatible with WARP_X=16 and the 8×8 cut decomposition: tid_xy=1 creates overlapping/out-of-range unrolled reads, while only four of eight tid_yy residues run. Launch the logical 16×64 shape.

  dim3 threads( hip::warp_size/2, hip::max_warps_per_thread_block * 2, 1 );

src/xc_integrator/local_work_driver/device/hip/kernels/hip_inc_potential.hip:212

• This asymmetric kernel has the same fixed 16×64 indexing contract, while the 32×32 HIP launch causes overlapping/out-of-bounds unrolled accesses and skips half of each cut block. Launch the dimensions represented by the kernel constants.

  dim3 threads( hip::warp_size/2, hip::max_warps_per_thread_block * 2, 1 );

src/xc_integrator/local_work_driver/device/hip/kernels/exx_ek_screening_bfn_stats.hip:405

• This duplicated conversion kernel has the same 64-vs-32 mismatch: the load loop leaves many collisions_buffer[threadIdx.y][threadIdx.x] rows uninitialized before they are consumed, yielding incorrect shell lists and nbe counts. Make its logical subgroup width and launch geometry consistent.

      for (int buffer_loop = 0; buffer_loop < warp_size; buffer_loop += warp_size/buffer_size) {
        const int t_id_x        = threadIdx.x % buffer_size;
        const int buffer_thread = threadIdx.x / buffer_size;
        const int buffer_idx    = buffer_thread + buffer_loop;
        if (j_block * buffer_size_bits + t_id_x * element_size < nshells && i_base + buffer_idx < ntasks) {

• Files reviewed: 25/25 changed files
• Comments generated: 10
• Review effort level: Balanced

---

💡 Add a code-review agent skill or configure MCP servers for context-aware, tailored reviews. Learn more in the docs.