PGI 编译器的 OpenACC 性能不佳

【问题标题】：OpenACC bad performance with PGI CompilerPGI 编译器的 OpenACC 性能不佳
【发布时间】：2020-10-22 14:03:57
【问题描述】：

我正在尝试使用 OpenACC 来完成我的期末论文，以提高代码的性能。但是，优化的代码比我尝试了很多东西的串行代码差，但没有一个有效。如果有人可以帮助我，我将不胜感激。

1332    void SOR()
1333    {
1334        iteracoes_solver = 0;
1335        
1336        
1337        {
1338    
1339            #pragma acc data copy(AC_lin[2:nx*ny], Po1_lin[2:nx*ny], AUXW1[2:nx*ny], Qo_lin[2:nx*ny], Tgx[2:nx*ny], Po3_lin[2:nx*ny], Tgy[2:nx*ny], P_auxi[2:nx*ny], vetErro[2:nx*ny]) \
1340            copyin(AC_lin[2:nx*ny], Tgx[2:nx*ny], Tgy[2:nx*ny])
1341            while (iteracoes_solver < 10001)
1342            {
1343    
1344                
1345                //#pragma acc kernels loop private(n) copyin(P_auxi[0:nxy]) vector(512)
1346                for (n = 1; n <= nxy; n++)    
1347                    P_auxi[n] = Po3_lin[n];
1348    
1349                
1350                {
1351    
1352                    erro = 0.0;
1353    
1354                    if (total_iter == 10000)
1355                    {
1356                        printf("---> TEMPO=%g\n---> Numero maximo de iteracoes atingido no solver\n---> Corrigir parametros\n", tempo_atual);//getch();
1357                        exit_tool();
1358                    }
1359    
1360                    n = 1;
1361                    auxf1 = (AC_lin[n] * Po1_lin[n]) + (AUXW1[n] * Qo_lin[n]) + (Tgx[n] * Po3_lin[n + 1]) + (Tgy[n] * Po3_lin[n + nx]) + ((1 - AUXW1[n]));
1362                    Paux = auxf1 / (AC_lin[n] + Tgx[n] + Tgy[n] + (1 - AUXW1[n]));
1363                    Paux = (omega * Paux) + ((1.0 - omega) * P_auxi[n]);
1364                    vetErro[n] = fabs(Po3_lin[n] - Paux);
1365                    Po3_lin[n] = Paux;
1366    
1367                    n = nx;
1368                    auxf1 = (AC_lin[n] * Po1_lin[n]) + (AUXW1[n] * Qo_lin[n]) + (Tgx[n - 1] * Po3_lin[n - 1]) + (Tgy[n] * Po3_lin[n + nx]) + ((1 - AUXW1[n]));
1369                    Paux = auxf1 / (AC_lin[n] + Tgx[n - 1] + Tgy[n] + ((1 - AUXW1[n])));
1370                    Paux = (omega * Paux) + ((1.0 - omega) * P_auxi[n]);
1371                    vetErro[n] = fabs(Po3_lin[n] - Paux);
1372                    Po3_lin[n] = Paux;
1373    
1374                    auxi11 = nxy - nx;
1375    
1376                    n = auxi11 + 1;
1377                    auxf1 = (AC_lin[n] * Po1_lin[n]) + (AUXW1[n] * Qo_lin[n]) + (Tgx[n] * Po3_lin[n + 1]) + (Tgy[n - nx] * Po3_lin[n - nx]) + ((1 - AUXW1[n]));
1378                    Paux = auxf1 / (AC_lin[n] + Tgx[n] + Tgy[n - nx] + ((1 - AUXW1[n])));
1379                    Paux = (omega * Paux) + ((1.0 - omega) * P_auxi[n]);
1380                    vetErro[n] = fabs(Po3_lin[n] - Paux);
1381                    Po3_lin[n] = Paux;
1382    
1383                    n = nxy;
1384                    auxf1 = (AC_lin[n] * Po1_lin[n]) + (AUXW1[n] * Qo_lin[n]) + (Tgx[n - 1] * Po3_lin[n - 1]) + (Tgy[n - nx] * Po3_lin[n - nx]) + ((1 - AUXW1[n]));
1385                    Paux = auxf1 / (AC_lin[n] + Tgx[n - 1] + Tgy[n - nx] + ((1 - AUXW1[n])));
1386                    Paux = (omega * Paux) + ((1.0 - omega) * P_auxi[n]);
1387                    vetErro[n] = fabs(Po3_lin[n] - Paux);
1388                    Po3_lin[n] = Paux;
1389    
1390    
1391                   
1392                    #pragma acc parallel loop private(n)
1393                    for (n = 2; n < nx; n++)
1394                    {
1395                        auxf1 = (AC_lin[n] * Po1_lin[n]) + (AUXW1[n] * Qo_lin[n]) + (Tgx[n] * Po3_lin[n + 1]) + (Tgx[n - 1] * Po3_lin[n - 1]) + (Tgy[n] * Po3_lin[n + nx]) + ((1 - AUXW1[n]));
1396                        Paux = auxf1 / (AC_lin[n] + Tgx[n] + Tgx[n - 1] + Tgy[n] + ((1 - AUXW1[n])));
1397                        Paux = (omega * Paux) + ((1.0 - omega) * P_auxi[n]);
1398                        vetErro[n] = fabs(Po3_lin[n] - Paux);
1399                        Po3_lin[n] = Paux;
1400                    }
1401    
1402    
1403                    
1404                    #pragma acc parallel loop private(n)
1405                    for (j = 2; j < ny; j++)
1406                    {
1407                        auxi11 = nx * (j - 1);
1408    
1409                        n = auxi11 + 1;
1410                        auxf1 = (AC_lin[n] * Po1_lin[n]) + (AUXW1[n] * Qo_lin[n]) + (Tgx[n] * Po3_lin[n + 1]) + (Tgy[n] * Po3_lin[n + nx]) + (Tgy[n - nx] * Po3_lin[n - nx]) + ((1 - AUXW1[n]));
1411                        Paux = auxf1 / (AC_lin[n] + Tgx[n] + Tgy[n] + Tgy[n - nx] + ((1 - AUXW1[n])));
1412                        Paux = (omega * Paux) + ((1.0 - omega) * P_auxi[n]);
1413                        vetErro[n] = fabs(Po3_lin[n] - Paux);
1414                        Po3_lin[n] = Paux;
1415                    }
1416    
1417    
1418                    
1419                    //#pragma acc kernels
1420                    for (j = 2; j < ny; j++)
1421                    {
1422                        auxi11 = nx * (j - 1);
1423    
1424                        n = auxi11 + nx;
1425                        auxf1 = (AC_lin[n] * Po1_lin[n]) + (AUXW1[n] * Qo_lin[n]) + (Tgx[n - 1] * Po3_lin[n - 1]) + (Tgy[n] * Po3_lin[n + nx]) + (Tgy[n - nx] * Po3_lin[n - nx]) + ((1 - AUXW1[n]));
1426                        Paux = auxf1 / (AC_lin[n] + Tgx[n - 1] + Tgy[n] + Tgy[n - nx] + ((1 - AUXW1[n])));
1427                        Paux = (omega * Paux) + ((1.0 - omega) * P_auxi[n]);
1428                        vetErro[n] = fabs(Po3_lin[n] - Paux);
1429                        Po3_lin[n] = Paux;
1430    
1431                    }
1432    
1433                   
1434                    auxi11 = nxy - nx;
1435    
1436    
1437                   
1438                    //#pragma acc kernels
1439                    for (i = 2; i < nx; i++)
1440                    {
1441                        n = auxi11 + i;
1442                        auxf1 = (AC_lin[n] * Po1_lin[n]) + (AUXW1[n] * Qo_lin[n]) + (Tgx[n] * Po3_lin[n + 1]) + (Tgx[n - 1] * Po3_lin[n - 1]) + (Tgy[n - nx] * Po3_lin[n - nx]) + ((1 - AUXW1[n]));
1443                        Paux = auxf1 / (AC_lin[n] + Tgx[n] + Tgx[n - 1] + Tgy[n - nx] + ((1 - AUXW1[n])));
1444                        Paux = (omega * Paux) + ((1.0 - omega) * P_auxi[n]);
1445                        vetErro[n] = fabs(Po3_lin[n] - Paux);
1446                        Po3_lin[n] = Paux;
1447                    }
1448                }
1449                
1450    
1451                
1452                
1453                //#pragma acc kernels loop private(auxf1,Paux,i,j,n) vector(512) independent
1454                for (j = 2; j < ny; j++)
1455                {
1456                    for (i = 2 + (j % 2); i < nx; i += 2)
1457                    {
1458                        n = (nx * (j - 1)) + i;
1459                        auxf1 = (AC_lin[n] * Po1_lin[n]) + (AUXW1[n] * Qo_lin[n]) + (Tgx[n] * Po3_lin[n + 1]) + (Tgx[n - 1] * Po3_lin[n - 1]) + (Tgy[n] * Po3_lin[n + nx]) + (Tgy[n - nx] * Po3_lin[n - nx]) + ((1 - AUXW1[n]));
1460                        Paux = auxf1 / (AC_lin[n] + Tgx[n] + Tgx[n - 1] + Tgy[n] + Tgy[n - nx] + ((1 - AUXW1[n])));
1461                        Paux = (omega * Paux) + ((1.0 - omega) * P_auxi[n]);
1462                        vetErro[n] = fabs(Po3_lin[n] - Paux);
1463                        Po3_lin[n] = Paux;
1464                    }
1465                }
1466                
1467                
1468                //#pragma acc kernels loop private(auxf1,Paux,i,j,n) vector(512) independent
1469                for (j = 2; j < ny; j++)
1470                {
1471                    for (i = 3 - (j % 2); i < nx; i += 2)
1472                    {
1473                        n = (nx * (j - 1)) + i;
1474                        auxf1 = (AC_lin[n] * Po1_lin[n]) + (AUXW1[n] * Qo_lin[n]) + (Tgx[n] * Po3_lin[n + 1]) + (Tgx[n - 1] * Po3_lin[n - 1]) + (Tgy[n] * Po3_lin[n + nx]) + (Tgy[n - nx] * Po3_lin[n - nx]) + ((1 - AUXW1[n]));
1475                        Paux = auxf1 / (AC_lin[n] + Tgx[n] + Tgx[n - 1] + Tgy[n] + Tgy[n - nx] + ((1 - AUXW1[n])));
1476                        Paux = (omega * Paux) + ((1.0 - omega) * P_auxi[n]);
1477                        vetErro[n] = fabs(Po3_lin[n] - Paux);
1478                        Po3_lin[n] = Paux;
1479                    }
1480                }
1481                
1482                
1483                #pragma acc parallel loop private(n)
1484                for (n = 0; n < nxy; n++) {
1485                    erro += vetErro[n];
1486                }
1487    
1488                
1489                if ((erro / nxy < tol)) {
1490                    //#pragma omp cancel parallel
1491                    break;
1492                }
1493                iteracoes_solver += 1;
1494            }
1495    
1496        }
1497    }

我放不下所有的decode，所以放了一部分代码。我评论了一些编译指示以供测试。当我的 nx=ny >=1000 时，串行代码优于优化代码。我不明白的另一件事是我删除了代码的所有 pragma，然后运行 ​​PGI 编译器来优化代码，串行代码更好，我不知道为什么。最后，我有两个问题：

1 - 我做错了什么？ 2 - 为什么没有编译指示的代码比串行代码差？

最好的问候，布雷诺。

更新： 我忘了把编译信息。我运行代码

pgcc -ta=multicore -Minfo=all file.c -o file.exe

并生成

 1360, FMA (fused multiply-add) instruction(s) generated
   1392, Generating Multicore code
       1393, #pragma acc loop gang
   1393, FMA (fused multiply-add) instruction(s) generated
   1404, Generating Multicore code
       1405, #pragma acc loop gang
   1405, FMA (fused multiply-add) instruction(s) generated
   1420, FMA (fused multiply-add) instruction(s) generated
   1439, FMA (fused multiply-add) instruction(s) generated
   1456, FMA (fused multiply-add) instruction(s) generated
   1471, FMA (fused multiply-add) instruction(s) generated
   1483, Generating Multicore code
       1484, #pragma acc loop gang
   1485, Generating implicit reduction(+:erro)

【问题讨论】：

  • 由于代码不完整、出错（如前所述，while 循环不会终止），并且您已注释掉所需的 pragma，因此很难提供任何具体建议。相反，编译器反馈消息 (-Minfo=accel) 告诉您什么？循环是成功并行化还是在设备上串行运行？探查器显示什么？您可以通过设置环境变量“PGI_ACC_TIME=1”来进行快速配置，或者根据您使用的编译器版本使用 pgprof 或 Nsight-Systems。
  • 正如 Mat 所说，-Minfo=accel 是你的朋友。另请注意，一种典型的方法是逐步并行化代码并检查性能和正确性，而不是进行大规模转换。如果您在并行化之前和之后对单个循环进行计时，那将为您提供很好的信息。
  • @MatColgrove，感谢您的回答。我用编译器消息更新了这个问题。我尝试更改环境变量 PGI_ACC_TIME，但我不能。我用“export PGI_ACC_TIME=1”替换了许多文件，但没有一个起作用：/
  • @Richard，也感谢您的回答。并感谢您的提示。我现在正在尝试这个。 :)
  • 好的，从更新中我看到您的目标是多核。 PGI_ACC_TIME 只会分析设备执行情况，因此解释了为什么没有给出输出。对于多核，请检查您的绑定以及如何使用核心。您是通过批处理作业还是在系统本地运行？如果在本地，打开一个新的 shell 并通过 'top' 监视运行以查看绑定。默认情况下，将使用所有内核，包括超线程，其中一种可能性是超额订阅。您可以通过环境变量 ACC_NUM_CORES 显式设置要使用的内核数。

标签： c hpc openacc


【解决方案1】：

通常，使用 OMP 功能会导致大量上下文切换，这实际上会减慢代码的执行速度。 SO OMP 非常适合 I/O 绑定代码，但不适用于 CPU 绑定代码。

重复计算应该移到循环之前：

什么是：2:nx*ny？

索引变量的“范围”应尽可能有限。因此，在编写for() 语句时，第一个参数还应声明索引变量，对于大多数for() 语句而言，索引变量类似于：

for( size_t i = 0; .,.. )

关于：1.0 这是 double 文字，但代码调用 fabs() 需要 float，建议文字为：1.0f

与几个计算一样，这个计算：(1.0 - omega) 不会改变，所以这个计算应该在循环之前只进行一次

如果我们没有看到串行代码，我们怎么能说为什么没有编译指示的代码比串行代码慢？

关于：

while (iteracoes_solver < 10001)

变量：iteracoes_solver 被初始化为 0，并且在发布的代码中的任何地方都没有递增。

【讨论】：

  • 嗨，@user3629249，感谢您的回答。在这种情况下，我需要指定数组的大小，大小为 [0:nx*ny] （我放错了，抱歉）。当我说串行代码时，我的意思是说代码是按顺序逐行执行的，因此串行代码将是上面在没有 pgi 编译器的情况下运行的代码。我用 iteracoes_solver 更新代码。
  • 自从 80 年代初/中期 C 首次面世以来，我就一直在编程。我不认识这个表达式：[0:nx*ny] 这是什么意思？
  • 你指的是第 1339 行，对吗？如果是，当我运行 openACC 时，它无法识别数组的大小，所以我找到的一种解决方案是在 pragma 中指定大小。我在 "Array Shape" 部分找到了这个问题的解决方案 [docs.computecanada.ca/wiki/OpenACC_Tutorial_-_Data_movement]。
猜你喜欢
  • 2019-06-10
  • 1970-01-01
  • 1970-01-01
  • 2015-11-03
  • 2023-03-24
  • 1970-01-01
  • 2016-04-10
  • 1970-01-01
  • 2014-03-17
相关资源
最近更新 更多
热门标签
Java Python linux javascript Mysql C# Docker 算法 前端 SpringBoot Redis Vue spring 设计模式 .net core .net kubernetes c++ 数据库 数据结构 大数据 js 机器学习 微服务 Android Go 程序员 面试 JVM ASP.net core 云原生 人工智能 后端 PHP git CSS golang k8s Nginx Django mybatis 深度学习 多线程 React 架构 devops 爬虫 云计算 Spring Boot LeetCode