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 = 1073741824;
    int sum = 0;

#pragma acc parallel loop copyout(sum)
    for (int i = 0; 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 << "
Count = " << count << ", duraion = " << time.count() << " s" << endl;        
    return 0;
}

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

D:CodeOpenACCOpenACCProjectOpenACCProject>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++11 -o acc.exe main.cpp -Minfo
cuan@CUAN:/mnt/d/Code/OpenACC/OpenACCProject/OpenACCProject$ ./acc.exe

Count = 1073741824, duraion = 0.483907 s

// 使用宏 ATOMIC，即使用原子操作
cuan@CUAN:/mnt/d/Code/OpenACC/OpenACCProject/OpenACCProject$ pgc++ -acc -std=c++11 -o acc.exe main.cpp -Minfo
main:
     15, Generating copyout(sum)
         Accelerator kernel generated
         Generating Tesla code
         18, #pragma acc loop gang, vector(128) /* blockIdx.x threadIdx.x */
cuan@CUAN:/mnt/d/Code/OpenACC/OpenACCProject/OpenACCProject$ ./acc.exe

Count = 1073741824, duraion = 0.248377 s

// 不用宏 ATOMIC，即不用原子操作
cuan@CUAN:/mnt/d/Code/OpenACC/OpenACCProject/OpenACCProject$ pgc++ -acc -std=c++11 -o acc.exe main.cpp -Minfo
main:
     15, Generating copyout(sum)
         Accelerator kernel generated
         Generating Tesla code
         18, #pragma acc loop seq
     23, Accelerator restriction: induction variable live-out from loop: sum    // 编译器提示强制原子操作
cuan@CUAN:/mnt/d/Code/OpenACC/OpenACCProject/OpenACCProject$ ./acc.exe

Count = 1073741824, 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 = 1073741824, length = count / 32;// 每一段的长度
    int sum = 0;

#pragma acc parallel loop copyout(sum)
    for (int start = 0; start < count; start+=length)                       // start 取每段的起点，共 count / length 段
    {
        const int end = (start + length < count) ? start + length : count;  // 每段的终点
        int subSum = 0;
#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 << "
Count = " << sum << ", duraion = " << time.count() << " s" << endl;                
    return 0;
}

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

cuan@CUAN:/mnt/d/Code/OpenACC/OpenACCProject/OpenACCProject$ pgc++ -acc -std=c++11 -o acc.exe main.cpp -Minfo
main:
     15, Generating copyout(sum)
         Accelerator kernel generated
         Generating Tesla code
         18, #pragma acc loop gang /* blockIdx.x */
         23, #pragma acc loop seq /* threadIdx.y */
             Generating reduction(+:subSum)
     23, Loop is parallelizable
cuan@CUAN:/mnt/d/Code/OpenACC/OpenACCProject/OpenACCProject$ ./acc.exe

Count = 1073741824, duraion = 0.246488 s