[Beignet] [PATCH V2 12/17] Utest: Add workgroup scan exclusive tests
Xiuli Pan
xiuli.pan at intel.com
Wed Apr 13 07:48:44 UTC 2016
char short are not required by the spec, they have very special max and
min value in the as EXP vlaue. We can just remove them.
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[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[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