Merge branch 'develop' into dsp-link

Author: Cstandardlib

.github/workflows/cuda.yml

@@ -3,6 +3,9 @@ name: CUDA Test
 on:
   workflow_dispatch:
   pull_request:
+  push:
+    branches:
+      - develop
 
 defaults:
   run:
@@ -19,10 +22,12 @@ jobs:
     if: github.repository_owner == 'deepmodeling'
     container:
       image: ghcr.io/deepmodeling/abacus-cuda
-      volumes:
-        - /tmp/ccache:/github/home/.ccache
       options: --gpus all
-      
+    env:
+      CCACHE_DIR: ${{ github.workspace }}/.ccache
+      CCACHE_BASEDIR: ${{ github.workspace }}
+      CCACHE_COMPRESS: "true"
+      CCACHE_MAXSIZE: 8G
     steps:
       - name: Checkout
         uses: actions/checkout@v7.0.1
@@ -34,6 +39,26 @@ jobs:
           sudo apt-get update
           sudo apt-get install -y ccache xz-utils ninja-build pkg-config
 
+      - name: Calculate compiler cache fingerprint
+        id: compiler
+        run: |
+          fingerprint="$({ c++ --version; nvcc --version; ccache --version; } | sha256sum | cut -d' ' -f1)"
+          echo "fingerprint=${fingerprint}" >> "${GITHUB_OUTPUT}"
+
+      - name: Restore compiler cache
+        uses: actions/cache@v4
+        with:
+          path: ${{ env.CCACHE_DIR }}
+          key: abacus-cuda-ccache-v1-${{ runner.os }}-${{ runner.arch }}-${{ steps.compiler.outputs.fingerprint }}-${{ github.sha }}
+          restore-keys: |
+            abacus-cuda-ccache-v1-${{ runner.os }}-${{ runner.arch }}-${{ steps.compiler.outputs.fingerprint }}-
+
+      - name: Prepare compiler cache
+        run: |
+          mkdir -p "${CCACHE_DIR}"
+          ccache --set-config="max_size=${CCACHE_MAXSIZE}"
+          ccache --zero-stats
+
       - name: Install external tools from toolchain
         run: |
           cd toolchain
@@ -51,6 +76,12 @@ jobs:
           cmake --build build -j "$(nproc)"
           cmake --install build
 
+      - name: Report compiler cache
+        if: always()
+        run: |
+          ccache --show-stats
+          ccache --show-stats >> "${GITHUB_STEP_SUMMARY}"
+
       - name: Module_LCAO CUDA Unittests
         env:
           GTEST_COLOR: 'yes'

.github/workflows/interface.yml

@@ -24,7 +24,7 @@ jobs:
             prefix: Bi2Se3_pw
             basis: pw
           - name: advance (LCAO)
-            script: example_advance.py
+            script: example_advanced.py
             prefix: Bi2Se3_advanced
             basis: lcao
 
@@ -58,12 +58,14 @@ jobs:
         env:
           SCRIPT: ${{ matrix.script }}
           PREFIX: ${{ matrix.prefix }}
+          BASIS: ${{ matrix.basis }}
         run: |
           python3 << 'PYEOF'
           import os, re, textwrap
 
           script = os.environ["SCRIPT"]
           prefix = os.environ["PREFIX"]
+          basis = os.environ["BASIS"]
 
           # ── 1. Patch the example script ────────────────────────
           with open(script) as f:
@@ -146,10 +148,12 @@ jobs:
               f"     {nk}",     "!     num_kpts",
               "     4     4     4", "!     mp_grid",
           ]
+          lines.append("begin kpoints")
           for ix in range(4):
               for iy in range(4):
                   for iz in range(4):
                       lines.append(f"  {ix/4:.15f}  {iy/4:.15f}  {iz/4:.15f}")
+          lines.append("end kpoints")
 
           lines.append(f"     {nntot}")
           lines.append("!     nntot")

AGENTS.md

@@ -136,6 +136,17 @@ The repository text files have been normalized to LF once. Day-to-day line
 ending enforcement should rely on staged/changed-file hooks and CI; rerun the
 full mixed-line-ending hook only for intentional repository-wide normalization.
 
+## Upstream Repository
+
+- Repository: https://github.com/deepmodeling/abacus-develop
+- Issues: https://github.com/deepmodeling/abacus-develop/issues
+- Pull requests: https://github.com/deepmodeling/abacus-develop/pulls
+- Upstream PRs are opened from personal fork branches
+  (`<fork-owner>:<branch>` into `develop`).
+- `workflow_dispatch`-only workflows (e.g. `.github/workflows/interface.yml`)
+  are not triggered by push/PR events; PR CI cannot verify such fixes, so
+  state "manual dispatch run required" in the PR verification notes.
+
 ## PR Self-Check
 
 - Confirm the PR body states exact commands run, whether they passed or failed,

source/CMakeLists.txt

@@ -705,7 +705,6 @@ target_link_libraries(
     psi_overall_init
     psi_init
     psi
-    dftu
     deltaspin
     container
     device
@@ -721,6 +720,7 @@ if(ENABLE_LCAO)
     PRIVATE
       hamilt_lcao
       tddft
+      dftu
       orb
       gint
       hcontainer

source/Makefile.Objects

@@ -840,8 +840,8 @@ OBJS_SRCPW=h_ewald_pw.o\
     update_cell_pw.o\
     dftu_base.o\
     dftu_base_io.o\
-    dftu_base_occ.o\
-    dftu_base_tools.o\
+    dftu_pw.o\
+    dftu_pw_tools.o\
     yukawa_screening.o\
     setup_dftu_pw.o\
     deltaspin_pw.o\

source/source_base/module_device/memory_op.cpp

@@ -1,6 +1,7 @@
 #include "memory_op.h"
 
 #include "source_base/memory_recorder.h"
+#include "source_base/tool_quit.h"
 #include "source_base/tool_threading.h"
 #ifdef __DSP
 #include "source_base/kernels/dsp/dsp_connector.h"
@@ -549,35 +550,52 @@ void set_memory(FPTYPE* arr, const int var, const size_t size, base_device::Abac
 
 template <typename FPTYPE>
 void synchronize_memory(FPTYPE* arr_out, const FPTYPE* arr_in, const size_t size, base_device::AbacusDevice_t device_type_out, base_device::AbacusDevice_t device_type_in){
-    if (device_type_out == base_device::AbacusDevice_t::CpuDevice || device_type_in == base_device::AbacusDevice_t::CpuDevice){
+    // The four source/destination combinations are mutually exclusive, so each
+    // branch must test BOTH devices. Using `||` here made the first branch match
+    // whenever either side was the CPU, which routed host<->device transfers to
+    // the host-to-host specialization.
+    if (device_type_out == base_device::AbacusDevice_t::CpuDevice && device_type_in == base_device::AbacusDevice_t::CpuDevice){
         synchronize_memory_op<FPTYPE, DEVICE_CPU, DEVICE_CPU>()(arr_out, arr_in, size);
     }
-    else if (device_type_out == base_device::AbacusDevice_t::CpuDevice || device_type_in == base_device::AbacusDevice_t::GpuDevice){
+#if __CUDA || __UT_USE_CUDA || __ROCM || __UT_USE_ROCM
+    else if (device_type_out == base_device::AbacusDevice_t::CpuDevice && device_type_in == base_device::AbacusDevice_t::GpuDevice){
         synchronize_memory_op<FPTYPE, DEVICE_CPU, DEVICE_GPU>()(arr_out, arr_in, size);
     }
-    else if (device_type_out == base_device::AbacusDevice_t::GpuDevice || device_type_in == base_device::AbacusDevice_t::CpuDevice){
+    else if (device_type_out == base_device::AbacusDevice_t::GpuDevice && device_type_in == base_device::AbacusDevice_t::CpuDevice){
         synchronize_memory_op<FPTYPE, DEVICE_GPU, DEVICE_CPU>()(arr_out, arr_in, size);
     }
-    else if (device_type_out == base_device::AbacusDevice_t::GpuDevice || device_type_in == base_device::AbacusDevice_t::GpuDevice){
+    else if (device_type_out == base_device::AbacusDevice_t::GpuDevice && device_type_in == base_device::AbacusDevice_t::GpuDevice){
         synchronize_memory_op<FPTYPE, DEVICE_GPU, DEVICE_GPU>()(arr_out, arr_in, size);
     }
+#endif
+    else {
+        ModuleBase::WARNING_QUIT("base_device::memory::synchronize_memory",
+                                 "unsupported source/destination device combination");
+    }
 }
 
 template <typename FPTYPE_out, typename FPTYPE_in>
 void cast_memory(FPTYPE_out* arr_out, const FPTYPE_in* arr_in, const size_t size, base_device::AbacusDevice_t device_type_out, base_device::AbacusDevice_t device_type_in)
 {
-    if (device_type_out == base_device::AbacusDevice_t::CpuDevice || device_type_in == base_device::AbacusDevice_t::CpuDevice){
+    // See synchronize_memory() above: dispatch on the exact (out, in) device pair.
+    if (device_type_out == base_device::AbacusDevice_t::CpuDevice && device_type_in == base_device::AbacusDevice_t::CpuDevice){
         cast_memory_op<FPTYPE_out, FPTYPE_in, DEVICE_CPU, DEVICE_CPU>()(arr_out, arr_in, size);
     }
-    else if (device_type_out == base_device::AbacusDevice_t::CpuDevice || device_type_in == base_device::AbacusDevice_t::GpuDevice){
+#if __CUDA || __UT_USE_CUDA || __ROCM || __UT_USE_ROCM
+    else if (device_type_out == base_device::AbacusDevice_t::CpuDevice && device_type_in == base_device::AbacusDevice_t::GpuDevice){
         cast_memory_op<FPTYPE_out, FPTYPE_in, DEVICE_CPU, DEVICE_GPU>()(arr_out, arr_in, size);
     }
-    else if (device_type_out == base_device::AbacusDevice_t::GpuDevice || device_type_in == base_device::AbacusDevice_t::CpuDevice){
+    else if (device_type_out == base_device::AbacusDevice_t::GpuDevice && device_type_in == base_device::AbacusDevice_t::CpuDevice){
         cast_memory_op<FPTYPE_out, FPTYPE_in, DEVICE_GPU, DEVICE_CPU>()(arr_out, arr_in, size);
     }
-    else if (device_type_out == base_device::AbacusDevice_t::GpuDevice || device_type_in == base_device::AbacusDevice_t::GpuDevice){
+    else if (device_type_out == base_device::AbacusDevice_t::GpuDevice && device_type_in == base_device::AbacusDevice_t::GpuDevice){
         cast_memory_op<FPTYPE_out, FPTYPE_in, DEVICE_GPU, DEVICE_GPU>()(arr_out, arr_in, size);
     }
+#endif
+    else {
+        ModuleBase::WARNING_QUIT("base_device::memory::cast_memory",
+                                 "unsupported source/destination device combination");
+    }
 }
 
 template <typename FPTYPE>
@@ -591,5 +609,24 @@ void delete_memory(FPTYPE* arr, base_device::AbacusDevice_t device_type)
     }
 }
 
+// Explicit instantiations of the runtime-dispatch wrappers, so that the
+// declarations in memory_op.h can actually be linked from another translation
+// unit (and covered by unit tests). cast_memory is instantiated only for the
+// type pairs that cast_memory_op provides for all four device combinations.
+template void synchronize_memory<int>(int*, const int*, const size_t, base_device::AbacusDevice_t, base_device::AbacusDevice_t);
+template void synchronize_memory<float>(float*, const float*, const size_t, base_device::AbacusDevice_t, base_device::AbacusDevice_t);
+template void synchronize_memory<double>(double*, const double*, const size_t, base_device::AbacusDevice_t, base_device::AbacusDevice_t);
+template void synchronize_memory<std::complex<float>>(std::complex<float>*, const std::complex<float>*, const size_t, base_device::AbacusDevice_t, base_device::AbacusDevice_t);
+template void synchronize_memory<std::complex<double>>(std::complex<double>*, const std::complex<double>*, const size_t, base_device::AbacusDevice_t, base_device::AbacusDevice_t);
+
+template void cast_memory<float, float>(float*, const float*, const size_t, base_device::AbacusDevice_t, base_device::AbacusDevice_t);
+template void cast_memory<double, double>(double*, const double*, const size_t, base_device::AbacusDevice_t, base_device::AbacusDevice_t);
+template void cast_memory<float, double>(float*, const double*, const size_t, base_device::AbacusDevice_t, base_device::AbacusDevice_t);
+template void cast_memory<double, float>(double*, const float*, const size_t, base_device::AbacusDevice_t, base_device::AbacusDevice_t);
+template void cast_memory<std::complex<float>, std::complex<float>>(std::complex<float>*, const std::complex<float>*, const size_t, base_device::AbacusDevice_t, base_device::AbacusDevice_t);
+template void cast_memory<std::complex<double>, std::complex<double>>(std::complex<double>*, const std::complex<double>*, const size_t, base_device::AbacusDevice_t, base_device::AbacusDevice_t);
+template void cast_memory<std::complex<float>, std::complex<double>>(std::complex<float>*, const std::complex<double>*, const size_t, base_device::AbacusDevice_t, base_device::AbacusDevice_t);
+template void cast_memory<std::complex<double>, std::complex<float>>(std::complex<double>*, const std::complex<float>*, const size_t, base_device::AbacusDevice_t, base_device::AbacusDevice_t);
+
 } // namespace memory
 } // namespace base_device

source/source_base/module_device/test/memory_test.cpp

@@ -170,6 +170,59 @@ TEST_F(TestModulePsiMemory, delete_memory_op_complex_double_cpu)
     delete_memory_complex_double_cpu_op()(hz_xx);
 }
 
+// ---------------------------------------------------------------------------
+// Runtime-dispatch wrappers (issue #7553).
+//
+// synchronize_memory()/cast_memory() pick a compile-time specialization from a
+// pair of runtime AbacusDevice_t values. The branches used to be joined with
+// `||`, so the first one matched whenever EITHER side was the CPU and every
+// host<->device transfer was served by the host-to-host specialization. The
+// checks below pin the exact-pair dispatch for all combinations available in
+// the current build.
+// ---------------------------------------------------------------------------
+
+TEST_F(TestModulePsiMemory, synchronize_memory_dispatch_cpu_to_cpu)
+{
+    std::vector<double> h_xx(xx.size(), 0);
+    base_device::memory::synchronize_memory(h_xx.data(),
+                                            xx.data(),
+                                            xx.size(),
+                                            base_device::AbacusDevice_t::CpuDevice,
+                                            base_device::AbacusDevice_t::CpuDevice);
+    for (int ii = 0; ii < xx.size(); ii++)
+    {
+        EXPECT_EQ(h_xx[ii], xx[ii]);
+    }
+}
+
+TEST_F(TestModulePsiMemory, synchronize_memory_dispatch_complex_cpu_to_cpu)
+{
+    std::vector<std::complex<double>> hz_xx(z_xx.size(), std::complex<double>(0, 0));
+    base_device::memory::synchronize_memory(hz_xx.data(),
+                                            z_xx.data(),
+                                            z_xx.size(),
+                                            base_device::AbacusDevice_t::CpuDevice,
+                                            base_device::AbacusDevice_t::CpuDevice);
+    for (int ii = 0; ii < z_xx.size(); ii++)
+    {
+        EXPECT_EQ(hz_xx[ii], z_xx[ii]);
+    }
+}
+
+TEST_F(TestModulePsiMemory, cast_memory_dispatch_cpu_to_cpu)
+{
+    std::vector<float> h_xx(xx.size(), 0);
+    base_device::memory::cast_memory(h_xx.data(),
+                                     xx.data(),
+                                     xx.size(),
+                                     base_device::AbacusDevice_t::CpuDevice,
+                                     base_device::AbacusDevice_t::CpuDevice);
+    for (int ii = 0; ii < xx.size(); ii++)
+    {
+        EXPECT_FLOAT_EQ(h_xx[ii], static_cast<float>(xx[ii]));
+    }
+}
+
 #if __UT_USE_CUDA || __UT_USE_ROCM
 TEST_F(TestModulePsiMemory, set_memory_op_double_gpu)
 {
@@ -347,4 +400,90 @@ TEST_F(TestModulePsiMemory, delete_memory_op_complex_double_gpu)
     delete_memory_complex_double_gpu_op()(thrust::raw_pointer_cast(dz_xx));
 }
 
+
+// Exact-pair dispatch across the host/device boundary (issue #7553). Before the
+// fix these three cases all reached synchronize_memory_op<..., CPU, CPU>, i.e. a
+// plain host memcpy on a device pointer.
+TEST_F(TestModulePsiMemory, synchronize_memory_dispatch_cpu_to_gpu)
+{
+    thrust::device_ptr<double> d_xx = thrust::device_malloc<double>(xx.size());
+    std::vector<double> hv_xx(xx.size(), 0);
+    thrust::copy(hv_xx.begin(), hv_xx.end(), d_xx);
+    base_device::memory::synchronize_memory(thrust::raw_pointer_cast(d_xx),
+                                            xx.data(),
+                                            xx.size(),
+                                            base_device::AbacusDevice_t::GpuDevice,
+                                            base_device::AbacusDevice_t::CpuDevice);
+
+    thrust::host_vector<double> h_xx(xx.size());
+    thrust::copy(d_xx, d_xx + xx.size(), h_xx.begin());
+    for (int ii = 0; ii < xx.size(); ii++)
+    {
+        EXPECT_EQ(h_xx[ii], xx[ii]);
+    }
+    thrust::device_free(d_xx);
+}
+
+TEST_F(TestModulePsiMemory, synchronize_memory_dispatch_gpu_to_cpu)
+{
+    thrust::device_ptr<double> d_xx = thrust::device_malloc<double>(xx.size());
+    thrust::copy(xx.begin(), xx.end(), d_xx);
+    thrust::host_vector<double> h_xx(xx.size());
+    base_device::memory::synchronize_memory(thrust::raw_pointer_cast(h_xx.data()),
+                                            thrust::raw_pointer_cast(d_xx),
+                                            xx.size(),
+                                            base_device::AbacusDevice_t::CpuDevice,
+                                            base_device::AbacusDevice_t::GpuDevice);
+
+    for (int ii = 0; ii < xx.size(); ii++)
+    {
+        EXPECT_EQ(h_xx[ii], xx[ii]);
+    }
+    thrust::device_free(d_xx);
+}
+
+TEST_F(TestModulePsiMemory, synchronize_memory_dispatch_gpu_to_gpu)
+{
+    thrust::device_ptr<double> d1_xx = thrust::device_malloc<double>(xx.size());
+    thrust::device_ptr<double> d2_xx = thrust::device_malloc<double>(xx.size());
+    thrust::copy(xx.begin(), xx.end(), d1_xx);
+    base_device::memory::synchronize_memory(thrust::raw_pointer_cast(d2_xx),
+                                            thrust::raw_pointer_cast(d1_xx),
+                                            xx.size(),
+                                            base_device::AbacusDevice_t::GpuDevice,
+                                            base_device::AbacusDevice_t::GpuDevice);
+
+    thrust::host_vector<double> h_xx(xx.size());
+    thrust::copy(d2_xx, d2_xx + xx.size(), h_xx.begin());
+    for (int ii = 0; ii < xx.size(); ii++)
+    {
+        EXPECT_EQ(h_xx[ii], xx[ii]);
+    }
+    thrust::device_free(d1_xx);
+    thrust::device_free(d2_xx);
+}
+
+TEST_F(TestModulePsiMemory, cast_memory_dispatch_cpu_to_gpu_and_back)
+{
+    thrust::device_ptr<float> d_xx = thrust::device_malloc<float>(xx.size());
+    base_device::memory::cast_memory(thrust::raw_pointer_cast(d_xx),
+                                     xx.data(),
+                                     xx.size(),
+                                     base_device::AbacusDevice_t::GpuDevice,
+                                     base_device::AbacusDevice_t::CpuDevice);
+
+    std::vector<double> h_xx(xx.size(), 0);
+    base_device::memory::cast_memory(h_xx.data(),
+                                     thrust::raw_pointer_cast(d_xx),
+                                     xx.size(),
+                                     base_device::AbacusDevice_t::CpuDevice,
+                                     base_device::AbacusDevice_t::GpuDevice);
+
+    for (int ii = 0; ii < xx.size(); ii++)
+    {
+        EXPECT_FLOAT_EQ(static_cast<float>(h_xx[ii]), static_cast<float>(xx[ii]));
+    }
+    thrust::device_free(d_xx);
+}
+
 #endif // __UT_USE_CUDA || __UT_USE_ROCM