OpenAcc的使用

安装cuda驱动

  1. 更新源
sudo add-apt-repository ppa:graphics-drivers/ppa
sudo apt update
  1. 用以下指令找到recommend的驱动
ubuntu-drivers devices
  1. 安装
sudo apt install nvidia-driver-XXX

安装后运行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}}

安装后重开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

然后在.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

重开terminal运行 nvc++有结果

安装gcc-11

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

编译OpenAcc

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

nvc++ testACC.cpp -o test -acc

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

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

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

#include <iostream>
#include <ctime>
#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

}

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

nvprof ./test

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

#include <iostream>
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;
}

openACC调用函数(gcc)

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

方法一:传形参

#include <iostream>
#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();
}

输出:1000 1000

方法二:直接调用

#include <iostream>
#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();
}

输出:1000 0

openACC调用math库(gcc)

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

#include <cmath>
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));
        }
    }
}

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

  • 1
    点赞
  • 7
    收藏
    觉得还不错? 一键收藏
  • 打赏
    打赏
  • 0
    评论

“相关推荐”对你有帮助么?

  • 非常没帮助
  • 没帮助
  • 一般
  • 有帮助
  • 非常有帮助
提交
评论
添加红包

请填写红包祝福语或标题

红包个数最小为10个

红包金额最低5元

当前余额3.43前往充值 >
需支付:10.00
成就一亿技术人!
领取后你会自动成为博主和红包主的粉丝 规则
hope_wisdom
发出的红包

打赏作者

萧易风船长

你的鼓励将是我创作的最大动力

¥1 ¥2 ¥4 ¥6 ¥10 ¥20
扫码支付:¥1
获取中
扫码支付

您的余额不足,请更换扫码支付或充值

打赏作者

实付
使用余额支付
点击重新获取
扫码支付
钱包余额 0

抵扣说明:

1.余额是钱包充值的虚拟货币,按照1:1的比例进行支付金额的抵扣。
2.余额无法直接购买下载,可以购买VIP、付费专栏及课程。

余额充值