opencl:异步复制函数的注意事项(async_work_group_copy/async_work_group_strided_copy)
2016-06-11 13:21
661 查看
OpenCL中的内置函数async_work_group_copy和async_work_group_strided_copy用于实现全局内存(global memory)和本地内存(local memory)之间的异步数据复制,在某些情况下,使用异步复制(async copy)的方式在全局内存和本地内存之间复制数据比直接赋值的方式要方便。
下面是async_work_group_copy的函数说明:
请注意用红线标注的两段话(async_work_group_strided_copy中的说明中也有同样的描述):
1:异步复制(async copy)会被工作组内的所有工作项执行,所以异步复制函数必须被所有工作项用同样的参数执行,否则结果是未定义的。
这句话有两个要点,a.异步复制(async copy)函数必须被所有的工作项执行,b.每个工作项执行异步复制(async copy)函数时所用的参数必须一样。a比较好理解,就是说不能有条件判断语句绕过异步复制(async copy)函数。b的意思就是所有的工作项在执行异步复制(async copy)函数时,复制的是同样的一段数据,从源地址和目标地址一样,复制的数量一样(对于async_work_group_strided_copy函数来说,每个数据的步长(stride)也必须一样),不能是每个工作项只复制自己所需要的一段数据。
2:异步复制(async copy)函数在执行复制之间不会执行任何隐式的源数据同步(比如用barrier函数进行同步)。
这一点就要求我们在执行多次异步复制的时候,要根据需要使用barrier函数进行进行数据同步。
比如两次调用异步复制函数复制的目的地址一样时,在这两次调用之间,就应该用barrier函数进行同步。
下面是我项目中一个实际的kernel函数,用于从积分图矩阵计算局部矩形区域的特征平均值。在这个kernel函数中同时展现符合了这两个要求的代码特特性。详见代码内的中文注释。
如果将上面的代码中async_work_group_strided_copy函数之间的barrier同步语句删除,有时也能正常执行,但并不是在所有的设备上都可以正常执行。
原来代码中async_work_group_strided_copy之间是没有用barrier做同步的,在AMD GPU上也能正常执行。
但当我把同样的代码在CPU上执行时,结果就不正确了。然后我在async_work_group_strided_copy之间加上barrier同步,执行结果才正确。这就印证了前面注意事项的第二条:异步复制函数本身是没有做数据同步的,必须根据需要做数据同步。
因为上面的代码中每次async_work_group_strided_copy函数的目标地址都是一样,如果没有barrier同步,有的工作项还没有来得及将数据从本地内存取走,异步复制就开始执行了会将本地内存中的结果冲掉。这时工作项从本地内存读取的数据就不对了。
下面是async_work_group_copy的函数说明:
请注意用红线标注的两段话(async_work_group_strided_copy中的说明中也有同样的描述):
1:异步复制(async copy)会被工作组内的所有工作项执行,所以异步复制函数必须被所有工作项用同样的参数执行,否则结果是未定义的。
这句话有两个要点,a.异步复制(async copy)函数必须被所有的工作项执行,b.每个工作项执行异步复制(async copy)函数时所用的参数必须一样。a比较好理解,就是说不能有条件判断语句绕过异步复制(async copy)函数。b的意思就是所有的工作项在执行异步复制(async copy)函数时,复制的是同样的一段数据,从源地址和目标地址一样,复制的数量一样(对于async_work_group_strided_copy函数来说,每个数据的步长(stride)也必须一样),不能是每个工作项只复制自己所需要的一段数据。
2:异步复制(async copy)函数在执行复制之间不会执行任何隐式的源数据同步(比如用barrier函数进行同步)。
这一点就要求我们在执行多次异步复制的时候,要根据需要使用barrier函数进行进行数据同步。
比如两次调用异步复制函数复制的目的地址一样时,在这两次调用之间,就应该用barrier函数进行同步。
下面是我项目中一个实际的kernel函数,用于从积分图矩阵计算局部矩形区域的特征平均值。在这个kernel函数中同时展现符合了这两个要求的代码特特性。详见代码内的中文注释。
__kernel void create_mean( matrix_info_cl im_info ,const __global float *integ_mat ,matrix_info_cl om_info ,__global float * mean_mat ,int win_size ,float inv_pixels_num ,int sample_step ){ __local float4 local_buf[MEAN_LOCAL_SIZE]; const int pixel_size=MEAN_LOCAL_SIZE<<2; const int group_y= (int)get_group_id(1); const int group_x= (int)get_group_id(0)*pixel_size; const int group_count=min(pixel_size,(int)(om_info.width-group_x)); const int local_id=get_local_id(0); float4 mean; // A1 #define INDEX_A1 (group_y*im_info.row_stride+group_x)*sample_step //A2 #define INDEX_A2 (group_y*im_info.row_stride+group_x)*sample_step+win_size //A3 #define INDEX_A3 (group_y*sample_step+win_size)*im_info.row_stride+group_x*sample_step //A4 #define INDEX_A4 (group_y*sample_step+win_size)*im_info.row_stride+group_x*sample_step+win_size local_buf[local_id]=0;// 本地内存初始化清零 barrier(CLK_LOCAL_MEM_FENCE);//同步函数 event_t evt; // 从全局内存到本地内存的异步复制,这里使用async_work_group_strided_copy做步长为sample_step的异步复制,将源数据中离散的数据复制到本地内存连续存储 // 注意:INDEX_A4的定义(下同),所有的工作项的原数据起始地址都是一样的。复制数量group_count也是一样的 evt=async_work_group_strided_copy((__local float*)local_buf,integ_mat+INDEX_A4,group_count,sample_step,0);wait_group_events(1,&evt); mean =local_buf[local_id]; barrier(CLK_LOCAL_MEM_FENCE);//因为第二次异步复制目标地址还是local_buf,所以这里要加上同步函数,以保证所有工作项在执行完前一条语句mean =local_buf[local_id];从local memory取走数据后同时开始执行异步复制 evt=async_work_group_strided_copy((__local float*)local_buf,integ_mat+INDEX_A1,group_count,sample_step,0);wait_group_events(1,&evt); mean+=local_buf[local_id]; barrier(CLK_LOCAL_MEM_FENCE); evt=async_work_group_strided_copy((__local float*)local_buf,integ_mat+INDEX_A2,group_count,sample_step,0);wait_group_events(1,&evt); mean-=local_buf[local_id]; barrier(CLK_LOCAL_MEM_FENCE); evt=async_work_group_strided_copy((__local float*)local_buf,integ_mat+INDEX_A3,group_count,sample_step,0);wait_group_events(1,&evt); mean-=local_buf[local_id]; local_buf[local_id] = mean*inv_pixels_num; barrier(CLK_LOCAL_MEM_FENCE);//同步函数,保证所有工作项都执行完上一条语句,将计算结果存入了local_buf后再同步执行异步复制 // 本地内存到全局内存的异步复制 evt=async_work_group_copy(mean_mat+group_y*om_info.row_stride+group_x,(__local float*)local_buf,group_count,0);wait_group_events(1,&evt); #undef INDEX_A1 #undef INDEX_A2 #undef INDEX_A3 #undef INDEX_A4 }
如果将上面的代码中async_work_group_strided_copy函数之间的barrier同步语句删除,有时也能正常执行,但并不是在所有的设备上都可以正常执行。
原来代码中async_work_group_strided_copy之间是没有用barrier做同步的,在AMD GPU上也能正常执行。
但当我把同样的代码在CPU上执行时,结果就不正确了。然后我在async_work_group_strided_copy之间加上barrier同步,执行结果才正确。这就印证了前面注意事项的第二条:异步复制函数本身是没有做数据同步的,必须根据需要做数据同步。
因为上面的代码中每次async_work_group_strided_copy函数的目标地址都是一样,如果没有barrier同步,有的工作项还没有来得及将数据从本地内存取走,异步复制就开始执行了会将本地内存中的结果冲掉。这时工作项从本地内存读取的数据就不对了。
相关文章推荐
- async.js 学习记录
- 批处理中Copy与Xcopy命令的区别小结
- dos 文件复制 copy命令
- async and await 的入门基础操作
- .NET中的async和await关键字使用及Task异步调用实例
- 谈谈xcopy中的排除copy
- php使用COPY函数更新配置文件的方法
- 详解C#中的Async和Await用法
- 关于async和await的一些误区实例详解
- jQuery中的ajax async同步和异步详解
- 浅谈Jquery中Ajax异步请求中的async参数的作用
- copy 将一个或多个文件从一个位置复制到其他位置
- cmd copy命令 文件复制
- ASP XML编程objXML.async = False第1/2页
- Ajax请求中async:false/true的作用分析
- .net4.5使用async和await异步编程实例
- async和DOM Script文件加载比较
- 浅谈node.js中async异步编程
- 浏览器环境下JavaScript脚本加载与执行探析之defer与async特性
- ajax中的async属性值之同步和异步及同步和异步区别