OpenMP 卸载到 Nvidia 错误减少

Question

我有兴趣使用 OpenMP 将工作卸载到 GPU。

下面的代码在 CPU

上给出了 sum 的正确值

//g++ -O3 -Wall foo.cpp -fopenmp
#pragma omp parallel for reduction(+:sum)                                                                                                                                    
for(int i = 0 ; i < 2000000000; i++) sum += i%11;

它也可以像这样在带有 OpenACC 的 GPU 上工作

//g++ -O3 -Wall foo.cpp -fopenacc   
#pragma acc parallel loop reduction(+:sum)                                                                                                                                    
for(int i = 0 ; i < 2000000000; i++) sum += i%11;

nvprof 表明它在 GPU 上运行，并且在 CPU.

上也比 OpenMP 快

然而，当我尝试像这样使用 OpenMP 卸载到 GPU 时

//g++ -O3 -Wall foo.cpp -fopenmp -fno-stack-protector
#pragma omp target teams distribute parallel for reduction(+:sum)
for(int i = 0 ; i < 2000000000; i++) sum += i%11;

sum 得到错误的结果（它只是 returns 零）。 nvprof 似乎表明它在 GPU 上运行，但它比 CPU 上的 OpenMP 慢得多。

为什么在 GPU 上使用 OpenMP 时缩减失败？

这是我用来测试这个的完整代码

#include <stdio.h>
//g++ -O3 -Wall acc2.cpp -fopenmp -fno-stack-protector                                                                                                                           
//sudo nvprof ./a.out                                                                                                                                                            
int main (void) {
  int sum = 0;
  //#pragma omp parallel for reduction(+:sum)                                                                                                                                    
  //#pragma acc parallel loop reduction(+:sum)                                                                                                                                   
  #pragma omp target teams distribute parallel for reduction(+:sum)
  for(int i = 0 ; i < 2000000000; i++) {
    sum += i%11;
  }
  printf("sum = %d\n",sum);
  return 0;
}

使用 GCC 7.2.0、Ubuntu 17.10 以及 gcc-offload-nvptx

Answer 1

解决方案是像这样添加子句 map(tofrom:sum)：

//g++ -O3 -Wall foo.cpp -fopenmp -fno-stack-protector
#pragma omp target teams distribute parallel for reduction(+:sum) map(tofrom:sum)
for(int i = 0 ; i < 2000000000; i++) sum += i%11;

这为 sum 获得了正确的结果，但是代码仍然比没有 target 的 OpenACC 或 OpenMP 慢得多。

更新： 速度的解决方案是添加 simd 子句。有关详细信息，请参阅此答案的末尾。

上面的解决方案在一行中有很多子句。可以这样分解：

#pragma omp target data map(tofrom: sum)
#pragma omp target teams distribute parallel for reduction(+:sum)
for(int i = 0 ; i < 2000000000; i++) sum += i%11;

另一种选择是使用 defaultmap(tofrom:scalar)

#pragma omp target teams distribute parallel for reduction(+:sum) defaultmap(tofrom:scalar)

显然，OpenMP 4.5 中的标量变量默认为 firstprivate。 https://developers.redhat.com/blog/2016/03/22/what-is-new-in-openmp-4-5-3/

如果您有多个要共享的标量值，

defaultmap(tofrom:scalar) 会很方便。

我也手动实现了减少，看看是否可以加快速度。我没有设法加快它的速度，但无论如何这是代码（我尝试了其他优化，但 none 其中有帮助）。

#include <omp.h>
#include <stdio.h>

//g++ -O3 -Wall acc2.cpp -fopenmp -fno-stack-protector
//sudo nvprof ./a.out

static inline int foo(int a, int b, int c) {
  return a > b ? (a/c)*b + (a%c)*b/c : (b/c)*a + (b%c)*a/c;
}

int main (void) {
  int nteams = 0, nthreads = 0;

  #pragma omp target teams map(tofrom: nteams) map(tofrom:nthreads)
  {
    nteams = omp_get_num_teams();
    #pragma omp parallel
    #pragma omp single
    nthreads = omp_get_num_threads();
  }
  int N = 2000000000;
  int sum = 0;

  #pragma omp declare target(foo)  

  #pragma omp target teams map(tofrom: sum)
  {
    int nteams = omp_get_num_teams();
    int iteam = omp_get_team_num();
    int start  = foo(iteam+0, N, nteams);
    int finish = foo(iteam+1, N, nteams);    
    int n2 = finish - start;
    #pragma omp parallel
    {
      int sum_team = 0;
      int ithread = omp_get_thread_num();
      int nthreads = omp_get_num_threads();
      int start2  = foo(ithread+0, n2, nthreads) + start;
      int finish2 = foo(ithread+1, n2, nthreads) + start;
      for(int i=start2; i<finish2; i++) sum_team += i%11;
      #pragma omp atomic
      sum += sum_team;
    }   
  }   

  printf("devices %d\n", omp_get_num_devices());
  printf("default device %d\n", omp_get_default_device());
  printf("device id %d\n", omp_get_initial_device());
  printf("nteams %d\n", nteams);
  printf("nthreads per team %d\n", nthreads);
  printf("total threads %d\n", nteams*nthreads);
  printf("sum %d\n", sum);
  return 0;
}

nvprof 表明大部分时间都花在 cuCtxSynchronize 上。使用 OpenACC，它大约是一半。

我终于设法大大加快了减少速度。解决方案是添加 simd 子句

#pragma omp target teams distribute parallel for simd reduction(+:sum) map(tofrom:sum).

一行九个子句。稍微短一点的解决方案是

#pragma omp target map(tofrom:sum)
#pragma omp teams distribute parallel for simd reduction(+:sum)

时间是

OMP_GPU    0.25 s
ACC        0.47 s
OMP_CPU    0.64 s

GPU 上的 OpenMP 现在比 CPU 上的 OpenACC 和 OpenMP 快得多。我不知道是否可以通过一些附加条款加快 OpenACC 的速度。

希望 Ubuntu 18.04 修复 gcc-offload-nvptx，这样它就不需要 -fno-stack-protector。

OpenMP 卸载到 Nvidia 错误减少

OpenMP offloading to Nvidia wrong reduction

c++

gcc

openmp

offloading

openacc