From: Grigore Lupescu <grigore.lupescu at intel.com>
Signed-off-by: Grigore Lupescu <grigore.lupescu at intel.com>
---
kernels/compiler_workgroup_scan_exclusive.cl | 170 ++++++++++++
utests/CMakeLists.txt | 1 +
utests/compiler_workgroup_scan_exclusive.cpp | 395 +++++++++++++++++++++++++++
3 files changed, 566 insertions(+)
create mode 100644 kernels/compiler_workgroup_scan_exclusive.cl
create mode 100644 utests/compiler_workgroup_scan_exclusive.cpp
diff --git a/kernels/compiler_workgroup_scan_exclusive.cl
b/kernels/compiler_workgroup_scan_exclusive.cl
new file mode 100644
index 0000000..fa0f81a
--- /dev/null
+++ b/kernels/compiler_workgroup_scan_exclusive.cl
@@ -0,0 +1,170 @@
+/*
+ * Workgroup scan exclusive add functions
+ */
+kernel void compiler_workgroup_scan_exclusive_add_char(global char *src,
global char *dst) {
+ char val = src[get_global_id(0)];
+ char sum = work_group_scan_exclusive_add(val);
+ dst[get_global_id(0)] = sum;
+}
+
+kernel void compiler_workgroup_scan_exclusive_add_uchar(global uchar *src,
global uchar *dst) {
+ uchar val = src[get_global_id(0)];
+ uchar sum = work_group_scan_exclusive_add(val);
+ dst[get_global_id(0)] = sum;
+}
+
+kernel void compiler_workgroup_scan_exclusive_add_short(global short *src,
global short *dst) {
+ short val = src[get_global_id(0)];
+ short sum = work_group_scan_exclusive_add(val);
+ dst[get_global_id(0)] = sum;
+}
+
+kernel void compiler_workgroup_scan_exclusive_add_ushort(global ushort *src,
global ushort *dst) {
+ ushort val = src[get_global_id(0)];
+ ushort sum = work_group_scan_exclusive_add(val);
+ dst[get_global_id(0)] = sum;
+}
+
+kernel void compiler_workgroup_scan_exclusive_add_int(global int *src, global
int *dst) {
+ int val = src[get_global_id(0)];
+ int sum = work_group_scan_exclusive_add(val);
+ dst[get_global_id(0)] = sum;
+}
+
+kernel void compiler_workgroup_scan_exclusive_add_uint(global uint *src,
global uint *dst) {
+ uint val = src[get_global_id(0)];
+ uint sum = work_group_scan_exclusive_add(val);
+ dst[get_global_id(0)] = sum;
+}
+
+kernel void compiler_workgroup_scan_exclusive_add_long(global long *src,
global long *dst) {
+ long val = src[get_global_id(0)];
+ long sum = work_group_scan_exclusive_add(val);
+ dst[get_global_id(0)] = sum;
+}
+
+kernel void compiler_workgroup_scan_exclusive_add_ulong(global ulong *src,
global ulong *dst) {
+ ulong val = src[get_global_id(0)];
+ ulong sum = work_group_scan_exclusive_add(val);
+ dst[get_global_id(0)] = sum;
+}
+
+kernel void compiler_workgroup_scan_exclusive_add_float(global float *src,
global float *dst) {
+ float val = src[get_global_id(0)];
+ float sum = work_group_scan_exclusive_add(val);
+ dst[get_global_id(0)] = sum;
+}
+
+/*
+ * Workgroup scan exclusive max functions
+ */
+kernel void compiler_workgroup_scan_exclusive_max_char(global char *src,
global char *dst) {
+ char val = src[get_global_id(0)];
+ char sum = work_group_scan_exclusive_max(val);
+ dst[get_global_id(0)] = sum;
+}
+
+kernel void compiler_workgroup_scan_exclusive_max_uchar(global uchar *src,
global uchar *dst) {
+ uchar val = src[get_global_id(0)];
+ uchar sum = work_group_scan_exclusive_max(val);
+ dst[get_global_id(0)] = sum;
+}
+
+kernel void compiler_workgroup_scan_exclusive_max_short(global short *src,
global short *dst) {
+ short val = src[get_global_id(0)];
+ short sum = work_group_scan_exclusive_max(val);
+ dst[get_global_id(0)] = sum;
+}
+
+kernel void compiler_workgroup_scan_exclusive_max_ushort(global ushort *src,
global ushort *dst) {
+ ushort val = src[get_global_id(0)];
+ ushort sum = work_group_scan_exclusive_max(val);
+ dst[get_global_id(0)] = sum;
+}
+
+kernel void compiler_workgroup_scan_exclusive_max_int(global int *src, global
int *dst) {
+ int val = src[get_global_id(0)];
+ int sum = work_group_scan_exclusive_max(val);
+ dst[get_global_id(0)] = sum;
+}
+
+kernel void compiler_workgroup_scan_exclusive_max_uint(global uint *src,
global uint *dst) {
+ uint val = src[get_global_id(0)];
+ uint sum = work_group_scan_exclusive_max(val);
+ dst[get_global_id(0)] = sum;
+}
+
+kernel void compiler_workgroup_scan_exclusive_max_long(global long *src,
global long *dst) {
+ long val = src[get_global_id(0)];
+ long sum = work_group_scan_exclusive_max(val);
+ dst[get_global_id(0)] = sum;
+}
+
+kernel void compiler_workgroup_scan_exclusive_max_ulong(global ulong *src,
global ulong *dst) {
+ ulong val = src[get_global_id(0)];
+ ulong sum = work_group_scan_exclusive_max(val);
+ dst[get_global_id(0)] = sum;
+}
+
+kernel void compiler_workgroup_scan_exclusive_max_float(global float *src,
global float *dst) {
+ float val = src[get_global_id(0)];
+ float sum = work_group_scan_exclusive_max(val);
+ dst[get_global_id(0)] = sum;
+}
+
+/*
+ * Workgroup scan exclusive min functions
+ */
+kernel void compiler_workgroup_scan_exclusive_min_char(global char *src,
global char *dst) {
+ char val = src[get_global_id(0)];
+ char sum = work_group_scan_exclusive_min(val);
+ dst[get_global_id(0)] = sum;
+}
+
+kernel void compiler_workgroup_scan_exclusive_min_uchar(global uchar *src,
global uchar *dst) {
+ uchar val = src[get_global_id(0)];
+ uchar sum = work_group_scan_exclusive_min(val);
+ dst[get_global_id(0)] = sum;
+}
+
+kernel void compiler_workgroup_scan_exclusive_min_short(global short *src,
global short *dst) {
+ short val = src[get_global_id(0)];
+ short sum = work_group_scan_exclusive_min(val);
+ dst[get_global_id(0)] = sum;
+}
+
+kernel void compiler_workgroup_scan_exclusive_min_ushort(global ushort *src,
global ushort *dst) {
+ ushort val = src[get_global_id(0)];
+ ushort sum = work_group_scan_exclusive_min(val);
+ dst[get_global_id(0)] = sum;
+}
+
+kernel void compiler_workgroup_scan_exclusive_min_int(global int *src, global
int *dst) {
+ int val = src[get_global_id(0)];
+ int sum = work_group_scan_exclusive_min(val);
+ dst[get_global_id(0)] = sum;
+}
+
+kernel void compiler_workgroup_scan_exclusive_min_uint(global uint *src,
global uint *dst) {
+ uint val = src[get_global_id(0)];
+ uint sum = work_group_scan_exclusive_min(val);
+ dst[get_global_id(0)] = sum;
+}
+
+kernel void compiler_workgroup_scan_exclusive_min_long(global long *src,
global long *dst) {
+ long val = src[get_global_id(0)];
+ long sum = work_group_scan_exclusive_min(val);
+ dst[get_global_id(0)] = sum;
+}
+
+kernel void compiler_workgroup_scan_exclusive_min_ulong(global ulong *src,
global ulong *dst) {
+ ulong val = src[get_global_id(0)];
+ ulong sum = work_group_scan_exclusive_min(val);
+ dst[get_global_id(0)] = sum;
+}
+
+kernel void compiler_workgroup_scan_exclusive_min_float(global float *src,
global float *dst) {
+ float val = src[get_global_id(0)];
+ float sum = work_group_scan_exclusive_min(val);
+ dst[get_global_id(0)] = sum;
+}
diff --git a/utests/CMakeLists.txt b/utests/CMakeLists.txt
index b849420..e58dea7 100644
--- a/utests/CMakeLists.txt
+++ b/utests/CMakeLists.txt
@@ -126,6 +126,7 @@ set (utests_sources
compiler_workgroup_broadcast.cpp
compiler_workgroup_predicate.cpp
compiler_workgroup_reduce.cpp
+ compiler_workgroup_scan_exclusive.cpp
compiler_async_stride_copy.cpp
compiler_insn_selection_min.cpp
compiler_insn_selection_max.cpp
diff --git a/utests/compiler_workgroup_scan_exclusive.cpp
b/utests/compiler_workgroup_scan_exclusive.cpp
new file mode 100644
index 0000000..e8644ee
--- /dev/null
+++ b/utests/compiler_workgroup_scan_exclusive.cpp
@@ -0,0 +1,395 @@
+#include <cstdint>
+#include <cstring>
+#include <iostream>
+#include <cstdlib>
+#include <iomanip>
+#include <algorithm>
+
+#include "utest_helper.hpp"
+
+using namespace std;
+
+/* NDRANGE */
+#define WG_GLOBAL_SIZE 64
+#define WG_LOCAL_SIZE 32
+
+enum WG_FUNCTION
+{
+ WG_SCAN_EXCLUSIVE_ADD,
+ WG_SCAN_EXCLUSIVE_MAX,
+ WG_SCAN_EXCLUSIVE_MIN
+};
+
+/*
+ * Generic compute-expected on CPU function for any workgroup type
+ * and any variable type
+ */
+template<class T>
+static void compute_expected(WG_FUNCTION wg_func,
+ T* input,
+ T* expected)
+{
+ if(wg_func == WG_SCAN_EXCLUSIVE_ADD)
+ {
+ expected[0] = 0;
+ expected[1] = input[0];
+ for(uint32_t i = 2; i < WG_LOCAL_SIZE; i++)
+ expected[i] = input[i] + expected[i - 1];
+ }
+ else if(wg_func == WG_SCAN_EXCLUSIVE_MAX)
+ {
+ expected[0] = 0;
+ expected[1] = input[0];
+ for(uint32_t i = 2; i < WG_LOCAL_SIZE; i++)
+ expected[i] = max(input[i], expected[i - 1]);
+ }
+ else if(wg_func == WG_SCAN_EXCLUSIVE_MIN)
+ {
+ expected[0] = 0;
+ expected[1] = input[0];
+ for(uint32_t i = 2; i < WG_LOCAL_SIZE; i++)
+ expected[i] = min(input[i], expected[i - 1]);
+ }
+}
+
+/*
+ * Generic input-expected generate function for any workgroup type
+ * and any variable type
+ */
+template<class T>
+static void generate_data(WG_FUNCTION wg_func,
+ T* &input,
+ T* &expected)
+{
+ input = new T[WG_GLOBAL_SIZE];
+ expected = new T[WG_GLOBAL_SIZE];
+
+ /* base value for all datatypes */
+ T base_val = (long)7 << (sizeof(T) * 5 - 3);
+
+ /* seed for random inputs */
+ srand (time(NULL));
+
+ /* generate inputs and expected values */
+ for(uint32_t gid = 0; gid < WG_GLOBAL_SIZE; gid += WG_LOCAL_SIZE)
+ {
+ /* input values */
+ cout << endl << "IN: " << endl;
+ for(uint32_t lid = 0; lid < WG_LOCAL_SIZE; lid++){
+ input[gid + lid] = (rand() % 2 - 1) * base_val + (rand() % 112);
+ cout << setw(4) << input[gid + lid] << ", " ;
+ if((lid + 1) % 8 == 0)
+ cout << endl;
+ }
+
+ /* expected values */
+ cout << endl << "EXP: " << endl;
+ compute_expected(wg_func, input + gid, expected + gid);
+ for(uint32_t lid = 0; lid < WG_LOCAL_SIZE; lid++){
+ cout << setw(4) << expected[gid + lid] << ", " ;
+ if((lid + 1) % 8 == 0)
+ cout << endl;
+ }
+ }
+}
+
+/*
+ * Generic workgroup utest function for any workgroup type
+ * and any variable type
+ */
+template<class T>
+static void workgroup_generic(WG_FUNCTION wg_func,
+ T* input,
+ T* expected)
+{
+ /* input and expected data */
+ generate_data(wg_func, input, expected);
+
+ /* prepare input for datatype */
+ 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]);
+ OCL_SET_ARG(1, sizeof(cl_mem), &buf[1]);
+
+ /* set input data for GPU */
+ OCL_MAP_BUFFER(0);
+ memcpy(buf_data[0], input, WG_GLOBAL_SIZE * sizeof(T));
+ OCL_UNMAP_BUFFER(0);
+
+ /* run the kernel on GPU */
+ globals[0] = WG_GLOBAL_SIZE;
+ locals[0] = WG_LOCAL_SIZE;
+ OCL_NDRANGE(1);
+
+ /* check if mistmatch */
+ OCL_MAP_BUFFER(1);
+ uint32_t mistmatches = 0;
+ cout << endl << endl << "CHECK" << endl;
+ for (uint32_t i = 0; i < WG_GLOBAL_SIZE; i++)
+ if(((T *)buf_data[1])[i] != *(expected + i)){
+ cout << "Err at " << i << ", " <<
+ ((T *)buf_data[1])[i] << " != " << *(expected + i) << endl;
+ mistmatches++;
+ }
+ cout << "MISTMATCHES " << mistmatches << endl;
+
+ cout << std::flush;
+ OCL_UNMAP_BUFFER(1);
+
+ OCL_ASSERT(mistmatches == 0);
+}
+
+/*
+ * Workgroup scan_exclusive add utest functions
+ */
+void compiler_workgroup_scan_exclusive_add_char(void)
+{
+ cl_char *input = NULL;
+ cl_char *expected = NULL;
+ OCL_CREATE_KERNEL_FROM_FILE("compiler_workgroup_scan_exclusive",
+ "compiler_workgroup_scan_exclusive_add_char");
+ workgroup_generic(WG_SCAN_EXCLUSIVE_ADD, input, expected);
+}
+MAKE_UTEST_FROM_FUNCTION(compiler_workgroup_scan_exclusive_add_char);
+void compiler_workgroup_scan_exclusive_add_uchar(void)
+{
+ cl_uchar *input = NULL;
+ cl_uchar *expected = NULL;
+ OCL_CREATE_KERNEL_FROM_FILE("compiler_workgroup_scan_exclusive",
+ "compiler_workgroup_scan_exclusive_add_uchar");
+ workgroup_generic(WG_SCAN_EXCLUSIVE_ADD, input, expected);
+}
+MAKE_UTEST_FROM_FUNCTION(compiler_workgroup_scan_exclusive_add_uchar);
+void compiler_workgroup_scan_exclusive_add_short(void)
+{
+ cl_short *input = NULL;
+ cl_short *expected = NULL;
+ OCL_CREATE_KERNEL_FROM_FILE("compiler_workgroup_scan_exclusive",
+ "compiler_workgroup_scan_exclusive_add_short");
+ workgroup_generic(WG_SCAN_EXCLUSIVE_ADD, input, expected);
+}
+MAKE_UTEST_FROM_FUNCTION(compiler_workgroup_scan_exclusive_add_short);
+void compiler_workgroup_scan_exclusive_add_ushort(void)
+{
+ cl_ushort *input = NULL;
+ cl_ushort *expected = NULL;
+ OCL_CREATE_KERNEL_FROM_FILE("compiler_workgroup_scan_exclusive",
+ "compiler_workgroup_scan_exclusive_add_ushort");
+ workgroup_generic(WG_SCAN_EXCLUSIVE_ADD, input, expected);
+}
+MAKE_UTEST_FROM_FUNCTION(compiler_workgroup_scan_exclusive_add_ushort);
+void compiler_workgroup_scan_exclusive_add_int(void)
+{
+ cl_int *input = NULL;
+ cl_int *expected = NULL;
+ OCL_CREATE_KERNEL_FROM_FILE("compiler_workgroup_scan_exclusive",
+ "compiler_workgroup_scan_exclusive_add_int");
+ workgroup_generic(WG_SCAN_EXCLUSIVE_ADD, input, expected);
+}
+MAKE_UTEST_FROM_FUNCTION(compiler_workgroup_scan_exclusive_add_int);
+void compiler_workgroup_scan_exclusive_add_uint(void)
+{
+ cl_uint *input = NULL;
+ cl_uint *expected = NULL;
+ OCL_CREATE_KERNEL_FROM_FILE("compiler_workgroup_scan_exclusive",
+ "compiler_workgroup_scan_exclusive_add_uint");
+ workgroup_generic(WG_SCAN_EXCLUSIVE_ADD, input, expected);
+}
+MAKE_UTEST_FROM_FUNCTION(compiler_workgroup_scan_exclusive_add_uint);
+void compiler_workgroup_scan_exclusive_add_long(void)
+{
+ cl_long *input = NULL;
+ cl_long *expected = NULL;
+ OCL_CREATE_KERNEL_FROM_FILE("compiler_workgroup_scan_exclusive",
+ "compiler_workgroup_scan_exclusive_add_long");
+ workgroup_generic(WG_SCAN_EXCLUSIVE_ADD, input, expected);
+}
+MAKE_UTEST_FROM_FUNCTION(compiler_workgroup_scan_exclusive_add_long);
+void compiler_workgroup_scan_exclusive_add_ulong(void)
+{
+ cl_ulong *input = NULL;
+ cl_ulong *expected = NULL;
+ OCL_CREATE_KERNEL_FROM_FILE("compiler_workgroup_scan_exclusive",
+ "compiler_workgroup_scan_exclusive_add_ulong");
+ workgroup_generic(WG_SCAN_EXCLUSIVE_ADD, input, expected);
+}
+MAKE_UTEST_FROM_FUNCTION(compiler_workgroup_scan_exclusive_add_ulong);
+void compiler_workgroup_scan_exclusive_add_float(void)
+{
+ cl_float *input = NULL;
+ cl_float *expected = NULL;
+ OCL_CREATE_KERNEL_FROM_FILE("compiler_workgroup_scan_exclusive",
+ "compiler_workgroup_scan_exclusive_add_float");
+ workgroup_generic(WG_SCAN_EXCLUSIVE_ADD, input, expected);
+}
+MAKE_UTEST_FROM_FUNCTION(compiler_workgroup_scan_exclusive_add_float);
+
+/*
+ * Workgroup scan_exclusive max utest functions
+ */
+void compiler_workgroup_scan_exclusive_max_char(void)
+{
+ cl_char *input = NULL;
+ cl_char *expected = NULL;
+ OCL_CREATE_KERNEL_FROM_FILE("compiler_workgroup_scan_exclusive",
+ "compiler_workgroup_scan_exclusive_max_char");
+ workgroup_generic(WG_SCAN_EXCLUSIVE_MAX, input, expected);
+}
+MAKE_UTEST_FROM_FUNCTION(compiler_workgroup_scan_exclusive_max_char);
+void compiler_workgroup_scan_exclusive_max_uchar(void)
+{
+ cl_uchar *input = NULL;
+ cl_uchar *expected = NULL;
+ OCL_CREATE_KERNEL_FROM_FILE("compiler_workgroup_scan_exclusive",
+ "compiler_workgroup_scan_exclusive_max_uchar");
+ workgroup_generic(WG_SCAN_EXCLUSIVE_MAX, input, expected);
+}
+MAKE_UTEST_FROM_FUNCTION(compiler_workgroup_scan_exclusive_max_uchar);
+void compiler_workgroup_scan_exclusive_max_short(void)
+{
+ cl_short *input = NULL;
+ cl_short *expected = NULL;
+ OCL_CREATE_KERNEL_FROM_FILE("compiler_workgroup_scan_exclusive",
+ "compiler_workgroup_scan_exclusive_max_short");
+ workgroup_generic(WG_SCAN_EXCLUSIVE_MAX, input, expected);
+}
+MAKE_UTEST_FROM_FUNCTION(compiler_workgroup_scan_exclusive_max_short);
+void compiler_workgroup_scan_exclusive_max_ushort(void)
+{
+ cl_ushort *input = NULL;
+ cl_ushort *expected = NULL;
+ OCL_CREATE_KERNEL_FROM_FILE("compiler_workgroup_scan_exclusive",
+ "compiler_workgroup_scan_exclusive_max_ushort");
+ workgroup_generic(WG_SCAN_EXCLUSIVE_MAX, input, expected);
+}
+MAKE_UTEST_FROM_FUNCTION(compiler_workgroup_scan_exclusive_max_ushort);
+void compiler_workgroup_scan_exclusive_max_int(void)
+{
+ cl_int *input = NULL;
+ cl_int *expected = NULL;
+ OCL_CREATE_KERNEL_FROM_FILE("compiler_workgroup_scan_exclusive",
+ "compiler_workgroup_scan_exclusive_max_int");
+ workgroup_generic(WG_SCAN_EXCLUSIVE_MAX, input, expected);
+}
+MAKE_UTEST_FROM_FUNCTION(compiler_workgroup_scan_exclusive_max_int);
+void compiler_workgroup_scan_exclusive_max_uint(void)
+{
+ cl_uint *input = NULL;
+ cl_uint *expected = NULL;
+ OCL_CREATE_KERNEL_FROM_FILE("compiler_workgroup_scan_exclusive",
+ "compiler_workgroup_scan_exclusive_max_uint");
+ workgroup_generic(WG_SCAN_EXCLUSIVE_MAX, input, expected);
+}
+MAKE_UTEST_FROM_FUNCTION(compiler_workgroup_scan_exclusive_max_uint);
+void compiler_workgroup_scan_exclusive_max_long(void)
+{
+ cl_long *input = NULL;
+ cl_long *expected = NULL;
+ OCL_CREATE_KERNEL_FROM_FILE("compiler_workgroup_scan_exclusive",
+ "compiler_workgroup_scan_exclusive_max_long");
+ workgroup_generic(WG_SCAN_EXCLUSIVE_MAX, input, expected);
+}
+MAKE_UTEST_FROM_FUNCTION(compiler_workgroup_scan_exclusive_max_long);
+void compiler_workgroup_scan_exclusive_max_ulong(void)
+{
+ cl_ulong *input = NULL;
+ cl_ulong *expected = NULL;
+ OCL_CREATE_KERNEL_FROM_FILE("compiler_workgroup_scan_exclusive",
+ "compiler_workgroup_scan_exclusive_max_ulong");
+ workgroup_generic(WG_SCAN_EXCLUSIVE_MAX, input, expected);
+}
+MAKE_UTEST_FROM_FUNCTION(compiler_workgroup_scan_exclusive_max_ulong);
+void compiler_workgroup_scan_exclusive_max_float(void)
+{
+ cl_float *input = NULL;
+ cl_float *expected = NULL;
+ OCL_CREATE_KERNEL_FROM_FILE("compiler_workgroup_scan_exclusive",
+ "compiler_workgroup_scan_exclusive_max_float");
+ workgroup_generic(WG_SCAN_EXCLUSIVE_MAX, input, expected);
+}
+MAKE_UTEST_FROM_FUNCTION(compiler_workgroup_scan_exclusive_max_float);
+
+/*
+ * Workgroup scan_exclusive min utest functions
+ */
+void compiler_workgroup_scan_exclusive_min_char(void)
+{
+ cl_char *input = NULL;
+ cl_char *expected = NULL;
+ OCL_CREATE_KERNEL_FROM_FILE("compiler_workgroup_scan_exclusive",
+ "compiler_workgroup_scan_exclusive_min_char");
+ workgroup_generic(WG_SCAN_EXCLUSIVE_MIN, input, expected);
+}
+MAKE_UTEST_FROM_FUNCTION(compiler_workgroup_scan_exclusive_min_char);
+void compiler_workgroup_scan_exclusive_min_uchar(void)
+{
+ cl_uchar *input = NULL;
+ cl_uchar *expected = NULL;
+ OCL_CREATE_KERNEL_FROM_FILE("compiler_workgroup_scan_exclusive",
+ "compiler_workgroup_scan_exclusive_min_uchar");
+ workgroup_generic(WG_SCAN_EXCLUSIVE_MIN, input, expected);
+}
+MAKE_UTEST_FROM_FUNCTION(compiler_workgroup_scan_exclusive_min_uchar);
+void compiler_workgroup_scan_exclusive_min_short(void)
+{
+ cl_short *input = NULL;
+ cl_short *expected = NULL;
+ OCL_CREATE_KERNEL_FROM_FILE("compiler_workgroup_scan_exclusive",
+ "compiler_workgroup_scan_exclusive_min_short");
+ workgroup_generic(WG_SCAN_EXCLUSIVE_MIN, input, expected);
+}
+MAKE_UTEST_FROM_FUNCTION(compiler_workgroup_scan_exclusive_min_short);
+void compiler_workgroup_scan_exclusive_min_ushort(void)
+{
+ cl_ushort *input = NULL;
+ cl_ushort *expected = NULL;
+ OCL_CREATE_KERNEL_FROM_FILE("compiler_workgroup_scan_exclusive",
+ "compiler_workgroup_scan_exclusive_min_ushort");
+ workgroup_generic(WG_SCAN_EXCLUSIVE_MIN, input, expected);
+}
+MAKE_UTEST_FROM_FUNCTION(compiler_workgroup_scan_exclusive_min_ushort);
+void compiler_workgroup_scan_exclusive_min_int(void)
+{
+ cl_int *input = NULL;
+ cl_int *expected = NULL;
+ OCL_CREATE_KERNEL_FROM_FILE("compiler_workgroup_scan_exclusive",
+ "compiler_workgroup_scan_exclusive_min_int");
+ workgroup_generic(WG_SCAN_EXCLUSIVE_MIN, input, expected);
+}
+MAKE_UTEST_FROM_FUNCTION(compiler_workgroup_scan_exclusive_min_int);
+void compiler_workgroup_scan_exclusive_min_uint(void)
+{
+ cl_uint *input = NULL;
+ cl_uint *expected = NULL;
+ OCL_CREATE_KERNEL_FROM_FILE("compiler_workgroup_scan_exclusive",
+ "compiler_workgroup_scan_exclusive_min_uint");
+ workgroup_generic(WG_SCAN_EXCLUSIVE_MIN, input, expected);
+}
+MAKE_UTEST_FROM_FUNCTION(compiler_workgroup_scan_exclusive_min_uint);
+void compiler_workgroup_scan_exclusive_min_long(void)
+{
+ cl_long *input = NULL;
+ cl_long *expected = NULL;
+ OCL_CREATE_KERNEL_FROM_FILE("compiler_workgroup_scan_exclusive",
+ "compiler_workgroup_scan_exclusive_min_long");
+ workgroup_generic(WG_SCAN_EXCLUSIVE_MIN, input, expected);
+}
+MAKE_UTEST_FROM_FUNCTION(compiler_workgroup_scan_exclusive_min_long);
+void compiler_workgroup_scan_exclusive_min_ulong(void)
+{
+ cl_ulong *input = NULL;
+ cl_ulong *expected = NULL;
+ OCL_CREATE_KERNEL_FROM_FILE("compiler_workgroup_scan_exclusive",
+ "compiler_workgroup_scan_exclusive_min_ulong");
+ workgroup_generic(WG_SCAN_EXCLUSIVE_MIN, input, expected);
+}
+MAKE_UTEST_FROM_FUNCTION(compiler_workgroup_scan_exclusive_min_ulong);
+void compiler_workgroup_scan_exclusive_min_float(void)
+{
+ cl_float *input = NULL;
+ cl_float *expected = NULL;
+ OCL_CREATE_KERNEL_FROM_FILE("compiler_workgroup_scan_exclusive",
+ "compiler_workgroup_scan_exclusive_min_float");
+ workgroup_generic(WG_SCAN_EXCLUSIVE_MIN, input, expected);
+}
+MAKE_UTEST_FROM_FUNCTION(compiler_workgroup_scan_exclusive_min_float);
--
2.5.0
_______________________________________________
Beignet mailing list
Beignet@lists.freedesktop.org
https://lists.freedesktop.org/mailman/listinfo/beignet
