OpenACC 简单的原子操作

▶ OpenACC 的原子操作，用到了 C++ 的一个高精度计时器

● 代码，直接的原子操作

 #include <iostream>

 #include <cstdlib>

 #include <chrono>

 #define ATOMIC

 using namespace std;

 using namespace std::chrono;

 int main()

 {

     high_resolution_clock::time_point t1 = high_resolution_clock::now();// 高精度计时器

     const int count = ;

     int sum = ;

 #pragma acc parallel loop copyout(sum)

     for (int i = ; i < count; i++)

     {

 #ifdef ATOMIC

     #pragma acc atomic update

 #endif

         sum++;

     }

     high_resolution_clock::time_point t2 = high_resolution_clock::now();

     duration<double> time = duration_cast<duration<double>>(t2 - t1);

     cout << "\nCount = " << count << ", duraion = " << time.count() << " s" << endl;

     return ;

 }

● 输出结果，不知道为什么，win10中的 pgCC 不能用

D:\Code\OpenACC\OpenACCProject\OpenACCProject>pgCC -acc -o acc_win10.exe main.cpp -Minfo

pgCC-Warning-C++ compilation is not supported: main.cpp

● 输出结果，WSL 中

// 不使用 OpenACC

cuan@CUAN:/mnt/d/Code/OpenACC/OpenACCProject/OpenACCProject$ pgc++ -std=c++ -o acc.exe main.cpp -Minfo

cuan@CUAN:/mnt/d/Code/OpenACC/OpenACCProject/OpenACCProject$ ./acc.exe

Count = , duraion = 0.483907 s

// 使用宏 ATOMIC，即使用原子操作

cuan@CUAN:/mnt/d/Code/OpenACC/OpenACCProject/OpenACCProject$ pgc++ -acc -std=c++ -o acc.exe main.cpp -Minfo

main:

     , Generating copyout(sum)

         Accelerator kernel generated

         Generating Tesla code

         , #pragma acc loop gang, vector(128) /* blockIdx.x threadIdx.x */

cuan@CUAN:/mnt/d/Code/OpenACC/OpenACCProject/OpenACCProject$ ./acc.exe

Count = , duraion = 0.248377 s

// 不用宏 ATOMIC，即不用原子操作

cuan@CUAN:/mnt/d/Code/OpenACC/OpenACCProject/OpenACCProject$ pgc++ -acc -std=c++ -o acc.exe main.cpp -Minfo

main:

     , Generating copyout(sum)

         Accelerator kernel generated

         Generating Tesla code

         , #pragma acc loop seq

     , Accelerator restriction: induction variable live-out from loop: sum    // 编译器提示强制原子操作

cuan@CUAN:/mnt/d/Code/OpenACC/OpenACCProject/OpenACCProject$ ./acc.exe

Count = , duraion = 0.247399 s

● 优化一下，使用分段计数

 #include <iostream>

 #include <cstdlib>

 #include <chrono>

 using namespace std;

 using namespace std::chrono;

 int main()

 {

     high_resolution_clock::time_point t1 = high_resolution_clock::now();

     const int count = , length = count / ;// 每一段的长度

     int sum = ;

 #pragma acc parallel loop copyout(sum)

     for (int start = ; start < count; start+=length)                       // start 取每段的起点，共 count / length 段

     {

         const int end = (start + length < count) ? start + length : count;  // 每段的终点

         int subSum = ;

 #pragma acc loop worker reduction(+:subSum)

         for (int j = start; j < end; j++)                                   // 每段从 start 加到 end

             subSum ++;

 #pragma acc atomic update

         sum += subSum;                                                      // 规约结果加到 sum 中来

     }

     high_resolution_clock::time_point t2 = high_resolution_clock::now();

     duration<double> time = duration_cast<duration<double>>(t2 - t1);

     cout << "\nCount = " << sum << ", duraion = " << time.count() << " s" << endl;

     return ;

 }

● 输出结果，好像好一点点

cuan@CUAN:/mnt/d/Code/OpenACC/OpenACCProject/OpenACCProject$ pgc++ -acc -std=c++ -o acc.exe main.cpp -Minfo

main:

     , Generating copyout(sum)

         Accelerator kernel generated

         Generating Tesla code

         , #pragma acc loop gang /* blockIdx.x */

         , #pragma acc loop seq /* threadIdx.y */

             Generating reduction(+:subSum)

     , Loop is parallelizable

cuan@CUAN:/mnt/d/Code/OpenACC/OpenACCProject/OpenACCProject$ ./acc.exe

Count = , duraion = 0.246488 s