[Beignet] [PATCH V2 14/17] Utest: Add workgroup all/any tests
Grigore Lupescu
grigore.lupescu at intel.com
Mon Apr 11 14:40:37 UTC 2016
From: Grigore Lupescu <grigore.lupescu at intel.com>
Added the following unit tests:
compiler_workgroup_any
compiler_workgroup_all
Signed-off-by: Grigore Lupescu <grigore.lupescu at intel.com>
---
kernels/compiler_workgroup_reduce.cl | 14 +++++
utests/compiler_workgroup_reduce.cpp | 114 ++++++++++++++++++++++++++++-------
2 files changed, 107 insertions(+), 21 deletions(-)
diff --git a/kernels/compiler_workgroup_reduce.cl b/kernels/compiler_workgroup_reduce.cl
index 34fd371..e5e53d8 100644
--- a/kernels/compiler_workgroup_reduce.cl
+++ b/kernels/compiler_workgroup_reduce.cl
@@ -1,6 +1,20 @@
/*
* Workgroup reduce add functions
*/
+kernel void compiler_workgroup_any(global int *src, global int *dst) {
+ int val = src[get_global_id(0)];
+ int predicate = work_group_any(val);
+ dst[get_global_id(0)] = predicate;
+}
+kernel void compiler_workgroup_all(global int *src, global int *dst) {
+ char val = src[get_global_id(0)];
+ int predicate = work_group_all(val);
+ dst[get_global_id(0)] = predicate;
+}
+
+/*
+ * Workgroup reduce add functions
+ */
kernel void compiler_workgroup_reduce_add_char(global char *src, global char *dst) {
char val = src[get_global_id(0)];
char sum = work_group_reduce_add(val);
diff --git a/utests/compiler_workgroup_reduce.cpp b/utests/compiler_workgroup_reduce.cpp
index 7fced74..f796d5c 100644
--- a/utests/compiler_workgroup_reduce.cpp
+++ b/utests/compiler_workgroup_reduce.cpp
@@ -9,19 +9,24 @@
using namespace std;
+/* set to 1 for debug, output of input-expected data */
+#define DEBUG_STDOUT 0
+
/* NDRANGE */
-#define WG_GLOBAL_SIZE 64
-#define WG_LOCAL_SIZE 32
+#define WG_GLOBAL_SIZE 64
+#define WG_LOCAL_SIZE 32
enum WG_FUNCTION
{
+ WG_ANY,
+ WG_ALL,
WG_REDUCE_ADD,
WG_REDUCE_MIN,
WG_REDUCE_MAX
};
/*
- * Generic compute-expected on CPU function for any workgroup type
+ * Generic compute-expected function for op REDUCE/ANY/ALL
* and any variable type
*/
template<class T>
@@ -29,7 +34,23 @@ static void compute_expected(WG_FUNCTION wg_func,
T* input,
T* expected)
{
- if(wg_func == WG_REDUCE_ADD)
+ if(wg_func == WG_ANY)
+ {
+ T wg_predicate = input[0];
+ for(uint32_t i = 1; i < WG_LOCAL_SIZE; i++)
+ wg_predicate = (int)wg_predicate | (int)input[i];
+ for(uint32_t i = 0; i < WG_LOCAL_SIZE; i++)
+ expected[i] = wg_predicate;
+ }
+ else if(wg_func == WG_ALL)
+ {
+ T wg_predicate = input[0];
+ for(uint32_t i = 1; i < WG_LOCAL_SIZE; i++)
+ wg_predicate = (int)wg_predicate & (int)input[i];
+ for(uint32_t i = 0; i < WG_LOCAL_SIZE; i++)
+ expected[i] = wg_predicate;
+ }
+ else if(wg_func == WG_REDUCE_ADD)
{
T wg_sum = input[0];
for(uint32_t i = 1; i < WG_LOCAL_SIZE; i++)
@@ -56,7 +77,7 @@ static void compute_expected(WG_FUNCTION wg_func,
}
/*
- * Generic input-expected generate function for any workgroup type
+ * Generic input-expected generate function for op REDUCE/ANY/ALL
* and any variable type
*/
template<class T>
@@ -67,7 +88,7 @@ static void generate_data(WG_FUNCTION wg_func,
input = new T[WG_GLOBAL_SIZE];
expected = new T[WG_GLOBAL_SIZE];
- /* base value for all datatypes */
+ /* base value for all data types */
T base_val = (long)7 << (sizeof(T) * 5 - 3);
/* seed for random inputs */
@@ -76,28 +97,49 @@ static void generate_data(WG_FUNCTION wg_func,
/* generate inputs and expected values */
for(uint32_t gid = 0; gid < WG_GLOBAL_SIZE; gid += WG_LOCAL_SIZE)
{
- /* input values */
+#if DEBUG_STDOUT
cout << endl << "IN: " << endl;
- for(uint32_t lid = 0; lid < WG_LOCAL_SIZE; lid++){
- input[gid + lid] = (rand() % 2 - 1) * base_val + (rand() % 112);
+#endif
+
+ /* input values */
+ for(uint32_t lid = 0; lid < WG_LOCAL_SIZE; lid++)
+ {
+ /* initially 0, augment after */
+ input[gid + lid] = 0;
+
+ /* check all data types, test ideal for QWORD types */
+ input[gid + lid] += ((rand() % 2 - 1) * base_val);
+ /* add trailing random bits, tests GENERAL cases */
+ input[gid + lid] += (rand() % 112);
+ /* always last bit is 1, ideal test ALL/ANY */
+ input[gid + lid] = (T)((long)input[gid + lid] | (long)1);
+
+#if DEBUG_STDOUT
+ /* output generated input */
cout << setw(4) << input[gid + lid] << ", " ;
if((lid + 1) % 8 == 0)
cout << endl;
+#endif
}
/* expected values */
- cout << endl << "EXP: " << endl;
compute_expected(wg_func, input + gid, expected + gid);
- for(uint32_t lid = 0; lid < WG_LOCAL_SIZE; lid++){
+
+#if DEBUG_STDOUT
+ /* output expected input */
+ cout << endl << "EXP: " << endl;
+ for(uint32_t lid = 0; lid < WG_LOCAL_SIZE; lid++) {
cout << setw(4) << expected[gid + lid] << ", " ;
if((lid + 1) % 8 == 0)
cout << endl;
}
+#endif
+
}
}
/*
- * Generic workgroup utest function for any workgroup type
+ * Generic workgroup utest function for op REDUCE/ANY/ALL
* and any variable type
*/
template<class T>
@@ -108,7 +150,7 @@ static void workgroup_generic(WG_FUNCTION wg_func,
/* input and expected data */
generate_data(wg_func, input, expected);
- /* prepare input for datatype */
+ /* prepare input for data type */
OCL_CREATE_BUFFER(buf[0], 0, WG_GLOBAL_SIZE * sizeof(T), NULL);
OCL_CREATE_BUFFER(buf[1], 0, WG_GLOBAL_SIZE * sizeof(T), NULL);
OCL_SET_ARG(0, sizeof(cl_mem), &buf[0]);
@@ -124,25 +166,55 @@ static void workgroup_generic(WG_FUNCTION wg_func,
locals[0] = WG_LOCAL_SIZE;
OCL_NDRANGE(1);
- /* check if mistmatch */
+ /* check if mismatch */
OCL_MAP_BUFFER(1);
- uint32_t mistmatches = 0;
- cout << endl << endl << "CHECK" << endl;
+ uint32_t mismatches = 0;
+
for (uint32_t i = 0; i < WG_GLOBAL_SIZE; i++)
- if(((T *)buf_data[1])[i] != *(expected + i)){
+ if(((T *)buf_data[1])[i] != *(expected + i))
+ {
+ /* found mismatch, increment */
+ mismatches++;
+
+#if DEBUG_STDOUT
+ /* output mismatch */
cout << "Err at " << i << ", " <<
((T *)buf_data[1])[i] << " != " << *(expected + i) << endl;
- mistmatches++;
+#endif
}
- cout << "MISTMATCHES " << mistmatches << endl;
- cout << std::flush;
+#if DEBUG_STDOUT
+ /* output mismatch count */
+ cout << "mismatches " << mismatches << endl;
+#endif
+
OCL_UNMAP_BUFFER(1);
- OCL_ASSERT(mistmatches == 0);
+ OCL_ASSERT(mismatches == 0);
}
/*
+ * Workgroup any/all utest functions
+ */
+void compiler_workgroup_any(void)
+{
+ cl_int *input = NULL;
+ cl_int *expected = NULL;
+ OCL_CREATE_KERNEL_FROM_FILE("compiler_workgroup_reduce",
+ "compiler_workgroup_any");
+ workgroup_generic(WG_ANY, input, expected);
+}
+MAKE_UTEST_FROM_FUNCTION(compiler_workgroup_any);
+void compiler_workgroup_all(void)
+{
+ cl_int *input = NULL;
+ cl_int *expected = NULL;
+ OCL_CREATE_KERNEL_FROM_FILE("compiler_workgroup_reduce",
+ "compiler_workgroup_all");
+ workgroup_generic(WG_ALL, input, expected);
+}
+MAKE_UTEST_FROM_FUNCTION(compiler_workgroup_all);
+/*
* Workgroup reduce add utest functions
*/
void compiler_workgroup_reduce_add_char(void)
--
2.5.0
More information about the Beignet
mailing list