为何OpenMP Offloading向量加法耗时比主机端OpenMP更长？

我分别编写了基于主机端OpenMP和OpenMP Offloading的向量加法代码，使用Intel oneAPI工具链编译后运行，发现OpenMP Offloading版本的运行耗时反而比主机端OpenMP版本更长，请问这是什么原因？

主机端代码（openmp-host.c） #

#include <assert.h>
#include <math.h>
#include <omp.h>
#include <stdio.h>
#include <stdlib.h>

int main(int argc, char *argv[]) {
  unsigned N = (argc > 1 ? atoi(argv[1]) : 1000000);

  float *a = (float *)calloc(N, sizeof(float));
  float *b = (float *)calloc(N, sizeof(float));
  float *c = (float *)calloc(N, sizeof(float));
  for (int i = 0; i < N; i++)
    a[i] = i, b[i] = N - i;

#pragma omp parallel
  {
    unsigned thrds = omp_get_num_threads(), tid = omp_get_thread_num();
    unsigned size = N / thrds, rem = N - size * thrds;
    size += (tid < rem);
    unsigned s = (tid < rem ? size * tid : (tid * size + rem)), e = s + size;

    double t = omp_get_wtime();
    
      for (unsigned i = s; i < e; i++){
        c[i] = a[i] + b[i];
      }
        
    t = omp_get_wtime() - t;

    if (tid == 0)
      printf("N: %u # threads: %u time: %e\n", N, thrds, t);
  }

  for (unsigned i = 0; i < N; i++)
    assert(fabs(c[i] - N) < 1e-8);

  free(a);

  return 0;
}

设备端代码（openmp-device.c） #

#include <assert.h>
#include <math.h>
#include <omp.h>
#include <stdio.h>
#include <stdlib.h>

int main(int argc, char *argv[]) {
  int N = (argc > 1 ? atoi(argv[1]) : 1000000);
  double start, end;
  int *a = (int *)calloc(N, sizeof(int));
  int *b = (int *)calloc(N, sizeof(int));
  int *c = (int *)calloc(N, sizeof(int));
  double t;
  for (int i = 0; i < N; i++) {
    a[i] = i;
    b[i] = N - i;
  }

#pragma omp target enter data map(to:a[0:N],b[0:N], c[0:N])

t= omp_get_wtime();
#pragma omp target teams distribute parallel for simd
for(int i=0; i<N; i++){
  c[i] = a[i] + b[i];
}

t = omp_get_wtime() - t;

#pragma omp target exit data map(from: c[0:N])

  
  printf("time: %e \n", t);

  for (int i = 0; i < N; i++)
    assert(abs(c[i] - N) < 1e-8);

  free(a);
  free(b);
  free(c);

  return 0;
}

编译命令 #

icx -qopenmp -fopenmp-targets=spir64 openmp-device.c -o omp_device
icx -qopenmp openmp-host.c -o omp_host

核心原因分析 #

1. 数据传输开销占比过高 #

OpenMP Offloading需要将主机端的数组a、b、c拷贝到加速设备（如GPU），计算完成后还要把结果数组c拷贝回主机。向量加法是计算强度极低的任务（每个元素仅一次加法操作），数据在主机和设备间传输的耗时远超过设备上的计算耗时，整体开销反而比主机端多线程执行更大。

你当前设置的N=1e6对应的数据量仅12MB（三个int数组：31e64字节），即使数据量不大，但对于低计算强度任务，传输延迟的影响依然非常显著。

2. 任务粒度与设备特性不匹配 #

GPU等加速设备擅长处理大规模、高并行度的计算任务，而向量加法的单操作计算量太小：

• 主机端OpenMP利用CPU多线程，且CPU缓存系统对连续内存访问的优化非常成熟，能高效处理这类简单任务；
• 设备端虽使用了teams distribute parallel for simd，但启动线程/团队的开销占比远高于计算本身，设备核心无法被充分利用。

3. 设备初始化与启动开销 #

第一次执行OpenMP Offloading任务时，加速设备需要完成初始化、内核加载等操作，这部分额外开销会被计入首次运行时间。而主机端OpenMP的线程启动开销相对极小，不会对结果产生明显影响。

4. 计时范围的隐性差异 #

两段代码的计时逻辑存在细节差异：

• 主机端仅计时并行计算部分；
• 设备端虽然只计时了计算阶段，但完整的Offloading流程包含前后的数据传输，若把传输时间计入，整体耗时会更高。

优化建议 #

• 增大计算规模：将N设置为1e8或更大，让计算耗时超过数据传输耗时，此时Offloading的并行优势才会体现；
• 优化数据传输：合并map子句减少拷贝次数，或使用use_device_ptr等特性避免不必要的数据拷贝；
• 匹配设备类型：若使用集成GPU，其计算能力可能不如多核CPU，建议切换到独立GPU测试；
• 调整并行粒度：通过num_teams、thread_limit等参数优化线程块大小，提升设备核心利用率；
• 多次运行取平均：排除首次设备初始化的额外开销，取多次运行的平均时间进行对比。

内容的提问来源于stack exchange，提问作者Ravindu Hirimuthugoda