summaryrefslogtreecommitdiff
path: root/kernels
diff options
context:
space:
mode:
authorGuo Yejun <yejun.guo@intel.com>2015-05-12 20:44:08 +0800
committerZhigang Gong <zhigang.gong@intel.com>2015-05-13 07:43:29 +0800
commit1257a871669ea3e5ab9f3bdc885d69d41c6a522c (patch)
tree13fbf225ef244da9bf4c86b7d3e3114b9d11de2b /kernels
parente8a7f0e248bdb0a2e579535fc8fd09ae2cc3c3a5 (diff)
add utest for intel_sub_group_shuffle
v2: correct kernel to be suitable for simd_width both 8 and 16 Signed-off-by: Guo Yejun <yejun.guo@intel.com> Reviewed-by: Zhigang Gong <zhigang.gong@intel.com>
Diffstat (limited to 'kernels')
-rw-r--r--kernels/compiler_sub_group_shuffle.cl18
1 files changed, 18 insertions, 0 deletions
diff --git a/kernels/compiler_sub_group_shuffle.cl b/kernels/compiler_sub_group_shuffle.cl
new file mode 100644
index 00000000..75adde3c
--- /dev/null
+++ b/kernels/compiler_sub_group_shuffle.cl
@@ -0,0 +1,18 @@
+__kernel void compiler_sub_group_shuffle(global int *dst, int c)
+{
+ int i = get_global_id(0);
+ if (i == 0)
+ dst[0] = get_sub_group_size();
+ dst++;
+
+ int from = i;
+ int j = get_sub_group_size() - get_sub_group_id() - 1;
+ int o0 = get_sub_group_id();
+ int o1 = intel_sub_group_shuffle(from, c);
+ int o2 = intel_sub_group_shuffle(from, 5);
+ int o3 = intel_sub_group_shuffle(from, j);
+ dst[i*4] = o0;
+ dst[i*4+1] = o1;
+ dst[i*4+2] = o2;
+ dst[i*4+3] = o3;
+}