| 场景 | 体系 | 原子数 | 系综 | HIP 2026.3 (ns/day) | CUDA 2023.2 (ns/day) | HIP / 老版本 |
|---|---|---|---|---|---|---|
| villin | 小蛋白 | ~40K | NVT | 323.49 | 101.18 | 3.20× |
| rnase_cubic | 可溶蛋白 | 24,040 | NVT | 147.94 | 66.93 | 2.21× |
| waterbox8 | 纯水 | 50,859 | NPT | 67.64 | 41.66 | 1.62× |
| adh_dodec | 膜蛋白 | 95,561 | NVT | 26.09 | 18.89 | 1.38× |
| aqp_ensemble | 水通道蛋白 | ~80K | NVT | 13.58 | 10.62 | 1.28× |
| ion_channel | 离子通道 | ~50K | NVT | 20.17 | 15.87 | 1.27× |
| stmv | 大病毒衣壳 | ~百万 | NVT | 2.30 | 1.83 | 1.26× |
测量方法：gmx_bench.slurm（每个 backend×system 各 3 次重复，
-notunepme -resetstep N/2 保证计时口径一致，报告 mean±sd；所有重复 sd < 2%）。
规律：体系越小、加速比越大——小体系 GPU 吃不饱，原生 HIP 内核更低的启动/调度开销优势更明显。
上一版报告里 waterbox8 NPT 是 15.31 vs 41.22（慢 2.7×），这是唯一一个 「HIP 变慢」的证据，也是整个 NPT 优化课题的起点。干净节点重测后：
| 污染节点 | 干净节点 | |
|---|---|---|
| HIP 2026.3 | 15.31（假象） | 67.64 |
| CUDA 2023.2 | 41.22 | 41.66 |
根因：HIP 后端对节点 CPU 争抢极敏感。同一争抢节点上，HIP 从 148 掉到 9.4 （15× 恶化），而 CUDA 翻译版几乎不动（稳定在 66.8）。上一版的 waterbox8 HIP 数字恰好落在 争抢节点上。由此，「waterbox8 NPT 慢」这一整条优化课题（virial 回退、DPP 归约、寄存器/占用率） 都是在追这个测量假象，详见第四节死胡同记录。
对每个场景对比 run.log 的 Energies 段，主能量项（LJ SR / Coulomb SR / Potential / 总能量） 全部一致到 ~0.01–0.05%，温度均落在目标值（~300 K），压强在系综波动范围内， 无 NaN/Inf。这个量级差异是「不同舍入顺序 → 数千步轨迹微发散」的正常水平，不是物理错误。
waterbox8 NPT 100 ps（50000 步）长跑：总能量 -5.718×10⁵ kJ/mol、 温度 299.88 K、压强 -1.8 bar，全程无漂移、无异常终止，与短跑一致的 67.5 ns/day。
gfx90a→gfx906 探针坑：GROMACS 的 gmxManageHipccConfig.cmake 用
--offload-arch=gfx90a 去探测 hipcc 是否支持各优化 flag，但 DTK 26.04 没有 gfx90a
设备库——每个探测都失败，于是所有优化 flag 被静默丢弃，性能掉到 ~17 ns/day 且
不报任何错误。修复：把探针架构 gfx90a 改成 gfx906（并补
-mcpu=gfx906、去掉 gfx906 不认的 -fno-gpu-rdc / -fno-slp-vectorize）。
这是影响最大的一处编译改动。
| 假线索 | 实验 | 结论 |
|---|---|---|
deviceStream.synchronize() 是同步开销大头 | 移除同步点重建对照 | 零变化，不是根因 |
virial 步 CPU 力归约回退（useGpuFBufferOps 在 computeVirial 步退回 CPU） | 源码坐实该回退存在 | 机制真实，但干净节点下不是净瓶颈（被 GPU 整体更快盖过） |
DPP 归约的动态 lane 转置 → ds_bpermute → 12.8× LDS | 改成原子归约对照 | 零变化，归约不是瓶颈 |
minBlocksPerMp=8→14 压 vgpr 提占用率 | 实测 vgpr 不降反升（32→36） | __launch_bounds__ 在这版 hipcc 上不按预期生效，收益 ~1%，已回退 |
squeue 确认节点无其他作业或 --exclusive 独占；
② 每点 3 次重复，sd > 2% 就判污染、重跑（gmx_bench.slurm 已内置）；
③ 跨节点对比时务必在同一批干净节点上对跑两个后端。
工具链：DTK 26.04（HIP 6.3.26113 / clang 17）+ GCC 12.2.0 + CMake 3.28.6 + OpenMPI 4.1.5。
# 0. 恢复 cmake_minimum_required 到 3.28（撤销任何 3.25 降级补丁）
cd $SRC
sed -i 's/cmake_minimum_required(VERSION 3\.25)/cmake_minimum_required(VERSION 3.28)/g' CMakeLists.txt
find . -name CMakeLists.txt -exec sed -i 's/cmake_minimum_required *(VERSION 3\.25)/cmake_minimum_required(VERSION 3.28)/g' {} \;
# 1. gfx906 架构探针补丁（关键，见 4.1）
cd $SRC/cmake
sed -i 's/--offload-arch=gfx90a/--offload-arch=gfx906/g' gmxManageHipccConfig.cmake
sed -i 's/gmx_hip_check_single_flag("-fno-gpu-rdc")/# REMOVED for gfx906: gmx_hip_check_single_flag("-fno-gpu-rdc")/' gmxManageHipccConfig.cmake
sed -i 's/gmx_hip_check_single_flag("-fno-slp-vectorize")/# REMOVED for gfx906: gmx_hip_check_single_flag("-fno-slp-vectorize")/' gmxManageHipccConfig.cmake
sed -i '/gmx_hip_check_single_flag("-ffast-math")/a\ gmx_hip_check_single_flag("-mcpu=gfx906")' gmxManageHipccConfig.cmake
# 2. 环境
module purge
module load mpi/openmpi/gcc-9.3.0/4.1.5
source <DTK_ROOT>/env.sh
export PATH=<GCC_ROOT>/bin:$PATH
export LD_LIBRARY_PATH=<GCC_ROOT>/lib64:$LD_LIBRARY_PATH
export amd_comgr_DIR=<DTK_ROOT>/lib64/cmake/amd_comgr
# 3. HIP clang wrapper（补 AMD GPU 优化标志）
cat > $BUILD/hip_wrapper.sh << 'EOF'
#!/bin/bash
exec <DTK_ROOT>/llvm/bin/clang++ \
--gcc-toolchain=<GCC_ROOT> \
-mllvm -amdgpu-early-inline-all=true \
-mllvm -amdgpu-function-calls=false \
"$@"
EOF
chmod +x $BUILD/hip_wrapper.sh
# 4. CMake 配置 + 编译
$CMAKE $SRC \
-DGMX_BUILD_OWN_FFTW=ON -DGMX_GPU=HIP -DGMX_MPI=ON \
-DCMAKE_C_COMPILER=$GCC_HOME/bin/gcc -DCMAKE_CXX_COMPILER=$GCC_HOME/bin/g++ \
-DCMAKE_HIP_COMPILER=$BUILD/hip_wrapper.sh \
-DCMAKE_PREFIX_PATH=<DTK_ROOT> \
-Damd_comgr_DIR=<DTK_ROOT>/lib64/cmake/amd_comgr \
-DGMX_SIMD=AVX2_256 -DGMX_GPU_FFT_LIBRARY=rocFFT \
-DGMX_GPU_NB_DISABLE_CLUSTER_PAIR_SPLIT=ON \
-DGMX_GPU_NB_NUM_CLUSTER_PER_CELL_X=1 -DGMX_GPU_NB_NUM_CLUSTER_PER_CELL_Y=1 -DGMX_GPU_NB_NUM_CLUSTER_PER_CELL_Z=1 \
-DGMX_HIP_TARGET_ARCH=gfx906 -DAMDGPU_TARGETS=gfx906 -DGPU_TARGETS=gfx906 \
-DCMAKE_BUILD_TYPE=Release
make -j8完整脚本：<USER_HOME>/scripts/gromacs_cmake28_build.slurm（可直接 sbatch 复现）。
当前最优产物在 <USER_HOME>/softwares/gromacs/build_cmake28/：
| 文件 | 大小 | 说明 |
|---|---|---|
bin/gmx_mpi | 113 KB | 主可执行文件（薄封装，动态链接库） |
lib/libgromacs_mpi.so.11.0.0 | 45 MB | 真正的计算库（含所有 HIP kernel） |
运行方式（需与编译时一致的 DTK 运行时）：
module purge
module load mpi/openmpi/gcc-9.3.0/4.1.5
source <DTK_ROOT>/env.sh
export LD_LIBRARY_PATH=<GCC_ROOT>/lib64:$LD_LIBRARY_PATH
/path/to/build_cmake28/bin/gmx_mpi mdrun -s case.tpr -nb gpu -pme gpu ...注意：GROMACS 2026.3 写的 tpr 是 version 138，集群预编译的 2023.2 读不了；反过来 2026.3 能读旧 tpr。跨版本对跑时各自用各自版本生成的 tpr。