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)