分类: 高性能计算
2016-02-11 16:07:57
MIC C编程(offload模式)
编程特点
简单---隐藏大量细节,语法与OpenMPI类似(不需要开辟空间)
灵活---OpenMP MPI(但是用的不多)pThread等多种方式
传统---与CPU编程一脉相承
MIC C扩展语言结构
编译指导方式(#pragma)
offload
--表示之后的代码段将使用offload模式运行
运行在其他设备上(MIC)
--标识内存传递参数,传递过程对用户透明
不需要手动写代码控制何时传入、何时传出
不需要手动申请卡上内存空间
不需要讲卡上内存与主机端内存空间手动对应
简单示例
#pragma offload target(mic)
for(i=0;i
{
printf(“Index:%d\n”,i);
}
循环和printf语句在MIC上执行,没有数据传输
此时循环是串行的,因为没有并行引语
offload语法
#pragma offload specifier[,specifier...]
specifier
示例
含义
target
target(mic:0)
使用什么设备运行
if
if(N>1000)
是否使用该设备
in
in(p:length(LEN) alloc_if(1))
数据传入device
out
out(p:length(LEN))
数据从device传出
inout
inout(p:length(LEN) align(8))
数据既传入又传出
nocopy
nocopy(p)
只开辟空间不传递
signal
signal(tag)
发送信号
wait
wait(tag1,tag2)
等待信号
mandatory
mandatory
必须用MIC运行
其中in/out/inout/nocopy可用的属性有
属性
示例
含义
length
length(LEN)
元素个数
alloc_if
alloc_if(1)
是否开辟内存
free_if
free_if(N>0)
offload后是否释放
align
align(8)
对齐
alloc
alloc(p[10:100])
开辟数组一部分
into
into(p[10:100])
传递数组一部分
语法详解target
#pragma offload target(mic) mic这里目前只支持MIC卡
--这是必须使用的属性
--表示下面的代码段使用MIC运行
--现阶段,target的参数只支持“mic”
--使用哪块卡,由驱动决定,如果系统没有MIC卡,则使用CPU运行
--num>=-1;当为-1时,系统自动选择一块mic卡,如果1号卡不存在或者有故障,则程序报错退出
--设备号=num%mic数量(也就是说num可以大于mic卡数,循环的)
语法详解in
#pragma offload target(mic) in(array) --传入CPU中的array,不负责传出
--可选属性
--可附加length、alloc_if、free_if、align、alloc、into属性
--表示进入offload区域时,将数组array从CPU内存中传递到MIC内存中,只传入不传出
--只需要传递数组,标量会自动以inout方式传递
语法详解out/inout/nocopy
可选属性
与in基本相似
out表示离开offload区域时,将数组array从MIC内存传递到CPU内存,只传出不传入
inout表示进入offload区域时,将数组从CPU内存传输到MIC内存,离开时将数组传回
nocopy表示仅开辟或保留空间,不传输数据
注意事项:
array是静态声明时,不用加长度参数
array是动态开辟时(堆),需要加长度参数length
array长度相同,或静态声明时,可以在一个传输区域内传输多个。例:in(a,b)
默认进入offload时开辟内存,离开时释放
in/out/inout/nocopy在一句offload中可以出现多次,但每个数组名只能出现一次
传输属性详解length
传输数组的元素个数,非数组长度!
只能使用在动态开辟的数组中
数组长度相同时可以共用
示例:
float *pArray,*pArray1;
pArray=(float*)malloc(LEN*sizeof(float));
pArray1=(float*)malloc(LEN*sizeof(float));
#pragma offload target \
in(pArray,pArray1:length(LEN)) 共用啦
{……}
传输属性详解alloc_if,free_if
控制是否在传输前开辟卡上内存空间和传完是否释放
常结合nocopy属性应用
例:
#pragma offload target \
in(pArray:length(LEN) free_if(0)) 里面是boolean型,“0”表示不释放空间
//进入下一个offload时,pArray的数据仍然存在(计算完后pArray存在于MIC中,不需要再从CPU端传输多MIC端了)
#pragma offload target \
nocopy(pArray:length(LEN) alloc_if(0)) 不用再释放空间了,也不用in再次传输了
//此时不需要开辟空间,否则会覆盖原有数据
(这里优化点:1:传输 2:开辟 3:释放)
数据传输优化Nocopy
p_c=...//p_c在每次迭代中值不变
for(i=0;i
{
p_in=...;//每次迭代计算时,p_in的值变化
#pragma offload target(mic) \
in(p_in:length(...)) in(p_c:length(...)) out(p_out:lenth(...))
{
kernel(p_in,p_c,p_out);
}
}
...=p_out;//CPU端在所有迭代完成之后才用到p_out的值
这里的冗余是p_c冗余传入p_out冗余输出
优化方案:
p_c=...;//p_c在每次迭代中值不改变
#pragma offload target(mic) \
in(p_c:length(...) alloc_if(1) free_if(0)) \ //开辟不释放
nocopy(p_in:length(...) alloc_if(1) free_if(0)) \
nocopy(p_out:length(...) alloc_if(1) free_if(0))
{
}//申请空间,并且不释放,传递p_c的值
for(i=0;i
p_in=...;
#pragma offload target(mic) \
in(p_in:length(...) alloc_if(0) free_if(0)) \ //p_in省去了steps次的开辟释放空间
nocopy(p_c) nocopy(p_out) //这里p_c省去了输入,p_out省去了输出
{
kernel(p_in,p_c,p_out);
}}
//以下是把最后的p_out输出来
#pragma offload target(mic) \
nocopy(p_c:length(...) alloc_if(0) free_if(1)) \
nocopy(p_in:length(...) alloc_if(0) free_if(1)) \
out(p_out:length(...) alloc_if(0) free_if(1))
{...}
...=p_out
传输属性详解align
控制卡上内存开辟的对齐大小
单位为字节
内存不对齐时会影响性能
一般用于传输不规则的结构体(定义顺序)
传输属性详解alloc/into
alloc开辟卡上内存,大小为数组的一部分
into传输数据,大小为数组的一部分
不能与inout/nocopy同用
传输属性特殊用法
#pragma offload target \
in(p[10:100]:alloc(p[5:1000]))
在设备开辟了1000个元素的数组p,数组下表的可用范围是从5开始,即5-10004.然后将主机端从p[10]开始的100个元素传到设备端的p[10]-p[109]的位置(mic卡上)
需要程序员保证不会越界
#pragma offload ... \
in(p[0:500]:into(p1[500:500]))
将主机端p[0]开始的500个元素的值,复制到设备端p1[500]-p1[999]的响应位置
如果同一offload语句有多个into,执行顺序未知,因此需要注意目的数组范围不能重叠
into不是简单拷贝,维度不同不能传输
语法详解signal/wait
#pragma offload target(mic) signal(tag)
可选属性,异步传输
等同于offload_transfer/offload_wait,但后者为单一语句,后面不跟代码段
在传输语句时加signal,CPU不等待MIC运行结束即继续运行。知道遇到wait语句
wait语句等待signal语句传输完成再继续运行
参数为一个变量或者数组名(即使同事传递多个)
示例:
float *in1;
#pragma offload target(mic:0) signal(in1)
{
mic_compute();
}
cpu_compute(); //本函数与上面的MIC函数并行执行
#pragma offload_wait target(mic:0) wait(in1) //这里与MIC合流
offload+OpenMP
在offload语句作用范围内使用OpenMp,即可在MIC上并行计算
例:
#pragma offload target(mic) out(arr:length(LEN))
#pragma omp parallel for
for(i=0;i
{
arr[i]=i*3.0/x;
}
变量和函数声明
_declspec(target(mic))
_attribute_((target(mic)))
修饰后,变量和函数可同时用于CPU和MIC,二者等价,attribute时注意有两层括号
例:
_attribute_((target(mic))) int a;
-attribute_((target(mic))) void fun();
变量和函数声明(批量)
同时声明多个变量和函数
两种形式
#pragma offload_attribute(push,target(mic))
//函数或变量声明
#pragma offload_attribute(pop)
或者:
#pragma offload_attribute(target(mic))
//函数或变量声明
#pragma offload_attribute(target(none))
判断代码段是否运行在MIC上
检查是否定义了宏_MIC_
#ifdef_MIC_
...
#else
...
#endif
需要封装在函数中,而不能直接用于offload代码段
_attribute_((target(mic))) void offload_check(void)
{
#ifdef _MIC_
printf(“check func:Run on the mic!\n”);
#else
printf(“check func:Run on the cpu!\n”);
#endif
}
语法详解if
#pragma offload target(mic) if(flag) //flag 是boolean型的
可选属性
如果flag为真,则使用MIC执行代码段,否则使用CPU执行
例:#pragma offload target(mic) if(N>100)
表示如果N大于100就在MIC卡上执行
可用于不同规模数据,或者CPU、MIC协同计算的情况
语法详解mandatory
可选属性
没有参数
强制下面的代码段在MIC上运行,如果MIC设备不可用,则报错退出--不去CPU上跑
多用于必须在MIC上运行的代码,例如使用SIMD写的代码段
程序编译
使用Intel编译器
编译选项与原CPU程序完全一样
如果只想使用CPU,添加选项”-no-offload”
如果使用了OpenMP,添加选项”-openmp” (如果不加此选项,就在mic的单核上运行)
target(mic:num)