From: Pan Xiuli <xiuli....@intel.com>
Signed-off-by: Pan Xiuli <xiuli....@intel.com>
---
kernels/compiler_subgroup_broadcast.cl | 10 ++++++++++
utests/compiler_subgroup_broadcast.cpp | 11 +++++++++++
2 files changed, 21 insertions(+)
diff --git a/kernels/compiler_subgroup_broadcast.cl
b/kernels/compiler_subgroup_broadcast.cl
index 63e9568..3d16d67 100644
--- a/kernels/compiler_subgroup_broadcast.cl
+++ b/kernels/compiler_subgroup_broadcast.cl
@@ -32,6 +32,16 @@ kernel void compiler_subgroup_broadcast_long(global long
*src,
long broadcast_val = sub_group_broadcast(val, simd_id);
dst[index] = broadcast_val;
}
+kernel void compiler_subgroup_broadcast_short(global short *src,
+ global short *dst,
+ uint simd_id)
+{
+ uint index = get_global_id(0);
+
+ short val = src[index];
+ short broadcast_val = sub_group_broadcast(val, simd_id);
+ dst[index] = broadcast_val;
+}
#else
#pragma OPENCL EXTENSION cl_khr_fp16 : enable
kernel void compiler_subgroup_broadcast_half(global half *src,
diff --git a/utests/compiler_subgroup_broadcast.cpp
b/utests/compiler_subgroup_broadcast.cpp
index 5aa749c..33ec43c 100644
--- a/utests/compiler_subgroup_broadcast.cpp
+++ b/utests/compiler_subgroup_broadcast.cpp
@@ -190,6 +190,17 @@ void compiler_subgroup_broadcast_long(void)
subgroup_generic(input, expected);
}
MAKE_UTEST_FROM_FUNCTION_WITH_ISSUE(compiler_subgroup_broadcast_long);
+void compiler_subgroup_broadcast_short(void)
+{
+ if(!cl_check_subgroups_short())
+ return;
+ cl_short *input = NULL;
+ cl_short *expected = NULL;
+ OCL_CREATE_KERNEL_FROM_FILE("compiler_subgroup_broadcast",
+ "compiler_subgroup_broadcast_short");
+ subgroup_generic(input, expected);
+}
+MAKE_UTEST_FROM_FUNCTION(compiler_subgroup_broadcast_short);
void compiler_subgroup_broadcast_half(void)
{
if(!cl_check_subgroups())
--
2.7.4
_______________________________________________
Beignet mailing list
Beignet@lists.freedesktop.org
https://lists.freedesktop.org/mailman/listinfo/beignet
