[Beignet] [PATCH V2 12/17] Utest: Add workgroup scan exclusive tests
Xiuli Pan
xiuli.pan at intel.com
Wed Apr 13 07:02:58 UTC 2016
The numeric_limits min, max for float could not give inf and -inf, it
will only return the minium and maximum positive finite. But the device
will return inf and -inf.
So the EXP value for the index 0 float max and min should refined with
has_infinity and infinity.
On Mon, Apr 11, 2016 at 05:40:00PM +0300, Grigore Lupescu wrote:
> From: Grigore Lupescu <grigore.lupescu at intel.com>
>
> Added the following unit tests:
> compiler_workgroup_scan_exclusive_add_char
> compiler_workgroup_scan_exclusive_add_uchar
> compiler_workgroup_scan_exclusive_add_short
> compiler_workgroup_scan_exclusive_add_ushort
> compiler_workgroup_scan_exclusive_add_int
> compiler_workgroup_scan_exclusive_add_uint
> compiler_workgroup_scan_exclusive_add_long
> compiler_workgroup_scan_exclusive_add_ulong
> compiler_workgroup_scan_exclusive_add_float
> compiler_workgroup_scan_exclusive_max_char
> compiler_workgroup_scan_exclusive_max_uchar
> compiler_workgroup_scan_exclusive_max_short
> compiler_workgroup_scan_exclusive_max_ushort
> compiler_workgroup_scan_exclusive_max_int
> compiler_workgroup_scan_exclusive_max_uint
> compiler_workgroup_scan_exclusive_max_long
> compiler_workgroup_scan_exclusive_max_ulong
> compiler_workgroup_scan_exclusive_max_float
> compiler_workgroup_scan_exclusive_min_char
> compiler_workgroup_scan_exclusive_min_uchar
> compiler_workgroup_scan_exclusive_min_short
> compiler_workgroup_scan_exclusive_min_ushort
> compiler_workgroup_scan_exclusive_min_int
> compiler_workgroup_scan_exclusive_min_uint
> compiler_workgroup_scan_exclusive_min_long
> compiler_workgroup_scan_exclusive_min_ulong
> compiler_workgroup_scan_exclusive_min_float
>
> 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 | 425 +++++++++++++++++++++++++++
> 3 files changed, 596 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 c75e6b7..0cd400c 100644
> --- a/utests/CMakeLists.txt
> +++ b/utests/CMakeLists.txt
> @@ -125,6 +125,7 @@ set (utests_sources
> compiler_async_copy.cpp
> compiler_workgroup_broadcast.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..9613140
> --- /dev/null
> +++ b/utests/compiler_workgroup_scan_exclusive.cpp
> @@ -0,0 +1,425 @@
> +#include <cstdint>
> +#include <cstring>
> +#include <iostream>
> +#include <cstdlib>
> +#include <iomanip>
> +#include <algorithm>
> +
> +#include "utest_helper.hpp"
> +
> +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
> +
> +enum WG_FUNCTION
> +{
> + WG_SCAN_EXCLUSIVE_ADD,
> + WG_SCAN_EXCLUSIVE_MAX,
> + WG_SCAN_EXCLUSIVE_MIN
> +};
> +
> +/*
> + * Generic compute-expected function for op SCAN EXCLUSIVE 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 - 1] + expected[i - 1];
> + }
> + else if(wg_func == WG_SCAN_EXCLUSIVE_MAX)
> + {
> + expected[0] = numeric_limits<T>::min();
expected[0] = numeric_limits<T>::has_infinity?-numeric_limits<T>::infinity(): numeric_limits<T>::min();
> + expected[1] = input[0];
> + for(uint32_t i = 2; i < WG_LOCAL_SIZE; i++)
> + expected[i] = max(input[i - 1], expected[i - 1]);
> + }
> + else if(wg_func == WG_SCAN_EXCLUSIVE_MIN)
> + {
> + expected[0] = numeric_limits<T>::max();
expected[0] = numeric_limits<T>::has_infinity?numeric_limits<T>::infinity(): numeric_limits<T>::max();
> + expected[1] = input[0];
> + for(uint32_t i = 2; i < WG_LOCAL_SIZE; i++)
> + expected[i] = min(input[i - 1], expected[i - 1]);
> + }
> +}
> +
> +/*
> + * Generic workgroup utest function for op SCAN EXCLUSIVE 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 data types */
> + 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)
> + {
> +#if DEBUG_STDOUT
> + cout << endl << "IN: " << endl;
> +#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);
> +
> +#if DEBUG_STDOUT
> + /* output generated input */
> + cout << setw(4) << input[gid + lid] << ", " ;
> + if((lid + 1) % 8 == 0)
> + cout << endl;
> +#endif
> + }
> +
> + /* expected values */
> + compute_expected(wg_func, input + gid, expected + gid);
> +
> +#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 op SCAN EXCLUSIVE 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 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]);
> + 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 mismatch */
> + OCL_MAP_BUFFER(1);
> + uint32_t mismatches = 0;
> +
> + for (uint32_t i = 0; i < WG_GLOBAL_SIZE; 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;
> +#endif
> + }
> +
> +#if DEBUG_STDOUT
> + /* output mismatch count */
> + cout << "mismatches " << mismatches << endl;
> +#endif
> +
> + OCL_UNMAP_BUFFER(1);
> +
> + OCL_ASSERT(mismatches == 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 at lists.freedesktop.org
> https://lists.freedesktop.org/mailman/listinfo/beignet
More information about the Beignet
mailing list