diff options
| author | Luo Xionghu <xionghu.luo@intel.com> | 2015-03-09 11:24:26 +0800 |
|---|---|---|
| committer | Zhigang Gong <zhigang.gong@intel.com> | 2015-03-09 16:30:45 +0800 |
| commit | 6e00c44f6bd4beacf2d9bf499faf7f9494e0d0be (patch) | |
| tree | 2d2eff8225ce4a736d5f263a3da2f9d1f0c44b2d /kernels | |
| parent | ce584ec4976e79a7635eaaafe4a035100397e732 (diff) | |
change the workitem related api to OVERLOABABLE.
the SPIR header file requirs these functions to be overlable.
(https://github.com/KhronosGroup/SPIR-Tools/blob/master/headers/opencl_spir.h)
Signed-off-by: Luo Xionghu <xionghu.luo@intel.com>
Reviewed-by: Zhigang Gong <zhigang.gong@linux.intel.com>
Diffstat (limited to 'kernels')
| -rw-r--r-- | kernels/compiler_async_copy.cl | 4 | ||||
| -rw-r--r-- | kernels/compiler_async_stride_copy.cl | 4 |
2 files changed, 4 insertions, 4 deletions
diff --git a/kernels/compiler_async_copy.cl b/kernels/compiler_async_copy.cl index dddde446..4beb4360 100644 --- a/kernels/compiler_async_copy.cl +++ b/kernels/compiler_async_copy.cl @@ -5,10 +5,10 @@ compiler_async_copy_##TYPE(__global TYPE *dst, __global TYPE *src, __local TYPE event_t event; \ int copiesPerWorkgroup = copiesPerWorkItem * get_local_size(0); \ int i; \ - event = async_work_group_copy((__local TYPE*)localBuffer, (__global const TYPE*)(src+copiesPerWorkgroup*get_group_id(0)), (size_t)copiesPerWorkgroup, (event_t)0 ); \ + event = async_work_group_copy((__local TYPE*)localBuffer, (__global const TYPE*)(src+copiesPerWorkgroup*get_group_id(0)), (size_t)copiesPerWorkgroup, 0 ); \ wait_group_events( 1, &event ); \ \ - event = async_work_group_copy((__global TYPE*)(dst+copiesPerWorkgroup*get_group_id(0)), (__local const TYPE*)localBuffer, (size_t)copiesPerWorkgroup, (event_t)0 ); \ + event = async_work_group_copy((__global TYPE*)(dst+copiesPerWorkgroup*get_group_id(0)), (__local const TYPE*)localBuffer, (size_t)copiesPerWorkgroup, 0 ); \ wait_group_events( 1, &event ); \ } diff --git a/kernels/compiler_async_stride_copy.cl b/kernels/compiler_async_stride_copy.cl index a9265888..5dbb5591 100644 --- a/kernels/compiler_async_stride_copy.cl +++ b/kernels/compiler_async_stride_copy.cl @@ -4,13 +4,13 @@ compiler_async_stride_copy(__global char4 *dst, __global char4 *src, __local cha event_t event; int copiesPerWorkgroup = copiesPerWorkItem * get_local_size(0); int i; - event = async_work_group_strided_copy( (__local char4*)localBuffer, (__global const char4*)(src+copiesPerWorkgroup*stride*get_group_id(0)), (size_t)copiesPerWorkgroup, (size_t)stride, (event_t)0 ); + event = async_work_group_strided_copy( (__local char4*)localBuffer, (__global const char4*)(src+copiesPerWorkgroup*stride*get_group_id(0)), (size_t)copiesPerWorkgroup, (size_t)stride, 0 ); wait_group_events( 1, &event ); for(i=0; i<copiesPerWorkItem; i++) localBuffer[ get_local_id( 0 )*copiesPerWorkItem+i ] = localBuffer[ get_local_id( 0 )*copiesPerWorkItem+i ] + (char4)(3); barrier(CLK_LOCAL_MEM_FENCE); - event = async_work_group_strided_copy((__global char4*)(dst+copiesPerWorkgroup*stride*get_group_id(0)), (__local const char4*)localBuffer, (size_t)copiesPerWorkgroup, (size_t)stride, (event_t)0 ); + event = async_work_group_strided_copy((__global char4*)(dst+copiesPerWorkgroup*stride*get_group_id(0)), (__local const char4*)localBuffer, (size_t)copiesPerWorkgroup, (size_t)stride, 0 ); wait_group_events( 1, &event ); } |