OpenACC 编程例子
小小河
2013 年 1 月 13 日
1. 条件编译宏变量_OPENACC
这个宏是可以当做整型变量输出，无法当作字符串输出
/* macro.c */
#include
int main()
{
printf("_OPENACC = %d\n", _OPENACC);
return 0;
}
程序输出为
_OPENACC = 20111111
这个宏的定义方式等价为
#define _OPENACC 20111111
2. 有数据依赖的 kernels loop 导语
#include 1
// export PGI_ACC_NOTIFY=1 2
// export PGI_ACC_TIME=1 3
// compile: pgcc -acc -Minfo t3-1.c 4
// run: ./t3-1.exe 5
6
#define n 1000 7
8
int main() 9
{ 10
int i; 11
int *a, *b, *c; 12
a = (int*)malloc(n*sizeof(int)); 13
b = (int*)malloc(n*sizeof(int)); 14
c = (int*)malloc(n*sizeof(int)); 15
for(i = 0; i < n; i++) 16
{ 17
a[i] = 0; 18
b[i] = c[i] = i; 19
} 20
#pragma acc kernels loop 21
for(i=0; i 2
#define n 1000 3
int main() 4
{ 5
int i; 6
int *restrict a, *restrict b, *restrict c; 7
a = (int*)malloc(n*sizeof(int)); 8
b = (int*)malloc(n*sizeof(int)); 9
c = (int*)malloc(n*sizeof(int)); 10
for(i = 0; i < n; i++) 11
{ 12
a[i] = 0; 13
b[i] = c[i] = i; 14
} 15
#pragma acc kernels loop 16
for(i=0; i
#define n 1000
int main()
{
int i;
int *a, *b, *c;
a = (int*)malloc(n*sizeof(int));
b = (int*)malloc(n*sizeof(int));
c = (int*)malloc(n*sizeof(int));
for(i = 0; i < n; i++)
{
a[i] = 0;
b[i] = c[i] = i;
}
#pragma acc kernels loop independent
for(i=0; i
#define n 1000
int main()
{
int i, a[n], b[n], c[n];
for(i = 0; i < n; i++)
{
a[i] = 0;
b[i] = c[i] = i;
}
#pragma acc kernels loop
for(i=0; i 1
#define N 1000 2
3
int main() 4
{ 5
int a[N],i; 6
for(i=0; i< N; i++) a[i] = -1; 7
8
#pragma acc kernels loop 9
for(i=0; i
#define N 1000
int main()
{
int a[N];
#pragma acc declare copy(a)
int i;
for(i=0; i< N; i++) a[i] = -1;
#pragma acc kernels loop
for(i=0; i
说明
关于失联党员情况说明岗位说明总经理岗位说明书会计岗位说明书行政主管岗位说明书
设备上的数据没有复制到主机端，需要使用 update 导语同步一下数据。将代码中
//#pragma acc update host(a)
的注释去掉，就以得到正确的结果
PGI$ ./datreg2.exe
a[1] = 2
8. Declare 导语执行时机
#include
#define N 1000
int main()
{
int a[N];
#pragma acc declare copy(a)
int i;
for(i=0; i< N; i++) a[i] = -10;
#pragma acc kernels loop
for(i=0; i
检测
工程第三方检测合同工程防雷检测合同植筋拉拔检测方案传感器技术课后答案检测机构通用要求培训
declare 导语复制数据的时机。如果数据复制发生在可执行语句之后，
那么最终的输出是 a[1] = 10；如果数据复制发生在可执行语句之前，那么 declare copy(a)的
执行结果就是设备端 a 所有元素的值皆为 0，经过 kernels 构件，a 所有元素的值被更新为
20，程序的最终输出是 a[1] = 20。
编译运行此程序，得到输出
PGI$ ./datreg2.exe
a[1] = 20
从而得知，declare 导语的执行时机为可执行语之前。
9. acc_device_t 类型枚举变量
在 openacc.h 中有如下定义
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_t;
使用运行时库函数 acc_get_num_devices 时，可以设定的参数有 5 个，查询 nVidia 设备时，
使用 acc_get_num_devices(acc_device_nvidia)
比 openacc
标准
excel标准偏差excel标准偏差函数exl标准差函数国标检验抽样标准表免费下载红头文件格式标准下载
多定义了一个 acc_device_nvidia，其它 PGI 编译器目前仅支持 nVidia
编译器。
10. Parallel 区域不支持 IO 语句
#include 1
int main() 2
{ 3
int a[10], i; 4
for(i=0; i<10; i++) 5
a[i] = i+10; 6
#pragma acc parallel 7
printf("a[0]=%d\n",a[0]); 8
return 0; 9
} 10
编译时将会报错
8, Accelerator restriction: unsupported call to 'printf'
这是因为不能在加速器上直接输入输出，所有输出操作必须绕道主机。
11. 在 kernels 区域中使用分支语句
program t1
implicit none
integer a(1000), b(1000)
integer i
a(:) = 1
do i = 1,1000
b(i)= i
enddo
!$acc kernels
do i = 1,1000
if(b(i)== 1) then
a(i) = 100
else
a(i) = 0
endif
enddo
!$acc end kernels
print*, "a = ", a(1:4)
endprogram
Kernels 区域的 if-else-endif 语句导致两个执行路径，在设备上运行时，会破坏 GPU 的合并访
问，效率降低至原来的 50%，计算速度与 IO 速度都下降。如果循环次数特别多（远超 1000
次），那么加速比仍然很可观。
12. 数据子语
1. 子数组范围指定方式
子数组指定下标范的
方法
快递客服问题件处理详细方法山木方法pdf计算方法pdf华与华方法下载八字理论方法下载
有多种形式。
对fortran, 假设定义了
integer A(100,10);
那么可以下列多种形式：
arr(下界 1:上界 1,下界 2：上界 2)，标准形式，同时给出所有维度的上界和下界；
A(:,:) 包含 A的所有元素，与 A等价；
A(:,1) 包含 A的第 1列中所有元素；
A(1:,1)包含 A 的第 1 列中所有元素，因为省略了第一个维度的上界，同时知道这个数组
声明时的上界为 100，所以认为上界就是 100；它等价于 A(1:100,1)和 A(:,1);
A(:100,1)包含 A的第 1列中所有元素，这里省略了第 1个维度的下界，编译器就认为下
界是 1；它等价于 A(1:100,1)和 A(:,1);
A(:,:2)等价于 A(:,1:2)；
program t2 1
implicit none 2
3
integer A(100,10) 4
integer lin, col 5
A = 0 6
print*, "A(1:2,1:2)=" , A(1:2,1:2) 7
8
!$acc kernels loop copy(A) 9
do lin = 1, 100 10
A(lin,1) = 1 11
enddo 12
print*, "A(1:2,1:2)=" , A(1:2,1:2) 13
14
!$acc kernels loop copy(A(:,:)) 15
do lin = 1, 100 16
A(lin,1) = 2 17
enddo 18
print*, "A(1:2,1:2)=" , A(1:2,1:2) 19
20
! only copy the first column 21
!$acc kernels loop copy(A(:,1)) 22
do lin = 1, 100 23
A(lin,1) = 3 24
enddo 25
print*, "A(1:2,1:2)=" , A(1:2,1:2) 26
27
! only copy the first column 28
!$acc kernels loop copy(A(1:,1)) 29
do lin = 1, 100 30
A(lin,1) = 4 31
enddo 32
print*, "A(1:2,1:2)=" , A(1:2,1:2) 33
34
! only copy the first column 35
!$acc kernels loop copy(A(:100,1)) 36
do lin = 1, 100 37
A(lin,1) = 5 38
enddo 39
print*, "A(1:2,1:2)=" , A(1:2,1:2) 40
41
! only copy the first column 42
!$acc kernels loop copy(A(:,:2)) 43
do lin = 1, 100 44
A(lin,1) = 6 45
enddo kernels loop copy(A(:,:2)) 46
print*, "A(1:2,1:2)=" , A(1:2,1:2) 47
48
endprogram 49
编译信息
PGI$ pgf90 -acc -Minfo d2.f90
t2:
6, Memory zero idiom, array assignment replaced by call to pgf90_mzero4
9, Generating present_or_copy(a(:,:))
Generating compute capability 1.0 binary
Generating compute capability 2.0 binary
10, Loop is parallelizable
Accelerator kernel generated
10, !$acc loop gang, vector(128) ! blockidx%x threadidx%x
CC 1.0 : 7 registers; 24 shared, 4 constant, 0 local memory bytes
CC 2.0 : 9 registers; 0 shared, 40 constant, 0 local memory bytes
15, Generating present_or_copy(a(:,:))
Generating compute capability 1.0 binary
Generating compute capability 2.0 binary
16, Loop is parallelizable
Accelerator kernel generated
16, !$acc loop gang, vector(128) ! blockidx%x threadidx%x
CC 1.0 : 7 registers; 24 shared, 4 constant, 0 local memory bytes
CC 2.0 : 9 registers; 0 shared, 40 constant, 0 local memory bytes
22, Generating present_or_copy(a(:,:1))
Generating compute capability 1.0 binary
Generating compute capability 2.0 binary
23, Loop is parallelizable
Accelerator kernel generated
23, !$acc loop gang, vector(128) ! blockidx%x threadidx%x
CC 1.0 : 7 registers; 24 shared, 4 constant, 0 local memory bytes
CC 2.0 : 9 registers; 0 shared, 40 constant, 0 local memory bytes
29, Generating present_or_copy(a(0:,:1))
Generating compute capability 1.0 binary
Generating compute capability 2.0 binary
30, Loop is parallelizable
Accelerator kernel generated
30, !$acc loop gang, vector(128) ! blockidx%x threadidx%x
CC 1.0 : 7 registers; 24 shared, 4 constant, 0 local memory bytes
CC 2.0 : 9 registers; 0 shared, 40 constant, 0 local memory bytes
36, Generating present_or_copy(a(:,:1))
Generating compute capability 1.0 binary
Generating compute capability 2.0 binary
37, Loop is parallelizable
Accelerator kernel generated
37, !$acc loop gang, vector(128) ! blockidx%x threadidx%x
CC 1.0 : 7 registers; 24 shared, 4 constant, 0 local memory bytes
CC 2.0 : 9 registers; 0 shared, 40 constant, 0 local memory bytes
43, Generating present_or_copy(a(:,:2))
Generating compute capability 1.0 binary
Generating compute capability 2.0 binary
44, Loop is parallelizable
Accelerator kernel generated
44, !$acc loop gang, vector(128) ! blockidx%x threadidx%x
CC 1.0 : 7 registers; 24 shared, 4 constant, 0 local memory bytes
CC 2.0 : 9 registers; 0 shared, 40 constant, 0 local memory bytes
运行结果
PGI$ ./d2.exe
A(1:2,1:2)= 0 0 0 0
A(1:2,1:2)= 1 1 0 0
A(1:2,1:2)= 2 2 0 0
A(1:2,1:2)= 3 3 0 0
A(1:2,1:2)= 4 4 0 0
A(1:2,1:2)= 5 5 0 0
A(1:2,1:2)= 6 6 0 0
由运行结果可以看出，设上的赋值语都被正确执行。
对 C 和 C++，标准方法为 arr[下界:长度]
如果声明
int A[100];
那么，A[:2]等价于 A[0], A[1]，因为下界省略，默认为 0；A[2:]等价于 A[2:98]，表示元素 A[2],
A[3],...,A[99]，因长度省略了，默认长度为声明的上界减去下界，即 100-2.
验证小程序
#include
int main()
{
int A[100];
int i;
#pragma acc kernels loop copy(A)
for(i=0; i<100; i++)
A[i] = 1;
printf("A[0:1]=%d\t%d\n",A[0],A[1]);
#pragma acc kernels loop copy(A[:])
for(i=0; i<100; i++)
A[i] = 2;
printf("A[0:1]=%d\t%d\n",A[0],A[1]);
#pragma acc kernels loop copy(A[:2])
for(i=0; i<2; i++)
A[i] = 3;
printf("A[0:2]=%d\t%d\t%d\n",A[0],A[1],A[2]);
}
2. C 中动态分配数组中指定子数组
对动态分配的数组，需要显式地指明数组的长度。因为数组是动态分配的，编译器难以确定
数组的上界，从而不能算出默认长度。
int N = 1000;
int restrict* p = (int*)malloc(N*sizeof(int));
如果在数据子语中使用 copy(p)、copy(p[1:])、copy(p[:])等，编译器都会报错，错误信
息类似为 Cannot determine bounds for array p；copy(p[:N/2])无报错。
总结起来，动态分配的数组，可以省略下界，但不能省略长度。
验证小程序
#include 1
int main() 2
{ 3
int N = 1000; 4
int restrict* p = (int*)malloc(N*sizeof(int)); 5
int i; 6
7
for(i=0; i
书
关于书的成语关于读书的排比句社区图书漂流公约怎么写关于读书的小报汉书pdf
面同意，禁止任何形式的商业使用。商业使用形式包括但不限于出版、复制、
传播、展示、引用、编辑。
4. 本文档允许以学术研究、技术交流为目
浙公网安备 33021202002483