我有兴趣使用 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
.