• OpenAcc的使用

    安装cuda驱动

    1. 更新源
    sudo add-apt-repository ppa:graphics-drivers/ppa
sudo apt update

    • 1
    • 2
    1. 用以下指令找到recommend的驱动
    ubuntu-drivers devices

    • 1
    1. 安装
    sudo apt install nvidia-driver-XXX

    • 1

    安装后运行nvidia-smi有结果

    安装cuda toolkit

    （其实不用安装，这是nvcc编译.cu的， 单纯记录一下）
    然后去cuda官网找相应版本的toolkit，不要用Ubuntu的源，会对不上版本而冲突

    需要在.bashrc中加入如下命令

    export CUDA_HOME=/usr/local/cuda 
export PATH=$PATH:$CUDA_HOME/bin 
export LD_LIBRARY_PATH=/usr/local/cuda-11.7/lib64${LD_LIBRARY_PATH:+:${LD_LIBRARY_PATH}}

    • 1
    • 2
    • 3

    安装后重开terminal运行nvcc -V有结果，且不会因为冲突导致运行nvidia-smi无结果

    安装nvhpc

    理论上g++-12（Ubuntu22的源里是带g++-12的，所以会很方便）可以编译OpenAcc的代码，但是结果没有达到我要的效果，遂采用nvhpc。

    NVIDIA HPC，原来好像是PGI
    根据官网，Ubuntu上安装运行如下命令：

    $ echo 'deb [trusted=yes] https://developer.download.nvidia.com/hpc-sdk/ubuntu/amd64 /' | sudo tee /etc/apt/sources.list.d/nvhpc.list
$ sudo apt-get update -y
$ sudo apt-get install -y nvhpc-22-7

    • 1
    • 2
    • 3

    然后在.bashrc中加入如下命令

    export NVARCH=`uname -s`_`uname -m`
export NVCOMPILERS=/opt/nvidia/hpc_sdk
export PATH=$NVCOMPILERS/$NVARCH/22.7/compilers/bin:$PATH
export MANPATH=$MANPATH:$NVCOMPILERS/$NVARCH/22.7/compilers/man

    • 1
    • 2
    • 3
    • 4

    重开terminal运行 nvc++有结果

    安装gcc-11

    sudo apt install nvptx-tools gcc-11-offload-nvptx

    • 1

    编译OpenAcc

    可以尝试编译以下代码 看运行时间

    nvc++ testACC.cpp -o test -acc

    • 1

    也可以用g++11，更方便，如果用QT 不用在QT折腾多编译器（QT尝试nvc++混合编译，但是失败了）

    g++-11 testACC.cpp -fopenacc -fcf-protection=none -fopt-info-optimized-omp -o test -fno-stack-protector```

    • 1

    其中-fopt-info-optimized-omp可选，用于输出OpenAcc优化的地方。
    在终端输入export GOMP_DEBUG=1 后可以看汇编输出

    #include 
#include 
#define N 1000
using namespace std;
int sum = 0;
int a[1010],b[1010], c[1010];

void init(){
        
    //[start:count]
    // 用copyin/copy/copyout 看是不是真的在gpu上跑
    #pragma acc enter data copyin(sum) copyin(a[1:N], b[1:N], c[1:N])
    ;
}

void update_host()
{
    #pragma acc update self(a[1:N], b[1:N], c[1:N], sum)
    ;
}

void update_device()
{
    #pragma acc update device(a[1:N], b[1:N], c[1:N], sum)
    ;
}

void exit(){
    #pragma acc exit data copyout(sum) copyout(a[1:N], b[1:N], c[1:N])
    ;
}

int main(){
    
    clock_t start, end;

    //忽略GPU启动时间
    init();

    start = clock();
    
    update_device();

    //尽量别把parallel围多个循环，因为loop会把下面的代码可能都loop一遍
    //不能单独用loop seq，loop seq只能在loop independent下面用，不然会被执行gang遍
    #pragma acc parallel present(a[1:N], b[1:N], c[1:N], sum) num_gangs(1024) // kernels 中多块循环是依次执行， 但是parallel中 多块循环是同时执行 ，这样会导致 a还没赋值，b已经被赋值了
    {
        #pragma acc loop independent
        for (int i = 1; i <= N; i++){
            
            #pragma acc loop independent
            for (int j = 1; j <= N; j++){
                #pragma acc loop independent
                for (int k = 1; k <= N; k++){
                    a[i] = i ;
                }
            }
        }

        #pragma acc loop independent
        for (int i = 1; i <= N; i++){
            b[i] = b[i] + i*i; //用加法，检测执行了多少遍
        }

        #pragma acc loop independent
        for(int i = 1; i <= N; i++){
            #pragma acc atomic
            sum = sum + b[i]; //用加法，检测执行了多少遍
        }

    }

    update_host();    
    exit();
    end = clock();
    
    cout<< b[90] <<" "<<sum <<" :"<< (double)(end - start)/CLOCKS_PER_SEC * 1000 <<"ms"<<endl; //N=1000 : 8100 333833500

}

    • 1
    • 2
    • 3
    • 4
    • 5
    • 6
    • 7
    • 8
    • 9
    • 10
    • 11
    • 12
    • 13
    • 14
    • 15
    • 16
    • 17
    • 18
    • 19
    • 20
    • 21
    • 22
    • 23
    • 24
    • 25
    • 26
    • 27
    • 28
    • 29
    • 30
    • 31
    • 32
    • 33
    • 34
    • 35
    • 36
    • 37
    • 38
    • 39
    • 40
    • 41
    • 42
    • 43
    • 44
    • 45
    • 46
    • 47
    • 48
    • 49
    • 50
    • 51
    • 52
    • 53
    • 54
    • 55
    • 56
    • 57
    • 58
    • 59
    • 60
    • 61
    • 62
    • 63
    • 64
    • 65
    • 66
    • 67
    • 68
    • 69
    • 70
    • 71
    • 72
    • 73
    • 74
    • 75
    • 76
    • 77
    • 78
    • 79

    可以用下面的语句看每一个模块运行的时间：

    nvprof ./test

    • 1

    还可以看一下并行效果
    理论上不开并行 结果是 0 0，开了结果是1 2，不过这在不同编译器是不一定的

    #include 
int main(){
    int a[100];
    #pragma acc enter data create(a[0:100])

    #pragma acc kernels present(a[0:100])
    {
        #pragma acc loop
        for (int i=0; i<100; i++){
            a[i] = i;
        }
        
        #pragma acc loop independent
        for (int i=1; i<100; i++){
            a[i] = a[i-1];
        }
    }
    #pragma acc exit data copyout(a[0:100]) 
    std::cout<< a[2] <<" "<< a[3] << std::endl;
    return 0;
}

    • 1
    • 2
    • 3
    • 4
    • 5
    • 6
    • 7
    • 8
    • 9
    • 10
    • 11
    • 12
    • 13
    • 14
    • 15
    • 16
    • 17
    • 18
    • 19
    • 20
    • 21

    openACC调用函数（gcc）

    如果调用一个通过形参的函数，直接在函数的定义前加上#pragma acc routine seq，但是如果不通过形参直接调用GPU上的参数，需要在参数声明的时候declare一下，但好像会导致这个数据传不回来，直接为0（应该是我的问题，希望大佬告诉我）（但是如果用nvc++编译下面两段程序，结果都是1000 1000），但可以通过其他变量间接传回来。

    方法一：传形参

    #include 
#define N 1000
using namespace std;

int k = 0;
// #pragma acc declare create(k)
int g = 0;
class C{
    public:
        C();
        ~C();
        void fun();
        void fun2(int & x);
        void fun1();
        void update_host();
        void update_device();
        void init();
        void exit();
    private:
        int *arr;

};


C::C(){

}

C::~C(){

}

void C::update_host(){
    #pragma acc update self(k,g)
    ;
}

void C::update_device(){
    #pragma acc update device(k,g)
    ;
}

void C::init(){
    #pragma acc enter data copyin(k,g)
    ;
}

void C::exit(){
    #pragma acc exit data copyout(k,g)
    ;
}

#pragma acc routine seq
void C::fun1(){
    // {
    //     #pragma acc atomic
    //     k = k + 1;
    // }
}

#pragma acc routine seq
void C::fun2(int & x){
    #pragma acc atomic
    x = x + 1;
}


void C::fun(){

    init();    
    update_device();
    #pragma acc parallel present(k,g)
    {
        #pragma acc loop independent
        for (int i = 1; i <= N; i++){
            // fun1();
            fun2(k);
        }
    }
    #pragma acc kernels present(k,g)
    {   
        g = k;
    }

    update_host();    
    exit();
    std::cout<<g<<" "<<k<<std::endl;

}

int main(){
    C cl;
    cl.fun();
}

    • 1
    • 2
    • 3
    • 4
    • 5
    • 6
    • 7
    • 8
    • 9
    • 10
    • 11
    • 12
    • 13
    • 14
    • 15
    • 16
    • 17
    • 18
    • 19
    • 20
    • 21
    • 22
    • 23
    • 24
    • 25
    • 26
    • 27
    • 28
    • 29
    • 30
    • 31
    • 32
    • 33
    • 34
    • 35
    • 36
    • 37
    • 38
    • 39
    • 40
    • 41
    • 42
    • 43
    • 44
    • 45
    • 46
    • 47
    • 48
    • 49
    • 50
    • 51
    • 52
    • 53
    • 54
    • 55
    • 56
    • 57
    • 58
    • 59
    • 60
    • 61
    • 62
    • 63
    • 64
    • 65
    • 66
    • 67
    • 68
    • 69
    • 70
    • 71
    • 72
    • 73
    • 74
    • 75
    • 76
    • 77
    • 78
    • 79
    • 80
    • 81
    • 82
    • 83
    • 84
    • 85
    • 86
    • 87
    • 88
    • 89
    • 90
    • 91
    • 92
    • 93
    • 94

    输出：1000 1000

    方法二：直接调用

    #include 
#define N 1000
using namespace std;

int k = 0;
#pragma acc declare create(k)
int g = 0;
class C{
    public:
        C();
        ~C();
        void fun();
        void fun2(int & x);
        void fun1();
        void update_host();
        void update_device();
        void init();
        void exit();
    private:
        int *arr;

};


C::C(){

}

C::~C(){

}

void C::update_host(){
    #pragma acc update self(k,g)
    ;
}

void C::update_device(){
    #pragma acc update device(k,g)
    ;
}

void C::init(){
    #pragma acc enter data copyin(k,g)
    ;
}

void C::exit(){
    #pragma acc exit data copyout(k,g)
    ;
}

#pragma acc routine seq
void C::fun1(){
    {
        #pragma acc atomic
        k = k + 1;
    }
}

#pragma acc routine seq
void C::fun2(int & x){
    #pragma acc atomic
    x = x + 1;
}


void C::fun(){

    init();    
    update_device();
    #pragma acc parallel present(k,g)
    {
        #pragma acc loop independent
        for (int i = 1; i <= N; i++){
            fun1();
            // fun2(k);
        }
    }
    #pragma acc kernels present(k,g)
    {   
        g = k;
    }

    update_host();    
    exit();
    std::cout<<g<<" "<<k<<std::endl;

}

int main(){
    C cl;
    cl.fun();
}

    • 1
    • 2
    • 3
    • 4
    • 5
    • 6
    • 7
    • 8
    • 9
    • 10
    • 11
    • 12
    • 13
    • 14
    • 15
    • 16
    • 17
    • 18
    • 19
    • 20
    • 21
    • 22
    • 23
    • 24
    • 25
    • 26
    • 27
    • 28
    • 29
    • 30
    • 31
    • 32
    • 33
    • 34
    • 35
    • 36
    • 37
    • 38
    • 39
    • 40
    • 41
    • 42
    • 43
    • 44
    • 45
    • 46
    • 47
    • 48
    • 49
    • 50
    • 51
    • 52
    • 53
    • 54
    • 55
    • 56
    • 57
    • 58
    • 59
    • 60
    • 61
    • 62
    • 63
    • 64
    • 65
    • 66
    • 67
    • 68
    • 69
    • 70
    • 71
    • 72
    • 73
    • 74
    • 75
    • 76
    • 77
    • 78
    • 79
    • 80
    • 81
    • 82
    • 83
    • 84
    • 85
    • 86
    • 87
    • 88
    • 89
    • 90
    • 91
    • 92
    • 93
    • 94

    输出：1000 0

    openACC调用math库（gcc）

    调用math的库在nvc++里没什么问题，gcc里可能会有问题。本人测试如下程序未发现报错：

    #include 
int main(){
    float f[100];
    #pragma acc parallel
    {
        #pragma acc loop
        for (int i = 0; i < 100; i++){
            float k = 1;
            f[i] = std::sin(double(k));
        }
    }
}

    • 1
    • 2
    • 3
    • 4
    • 5
    • 6
    • 7
    • 8
    • 9
    • 10
    • 11
    • 12

    编译指令：g++ -fopenacc -fcf-protection=none -fopt-info-optimized-omp -foffload=-lm test.cpp -o test

  • 相关阅读:
    java VisualVM工具连接远程服务和实践
    [JS]迭代器、可迭代对象、生成器详解
    docker镜像升级python3.5→python3.9（Ubuntu）
    java Optional操作
    网上最全的套接字socket
    情感丰富的文字
    WPF自定义控件与样式(3)-TextBox & RichTextBox & PasswordBox样式、水印、Label标签、功能扩展...
    特殊的转义字符—— \b 退格字符 ASCII 08
    如何远程控制别人电脑进行技术支持？
    web前端期末大作业——HTML+CSS+JavaScript仿王者荣耀游戏网站设计与制作
  • 原文地址：https://blog.csdn.net/qq_35975855/article/details/126978969