CUDA 学习笔记(六)

前言

本小节介绍的是OpenACC,原书中写到的介绍是:OpenACC is a complementary programming model to CUDA that uses a compiler directives-based API designed for high performance, programmability, and portability across many platforms. 通篇读下来后,博主认为该工具主要的作用在于使用简单的语法实现CUDA底层调用gpu硬件的加速,可以理解为对CUDA编程的封装(?

想要具体学习OpenACC可以去官网查看他的文档及API说明:https://www.openacc.org/resources,博主仅针对原书的介绍进行简单的总结

所有涉及到的代码可从博主的代码仓库中查看:https://github.com/palLiHua/cuda

第八章 CUDA相关库及OpenACC

8.4 OpenACC介绍

概述

OpenACC的线程模型包括gang,worker,vector,类比于CUDA编程中的blocks,warp,thread.与CUDA最大的区别是OpenACC显示暴露了warp,而在CUDA编程中代码逻辑层面是看不到warp的。博主这里说明得可能不够清晰,但学过前面前三章的童鞋应该能明白我在说什么,原书中的介绍是:OpenACC exposes the concept of workers (that is, warps) directly in the programming model, whereas in CUDA you do not explicitly create warps of threads.

OpenACC执行的模式主要有三种,假设有G个gangs,W个workers,每个vector有V个elements,三种模式及其运行情况如下:

1. gang-redundant mode

这种模式下每一个gang中只有一个worker工作,每一个worker又只有一个element工作,即共有G个线程在同时运行。类比于CUDA编程的话,就是创建了G个block,但每个block只有一个thread工作,等价于如下代码:

__global__ void kernel(...) { 
    if (threadIdx.x == 0) 
    {
        foo(); 
    }
}

2. gang-partitioned mode

这种模式下线程工作模式同gang-redundant相同,即同时有G个线程并行运行,那就有疑问了,和gang-redundant有什么区别呢?原文解释是:In gang-partitioned mode, there is still only one active vector element and one active worker in each gang, but each active vector element is executing a different piece of the parallel region.

给出的代码示例是:

__global__ void kernel(int *in1, int *in2, int *out, int N) 
{ 
    if (threadIdx.x == 0) 
    {
        int i;
        for (i = blockIdx.x; i < N; i += gridDim.x) 
        {
            out[i] = in1[i] + in2[i]; 
        }
    } 
}


查了一下资料以及问询GPT,个人认为和gang-redundant的区别体现在gang-redundant是所有线程做着同样的工作及处理同样的数据,换句话说就是在重复性工作,没有任何意义(可能?),而gang-partitioned则是每个gang在处理着不同的数据。

3. worker-single mode

在该模式下又可以分为worker-partitioned mode和vector-partitioned mode。有了上面的了解,不难理解worker-partitioned mode就是每个gang下所有worker都开始并行运行了,因此总共会有G*W个线程同时工作,等价于:

__global__ void kernel(int *in1, int *in2, int *out, int N) 
{ 
    if (threadIdx.x % warpSize == 0) 
    {
        int warpId = threadIdx.x / warpSize;
        int warpsPerBlock = blockDim.x / warpSize;
        int i;
        for (i = blockIdx.x * warpsPerBlock + warpId; i < N; 
                            i += gridDim.x * warpsPerBlock)
        { 
            out[i] = in1[i] + in2[i];
        } 
    }
}

那么在vector-partitioned mode下就会有G*W*V个线程同时执行。

使用OpenACC计算指令

主要可以使用#pragma acc kernels和#pragma acc parallel两种方式以启用gpu加速。

#pragma acc kernels

使用该指令通常会使coder更轻松,包括数据传输、并行程度均有编译器自行决定,但代价自然就是性能的下降。具体的使用方法可以参考acc.c,相较于CUDA需要自己手写核函数,OpenACC仅需要一行代码就可以实现gpu加速,非常便捷。使用的编译器是PGI,博主下载的是英伟达提供的HPC,可从https://developer.nvidia.com/hpc-sdk处下载,使用指令pgcc -acc -Minfo=accel acc.c -o test进行编译,其中-Minfo=accel可以用于查看编译的一些细节,顺利的话你会得到类似于如下的输出:

从红框圈出部分不难看出,两处循环使用了gpu进行加速,前两行代码也不难猜出,要在gpu上运行的话,它自动帮我们完成了数据host to device以及device to host传输的工作。

使用这种方式还可以加上if条件语句以及async(id)用于异步运行,具体用法参考官方文档及例子。需要值得注意的一点是,if条件不满足情况下,使用pgcc -acc -Minfo=accel编译仍会有上图中的输出,但具体有没有使用gpu加速其实是在运行过程中才决定的。

#pragma acc parallel

使用parallel的好处是你可以决定程序运行的并发程度,例如使用num_gangs(int)指定gang的数量,同样也可以加入if、async、wait等指令,与kernels最大的区别就是你能控制的部分更多了。在执行循环时需要加上#pragma acc loop,默认情况下是gang-redundant mode,如果要增加并发度,那么就要再加上gang用以表明使用gang-partitioned mode,代码如下:

#pragma acc parallel {
    int b = a + c; 
#pragma acc loop gang
    for (i = 0; i < N; i++) 
    { 
        ...
    } 
}

那么同理加上worker/vector就能切换至worker-partitioned mode或vector-partitioned mode

使用OpenACC数据指令

上面介绍的过程中,我们都默认数据传输是自动进行的,该小节介绍的就是我们如何手工控制数据的传输。有两种方式,第一种是使用data copyin/ data copyout,具体如何使用可参考data.c,这里仅贴出核心代码

#pragma acc data copyin(A[0:N], B[0:N]) copyout(C[0:N], D[0:N])
    {
#pragma acc parallel
        {
#pragma acc loop
            for (i = 0; i < N; i++)
            {
                C[i] = A[i] + B[i];
            }

#pragma acc loop
            for (i = 0; i < N; i++)
            {
                D[i] = C[i] * A[i];
            }
        }
    }

这里要提到的一点是,为写代码方便,上述代码等价于:

#pragma acc parallel copyin(A[0:N], B[0:N]) copyout(C[0:N], D[0:N])
    {
#pragma acc loop
        for (i = 0; i < N; i++)
        {
            C[i] = A[i] + B[i];
        }

#pragma acc loop
        for (i = 0; i < N; i++)
        {
            D[i] = C[i] * A[i];
        }
    }

这种方式存在的一个问题是,数据传输会block程序的进行,即不能异步进行,那么这时候就需要第二种方式闪亮登场。

在第二种方式下,仅需要把指令改为#pragma acc enter data和#pragma acc exit data,如果需要异步进行,则仅需要加上async(int),可以参考原书中的这段代码:

#pragma acc enter data copyin(B[0:N]) async(0)

do_some_heavy_work(C);

#pragma acc kernels async(1) wait(0)
{
    for (i = 0; i < N; i++) 
    {
        A[i] = do_work(B[i]);
    }
}

#pragma acc exit data copyout(A[0:N]) async(2) wait(1)

#pragma wait(2)

OpenACC Runtime API

除了上述介绍的compiler directives,OpenACC还提供了一系列函数指令,包含在openacc.h头文件中,这里不作解释,可自行查看文档

OpenACC和CUDA一同使用

CUDA编程的优势在于能自主控制的部分更多,而OpenACC自然在于其编程的便利性,这里就以一个例子来介绍下如何同时使用这两个工具,代码可参考cuda_openacc.cu,这个例子中使用curand产生随机的两个矩阵,然后使用3个for循环进行矩阵相乘操作(使用OpenACC),最后将生成的矩阵进行每个元素加和的操作(使用cublas库的cublasSasum先进行每行求和,最后再将vector加和)

在编译的时候,原书给出的指令是pgcpp -acc cuda-openacc.cu -o cuda-openacc -Minfo=accel  -L${CUDA_HOME}/lib64 -lcurand -lcublas –lcudart,但博主在执行过程中发现根本没有pgcpp,查看了一下英伟达提供的指令:

那就尝试一下pgc++吧,执行后的结果报错:

显示找不到curand头文件,之前使用nvcc时都可以链接,怎么到这就不行了?只能google一下,有个2015的回答是这样:

ok,那就加上-Mcuda试试:

感觉可行的样子,换成-cuda结果如下:

终于成功,看下输出,有个比较有意思的点:

43, #pragma acc loop seq

        Generating implicit reduction(+:sum)

如果看过代码的童鞋会发现我在第43行的for循环中并没有加上编译指令,编译器自动帮我选择了合适的编译指令,reduction什么作用可以查看文档或者原书371页

小结

这一小节也只是简单介绍了下OpenACC,总体来说整个第八章都是概述性的介绍,所以博主也只是挑了些重点写,实际需要哪个库还是需要查看文档。下一章进入多gpu编程,希望自己还能坚持下去,也请各位点个赞叭

评论
添加红包

请填写红包祝福语或标题

红包个数最小为10个

红包金额最低5元

当前余额3.43前往充值 >
需支付:10.00
成就一亿技术人!
领取后你会自动成为博主和红包主的粉丝 规则
hope_wisdom
发出的红包
实付
使用余额支付
点击重新获取
扫码支付
钱包余额 0

抵扣说明:

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

余额充值