MIC C编程（offload模式）
编程特点
简单---隐藏大量细节，语法与OpenMPI类似（不需要开辟空间）
灵活---OpenMP MPI（但是用的不多）pThread等多种方式
传统---与CPU编程一脉相承
MIC C扩展语言结构
编译指导方式（#pragma）
offload
--表示之后的代码段将使用offload模式运行
运行在其他设备上（MIC）
--标识内存传递参数，传递过程对用户透明
不需要手动写代码控制何时传入、何时传出
不需要手动申请卡上内存空间
不需要讲卡上内存与主机端内存空间手动对应
简单示例
#pragma offload target(mic)
for(i=0;i<LEN;i++)
{
printf(“Index:%d ”,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运行
target(mic:num)
--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<steps;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<steps;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<LEN;++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! ”);
#else
printf(“check func:Run on the cpu! ”);
#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的单核上运行)