OpenACC 数据管理语句

▶ 书中第4章，数据管理部分的代码和说明

● 代码，关于 copy，copyin，copyout，create

#include <stdio.h>
#include <openacc.h>

int main()
{
    const int length = 1024;
    int a[length], b[length], c[length],d[length];

    for (int i = 0; i < length; a[i] = b[i] = c[i] = 1);
    {
#pragma acc kernels create(d)
        for (int i = 0; i < length; i++)
        {
            a[i] ++;
            c[i] = a[i] + b[i];
            d[i] = 0;
        }
    }
    for (int i = 0; i < 10; i++)
        printf("a[%d] = %d, c[%d] = %d
", i, a[i], i, c[i]);
    getchar();
    return 0;
}

● 输出结果，显式创建了中间变量 d，隐式创建了 a，b，c，并具有不同的拷贝属性

D:CodeOpenACCOpenACCProjectOpenACCProject>pgcc -acc -Minfo main.c -o main_acc.exe
main:
     11, Generating create(d[:])
         Generating implicit copyout(c[:1024])
         Generating implicit copyin(b[:1024])
         Generating implicit copy(a[:1024])
     12, Loop is parallelizable
         Accelerator kernel generated
         Generating Tesla code
         12, #pragma acc loop gang, vector(128) /* blockIdx.x threadIdx.x */

● 在 kernels 里单独使用 copyout 时报警告：PGC-W-0996-The directive #pragma acc copyout is deprecated; use #pragma acc declare copyout instead (main.c: XX)

● enter data 和 exit data 用于 C++。

■ 首先，windows 中 pgi 不支持 C++ 编译，只有 pgcc.exe 而没有 pgc++*.exe，只能乖乖到 Linux 下去写！

■ 书上的代码有点问题，大意是 OpenACC 的 copy 是浅拷贝，对于内含指针的数据结构（如 vector，class）不会连着指针指向的对象一起拷。这里有两种解决办法，一种是去结构化，将 class 中的数据集中成简单数组来进行拷贝；另一种是使用 Managed 内存，也就不存在显式拷贝的问题了。【https://stackoverflow.com/questions/53860467/how-to-copy-on-gpu-a-vector-of-vector-pointer-memory-allocated-in-openacc】

■ 书上的代码没有采用这两种解决方案，会报错 “call to cuStreamSynchronize returned error 700: Illegal address during kernel execution” 以及 “call to cuStreamSynchronize returned error 700: Illegal address during kernel execution”，这个问题还蛮常见的【https://stackoverflow.com/search?q=call+to+returned+error+700%3A+Illegal+address+during+kernel+execution】

● 使用去结构化来使用数组

#include <iostream>
#include <vector>
#include <cstdint>

using namespace std;

int main()
{
    const int vectorCount = 1024, vectorLength = 20;
    long sum = 0;

    vector<int32_t> *vectorTable = new vector<int32_t>[vectorCount]; // 1024 个向量，每个向量放入 20 个元素
    for (int i = 0; i < vectorCount; i++)
    {
        for (int j = 0; j < vectorLength; j++)
            vectorTable[i].push_back(i);
    }
    int32_t **arrayTable = new int32_t *[vectorCount]; // 仅包含向量数据的数组，与 vectorTable 对应
    int *vectorSize = new int[vectorCount];            // 每个向量的尺寸

#pragma acc enter data create(arrayTable [0:vectorCount] [0:0]) // 设备中创建 arryTable，注意维度
    for (int i = 0; i < vectorCount; i++)
    {
        int sze = vectorTable[i].size();
        vectorSize[i] = sze;
        arrayTable[i] = vectorTable[i].data();        // 把每个向量数据的指针赋给 arrayTable
#pragma acc enter data copyin(arrayTable [i:1][:sze]) // 把每个向量的数据拷贝进设备
    }
#pragma acc enter data copyin(vectorSize[:vectorCount]) // 向量尺寸也放进设备

#pragma acc parallel loop gang vector reduction(+: sum) present(arrayTable, vectorSize) // 规约计算
    for (int i = 0; i < vectorCount; i++)
    {
        for (int j = 0; j < vectorSize[i]; ++j)
            sum += arrayTable[i][j];
    }
    cout << "Sum: " << sum << endl;

#pragma acc exit data delete (vectorSize)
#pragma acc exit data delete (arrayTable)
    delete[] vectorSize;
    delete[] vectorTable;
    return 0;
}

● 输出结果

cuan@CUAN:~$ pgc++ main.cpp -o main.exe --c++11 -ta=tesla -Minfo -acc
main:
     95, Generating enter data create(arrayTable[:1024][:0])
     32, Generating enter data copyin(arrayTable[i][:sze+1],vectorSize[:1024])
         Generating implicit copy(sum)
         Generating present(vectorSize[:])
         Generating Tesla code
         32, Generating reduction(+:sum)
         32, #pragma acc loop gang, vector(128) /* blockIdx.x threadIdx.x */
         34, #pragma acc loop seq
     32, Generating present(arrayTable[:][:])
     34, Loop is parallelizable
     41, Generating exit data delete(vectorSize[:1],arrayTable[:1][:1])
cuan@CUAN:~$ ./main.exe
Sum: 10475520

● 使用 Managed 内存

#include <iostream>

using namespace std;

class ivector
{
public:
    int len;
    int *arr;
    ivector(int length)
    {
        len = length;
        arr = new int[len];
#pragma acc enter data copyin(this)
#pragma acc enter data create(arr [0:len])
#pragma acc parallel loop present(arr [0:len])
        for (int iend = len, i = 0; i < iend; i++)              // 使用临时变量 iend，防止编译器认为 len 值在循环中会改变，从而拒绝并行化
            arr[i] = i;
    }

    ivector(const ivector &s)
    {
        len = s.len;
        arr = new int[len];
#pragma acc enter data copyin(this)
#pragma acc enter data create(arr [0:len])
#pragma acc parallel loop present(arr [0:len], s.arr [0:len])   // s 也已经在设备上了
        for (int iend = len, i = 0; i < iend; i++)
            arr[i] = s.arr[i];
    }

    ~ivector()
    {
#pragma acc exit data delete (arr)                              // 销毁对象时依次销毁设备上的 arr 和 this
#pragma acc exit data delete (this)
        cout << "deconstruction!" << endl;
        delete[] arr;
        len = 0;
    }

    int &operator[](int i)
    {
        if (i < 0 || i >= this->len)
            return arr[0];
        return arr[i];
    }

    void add(int c)
    {
#pragma acc kernels loop present(arr [0:len])                   // 每次涉及修改 arr 的操作都要注明 present
        for (int iend = len, i = 0; i < iend; i++)
            arr[i] += c;
    }

    void updateHost()                                           // 手动更新主机端数据
    {
#pragma acc update host(arr [0:len])
    }
};

int main()
{
    ivector s1(20);
    s1.add(10);
    s1.updateHost();
    cout << "s1[1] = " << s1[1] << endl;

    ivector s2(s1);
    s2.updateHost();
    cout << "s2[1] = " << s2[1] << endl;

    return 0;
}

● 输出结果，不加 -ta=tesla:managed 会报错【填坑】

cuan@CUAN:~$ pgc++ main.cpp -o main.exe --c++11 -ta=tesla:managed -Minfo -acc
ivector::ivector(int):
     13, Generating enter data copyin(this[:1])
         Generating enter data create(arr[:len])
         Generating Tesla code
         17, #pragma acc loop gang, vector(128) /* blockIdx.x threadIdx.x */
     13, Generating implicit copy(this[:])
         Generating present(arr[:len])
ivector::ivector(const ivector&):
     24, Generating enter data create(arr[:len])
         Generating enter data copyin(this[:1])
         Generating present(arr[:len])
         Generating Tesla code
         28, #pragma acc loop gang, vector(128) /* blockIdx.x threadIdx.x */
     24, Generating implicit copyin(s[:])
         Generating implicit copy(this[:])
         Generating present(s->arr[:len])
ivector::~ivector():
     36, Generating exit data delete(this[:1],arr[:1])
ivector::add(int):
      0, Generating Tesla code
     49, Accelerator serial kernel generated
         Generating implicit copy(this[:])
         Generating present(arr[:len])
     51, Loop is parallelizable
         Generating Tesla code
         51, #pragma acc loop gang, vector(128) /* blockIdx.x threadIdx.x */
ivector::updateHost():
     58, Generating update self(arr[:len])
cuan@CUAN:~$ ./main.exe
launch CUDA kernel  file=/home/cuan/main.cpp function=_ZN7ivectorC1Ei line=13 device=0 threadid=1 num_gangs=1024 num_workers=1 vector_length=128 grid=1024 block=128
launch CUDA kernel  file=/home/cuan/main.cpp function=_ZN7ivector3addEi line=49 device=0 threadid=1 num_gangs=1 num_workers=1 vector_length=1 grid=1 block=1
launch CUDA kernel  file=/home/cuan/main.cpp function=_ZN7ivector3addEi line=51 device=0 threadid=1 num_gangs=1024 num_workers=1 vector_length=128 grid=1024 block=128
s1[1] = 11
launch CUDA kernel  file=/home/cuan/main.cpp function=_ZN7ivectorC1ERKS_ line=24 device=0 threadid=1 num_gangs=1024 num_workers=1 vector_length=128 grid=1024 block=128
s2[1] = 11
deconstruction!
deconstruction!

 ● 在这本书上找到了 C++ 中使用 OpenACC 的办法 【https://www.elsevier.com/books/parallel-programming-with-openacc/farber/978-0-12-410397-9】，代码是【https://github.com/rmfarber/ParallelProgrammingWithOpenACC/tree/master/Chapter05】下的 accList.double.cpp

// accList.h
#ifndef ACC_LIST_H
#define ACC_LIST_H

#include <cstdlib>
#include <cassert>
#ifdef _OPENACC
#include <openacc.h>
#endif

template<typename T>
class accList
{
public:
    explicit accList() {}
    explicit accList(size_t size)           // 构造函数把 this 指针拷进设备，然后创建内存
    {
#pragma acc enter data copyin(this)         
        allocate(size);
    }

    ~accList()                              // 析构时释放内存，再删除 this 指针
    {
        release();
#pragma acc exit data delete(this)          
    }

#pragma acc routine seq
    T& operator[](size_t idx)
    {
        return _A[idx];
    }

#pragma acc routine seq
    const T& operator[](size_t idx) const
    {
        return _A[idx];
    }

    size_t size() const
    {
        return _size;
    }

    accList& operator=(const accList& B)
    {
        allocate(B.size());
        for (size_t j = 0; j < _size; ++j)
        {
            _A[j] = B[j];
        }
        accUpdateDevice();
        return *this;
    }

    void insert(size_t idx, const T& val)
    {
        _A[idx] = val;
    }
    void insert(size_t idx, const T* val)
    {
        _A[idx] = *val;
    }

    void accUpdateSelf()
    {
        accUpdateSelfT(_A, 0);
    }
    void accUpdateDevice()
    {
        accUpdateDeviceT(_A, 0);
    }

private:
    T * _A{ nullptr };                      // 数据成员只有指针和长度
    size_t _size{ 0 };

    void release()
    {
        if (_size > 0)
        {
#pragma acc exit data delete(_A[0:_size])   // 释放内存时删除设备内存
            delete[] _A;
            _A = nullptr;
            _size = 0;
        }
    }

    void allocate(size_t size)
    {
        if (_size != size)                  // 申请内存尺寸与当前尺寸不一致时重新开辟一块
        {
            release();
            _size = size;
#pragma acc update device(_size)
            if (_size > 0)
            {
                _A = new T[_size];
#ifdef _OPENACC                             // 有 OpenACC 的话检查 _A 是否已经在设备上了
                assert(!acc_is_present(&_A[0], sizeof(T)));
#endif
#pragma acc enter data create(_A[0:_size])  // 在设备上申请新内存
            }
        }
    }

    template<typename U>
    void accUpdateSelfT(U *p, long)
    {
#pragma acc update self(p[0:_size])
    }

    template<typename U>
    auto accUpdateSelfT(U *p, int) -> decltype(p->accUpdateSelf())
    {
        for (size_t j = 0; j < _size; ++j)
        {
            p[j].accUpdateSelf();
        }
    }

    template<typename U>
    void accUpdateDeviceT(U *p, long)
    {
#pragma acc update device(p[0:_size])
    }

    template<typename U>
    auto accUpdateDeviceT(U *p, int) -> decltype(p->accUpdateDevice())
    {
        for (size_t j = 0; j < _size; ++j)
        {
            p[j].accUpdateDevice();
        }
    }
};
#endif

// main.cpp
#include <iostream>
#include <cstdlib>
#include <cstdint>
#include "accList.h" 
using namespace std;
#ifndef N
#define N 1024
#endif

int main()
{
    accList<double> A(N), B(N);
    for (int i = 0; i < B.size(); ++i)
        B[i] = 2.5;
    B.accUpdateDevice();                                // 手动更新设备内存
#pragma acc parallel loop gang vector present(A,B)
    for (int i = 0; i < A.size(); ++i)
        A[i] = B[i] + i;
    A.accUpdateSelf();                                  // 手动更新主机内存
    for (int i = 0; i<10; ++i)
        cout << "A[" << i << "]: " << A[i] << endl;
    cout << "......" << endl;
    for (int i = N - 10; i<N; ++i)
        cout << "A[" << i << "]: " << A[i] << endl;
    return 0;
}

● 运行结果

cuan@CUAN:~/acc$ pgc++ main.cpp -o main.exe -Minfo -acc
main:
     16, Generating present(B,A)
         Generating Tesla code
         18, #pragma acc loop gang, vector(128) /* blockIdx.x threadIdx.x */
accList<double>::accList(unsigned long):
      4, include "accList.h"
          18, Generating enter data copyin(this[:1])
accList<double>::~accList():
      4, include "accList.h"
          25, Generating exit data delete(this[:1])
accList<double>::operator [](unsigned long):
      4, include "accList.h"
          29, Generating acc routine seq
              Generating Tesla code
accList<double>::size() const:
      4, include "accList.h"
          40, Generating implicit acc routine seq
              Generating acc routine seq
              Generating Tesla code
accList<double>::release():
      4, include "accList.h"
          82, Generating exit data delete(_A[:_size])
accList<double>::allocate(unsigned long):
      4, include "accList.h"
          95, Generating update device(_size)
         102, Generating enter data create(_A[:_size])
void accList<double>::accUpdateSelfT<double>(T1 *, long):
      4, include "accList.h"
         110, Generating update self(p[:_size])
void accList<double>::accUpdateDeviceT<double>(T1 *, long):
      4, include "accList.h"
         125, Generating update device(p[:_size])
cuan@CUAN:~/acc$ ./main.exe 
launch CUDA kernel  file=/home/cuan/acc/main.cpp function=main line=16 device=0 threadid=1 num_gangs=1024 num_workers=1 vector_length=128 grid=1024 block=128
A[0]: 2.5
A[1]: 3.5
A[2]: 4.5
A[3]: 5.5
A[4]: 6.5
A[5]: 7.5
A[6]: 8.5
A[7]: 9.5
A[8]: 10.5
A[9]: 11.5
......
A[1014]: 1016.5
A[1015]: 1017.5
A[1016]: 1018.5
A[1017]: 1019.5
A[1018]: 1020.5
A[1019]: 1021.5
A[1020]: 1022.5
A[1021]: 1023.5
A[1022]: 1024.5
A[1023]: 1025.5

Accelerator Kernel Timing data
/home/cuan/acc/main.cpp
  main  NVIDIA  devicenum=0
    time(us): 13
    16: compute region reached 1 time
        16: kernel launched 1 time
            grid: [1024]  block: [128]
             device time(us): total=13 max=13 min=13 avg=13
            elapsed time(us): total=329 max=329 min=329 avg=329
    16: data region reached 2 times
/home/cuan/acc/main.cpp
  _ZN7accListIdEC1Em  NVIDIA  devicenum=0
    time(us): 12
    18: data region reached 2 times
        18: data copyin transfers: 2
             device time(us): total=12 max=9 min=3 avg=6
/home/cuan/acc/main.cpp
  _ZN7accListIdED1Ev  NVIDIA  devicenum=0
    time(us): 0
    25: data region reached 2 times
/home/cuan/acc/main.cpp
  _ZN7accListIdE7releaseEv  NVIDIA  devicenum=0
    time(us): 7
    82: data region reached 2 times
        82: data copyin transfers: 2
             device time(us): total=7 max=4 min=3 avg=3
/home/cuan/acc/main.cpp
  _ZN7accListIdE8allocateEm  NVIDIA  devicenum=0
    time(us): 11
    95: update directive reached 2 times
        95: data copyin transfers: 2
             device time(us): total=5 max=3 min=2 avg=2
    102: data region reached 2 times
        102: data copyin transfers: 2
             device time(us): total=6 max=3 min=3 avg=3
/home/cuan/acc/main.cpp
  _ZN7accListIdE14accUpdateSelfTIdEEvPT_l  NVIDIA  devicenum=0
    time(us): 10
    110: update directive reached 1 time
        110: data copyout transfers: 1
             device time(us): total=10 max=10 min=10 avg=10
/home/cuan/acc/main.cpp
  _ZN7accListIdE16accUpdateDeviceTIdEEvPT_l  NVIDIA  devicenum=0
    time(us): 3
    125: update directive reached 1 time
        125: data copyin transfers: 1
             device time(us): total=3 max=3 min=3 avg=3