diff options
| author | Pan Xiuli <xiuli.pan@intel.com> | 2016-10-19 14:37:23 +0800 |
|---|---|---|
| committer | Yang Rong <rong.r.yang@intel.com> | 2016-11-03 12:24:10 +0800 |
| commit | 45544098d5253c37bca03aaf860c3e91816821dd (patch) | |
| tree | 65ff13b2b3098ae558fac575f4f867a1dd1173c3 /kernels/compiler_subgroup_buffer_block_read.cl | |
| parent | 24f8f5dcd6306b0662bc3d7fdd70988c35646748 (diff) | |
| download | beignet-45544098d5253c37bca03aaf860c3e91816821dd.tar.gz | |
Utest: Add subgroup block read/write ushort test case
Add ushort block read/write for buffer and image.
Refine uint block read/write with suffix _ui.
Signed-off-by: Pan Xiuli <xiuli.pan@intel.com>
Reviewed-by: Yang Rong <rong.r.yang@intel.com>
Diffstat (limited to 'kernels/compiler_subgroup_buffer_block_read.cl')
| -rw-r--r-- | kernels/compiler_subgroup_buffer_block_read.cl | 47 |
1 files changed, 39 insertions, 8 deletions
diff --git a/kernels/compiler_subgroup_buffer_block_read.cl b/kernels/compiler_subgroup_buffer_block_read.cl index 9edaa2ed..4cbf8945 100644 --- a/kernels/compiler_subgroup_buffer_block_read.cl +++ b/kernels/compiler_subgroup_buffer_block_read.cl @@ -1,31 +1,62 @@ -__kernel void compiler_subgroup_buffer_block_read1(global uint *src, global uint *dst) +__kernel void compiler_subgroup_buffer_block_read_ui1(global uint *src, global uint *dst) { int id = get_global_id(0); global uint * p = src + get_sub_group_id() * get_max_sub_group_size(); - uint tmp = intel_sub_group_block_read(p); + uint tmp = intel_sub_group_block_read_ui(p); dst[id] = tmp; } -__kernel void compiler_subgroup_buffer_block_read2(global uint *src, global uint2 *dst) +__kernel void compiler_subgroup_buffer_block_read_ui2(global uint *src, global uint2 *dst) { int id = get_global_id(0); global uint * p = src + get_sub_group_id() * get_max_sub_group_size()*2; - uint2 tmp = intel_sub_group_block_read2(p); + uint2 tmp = intel_sub_group_block_read_ui2(p); dst[id] = tmp; } -__kernel void compiler_subgroup_buffer_block_read4(global uint *src, global uint4 *dst) +__kernel void compiler_subgroup_buffer_block_read_ui4(global uint *src, global uint4 *dst) { int id = get_global_id(0); global uint * p = src + get_sub_group_id() * get_max_sub_group_size()*4; - uint4 tmp = intel_sub_group_block_read4(p); + uint4 tmp = intel_sub_group_block_read_ui4(p); dst[id] = tmp; } -__kernel void compiler_subgroup_buffer_block_read8(global uint *src, global uint8 *dst) +__kernel void compiler_subgroup_buffer_block_read_ui8(global uint *src, global uint8 *dst) { int id = get_global_id(0); global uint * p = src + get_sub_group_id() * get_max_sub_group_size()*8; - uint8 tmp = intel_sub_group_block_read8(p); + uint8 tmp = intel_sub_group_block_read_ui8(p); dst[id] = tmp; } +#ifdef SHORT +__kernel void compiler_subgroup_buffer_block_read_us1(global ushort *src, global ushort *dst) +{ + int id = get_global_id(0); + global ushort * p = src + get_sub_group_id() * get_max_sub_group_size(); + ushort tmp = intel_sub_group_block_read_us(p); + dst[id] = tmp; +} +__kernel void compiler_subgroup_buffer_block_read_us2(global ushort *src, global ushort2 *dst) +{ + int id = get_global_id(0); + global ushort * p = src + get_sub_group_id() * get_max_sub_group_size()*2; + ushort2 tmp = intel_sub_group_block_read_us2(p); + dst[id] = tmp; +} +__kernel void compiler_subgroup_buffer_block_read_us4(global ushort *src, global ushort4 *dst) +{ + int id = get_global_id(0); + global ushort * p = src + get_sub_group_id() * get_max_sub_group_size()*4; + ushort4 tmp = intel_sub_group_block_read_us4(p); + dst[id] = tmp; +} + +__kernel void compiler_subgroup_buffer_block_read_us8(global ushort *src, global ushort8 *dst) +{ + int id = get_global_id(0); + global ushort * p = src + get_sub_group_id() * get_max_sub_group_size()*8; + ushort8 tmp = intel_sub_group_block_read_us8(p); + dst[id] = tmp; +} +#endif |