OpenACC GPU版点积性能劣于CPU串行版，求问题排查

问题描述 #

在C程序中实现点积计算，目标对比CPU串行版本与基于OpenACC的GPU并行版本性能，但GPU版本速度始终慢于CPU，测试了n从10¹到10⁷的所有场景均不符合预期，运行环境为集群。

串行版本代码 #

double dot_product_cpu(int n, double* x, double* y)
{
    double result = 0.0;
    for (int i = 0; i < n; i++)
    {
        result += x[i] * y[i];
    }
    return result;
}

OpenACC GPU并行版本代码 #

double dot_product_gpu(int n, double* x, double* y)
{
    double result = 0.0;

    #pragma acc data copyin(x[0:n], y[0:n]) copy(result)
    {
        #pragma acc parallel loop reduction(+:result)
        for (int i = 0; i < n; i++)
        {
            result += x[i] * y[i];
        }
    }

    return result;
}

核心原因与优化方案 #

1. 数据传输开销主导性能

GPU加速的最大瓶颈往往是CPU与GPU间的数据传输，原代码每次调用dot_product_gpu都会执行copyin（CPU→GPU）和copy（GPU→CPU）操作，小数据量下传输耗时远超过计算耗时；即使n=10⁷，单次传输的开销也可能抵消并行计算的收益。

优化：复用GPU内存
将数据拷贝与计算分离，仅在初始化和数据更新时传输，计算阶段直接使用GPU上的内存：

// 初始化GPU内存（仅调用一次）
void init_gpu_buffers(int n, double* x, double* y) {
    #pragma acc enter data create(x[0:n], y[0:n])
}

// 更新GPU内存中的数据（仅当x/y变化时调用）
void update_gpu_data(int n, double* x, double* y) {
    #pragma acc update device(x[0:n], y[0:n])
}

// 优化后的GPU点积计算，无需重复拷贝
double dot_product_gpu_optimized(int n, double* x, double* y) {
    double result = 0.0;
    #pragma acc parallel loop reduction(+:result) gang vector(1024) present(x[0:n], y[0:n])
    for (int i = 0; i < n; i++) {
        result += x[i] * y[i];
    }
    return result;
}

// 释放GPU内存（程序结束前调用）
void free_gpu_buffers(int n, double* x, double* y) {
    #pragma acc exit data delete(x[0:n], y[0:n])
}

2. 并行粒度未适配GPU架构

OpenACC默认的并行分区策略可能未充分利用GPU的硬件资源（如warp/wavefront大小），手动指定gang和vector层级可以提升并行效率。例如vector(1024)适配多数NVIDIA GPU的warp大小（32的倍数），AMD GPU可设为64。

3. 编译器优化未拉满

确保编译时开启最高级优化，并指定正确的GPU算力版本：

• PGI/NVHPC编译器：pgcc -O3 -acc -gpu=cc70 -o dot_product dot_product.c（cc70对应Volta架构，根据集群GPU型号调整）
• GCC编译器：gcc -O3 -fopenacc -o dot_product dot_product.c

4. CPU版本已被自动向量化

默认情况下，开启-O3优化后，编译器会自动对CPU串行代码进行向量化（SIMD指令），此时CPU性能已大幅提升。若要对比纯串行性能，可禁用CPU向量化：-fno-tree-vectorize（GCC）或-no-vec（PGI）。

5. 数据对齐优化

GPU对内存访问的对齐要求较高，声明数组时添加__attribute__((aligned(64)))可避免非对齐访问带来的性能损耗：

double x[n] __attribute__((aligned(64)));
double y[n] __attribute__((aligned(64)));

内容的提问来源于stack exchange，提问作者MARTÍ ARMENGOL AYALA