[Beignet] [PATCH 10/14] Utest: Add test case for sub group short builtin functions
Xiuli Pan
xiuli.pan at intel.com
Wed Oct 12 08:56:40 UTC 2016
From: Pan Xiuli <xiuli.pan at intel.com>
Signed-off-by: Pan Xiuli <xiuli.pan at intel.com>
---
kernels/compiler_subgroup_reduce.cl | 22 ++++++++++
kernels/compiler_subgroup_scan_exclusive.cl | 36 ++++++++++++++++
kernels/compiler_subgroup_scan_inclusive.cl | 36 ++++++++++++++++
utests/compiler_subgroup_reduce.cpp | 66 +++++++++++++++++++++++++++++
utests/compiler_subgroup_scan_exclusive.cpp | 66 +++++++++++++++++++++++++++++
utests/compiler_subgroup_scan_inclusive.cpp | 66 +++++++++++++++++++++++++++++
6 files changed, 292 insertions(+)
diff --git a/kernels/compiler_subgroup_reduce.cl b/kernels/compiler_subgroup_reduce.cl
index 6d7ecfd..79d8e7d 100644
--- a/kernels/compiler_subgroup_reduce.cl
+++ b/kernels/compiler_subgroup_reduce.cl
@@ -73,6 +73,17 @@ kernel void compiler_subgroup_reduce_add_float(global float *src, global float *
/*
* Subgroup reduce max functions
*/
+kernel void compiler_subgroup_reduce_max_short(global short *src, global short *dst) {
+ short val = src[get_global_id(0)];
+ short sum = sub_group_reduce_max(val);
+ dst[get_global_id(0)] = sum;
+}
+kernel void compiler_subgroup_reduce_max_ushort(global ushort *src, global ushort *dst) {
+ ushort val = src[get_global_id(0)];
+ //printf("src is %d\n",val);
+ ushort sum = sub_group_reduce_max(val);
+ dst[get_global_id(0)] = sum;
+}
kernel void compiler_subgroup_reduce_max_int(global int *src, global int *dst) {
int val = src[get_global_id(0)];
int sum = sub_group_reduce_max(val);
@@ -106,6 +117,17 @@ kernel void compiler_subgroup_reduce_max_float(global float *src, global float *
/*
* Subgroup reduce min functions
*/
+kernel void compiler_subgroup_reduce_min_short(global short *src, global short *dst) {
+ short val = src[get_global_id(0)];
+ short sum = sub_group_reduce_min(val);
+ dst[get_global_id(0)] = sum;
+}
+kernel void compiler_subgroup_reduce_min_ushort(global ushort *src, global ushort *dst) {
+ ushort val = src[get_global_id(0)];
+ //printf("src is %d\n",val);
+ ushort sum = sub_group_reduce_min(val);
+ dst[get_global_id(0)] = sum;
+}
kernel void compiler_subgroup_reduce_min_int(global int *src, global int *dst) {
int val = src[get_global_id(0)];
int sum = sub_group_reduce_min(val);
diff --git a/kernels/compiler_subgroup_scan_exclusive.cl b/kernels/compiler_subgroup_scan_exclusive.cl
index ca0ada2..2c4b928 100644
--- a/kernels/compiler_subgroup_scan_exclusive.cl
+++ b/kernels/compiler_subgroup_scan_exclusive.cl
@@ -2,6 +2,18 @@
* Subgroup scan exclusive add functions
*/
#ifndef HALF
+kernel void compiler_subgroup_scan_exclusive_add_short(global short *src, global short *dst) {
+ short val = src[get_global_id(0)];
+ short sum = sub_group_scan_exclusive_add(val);
+ dst[get_global_id(0)] = sum;
+}
+
+kernel void compiler_subgroup_scan_exclusive_add_ushort(global ushort *src, global ushort *dst) {
+ ushort val = src[get_global_id(0)];
+ ushort sum = sub_group_scan_exclusive_add(val);
+ dst[get_global_id(0)] = sum;
+}
+
kernel void compiler_subgroup_scan_exclusive_add_int(global int *src, global int *dst) {
int val = src[get_global_id(0)];
int sum = sub_group_scan_exclusive_add(val);
@@ -35,6 +47,18 @@ kernel void compiler_subgroup_scan_exclusive_add_float(global float *src, global
/*
* Subgroup scan exclusive max functions
*/
+kernel void compiler_subgroup_scan_exclusive_max_short(global short *src, global short *dst) {
+ short val = src[get_global_id(0)];
+ short sum = sub_group_scan_exclusive_max(val);
+ dst[get_global_id(0)] = sum;
+}
+
+kernel void compiler_subgroup_scan_exclusive_max_ushort(global ushort *src, global ushort *dst) {
+ ushort val = src[get_global_id(0)];
+ ushort sum = sub_group_scan_exclusive_max(val);
+ dst[get_global_id(0)] = sum;
+}
+
kernel void compiler_subgroup_scan_exclusive_max_int(global int *src, global int *dst) {
int val = src[get_global_id(0)];
int sum = sub_group_scan_exclusive_max(val);
@@ -68,6 +92,18 @@ kernel void compiler_subgroup_scan_exclusive_max_float(global float *src, global
/*
* Subgroup scan exclusive min functions
*/
+kernel void compiler_subgroup_scan_exclusive_min_short(global short *src, global short *dst) {
+ short val = src[get_global_id(0)];
+ short sum = sub_group_scan_exclusive_min(val);
+ dst[get_global_id(0)] = sum;
+}
+
+kernel void compiler_subgroup_scan_exclusive_min_ushort(global ushort *src, global ushort *dst) {
+ ushort val = src[get_global_id(0)];
+ ushort sum = sub_group_scan_exclusive_min(val);
+ dst[get_global_id(0)] = sum;
+}
+
kernel void compiler_subgroup_scan_exclusive_min_int(global int *src, global int *dst) {
int val = src[get_global_id(0)];
int sum = sub_group_scan_exclusive_min(val);
diff --git a/kernels/compiler_subgroup_scan_inclusive.cl b/kernels/compiler_subgroup_scan_inclusive.cl
index e97521c..def941c 100644
--- a/kernels/compiler_subgroup_scan_inclusive.cl
+++ b/kernels/compiler_subgroup_scan_inclusive.cl
@@ -2,6 +2,18 @@
* Subgroup scan inclusive add functions
*/
#ifndef HALF
+kernel void compiler_subgroup_scan_inclusive_add_short(global short *src, global short *dst) {
+ short val = src[get_global_id(0)];
+ short sum = sub_group_scan_inclusive_add(val);
+ dst[get_global_id(0)] = sum;
+}
+
+kernel void compiler_subgroup_scan_inclusive_add_ushort(global ushort *src, global ushort *dst) {
+ ushort val = src[get_global_id(0)];
+ ushort sum = sub_group_scan_inclusive_add(val);
+ dst[get_global_id(0)] = sum;
+}
+
kernel void compiler_subgroup_scan_inclusive_add_int(global int *src, global int *dst) {
int val = src[get_global_id(0)];
int sum = sub_group_scan_inclusive_add(val);
@@ -35,6 +47,18 @@ kernel void compiler_subgroup_scan_inclusive_add_float(global float *src, global
/*
* Subgroup scan inclusive max functions
*/
+kernel void compiler_subgroup_scan_inclusive_max_short(global short *src, global short *dst) {
+ short val = src[get_global_id(0)];
+ short sum = sub_group_scan_inclusive_max(val);
+ dst[get_global_id(0)] = sum;
+}
+
+kernel void compiler_subgroup_scan_inclusive_max_ushort(global ushort *src, global ushort *dst) {
+ ushort val = src[get_global_id(0)];
+ ushort sum = sub_group_scan_inclusive_max(val);
+ dst[get_global_id(0)] = sum;
+}
+
kernel void compiler_subgroup_scan_inclusive_max_int(global int *src, global int *dst) {
int val = src[get_global_id(0)];
int sum = sub_group_scan_inclusive_max(val);
@@ -68,6 +92,18 @@ kernel void compiler_subgroup_scan_inclusive_max_float(global float *src, global
/*
* Subgroup scan inclusive min functions
*/
+kernel void compiler_subgroup_scan_inclusive_min_short(global short *src, global short *dst) {
+ short val = src[get_global_id(0)];
+ short sum = sub_group_scan_inclusive_min(val);
+ dst[get_global_id(0)] = sum;
+}
+
+kernel void compiler_subgroup_scan_inclusive_min_ushort(global ushort *src, global ushort *dst) {
+ ushort val = src[get_global_id(0)];
+ ushort sum = sub_group_scan_inclusive_min(val);
+ dst[get_global_id(0)] = sum;
+}
+
kernel void compiler_subgroup_scan_inclusive_min_int(global int *src, global int *dst) {
int val = src[get_global_id(0)];
int sum = sub_group_scan_inclusive_min(val);
diff --git a/utests/compiler_subgroup_reduce.cpp b/utests/compiler_subgroup_reduce.cpp
index ff545c6..157086a 100644
--- a/utests/compiler_subgroup_reduce.cpp
+++ b/utests/compiler_subgroup_reduce.cpp
@@ -357,6 +357,28 @@ void compiler_subgroup_reduce_add_half(void)
subgroup_generic(WG_REDUCE_ADD, input, expected, true);
}
MAKE_UTEST_FROM_FUNCTION(compiler_subgroup_reduce_add_half);
+void compiler_subgroup_reduce_add_short(void)
+{
+ if(!cl_check_subgroups_short())
+ return;
+ cl_short *input = NULL;
+ cl_short *expected = NULL;
+ OCL_CREATE_KERNEL_FROM_FILE("compiler_subgroup_reduce",
+ "compiler_subgroup_reduce_add_short");
+ subgroup_generic(WG_REDUCE_ADD, input, expected);
+}
+MAKE_UTEST_FROM_FUNCTION(compiler_subgroup_reduce_add_short);
+void compiler_subgroup_reduce_add_ushort(void)
+{
+ if(!cl_check_subgroups_short())
+ return;
+ cl_ushort *input = NULL;
+ cl_ushort *expected = NULL;
+ OCL_CREATE_KERNEL_FROM_FILE("compiler_subgroup_reduce",
+ "compiler_subgroup_reduce_add_ushort");
+ subgroup_generic(WG_REDUCE_ADD, input, expected);
+}
+MAKE_UTEST_FROM_FUNCTION(compiler_subgroup_reduce_add_ushort);
/*
* Workgroup reduce max utest functions
@@ -430,6 +452,28 @@ void compiler_subgroup_reduce_max_half(void)
subgroup_generic(WG_REDUCE_MAX, input, expected, true);
}
MAKE_UTEST_FROM_FUNCTION(compiler_subgroup_reduce_max_half);
+void compiler_subgroup_reduce_max_short(void)
+{
+ if(!cl_check_subgroups_short())
+ return;
+ cl_short *input = NULL;
+ cl_short *expected = NULL;
+ OCL_CREATE_KERNEL_FROM_FILE("compiler_subgroup_reduce",
+ "compiler_subgroup_reduce_max_short");
+ subgroup_generic(WG_REDUCE_MAX, input, expected);
+}
+MAKE_UTEST_FROM_FUNCTION(compiler_subgroup_reduce_max_short);
+void compiler_subgroup_reduce_max_ushort(void)
+{
+ if(!cl_check_subgroups_short())
+ return;
+ cl_ushort *input = NULL;
+ cl_ushort *expected = NULL;
+ OCL_CREATE_KERNEL_FROM_FILE("compiler_subgroup_reduce",
+ "compiler_subgroup_reduce_max_ushort");
+ subgroup_generic(WG_REDUCE_MAX, input, expected);
+}
+MAKE_UTEST_FROM_FUNCTION(compiler_subgroup_reduce_max_ushort);
/*
* Workgroup reduce min utest functions
@@ -503,3 +547,25 @@ void compiler_subgroup_reduce_min_half(void)
subgroup_generic(WG_REDUCE_MIN, input, expected, true);
}
MAKE_UTEST_FROM_FUNCTION(compiler_subgroup_reduce_min_half);
+void compiler_subgroup_reduce_min_short(void)
+{
+ if(!cl_check_subgroups_short())
+ return;
+ cl_short *input = NULL;
+ cl_short *expected = NULL;
+ OCL_CREATE_KERNEL_FROM_FILE("compiler_subgroup_reduce",
+ "compiler_subgroup_reduce_min_short");
+ subgroup_generic(WG_REDUCE_MIN, input, expected);
+}
+MAKE_UTEST_FROM_FUNCTION(compiler_subgroup_reduce_min_short);
+void compiler_subgroup_reduce_min_ushort(void)
+{
+ if(!cl_check_subgroups_short())
+ return;
+ cl_ushort *input = NULL;
+ cl_ushort *expected = NULL;
+ OCL_CREATE_KERNEL_FROM_FILE("compiler_subgroup_reduce",
+ "compiler_subgroup_reduce_min_ushort");
+ subgroup_generic(WG_REDUCE_MIN, input, expected);
+}
+MAKE_UTEST_FROM_FUNCTION(compiler_subgroup_reduce_min_ushort);
diff --git a/utests/compiler_subgroup_scan_exclusive.cpp b/utests/compiler_subgroup_scan_exclusive.cpp
index e51b78d..4f3e5ea 100644
--- a/utests/compiler_subgroup_scan_exclusive.cpp
+++ b/utests/compiler_subgroup_scan_exclusive.cpp
@@ -312,6 +312,28 @@ void compiler_subgroup_scan_exclusive_add_half(void)
subgroup_generic(WG_SCAN_EXCLUSIVE_ADD, input, expected, true);
}
MAKE_UTEST_FROM_FUNCTION(compiler_subgroup_scan_exclusive_add_half);
+void compiler_subgroup_scan_exclusive_add_short(void)
+{
+ if(!cl_check_subgroups())
+ return;
+ cl_short *input = NULL;
+ cl_short *expected = NULL;
+ OCL_CREATE_KERNEL_FROM_FILE("compiler_subgroup_scan_exclusive",
+ "compiler_subgroup_scan_exclusive_add_short");
+ subgroup_generic(WG_SCAN_EXCLUSIVE_ADD, input, expected);
+}
+MAKE_UTEST_FROM_FUNCTION(compiler_subgroup_scan_exclusive_add_short);
+void compiler_subgroup_scan_exclusive_add_ushort(void)
+{
+ if(!cl_check_subgroups())
+ return;
+ cl_ushort *input = NULL;
+ cl_ushort *expected = NULL;
+ OCL_CREATE_KERNEL_FROM_FILE("compiler_subgroup_scan_exclusive",
+ "compiler_subgroup_scan_exclusive_add_ushort");
+ subgroup_generic(WG_SCAN_EXCLUSIVE_ADD, input, expected);
+}
+MAKE_UTEST_FROM_FUNCTION(compiler_subgroup_scan_exclusive_add_ushort);
/*
* Workgroup scan_exclusive max utest functions
@@ -385,6 +407,28 @@ void compiler_subgroup_scan_exclusive_max_half(void)
subgroup_generic(WG_SCAN_EXCLUSIVE_MAX, input, expected, true);
}
MAKE_UTEST_FROM_FUNCTION(compiler_subgroup_scan_exclusive_max_half);
+void compiler_subgroup_scan_exclusive_max_short(void)
+{
+ if(!cl_check_subgroups())
+ return;
+ cl_short *input = NULL;
+ cl_short *expected = NULL;
+ OCL_CREATE_KERNEL_FROM_FILE("compiler_subgroup_scan_exclusive",
+ "compiler_subgroup_scan_exclusive_max_short");
+ subgroup_generic(WG_SCAN_EXCLUSIVE_MAX, input, expected);
+}
+MAKE_UTEST_FROM_FUNCTION(compiler_subgroup_scan_exclusive_max_short);
+void compiler_subgroup_scan_exclusive_max_ushort(void)
+{
+ if(!cl_check_subgroups())
+ return;
+ cl_ushort *input = NULL;
+ cl_ushort *expected = NULL;
+ OCL_CREATE_KERNEL_FROM_FILE("compiler_subgroup_scan_exclusive",
+ "compiler_subgroup_scan_exclusive_max_ushort");
+ subgroup_generic(WG_SCAN_EXCLUSIVE_MAX, input, expected);
+}
+MAKE_UTEST_FROM_FUNCTION(compiler_subgroup_scan_exclusive_max_ushort);
/*
* Workgroup scan_exclusive min utest functions
@@ -458,3 +502,25 @@ void compiler_subgroup_scan_exclusive_min_half(void)
subgroup_generic(WG_SCAN_EXCLUSIVE_MIN, input, expected, true);
}
MAKE_UTEST_FROM_FUNCTION(compiler_subgroup_scan_exclusive_min_half);
+void compiler_subgroup_scan_exclusive_min_short(void)
+{
+ if(!cl_check_subgroups())
+ return;
+ cl_short *input = NULL;
+ cl_short *expected = NULL;
+ OCL_CREATE_KERNEL_FROM_FILE("compiler_subgroup_scan_exclusive",
+ "compiler_subgroup_scan_exclusive_min_short");
+ subgroup_generic(WG_SCAN_EXCLUSIVE_MIN, input, expected);
+}
+MAKE_UTEST_FROM_FUNCTION(compiler_subgroup_scan_exclusive_min_short);
+void compiler_subgroup_scan_exclusive_min_ushort(void)
+{
+ if(!cl_check_subgroups())
+ return;
+ cl_ushort *input = NULL;
+ cl_ushort *expected = NULL;
+ OCL_CREATE_KERNEL_FROM_FILE("compiler_subgroup_scan_exclusive",
+ "compiler_subgroup_scan_exclusive_min_ushort");
+ subgroup_generic(WG_SCAN_EXCLUSIVE_MIN, input, expected);
+}
+MAKE_UTEST_FROM_FUNCTION(compiler_subgroup_scan_exclusive_min_ushort);
diff --git a/utests/compiler_subgroup_scan_inclusive.cpp b/utests/compiler_subgroup_scan_inclusive.cpp
index 0f0df1c..8f8c264 100644
--- a/utests/compiler_subgroup_scan_inclusive.cpp
+++ b/utests/compiler_subgroup_scan_inclusive.cpp
@@ -298,6 +298,28 @@ void compiler_subgroup_scan_inclusive_add_half(void)
subgroup_generic(WG_SCAN_INCLUSIVE_ADD, input, expected, true);
}
MAKE_UTEST_FROM_FUNCTION(compiler_subgroup_scan_inclusive_add_half);
+void compiler_subgroup_scan_inclusive_add_short(void)
+{
+ if(!cl_check_subgroups())
+ return;
+ cl_short *input = NULL;
+ cl_short *expected = NULL;
+ OCL_CREATE_KERNEL_FROM_FILE("compiler_subgroup_scan_inclusive",
+ "compiler_subgroup_scan_inclusive_add_short");
+ subgroup_generic(WG_SCAN_INCLUSIVE_ADD, input, expected);
+}
+MAKE_UTEST_FROM_FUNCTION(compiler_subgroup_scan_inclusive_add_short);
+void compiler_subgroup_scan_inclusive_add_ushort(void)
+{
+ if(!cl_check_subgroups())
+ return;
+ cl_ushort *input = NULL;
+ cl_ushort *expected = NULL;
+ OCL_CREATE_KERNEL_FROM_FILE("compiler_subgroup_scan_inclusive",
+ "compiler_subgroup_scan_inclusive_add_ushort");
+ subgroup_generic(WG_SCAN_INCLUSIVE_ADD, input, expected);
+}
+MAKE_UTEST_FROM_FUNCTION(compiler_subgroup_scan_inclusive_add_ushort);
/*
* Workgroup scan_inclusive max utest functions
@@ -371,6 +393,28 @@ void compiler_subgroup_scan_inclusive_max_half(void)
subgroup_generic(WG_SCAN_INCLUSIVE_MAX, input, expected, true);
}
MAKE_UTEST_FROM_FUNCTION(compiler_subgroup_scan_inclusive_max_half);
+void compiler_subgroup_scan_inclusive_max_short(void)
+{
+ if(!cl_check_subgroups())
+ return;
+ cl_short *input = NULL;
+ cl_short *expected = NULL;
+ OCL_CREATE_KERNEL_FROM_FILE("compiler_subgroup_scan_inclusive",
+ "compiler_subgroup_scan_inclusive_max_short");
+ subgroup_generic(WG_SCAN_INCLUSIVE_MAX, input, expected);
+}
+MAKE_UTEST_FROM_FUNCTION(compiler_subgroup_scan_inclusive_max_short);
+void compiler_subgroup_scan_inclusive_max_ushort(void)
+{
+ if(!cl_check_subgroups())
+ return;
+ cl_ushort *input = NULL;
+ cl_ushort *expected = NULL;
+ OCL_CREATE_KERNEL_FROM_FILE("compiler_subgroup_scan_inclusive",
+ "compiler_subgroup_scan_inclusive_max_ushort");
+ subgroup_generic(WG_SCAN_INCLUSIVE_MAX, input, expected);
+}
+MAKE_UTEST_FROM_FUNCTION(compiler_subgroup_scan_inclusive_max_ushort);
/*
* Workgroup scan_inclusive min utest functions
@@ -444,3 +488,25 @@ void compiler_subgroup_scan_inclusive_min_half(void)
subgroup_generic(WG_SCAN_INCLUSIVE_MIN, input, expected, true);
}
MAKE_UTEST_FROM_FUNCTION(compiler_subgroup_scan_inclusive_min_half);
+void compiler_subgroup_scan_inclusive_min_short(void)
+{
+ if(!cl_check_subgroups())
+ return;
+ cl_short *input = NULL;
+ cl_short *expected = NULL;
+ OCL_CREATE_KERNEL_FROM_FILE("compiler_subgroup_scan_inclusive",
+ "compiler_subgroup_scan_inclusive_min_short");
+ subgroup_generic(WG_SCAN_INCLUSIVE_MIN, input, expected);
+}
+MAKE_UTEST_FROM_FUNCTION(compiler_subgroup_scan_inclusive_min_short);
+void compiler_subgroup_scan_inclusive_min_ushort(void)
+{
+ if(!cl_check_subgroups())
+ return;
+ cl_ushort *input = NULL;
+ cl_ushort *expected = NULL;
+ OCL_CREATE_KERNEL_FROM_FILE("compiler_subgroup_scan_inclusive",
+ "compiler_subgroup_scan_inclusive_min_ushort");
+ subgroup_generic(WG_SCAN_INCLUSIVE_MIN, input, expected);
+}
+MAKE_UTEST_FROM_FUNCTION(compiler_subgroup_scan_inclusive_min_ushort);
--
2.7.4
More information about the Beignet
mailing list