From 1257a871669ea3e5ab9f3bdc885d69d41c6a522c Mon Sep 17 00:00:00 2001 From: Guo Yejun Date: Tue, 12 May 2015 20:44:08 +0800 Subject: 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 Reviewed-by: Zhigang Gong --- kernels/compiler_sub_group_shuffle.cl | 18 ++++++++++++++++++ 1 file changed, 18 insertions(+) create mode 100644 kernels/compiler_sub_group_shuffle.cl (limited to 'kernels') 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; +} -- cgit v1.2.3