opencl:异步复制函数的注意事项(async_work_group_copy/async_work_group_strided_copy)

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函数中同时展现符合了这两个要求的代码特特性。详见代码内的中文注释。

__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同步，有的工作项还没有来得及将数据从本地内存取走,异步复制就开始执行了会将本地内存中的结果冲掉。这时工作项从本地内存读取的数据就不对了。