OpenMP卸载到Nvidia错误还原

OpenMP offloading to Nvidia wrong reduction

本文关键字：错误 还原 Nvidia 卸载 OpenMP      更新时间：2023-10-16

我有兴趣使用 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慢得多。

为什么 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 = %dn",sum);
return 0;
}

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

解决方案是添加子句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 %dn", omp_get_num_devices());
printf("default device %dn", omp_get_default_device());
printf("device id %dn", omp_get_initial_device());
printf("nteams %dn", nteams);
printf("nthreads per team %dn", nthreads);
printf("total threads %dn", nteams*nthreads);
printf("sum %dn", 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.

相关文章：