CUDA C编程（三十）OpenACC的使用

OpenACC是CUDA的一个补充编程模型，使用基于编译器指令的API，具有高性能、可编程性和跨平台可移植性。接下来将会介绍OpenACC的概念和方法，重点介绍CUDA和OpenACC之间的关系。
OpenACC的线程模型与CUDA的线程模型类似，但添加了一个并行的维度。OpenACC可以分为gang、worker和vector3个并行层次。在上层，gang类类似于CUDA线程块。一个gang可包含一个或多个执行的线程，在每个gang内部每个gang都包含一个或多个worker。在CUDA中，一个worker类似于线程中的一个线程束。在CUDA中，一个worker类似于线程中的一个线程束。每个worker都有一个向量宽度，由一个或多个同时执行指令的向量元素组成。每个向量元素都类似一个CUDA线程，因为它是一个单一的执行流。OpenACC和CUDA线程模型之间的主要区别在于，OpenACC在编程模型中直接指出了worker的概念(即线程束)，而在CUDA中并没有明确建立线程束。
OpenACC平台模型与CUDA类似，但它使用不同的术语和略有不同的抽象概念。OpenACC的目标是建立一个具有单线程的主机程序平台，在该主机程序中，将内核交付给多处理单元(PU),在此平台上，每个PU一次只运行一个gang。每个PU可以同时执行多个独立的并发执行线程(worker)。每个执行线程可以执行具有一定向量宽度的向量运行。每个worker里的并行以及一个跨向量操作的并行被称为向量并行。当在GPU上使用OpenACC时，一个PU就类似于一个SM。
根据任务是否通过gang、worker、vector并行执行，OpenACC执行被分成几种模式。现在，假设在一个OpenACC程序的并行计算区域中，创建了G个gang，其中每个gang包含W个worker，每个worker的向量宽度为V。那么，总共有G×W×V个执行线程处理这个并行区域。
当开始执行并行区域时，gang以gang冗余模式执行，这有利于在并行执行开始前对gang的状态进行初始化。在gang冗余模式中，每个gang的worker中只有一个活跃vector元素，其他worker和vector元素是闲置的，因此只有G个活跃的执行线程。此外，每个gang都执行相同的运算，所以在这个阶段没有通过gang并行任务。在CUDA中，gang冗余并行将作为执行相同计算的线程块里的一个线程来实现:

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

在OpenACC并行区域的某些地方，应用程序可能通过gang转换为并行执行。在这种情况下，程序以gang分裂模式执行。在分裂模式下，每个gang中仍然只有一个活跃地vector元素和一个活动的worker，但每个活跃的vector元素执行不同的并行区域。因此，该计算区域的任务被分散到各个gang中。在CUDA中，gang分裂模式将作为一个线程来实现，在这个线程里每个线程块处理分离的数据点。对于向量加法，在gang分裂模式下执行的CUDA内核如下所示：

__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];}
}

那么对于worker并行和vector并行呢？当在一个gang中只有一个活跃的worker时，程序处于单一worker模式。当worker中只有一个活跃的vector元素时，程序处于单一vector模式。因此，gang冗余模式和gang分裂模式也可以被称为单一worker模式和单一vector模式。
在worker分裂模式下，并行区域的工作被划分到多个gang和多个worker中。使用所有gang里的所有worker可以提供G×W路并行。在CUDA中，worker分裂模式通过每个线程束中的第一个线程来实现：

__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分裂模式下，工作任务在gang、worker和vector通道上进行划分，同时提供G×W×V路并行。这个模式与编写的CUDA内核模式最为相似。这些不同的OpenACC模式，使得一个应用程序的并行性可在代码的并行区域内进行动态调整。
当使用OpenACC时，由程序员用编译器指令指定并行区域，或是并行运行。编译器指令还可以指定使用何种类型的并行处理。编译器指令是一行源代码。用C/C++编写，开头为#pragma。OpenACC指令使用acc关键字作为唯一标识，这意味着所有OpenACC指令都是以#pragma acc开头的。
尽管编译器指令是程序源代码的一部分，但它们对编译器生成的可执行文件的影响是不一定的。如果编译器无法识别或不支持#pragma的这种类型，那么在编译时就会忽略#pragma的存在。目前，PGI、Cray和CAPS编译器都支持OpenACC指令。
使用OpenACC编译器能对程序的OpenACC指令做出解释，也能对源代码进行自动分析以自动生成加速器源代码。因此，只需要添加几行指令代码，就可以在GPU上自动执行应用程序了。如下是OpenACC向量加法的实现：

#pragma acc kernels
for(int i = 0; i < N; i++){C[i] = A[i] + B[i];
}

#pragma acc内核指令标志着下列顺序代码块在OpenACC加速器上符合执行的条件(如GPU)。如果构建源代码的编译器是支持OpenACC的，那么它会分析循环过程，决定在gang、worker和vector通道上并行执行的策略，然后自动生成内存拷贝、内核启动及在GPU上并行执行循环部分必需的内核代码。与此相反，在CUDA中的手动编码(没有CUDA统一内存)则要求：1)把循环体转换为CUDA_global_kernel，2)用cudaMalloc分配内存，3)用cudaMemcpy将数据拷贝到设备端，4)启动内核，5）将数据拷贝回主机端。有了OpenACC,所有这些都可以由#pragma来完成。
除了编译器指令，OpenACC也提供了一系列库函数。与OpenACC的编译器指令相比，这些函数可以实现与其相同的、互补的或独有的功能。

OpenACC 计算指令的使用
在OpenACC中，用计算编译程序指令来通知编译器是如何让一个代码块并行执行的。有两个相关的计算指令：#pragma acc kernels 和 #pragma acc parallel。

内核指令的使用
#pragma acc kernels采取了一种比#pragma acc parallel更自动化且编译器可驱动的方法。当在一个代码块中应用内核指令时。编译器会自动分析这个代码块中可并行的循环。当找到可并行的区域后，编译器可以在每个并行循环中使用任意配置的gang、worker和gang分裂宽度来对并行执行进行调度。也就是说，编译器会自动决定何时使用gang冗余模式、gang分裂模式、worker分裂模式等。在CUDA中，编译器会搜寻代码块，在CUDA内核中，这些代码块是由并行循环所执行的内核指令所修饰的。在内核块中不能并行的其他代码仍要执行，只是不以并行方式执行。
可以从Wrox.com下载的simple-kernels.c中的代码作为简单的示例。以下代码段是核心代码：

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

这个内核块包含两个可以并行的循环。如果有一个PGI编译器，那么就可以用以下命令进行编译：

    $ pgcc -acc simple-kernels.c -o simple-kernels

为PGI编译器添加-acc标志使其支持OpenACC,允许其在所提供的代码中识别任何带有 #pragma acc 的指令。同时强烈建议在PGI编译器中添加-Minfo = accel标志，以弄清楚自动并行化是如何实现的。在simple-kernels.c使用-Minfo = accel可得到以下输出：

现在，忽略那些有present_or_copyoyt和present_or_copy的语句。这些在后面会有所提及。使用36和39标记的行是源代码中每个循环开始的地方。在这两种情况下，OpenACC能自动找到可并行化的循环。它还为每个循环所用的并行化策略输出相关信息：#pragma acc loop gang, vector(128) /* blockIdx.x threadIdx.x */。信息表明，在一个有128个元素的向量宽度下，这两个循环都通过gang和vector完全并行化了。在CUDA中，这就对应一个128线程的线程块，需要尽可能多的块启动来并行执行循环迭代。
内核指令后面也可能会加一些有修饰其行为的选项，例如：#pragma acc kernels if(cond)，如果cond是false,那么应禁止代码块在OpenACC加速器上执行。如果在GPU上并行执行没有任何意义，在这种情况下，我们若想阻止，那么上述指令就非常有用，类似如下命令：#pragma acc kernels if( N < 128),在OpenACC中，对所有计算来说，在内核指令结束时有一个默认的等待命令。但是，如果使用async子句那么执行就不会被阻塞：#pragma acc kernels async(id),async子句接受一个可选的整型参数。这个整数给内核块指定的唯一ID，允许相同的整数ID在之后被用于测试或等待这个内核块的实现。如果没有给出这个ID,那么内核块仍异步执行，但是不会等待那个内核块执行完毕。例如，如果用如下命令创建一个内核块：#pragma acc kernels async(2),然后应用程序可以使用如下等待命令，等待完成与内核指令相关的计算：#pragma acc wait(3);或者通过调用库函数acc_async_wait:acc_async_wait(3);我们还可以使用一个空的等待指令来等待所有异步任务的完成：#pragma acc wait,或使用库函数acc_async_wait_all:acc_async_wait_all();
在CUDA中，使用整数ID的异步指令和函数，类似于使用cudaEvent_t来识别等待执行任务中的一个点。然后，对异步任务使用一个等待指令或阻塞函数，类似于使用cudaEventSynchronize函数阻塞某个事件。如果没有整数ID,等待行为类似于cudaDeviceSynchronize调用。
我们也可以在内核指令中添加一个wait子句，确保内核区域的执行在下列情形之前未启动：1)之前所有异步任务均已完成，2)与所提供的整数ID相关的任务都已经完成。将内核区域的异步和wait子句结合起来，可以链接到异步加速区域：

#pragma acc kernels async(0)
{...
}
#pragma acc kernels wait(0) async(1)
{...
}#pragma acc kernels wait(1) async(2)
{...
}
#pragma acc wait(2)

此外，OpenACC还支持检查异步计算是否是在没有阻塞的情况下完成的。这只能用库函数来实现。acc_async_test(int)检查所给出的ID内核是否已经执行完毕，而acc_async_test_all检测所有的异步命令是否都已经完成。如果所有的异步命令已经完成，就返回一个非零值；否则，返回零。
注意可以将多个子句进行组合。例如，下面的#pragma指令将一个内核区域标记为异步的，但只有在N>128的加速器上才能执行。对于这个例子，if子句应用于内核指令，但不是内核指令的异步性。#pragma acc kernels if(N>128) async
并行指令的使用
内核指令及其相关的子句对应用程序的加速来说是一个强大的工具，使用它们可以让我们对应用程序的实际执行有较少的控制。OpenACC编译器会自动分析代码并选择一个合适的并行策略，这只需要我们进行很少的参与。为了解决这个问题，OpenACC添加了另一个类似于内核的指令，但这个指令提供了更多的执行控制：#pragma acc并行指令。内核指令允许编译器将标记代码分组到尽可能多的加速器内核中，该内核中包含编译器认为所有必要的并行部分。在使用一个并行指令时，所有的gang和worker都在并行区域的开始位置启动，在末尾处停止执行。尽管编译器可以基于我们的指令在多种执行模式之间转换，但它不能调整在并行区域中间位置的并行维度。与CUDA中一样，这使我们可以完全掌握并行性的创建。
异步并行指令支持那些解释一部分内核指令的子句，如if子句、async子句和wait子句。此外，并行指令还支持以下子句：使用num_gangs(int)设置gang的数量，用num_workers(int)设置worker的数量，以及用vector_length(int)设置每个worker的向量宽度。我们应该熟悉在CUDA中配置线程块的数量和每个块中线程的数量，这里只不过多了一个维度。并行指令还支持reduction子句。一旦一个并行区域执行结束，reduction子句就会自动将每个gang的输出结合起来，处理成一个单一值进行输出，一旦并行指定完成就可以获得这个值。reduction子句需要一个运算和一个变量列表来实现，它们之间用冒号隔开：#pragma acc parallel reduction(op:var1,var2,...)。在并行区域的每个gang中每个变量都有一个变量副本var1,var2…，将其初始化为一个默认的、运算指定的初始值。当并行区域执行结束时，每个gang中的副本都执行运算op，并将运算结果作为最终结果进行输出。例如，如果我们想通过gang对变量result求和，可以使用如下命令：#pragma acc parallel reduction(+:result)。OpenACC支持各种简化运算符，包括+、*、max、min、&、|、^、&&以及||。
在CUDA中，reduction子句将通过在__shared__内存中存储一个标量来完成实现，从每个线程块不断更新它的值，并且在内核结束时使用原子操作将每个线程块写入的值进行结合。这比使用reduction子句需要更多的编程工作，但却使之更可控、更具有可自定义性(通过启动自定义原子操作)。
并行指令中也可以使用private和firstprivate子句。private和firstprivate子句运用变量列表。当使用private时，会为每个gang创建一个private型复制变量。只有该gang可以使用该变量的拷贝，因此该值的改变对其他gang或者主机应用程序是不可见的。例如，下面的代码段：

int a;
#pragma acc parallel private(a)
{a = ...;
}

并行区域中的每个gang将a的复制变量设为不同的值。这些值对其他gang或主机应用程序是不可见的。从概念上讲，我们可以认为它类似于CUDA中的__shared__内存变量。
firstprivate子句与private子句功能相同，但是要将每个gang中的private型变量的值初始化为主机上该变量的当前值。代码段如下：

int a = 5;
#pragma acc parallel firstprivate(a)
{...
}

并行区域中的每个gang会以a的值被设置为5的复制变量开始。对于每个gang来说，对a的任何更改都是private型。
循环指令的使用
并行指令的挑战是，需要我们为要加速的编译器明确标注并行性。并行区域总是以gang冗余模式开始的。执行并行模式之间的转换(如gang分裂模式或work分裂模式)需要对有更高并行期待水平的编译器有明确的指示。这是通过使用#pragma acc循环指令标记并行循环来完成的，我们可以直接对该循环所使用的执行模式进行操作。例如，我们可以用之前使用的并行指令和循环指令实现较早的simple-kerhcls示例，目的是将包含的循环标记为可并行的(如下所示)：

#pragma acc parallel{#pragma acc loopfor(i = 0; i < N; i++){C[i] = A[i] + B[i];}
#pragma acc loopfor(i = 0; i < N; i++){D[i] = C[i] + A[i];}
}

由于在这个例子中没有为循环指令添加子句，所以编译器可以自由使用它认为最优的任何循环调度。程序员也可以通过对循环指令添加gang、worker或vector子句显式地控制每一级的并行性。当在循环指令中添加了以上列出的一个或多个子句时，这个循环在各自的维度上就可以并行执行了。例如，考虑以下代码段：

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

这里，并行区域以gang冗余模式开始。当遇到循环指令时，由于gang子句的存在，执行会切换至gang分裂模式。
然而，循环指令不是仅在一个并行区域内有效。它也可以与内核指令相结合，为编译器标记并行循环，目的是将其变成加速器内核。然而，其子句的意义因上下文的不同而有所差别。下表列出了可应用于循环指令的子句，以及依赖于它们在并行内部或内核区域使用时含义的变化。

我们可以把并行或内核循环指令结合到一个pragma中：

#pragma acc parallel loop
for(i = 0; i < N; i++)
{...
}
#pragma acc kernels loop
for(i = 0; i < N; i++)
{...
}

这些只是扩展并行指令和内核指令的语法修改，后面紧跟着一个循环指令：

#pragma acc parallel loop
{#pragma acc loopfor(i = 0; i < N; i++){...}
}
#pragma acc kernels
{#pragma acc loopfor(i = 0; i < N; i++){...}
}

OpenACC 数据指令的使用
在编写OpenACC程序时，我们可能对数据转移问题毫不关心。但是，这样做会让OpenACC进行了很多不必要的通信从而使性能显著下降。接下来将介绍#pragma acc data指令是如何显式地在主机和OpenACC加速器之间进行通信的。

数据指令的使用
在OpenACC中，#pragma acc data被显式地用于在主机应用程序和加速器之间传输数据，类似于CUDA中的cudaMemcpy。与kernels和parallel指令类似，数据被应用到代码的某个区域。它定义了在该区域边界处必须进行的数据传输工作。例如，可以把一个变量标记为copyin。也就是说，可以将这个变量在该区域的起始位置传送给加速器，但最后不能传出。相反的，copyout是在数据区的末端将该变量传回主机端，但不能在该数据区域的起始位置将其传给加速器。
我们可以从Wrox.com中下载simple-data.c进行学习。simple-data.c是simple-parallel.c的扩展。核心代码如下：

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

所添加到#pragma acc data指令通知了编译器：只有A和B应该拷贝到设备端，只有C和D应该拷贝回来。这段代码还指明了传输数组的范围，在这种情况下，传输的应该是整个数组。在某些情况下，编译器能够推断出要复制的数组大小，这能略微简化代码：#pragma acc data copyin(A,B) copyout(C,D),以上修改的结果是，与没有使用数据指令的传输过程相比，要传输的字节数减少了一半。
除了数据指令，在执行过程中也可以用#pragma acc enter data和#pragma acc exit data来标记任意节点传入和传出加速器的数组。当编译器遇到enter data指令时，它会指明哪些数据应该复制到设备端。这些数据将继续留在设备端，直到编译器遇到将其传回的exit data指令或者程序终止执行。当在async子句和wait子句相结合时，enter data指令和exit data指令能够发挥最大的作用。注意，data指令不支持async子句和wait子句。
当把async子句应用到enter data和exit data指令中，它会创建数据传入或传出加速器的异步传输任务，类似于cudaMemcpyAsync。正如在CUDA中异步拷贝是很有用的一样，作为一种重叠计算和通信的方法，在OpenACC中他也是很有用的。当把wait子句应用到enter data和exit data指令中时，它的作用与在kernels指令或parallel指令中一样：通信指令要等待其他异步任务结束后再执行。需要注意的是通信指令(即 enter data和exit data指令)可以使用async和wait子句来交互异步计算任务(即kernels指令和parallel指令)，反之亦然。参考下面的代码：

int *A = init_data(N);
int *B = init_more_data(N);do_some_heavy_work(C);#pragma acc data copyin(B[0:N]) copyout(A[0:N])
{#pragma acc kernels{for(i = 0; i < N; i++){A[i] = do_work(B[i]);}}
}
do_lots_more_work(D);

这里，data指令将B传送到设备端，然后将A传输回来。但是，由于data指令使用的是同步传输，所以这个应用程序必须停止并且等待要传送的潜在的大数组。如果用async子句代替enter data和exit data指令，那么do_some_heavy_work和do_lots_more_work中的通信开销可以被隐藏：

int *A = init_data(N);
int *B = init_more_data(N);#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)do_lots_more_work(D);#pragma wait(2)

这里，enter data指令用于将B异步传输到设备端，与do_some_heavy_work的作用相同。然后，内核指令使用wait子句确保B的异步拷贝已经完成，并启动一个异步计算任务。随后，exit data指令用来将A异步地传回，与do_lots_more_work的作用相同，但要先等待内核区域执行结束。最后，必须使用一个wait指令确保A已经传送回主机端。
与kernels指令和parallel指令类似，data指令支持和共享多种子句。下表列出了data指令所支持的子句以及支持这些子句的指令：

为内核指令和并行指令添加data子句
通常情况下，由于输入和输出数据都是在计算区域之前或之后进行传输的，所以数据指令与计算指令紧密相关。尽管可以对每个任务使用独立的指令，但OpenACC还是支持使用计算指令上的data子句来简化代码。
例如，考虑前面提到的simple-data.c中的核心逻辑：

#pragma acc data copyin(A[0:N],B[0:N]) copyout(C[0:N], D[0:N]){#pragma acc parallel{...}

可以为并行指令添加一个data子句而不是保留两个分离的编译器指令：

#pragma acc parallel copyin(A[0:N],B[0:N]) copyout(C[0:N], D[0:N])

这种改变简化了源码，并且更容易看到的是，并行区域和数据传输是相关联的。内核指令和并行指令都支持上表中列举的copy,copy in,copy out,create,present,present_ or_copy,present_or_copyin,present_or_copyout,present_or_create和deviceptr子句。
OpenACC 运行时 API
除了编译器指令，OpenACC也提供了一个函数库。在介绍async和wait子句时已经有所提及：函数acc_async_wait、acc_async_wait_all、acc_async_test和acc_async_test_all都是OpenACC运行时API的一部分。使用OpenACC运行时API中的函数要求添加头文件openacc.h。
许多用OpenACC编写的程序完全可以不使用OpenACC运行时API，因为在许多情况下，编译器指令就可以提供所需要的功能。不过，仍然有一些由运行时API提供的操作是OpenACC编译器指令不能提供的。
OpenACC运行时API函数可分为4个部分：设备管理、异步控制、运行时初始化和内存管理。接下来着重介绍一些比较有用的函数。
设备管理函数允许我们显式控制使用哪个加速器或加速器类型来执行OpenACC计算区域。许多设备管理函数使用acc_device_t类型，这是一个枚举类型，它代表了由OpenACC实现所支持的不同设备类型。最低限度，所有的OpenACC实现必须支持acc_device_none、acc_device_default、acc_device_host和acc_device_not_host类型，当然也可支持其他类型，例如，PGI 14.4支持以下设备类型：

typedef enum{acc_device_none = 0,acc_device_default = 1,acc_device_host = 2,acc_device_not_host = 3,acc_device_nvidia = 4,acc_device_radeon = 5,acc_device_xeonphi = 6,acc_device_pgi_opencl = 7,acc_device_nvidia_opencl = 8,acc_device_opencl = 9
}acc_device_t

下表列出了一些设备管理函数：

异步控制函数允许我们检查或等待异步操作的执行状态。异步操作包括使用并行指令和内核指令创建的异步计算和使用OpenACC数据创建的异步通信。下表给出了一些异步控制函数：

运行初始化函数用来初始化或管理OpenACC的内部状态。下表给出了一些运行时初始化函数。如果acc_init没有被OpenACC的应用程序显式调用，那么运行时初始化作为应用程序的第一个OpenACC操作来执行。

内存管理函数用于管理加速器内存分配以及在主机和加速器之间的数据传输。因此，在许多情况下，他们的功能与OpenACC的数据指令和子句相同。下表中给出了一些内存管理函数：

OpenACC 和 CUDA 库的结合
尽管CUDA和OpenACC是相互独立的编程模型，但它们仍可在同一应用程序中使用。这样就需要更改应用程序的编程方式，同时必须使用deviceptr子句来实现CUDA和OpenACC之间的数据共享。示例程序可从Wrox.com上下载cuda-openacc.cu,cuda-openacc.cu可以分成以下几步：
1.使用cudaMalloc为矩阵分配设备内存，使用curandCreateGenerator和cublasCreate为cuRAND和CUBLAS库创建句柄；
2.cuRAND库中的curandGenerateUniform函数产生的随机数据对设备内存中的输入矩阵进行填充；
3.使用OpenACC指令，在GPU上并行执行两个矩阵间的乘法；
4.cublasSasum用于计算输出矩阵中所有元素的总和；
5.使用cudaFree释放设备内存。
用PGI OpenACC编译器对cuda-openacc.cu进行编译，命令如下：$ pgcpp -acc cuda-openacc.cu -o cuda-openacc -Minfo=accel \ -L${CUDA_HOME}/lib64 -lcurand -lcublas -lcudart,注意这里的C++编译器pgcpp与CUDA库是兼容的。保留-acc参数以支持OpenACC,-Minfo=accel用来表示在OpenACC并行计算区域内的诊断信息。此外，将CUDA库的路径应添加到编译器的库路径中，以便编译器可以找到cuRAND,cuBLAS和CUDA运行时函数的定义。此命令的输出如下：

需要注意的是内核是在70和72行的循环上产生的，包括使用gang并行的外层循环和使用worker及vector并行的内部循环。
接下来可能有两点陌生的知识：1.deviceptr子句是在并行指令中使用的。deviceptr允许一个应用程序显式地分配和管理自己的设备内存，然后直接将其传递给OpenACC计算区域。在这种情况下，在用cuRAND填充数据之前，cuda-openacc应用程序会用cudaMalloc显式的分配自己的设备内存。然后，使用deviceptr给OpenACC内核提供相同设备内存的直接访问权，而不必在使用copyin之前将其传回主机。deviceptr是允许OpenACC和其他GPU编程框架相结合的一个关键组成部分。2.示例cuda-openacc.cu中也使用了cublasSetPointerMode来衡量cuBLAS函数是否使用主机或设备的指针来返回标量结果。在这种情况下，cublasSasum返回一个标量结果作为它的最后一个参数。最初，在设备端对输出矩阵的行进行求和时，cublasSetPointerMode用来将模式设置为CUBLAS_POINTER_MODE_DEVICE。当执行最终的跨行求和时，模式切换为CUBLAS_POINTER_MODE_HOST,返回地址为主机应用程序地址空间中的一个变量。

OpenACC 小结
OpenACC是一个灵活的、易于使用的、高性能的编程模型，它在许多方面对CUDA和CUDA库进行了补充。与CUDA库i相比，OpenACC的使用更加灵活，允许我们使用C语言编写自己的计算函数。与CUDA相比，OpenACC的使用更为方便，在通信和计算方面比CUDA C需要更少的人为参与。

CUDA C编程（三十）OpenACC的使用相关推荐

5岁自学python编程-三十岁了还可以学编程吗？只要你想，只要你做，什么时候都不晚...
文章有点长,希望你可以看完.如果能给迷茫的朋友一点勇气,我不胜荣幸. 做一件事最好的时间是十年前,其次是现在. 如果想学编程却苦于不知道从何开始,也没有系统的学习路径,那就来加入我们一起学习在这里不 ...
CUDA学习（三十六）
内置的矢量类型: char, short, int, long, longlong, float, double: 这些是从基本整数和浮点类型派生的矢量类型. 它们是结构,第一,第二,第三和第四个组件 ...
java并发编程(三十五)——公平与非公平锁实战
前言在 java并发编程(十六)--锁的七大分类及特点一文中我们对锁有各个维度的分类,其中有一个维度是公平/非公平,本文我们来探讨下公平与非公平锁. 公平|非公平首先,我们来看下什么是公平锁和非 ...
cmdb python 采集虚拟机_Python编程(三十四)：CMDB后台管理、封装自定义JS组件、前端td标签定制显示内容及属性...
一. CMDB后台管理 CMDB管理主要分为采集资产.API接口.后台管理.这里主要介绍CMDB后台管理. - 采集资产 - API - 后台管理- 资产列表- 业务线列表- 用户列表- 组列表... ...
Python编程基础：第三十九节面向对象编程Object Oriented Programming
第三十九节面向对象编程Object Oriented Programming 前言实践前言到目前为止我们都是函数式编程,也即将每一个功能块写为一个函数.其实还有一种更常用的编程方式被称为面向对 ...
Python编程基础：第三十八节问答游戏Quiz Game
第三十八节问答游戏Quiz Game 前言实践前言我们这一节还是对之前学习内容的一个综合运用,主要涉及到函数编程.字典以及列表的使用.条件语句.循环结构等等.通过本节的学习读者可以检验之前内容 ...
Python编程基础：第三十六节模块Modules
第三十六节模块Modules 前言实践前言我们目前所有的代码都写在一个文档里面.如果你的项目比较大,那么把所有功能写在一个文件里就非常不便于后期维护.为了提高我们代码的可读性,降低后期维护的成 ...
实现TeX的算法：回首编程技术的过去三十年
Glenn Vanderburg是LivingSocial的工程主管,在最近的ClojureConj会议上,对他使用Clojure实现TeX算法的工作做了很有趣的叙述.在那个过程中,他发现在过去三十年 ...
三十天学不会TCP,UDP/IP网络编程-IP头格式祥述
我又来了,这篇文章还是来做(da)推(guang)介(gao)我自己的!俗话说事不过三,我觉得我下次得换个说法了,不然估计要被厌恶了,但是我是好心呐,一定要相信我纯洁的眼神.由于这两年接触到了比较多的 ...

CUDA C编程（三十）OpenACC的使用

CUDA C编程（三十）OpenACC的使用相关推荐

最新文章

热门文章