OpenMP向Nvidia的错误归约

Question

OpenMP向Nvidia的错误归约

8

我希望通过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，它返回的结果是错误的（仅返回零）。nvprof显示该代码在GPU上运行，但比CPU上的OpenMP慢得多。

为什么OpenMP在GPU上的reduction操作会失败？

以下是我用来测试的完整代码：

#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。

- Z boson

可能sum被映射到设备作为firstprivate（这是目标区域中标量变量的默认数据共享属性）。 - Ilya Verbin

1

是的，在OpenMP 4.0和OpenMP 4.5之间，默认映射发生了变化。 - Ilya Verbin

@IlyaVerbin，你有没有想法为什么在相同的硬件上，OpenMP比OpenACC慢得多？OpenMP大约慢了两倍。 - Z boson

2

不，我没有GPU编程经验。我建议比较生成的两种情况下的PTX代码。“gcc -foffload=-save-temps ...”应该会将其转储出来。 - Ilya Verbin

为什么要使用 sudo 运行 nvprof？ - cavalcantelucas

显示剩余2条评论

1个回答

网页内容由stack overflow 提供, 点击上面的

可以查看英文原文，
原文链接

- Z boson · Accepted Answer

解决方案是添加如下语句：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)很方便。

我还手动实现了归约以查看是否可以加速。我没有设法加速它，但这里是代码（我已经尝试了其他优化，但它们都没有帮助）。

#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。