feat: GPU-accelerated WT KEDF multi_kernel convolution

[SunsetStand]

Reminder

• Have you linked an issue with this pull request?
• Have you added adequate unit tests and/or case tests for your pull request?
• Have you noticed possible changes of behavior below or in the linked issue?
• Have you explained the changes of codes in core modules of ESolver, HSolver, ElecState, Hamilt, Operator or Psi? (ignore if not applicable)

Linked Issue

This is a new feature — no existing issue. A CPU-vs-GPU correctness and performance benchmark is provided in the PR description below. An issue can be opened for discussion if preferred.

Unit Tests and/or Case Tests for my changes

A standalone benchmark (ofdft_cuda/) was used to verify correctness (GPU vs FFTW3 CPU reference, error < 1e-7 for WT KEDF) and measure performance (14.2× speedup at 96³ grid on RTX 4060). Integration into ABACUS's existing GPU CI pipeline (e.g., tests/integrate/ GPU OFDFT cases) is planned as a follow-up once CI GPU runners are confirmed available for this module.

What's changed?

This PR adds GPU acceleration for the WT KEDF multi_kernel() function, which is the most expensive single operation in OFDFT Wang-Teter calculations (up to 40% of total SCF time). The implementation:

1. Uses ABACUS's existing GPU infrastructure — pw_rho->real2recip_gpu() / recip2real_gpu() for FFT and memory_op for device memory management. No new external dependencies.
2. Adds a single CUDA kernel (kedf_wt_recip_multiply) for element-wise G-space kernel multiplication, following the same pattern as existing GPU kernels in source_base/kernels/cuda/.
3. Persistent GPU buffers are lazily allocated on first call and reused across SCF iterations. The WT kernel array (kernel_) is copied to device once since it is constant throughout the SCF cycle.
4. Zero overhead when CUDA is disabled — all GPU code is guarded by #ifdef __CUDA and the CPU path is completely untouched.
   The GPU dispatch is a simple 5-line addition at the top of multi_kernel():

#ifdef __CUDA
    if (pw_rho->device == "gpu") {
        this->multi_kernel_gpu(prho, rkernel_rho, exponent, pw_rho);
        return;
    }
#endif

Performance: on an RTX 4060 Laptop GPU, the GPU path achieves 14.2× speedup for WT KEDF at typical OFDFT grid sizes (96³) compared to FFTW3 CPU, with correctness verified to < 1e-7 relative error. A full benchmark report is available in the standalone prototype (examples/ or as supplementary material upon request).

Any changes of core modules? (ignore if not applicable)

N/A — only modifies the OFDFT KEDF module (source_pw/module_ofdft/), which is not a core ESolver/Hamilt/Operator module.

[mohanchen]

Nice try, could you provide some tests/examples in the PR? you can check out /tests/07_OFDFT

[SunsetStand]

Thanks for the review! I've added a GPU WT KEDF test case:

Test location: tests/07_OFDFT/31_OF_KE_WT_GPU/

It mirrors 09_OF_KE_WT (Al FCC, WT KEDF, symmetry=on) with device gpu added to INPUT. The GPU path performs identical math—cuFFT replaces CPU FFT + GPU kernels for element-wise ops—so results should match the CPU reference within tolerance.

[SunsetStand]

Optimizations applied (commit 4129f9f)
Three targeted optimizations for the WT KEDF GPU kernels:
double2 replaces thrust::complex (AoS → native type)
thrust::complex uses AoS layout {real, imag} but real_to_complex only writes imag=0 and complex_to_real_norm only reads real part — wasting 50% memory bandwidth. Using native double2 (layout-compatible with cufftDoubleComplex) eliminates the abstraction overhead and enables better compiler vectorization. No API breakage, double2 is trivially castable to cufftDoubleComplex*.
Grid-stride loops
Replaced if (idx >= n) return with for (i = idx; i < n; i += stride) for flexible occupancy across arbitrary grid sizes.
GPU-side rho^exponent
Moved std::pow(rho[i], exponent) from CPU to a GPU kernel, eliminating one H→D transfer per SCF iteration. In the ABACUS WT KEDF call chain, prho needs one H→D copy regardless — the GPU pow kernel runs on the copied data directly, avoiding the CPU compute + second transfer.
Benchmark (RTX 4060 Laptop, standalone benchmark)

Grid	Original	Optimized	Speedup
32³	0.57 ms	0.32 ms	1.78×
64³	3.34 ms	1.13 ms	2.96×
96³	11.1 ms	3.42 ms	3.25×

Kernel-only (data already on GPU): 1.76× for 96³.
What was tested but NOT included
Thread coarsening (4 elements/thread) showed a ~4% regression on Ada Lovelace (SM 8.9) — fewer active warps reduced the GPU's ability to hide memory latency for these memory-bound kernels. Left for future architecture-specific tuning.
Correctness
Max relative error vs. original output: ~1e-16 for 32³/64³ grids. No existing test changes needed.