//
// Copyright (c) 2020 The Khronos Group Inc.
//
// Licensed under the Apache License, Version 2.0 (the "License");
// you may not use this file except in compliance with the License.
// You may obtain a copy of the License at
//
//    http://www.apache.org/licenses/LICENSE-2.0
//
// Unless required by applicable law or agreed to in writing, software
// distributed under the License is distributed on an "AS IS" BASIS,
// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
// See the License for the specific language governing permissions and
// limitations under the License.
//
#ifndef SUBGROUPCOMMONTEMPLATES_H
#define SUBGROUPCOMMONTEMPLATES_H

#include "typeWrappers.h"
#include "CL/cl_half.h"
#include "subhelpers.h"
#include <set>
#include <algorithm>

// DESCRIPTION :
// sub_group_broadcast - each work_item registers it's own value.
// All work_items in subgroup takes one value from only one (any) work_item
// sub_group_broadcast_first - same as type 0. All work_items in
// subgroup takes only one value from only one chosen (the smallest subgroup ID)
// work_item
// sub_group_non_uniform_broadcast - same as type 0 but
// only 4 work_items from subgroup enter the code (are active)
template <typename Ty, SubgroupsBroadcastOp operation> struct BC
{
    static void log_test(const WorkGroupParams &test_params,
                         const char *extra_text)
    {
        log_info("  sub_group_%s(%s)...%s\n", operation_names(operation),
                 TypeManager<Ty>::name(), extra_text);
    }

    static void gen(Ty *x, Ty *t, cl_int *m, const WorkGroupParams &test_params)
    {
        int i, ii, j, k, n;
        int ng = test_params.global_workgroup_size;
        int nw = test_params.local_workgroup_size;
        int ns = test_params.subgroup_size;
        int nj = (nw + ns - 1) / ns;
        int d = ns > 100 ? 100 : ns;
        int non_uniform_size = ng % nw;
        ng = ng / nw;
        int last_subgroup_size = 0;
        ii = 0;

        if (non_uniform_size)
        {
            ng++;
        }
        for (k = 0; k < ng; ++k)
        { // for each work_group
            if (non_uniform_size && k == ng - 1)
            {
                set_last_workgroup_params(non_uniform_size, nj, ns, nw,
                                          last_subgroup_size);
            }
            for (j = 0; j < nj; ++j)
            { // for each subgroup
                ii = j * ns;
                if (last_subgroup_size && j == nj - 1)
                {
                    n = last_subgroup_size;
                }
                else
                {
                    n = ii + ns > nw ? nw - ii : ns;
                }
                int bcast_if = 0;
                int bcast_elseif = 0;
                int bcast_index = (int)(genrand_int32(gMTdata) & 0x7fffffff)
                    % (d > n ? n : d);
                // l - calculate subgroup local id from which value will be
                // broadcasted (one the same value for whole subgroup)
                if (operation != SubgroupsBroadcastOp::broadcast)
                {
                    // reduce brodcasting index in case of non_uniform and
                    // last workgroup last subgroup
                    if (last_subgroup_size && j == nj - 1
                        && last_subgroup_size < NR_OF_ACTIVE_WORK_ITEMS)
                    {
                        bcast_if = bcast_index % last_subgroup_size;
                        bcast_elseif = bcast_if;
                    }
                    else
                    {
                        bcast_if = bcast_index % NR_OF_ACTIVE_WORK_ITEMS;
                        bcast_elseif = NR_OF_ACTIVE_WORK_ITEMS
                            + bcast_index % (n - NR_OF_ACTIVE_WORK_ITEMS);
                    }
                }

                for (i = 0; i < n; ++i)
                {
                    if (operation == SubgroupsBroadcastOp::broadcast)
                    {
                        int midx = 4 * ii + 4 * i + 2;
                        m[midx] = (cl_int)bcast_index;
                    }
                    else
                    {
                        if (i < NR_OF_ACTIVE_WORK_ITEMS)
                        {
                            // index of the third
                            // element int the vector.
                            int midx = 4 * ii + 4 * i + 2;
                            // storing information about
                            // broadcasting index -
                            // earlier calculated
                            m[midx] = (cl_int)bcast_if;
                        }
                        else
                        { // index of the third
                          // element int the vector.
                            int midx = 4 * ii + 4 * i + 3;
                            m[midx] = (cl_int)bcast_elseif;
                        }
                    }

                    // calculate value for broadcasting
                    cl_ulong number = genrand_int64(gMTdata);
                    set_value(t[ii + i], number);
                }
            }
            // Now map into work group using map from device
            for (j = 0; j < nw; ++j)
            { // for each element in work_group
                // calculate index as number of subgroup
                // plus subgroup local id
                x[j] = t[j];
            }
            x += nw;
            m += 4 * nw;
        }
    }

    static test_status chk(Ty *x, Ty *y, Ty *mx, Ty *my, cl_int *m,
                           const WorkGroupParams &test_params)
    {
        int ii, i, j, k, l, n;
        int ng = test_params.global_workgroup_size;
        int nw = test_params.local_workgroup_size;
        int ns = test_params.subgroup_size;
        int nj = (nw + ns - 1) / ns;
        Ty tr, rr;
        int non_uniform_size = ng % nw;
        ng = ng / nw;
        int last_subgroup_size = 0;
        if (non_uniform_size) ng++;

        for (k = 0; k < ng; ++k)
        { // for each work_group
            if (non_uniform_size && k == ng - 1)
            {
                set_last_workgroup_params(non_uniform_size, nj, ns, nw,
                                          last_subgroup_size);
            }
            for (j = 0; j < nw; ++j)
            { // inside the work_group
                mx[j] = x[j]; // read host inputs for work_group
                my[j] = y[j]; // read device outputs for work_group
            }

            for (j = 0; j < nj; ++j)
            { // for each subgroup
                ii = j * ns;
                if (last_subgroup_size && j == nj - 1)
                {
                    n = last_subgroup_size;
                }
                else
                {
                    n = ii + ns > nw ? nw - ii : ns;
                }

                // Check result
                if (operation == SubgroupsBroadcastOp::broadcast_first)
                {
                    int lowest_active_id = -1;
                    for (i = 0; i < n; ++i)
                    {

                        lowest_active_id = i < NR_OF_ACTIVE_WORK_ITEMS
                            ? 0
                            : NR_OF_ACTIVE_WORK_ITEMS;
                        //  findout if broadcasted
                        //  value is the same
                        tr = mx[ii + lowest_active_id];
                        //  findout if broadcasted to all
                        rr = my[ii + i];

                        if (!compare(rr, tr))
                        {
                            log_error(
                                "ERROR: sub_group_broadcast_first(%s) "
                                "mismatch "
                                "for local id %d in sub group %d in group "
                                "%d\n",
                                TypeManager<Ty>::name(), i, j, k);
                            return TEST_FAIL;
                        }
                    }
                }
                else
                {
                    for (i = 0; i < n; ++i)
                    {
                        if (operation == SubgroupsBroadcastOp::broadcast)
                        {
                            int midx = 4 * ii + 4 * i + 2;
                            l = (int)m[midx];
                            tr = mx[ii + l];
                        }
                        else
                        {
                            if (i < NR_OF_ACTIVE_WORK_ITEMS)
                            { // take index of array where info
                              // which work_item will be
                              // broadcast its value is stored
                                int midx = 4 * ii + 4 * i + 2;
                                // take subgroup local id of
                                // this work_item
                                l = (int)m[midx];
                                // take value generated on host
                                // for this work_item
                                tr = mx[ii + l];
                            }
                            else
                            {
                                int midx = 4 * ii + 4 * i + 3;
                                l = (int)m[midx];
                                tr = mx[ii + l];
                            }
                        }
                        rr = my[ii + i]; // read device outputs for
                                         // work_item in the subgroup

                        if (!compare(rr, tr))
                        {
                            log_error("ERROR: sub_group_%s(%s) "
                                      "mismatch for local id %d in sub "
                                      "group %d in group %d - %s\n",
                                      operation_names(operation),
                                      TypeManager<Ty>::name(), i, j, k,
                                      print_expected_obtained(tr, rr).c_str());
                            return TEST_FAIL;
                        }
                    }
                }
            }
            x += nw;
            y += nw;
            m += 4 * nw;
        }
        return TEST_PASS;
    }
};

static float to_float(subgroups::cl_half x) { return cl_half_to_float(x.data); }

static subgroups::cl_half to_half(float x)
{
    subgroups::cl_half value;
    value.data = cl_half_from_float(x, g_rounding_mode);
    return value;
}

// for integer types
template <typename Ty> inline Ty calculate(Ty a, Ty b, ArithmeticOp operation)
{
    switch (operation)
    {
        case ArithmeticOp::add_: return a + b;
        case ArithmeticOp::max_: return a > b ? a : b;
        case ArithmeticOp::min_: return a < b ? a : b;
        case ArithmeticOp::mul_: return a * b;
        case ArithmeticOp::and_: return a & b;
        case ArithmeticOp::or_: return a | b;
        case ArithmeticOp::xor_: return a ^ b;
        case ArithmeticOp::logical_and: return a && b;
        case ArithmeticOp::logical_or: return a || b;
        case ArithmeticOp::logical_xor: return !a ^ !b;
        default: log_error("Unknown operation request\n"); break;
    }
    return 0;
}
// Specialize for floating points.
template <>
inline cl_double calculate(cl_double a, cl_double b, ArithmeticOp operation)
{
    switch (operation)
    {
        case ArithmeticOp::add_: {
            return a + b;
        }
        case ArithmeticOp::max_: {
            return a > b ? a : b;
        }
        case ArithmeticOp::min_: {
            return a < b ? a : b;
        }
        case ArithmeticOp::mul_: {
            return a * b;
        }
        default: log_error("Unknown operation request\n"); break;
    }
    return 0;
}

template <>
inline cl_float calculate(cl_float a, cl_float b, ArithmeticOp operation)
{
    switch (operation)
    {
        case ArithmeticOp::add_: {
            return a + b;
        }
        case ArithmeticOp::max_: {
            return a > b ? a : b;
        }
        case ArithmeticOp::min_: {
            return a < b ? a : b;
        }
        case ArithmeticOp::mul_: {
            return a * b;
        }
        default: log_error("Unknown operation request\n"); break;
    }
    return 0;
}

template <>
inline subgroups::cl_half calculate(subgroups::cl_half a, subgroups::cl_half b,
                                    ArithmeticOp operation)
{
    switch (operation)
    {
        case ArithmeticOp::add_: return to_half(to_float(a) + to_float(b));
        case ArithmeticOp::max_:
            return to_float(a) > to_float(b) || is_half_nan(b.data) ? a : b;
        case ArithmeticOp::min_:
            return to_float(a) < to_float(b) || is_half_nan(b.data) ? a : b;
        case ArithmeticOp::mul_: return to_half(to_float(a) * to_float(b));
        default: log_error("Unknown operation request\n"); break;
    }
    return to_half(0);
}

template <typename Ty> bool is_floating_point()
{
    return std::is_floating_point<Ty>::value
        || std::is_same<Ty, subgroups::cl_half>::value;
}

template <typename Ty, ArithmeticOp operation>
void generate_inputs(Ty *x, Ty *t, cl_int *m, int ns, int nw, int ng)
{
    int nj = (nw + ns - 1) / ns;

    std::vector<cl_ulong> safe_values;
    if (operation == ArithmeticOp::mul_ || operation == ArithmeticOp::add_)
    {
        fill_and_shuffle_safe_values(safe_values, ns);
    }

    for (int k = 0; k < ng; ++k)
    {
        for (int j = 0; j < nj; ++j)
        {
            int ii = j * ns;
            int n = ii + ns > nw ? nw - ii : ns;

            for (int i = 0; i < n; ++i)
            {
                cl_ulong out_value;
                if (operation == ArithmeticOp::mul_
                    || operation == ArithmeticOp::add_)
                {
                    out_value = safe_values[i];
                }
                else
                {
                    out_value = genrand_int64(gMTdata) % (32 * n);
                    if ((operation == ArithmeticOp::logical_and
                         || operation == ArithmeticOp::logical_or
                         || operation == ArithmeticOp::logical_xor)
                        && ((out_value >> 32) & 1) == 0)
                        out_value = 0; // increase probability of false
                }
                set_value(t[ii + i], out_value);
            }
        }

        // Now map into work group using map from device
        for (int j = 0; j < nw; ++j)
        {
            x[j] = t[j];
        }

        x += nw;
        m += 4 * nw;
    }
}

template <typename Ty, ShuffleOp operation> struct SHF
{
    static void log_test(const WorkGroupParams &test_params,
                         const char *extra_text)
    {
        log_info("  sub_group_%s(%s)...%s\n", operation_names(operation),
                 TypeManager<Ty>::name(), extra_text);
    }

    static void gen(Ty *x, Ty *t, cl_int *m, const WorkGroupParams &test_params)
    {
        int i, ii, j, k, n;
        cl_uint l;
        int nw = test_params.local_workgroup_size;
        int ns = test_params.subgroup_size;
        int ng = test_params.global_workgroup_size;
        int nj = (nw + ns - 1) / ns;
        ii = 0;
        ng = ng / nw;
        for (k = 0; k < ng; ++k)
        { // for each work_group
            for (j = 0; j < nj; ++j)
            { // for each subgroup
                ii = j * ns;
                n = ii + ns > nw ? nw - ii : ns;
                for (i = 0; i < n; ++i)
                {
                    int midx = 4 * ii + 4 * i + 2;
                    l = (((cl_uint)(genrand_int32(gMTdata) & 0x7fffffff) + 1)
                         % (ns * 2 + 1))
                        - 1;
                    switch (operation)
                    {
                        case ShuffleOp::shuffle:
                        case ShuffleOp::shuffle_xor:
                        case ShuffleOp::shuffle_up:
                        case ShuffleOp::shuffle_down:
                            // storing information about shuffle index/delta
                            m[midx] = (cl_int)l;
                            break;
                        case ShuffleOp::rotate:
                        case ShuffleOp::clustered_rotate:
                            // Storing information about rotate delta.
                            // The delta must be the same for each thread in
                            // the subgroup.
                            if (i == 0)
                            {
                                m[midx] = (cl_int)l;
                            }
                            else
                            {
                                m[midx] = m[midx - 4];
                            }
                            break;
                        default: break;
                    }
                    cl_ulong number = genrand_int64(gMTdata);
                    set_value(t[ii + i], number);
                }
            }
            // Now map into work group using map from device
            for (j = 0; j < nw; ++j)
            { // for each element in work_group
                x[j] = t[j];
            }
            x += nw;
            m += 4 * nw;
        }
    }

    static test_status chk(Ty *x, Ty *y, Ty *mx, Ty *my, cl_int *m,
                           const WorkGroupParams &test_params)
    {
        int ii, k;
        size_t n;
        cl_uint l;
        size_t nw = test_params.local_workgroup_size;
        size_t ns = test_params.subgroup_size;
        int ng = test_params.global_workgroup_size;
        size_t nj = (nw + ns - 1) / ns;
        Ty tr, rr;
        ng = ng / nw;

        for (k = 0; k < ng; ++k)
        { // for each work_group
            for (size_t j = 0; j < nw; ++j)
            { // inside the work_group
                mx[j] = x[j]; // read host inputs for work_group
                my[j] = y[j]; // read device outputs for work_group
            }

            for (size_t j = 0; j < nj; ++j)
            { // for each subgroup
                ii = j * ns;
                n = ii + ns > nw ? nw - ii : ns;

                for (size_t i = 0; i < n; ++i)
                { // inside the subgroup
                  // shuffle index storage
                    int midx = 4 * ii + 4 * i + 2;
                    l = m[midx];
                    rr = my[ii + i];
                    cl_uint tr_idx;
                    bool skip = false;
                    switch (operation)
                    {
                        // shuffle basic - treat l as index
                        case ShuffleOp::shuffle: tr_idx = l; break;
                        // shuffle xor - treat l as mask
                        case ShuffleOp::shuffle_xor: tr_idx = i ^ l; break;
                        // shuffle up - treat l as delta
                        case ShuffleOp::shuffle_up:
                            if (l >= ns) skip = true;
                            tr_idx = i - l;
                            break;
                        // shuffle down - treat l as delta
                        case ShuffleOp::shuffle_down:
                            if (l >= ns) skip = true;
                            tr_idx = i + l;
                            break;
                        // rotate - treat l as delta
                        case ShuffleOp::rotate:
                            tr_idx = (i + l) % test_params.subgroup_size;
                            break;
                        case ShuffleOp::clustered_rotate: {
                            tr_idx = ((i & ~(test_params.cluster_size - 1))
                                      + ((i + l) % test_params.cluster_size));
                            break;
                        }
                        default: break;
                    }

                    if (!skip && tr_idx < n)
                    {
                        tr = mx[ii + tr_idx];

                        if (!compare(rr, tr))
                        {
                            log_error("ERROR: sub_group_%s(%s) mismatch for "
                                      "local id %d in sub group %d in group "
                                      "%d\n",
                                      operation_names(operation),
                                      TypeManager<Ty>::name(), i, j, k);
                            return TEST_FAIL;
                        }
                    }
                }
            }
            x += nw;
            y += nw;
            m += 4 * nw;
        }
        return TEST_PASS;
    }
};

template <typename Ty, ArithmeticOp operation> struct SCEX_NU
{
    static void log_test(const WorkGroupParams &test_params,
                         const char *extra_text)
    {
        std::string func_name = (test_params.all_work_item_masks.size() > 0
                                     ? "sub_group_non_uniform_scan_exclusive"
                                     : "sub_group_scan_exclusive");
        log_info("  %s_%s(%s)...%s\n", func_name.c_str(),
                 operation_names(operation), TypeManager<Ty>::name(),
                 extra_text);
    }

    static void gen(Ty *x, Ty *t, cl_int *m, const WorkGroupParams &test_params)
    {
        int nw = test_params.local_workgroup_size;
        int ns = test_params.subgroup_size;
        int ng = test_params.global_workgroup_size;
        ng = ng / nw;
        generate_inputs<Ty, operation>(x, t, m, ns, nw, ng);
    }

    static test_status chk(Ty *x, Ty *y, Ty *mx, Ty *my, cl_int *m,
                           const WorkGroupParams &test_params)
    {
        int ii, i, j, k, n;
        int nw = test_params.local_workgroup_size;
        int ns = test_params.subgroup_size;
        int ng = test_params.global_workgroup_size;
        bs128 work_items_mask = test_params.work_items_mask;
        int nj = (nw + ns - 1) / ns;
        Ty tr, rr;
        ng = ng / nw;

        std::string func_name = (test_params.all_work_item_masks.size() > 0
                                     ? "sub_group_non_uniform_scan_exclusive"
                                     : "sub_group_scan_exclusive");

        // for uniform case take into consideration all workitems
        if (!work_items_mask.any())
        {
            work_items_mask.set();
        }
        for (k = 0; k < ng; ++k)
        { // for each work_group
            // Map to array indexed to array indexed by local ID and sub group
            for (j = 0; j < nw; ++j)
            { // inside the work_group
                mx[j] = x[j]; // read host inputs for work_group
                my[j] = y[j]; // read device outputs for work_group
            }
            for (j = 0; j < nj; ++j)
            {
                ii = j * ns;
                n = ii + ns > nw ? nw - ii : ns;
                std::set<int> active_work_items;
                for (i = 0; i < n; ++i)
                {
                    if (work_items_mask.test(i))
                    {
                        active_work_items.insert(i);
                    }
                }
                if (active_work_items.empty())
                {
                    continue;
                }
                else
                {
                    tr = TypeManager<Ty>::identify_limits(operation);
                    for (const int &active_work_item : active_work_items)
                    {
                        rr = my[ii + active_work_item];
                        if (!compare_ordered(rr, tr))
                        {
                            log_error(
                                "ERROR: %s_%s(%s) "
                                "mismatch for local id %d in sub group %d in "
                                "group %d %s\n",
                                func_name.c_str(), operation_names(operation),
                                TypeManager<Ty>::name(), i, j, k,
                                print_expected_obtained(tr, rr).c_str());
                            return TEST_FAIL;
                        }
                        tr = calculate<Ty>(tr, mx[ii + active_work_item],
                                           operation);
                    }
                }
            }
            x += nw;
            y += nw;
            m += 4 * nw;
        }

        return TEST_PASS;
    }
};

// Test for scan inclusive non uniform functions
template <typename Ty, ArithmeticOp operation> struct SCIN_NU
{
    static void log_test(const WorkGroupParams &test_params,
                         const char *extra_text)
    {
        std::string func_name = (test_params.all_work_item_masks.size() > 0
                                     ? "sub_group_non_uniform_scan_inclusive"
                                     : "sub_group_scan_inclusive");
        log_info("  %s_%s(%s)...%s\n", func_name.c_str(),
                 operation_names(operation), TypeManager<Ty>::name(),
                 extra_text);
    }

    static void gen(Ty *x, Ty *t, cl_int *m, const WorkGroupParams &test_params)
    {
        int nw = test_params.local_workgroup_size;
        int ns = test_params.subgroup_size;
        int ng = test_params.global_workgroup_size;
        ng = ng / nw;
        generate_inputs<Ty, operation>(x, t, m, ns, nw, ng);
    }

    static test_status chk(Ty *x, Ty *y, Ty *mx, Ty *my, cl_int *m,
                           const WorkGroupParams &test_params)
    {
        int ii, i, j, k, n;
        int nw = test_params.local_workgroup_size;
        int ns = test_params.subgroup_size;
        int ng = test_params.global_workgroup_size;
        bs128 work_items_mask = test_params.work_items_mask;

        int nj = (nw + ns - 1) / ns;
        Ty tr, rr;
        ng = ng / nw;

        std::string func_name = (test_params.all_work_item_masks.size() > 0
                                     ? "sub_group_non_uniform_scan_inclusive"
                                     : "sub_group_scan_inclusive");

        // for uniform case take into consideration all workitems
        if (!work_items_mask.any())
        {
            work_items_mask.set();
        }
        // std::bitset<32> mask32(use_work_items_mask);
        // for (int k) mask32.count();
        for (k = 0; k < ng; ++k)
        { // for each work_group
            // Map to array indexed to array indexed by local ID and sub group
            for (j = 0; j < nw; ++j)
            { // inside the work_group
                mx[j] = x[j]; // read host inputs for work_group
                my[j] = y[j]; // read device outputs for work_group
            }
            for (j = 0; j < nj; ++j)
            {
                ii = j * ns;
                n = ii + ns > nw ? nw - ii : ns;
                std::set<int> active_work_items;
                int catch_frist_active = -1;

                for (i = 0; i < n; ++i)
                {
                    if (work_items_mask.test(i))
                    {
                        if (catch_frist_active == -1)
                        {
                            catch_frist_active = i;
                        }
                        active_work_items.insert(i);
                    }
                }
                if (active_work_items.empty())
                {
                    continue;
                }
                else
                {
                    tr = TypeManager<Ty>::identify_limits(operation);
                    for (const int &active_work_item : active_work_items)
                    {
                        rr = my[ii + active_work_item];
                        if (active_work_items.size() == 1)
                        {
                            tr = mx[ii + catch_frist_active];
                        }
                        else
                        {
                            tr = calculate<Ty>(tr, mx[ii + active_work_item],
                                               operation);
                        }
                        if (!compare_ordered<Ty>(rr, tr))
                        {
                            log_error(
                                "ERROR: %s_%s(%s) "
                                "mismatch for local id %d in sub group %d "
                                "in "
                                "group %d %s\n",
                                func_name.c_str(), operation_names(operation),
                                TypeManager<Ty>::name(), active_work_item, j, k,
                                print_expected_obtained(tr, rr).c_str());
                            return TEST_FAIL;
                        }
                    }
                }
            }
            x += nw;
            y += nw;
            m += 4 * nw;
        }

        return TEST_PASS;
    }
};

// Test for reduce non uniform functions
template <typename Ty, ArithmeticOp operation> struct RED_NU
{
    static void log_test(const WorkGroupParams &test_params,
                         const char *extra_text)
    {
        std::string func_name = (test_params.all_work_item_masks.size() > 0
                                     ? "sub_group_non_uniform_reduce"
                                     : "sub_group_reduce");
        log_info("  %s_%s(%s)...%s\n", func_name.c_str(),
                 operation_names(operation), TypeManager<Ty>::name(),
                 extra_text);
    }

    static void gen(Ty *x, Ty *t, cl_int *m, const WorkGroupParams &test_params)
    {
        int nw = test_params.local_workgroup_size;
        int ns = test_params.subgroup_size;
        int ng = test_params.global_workgroup_size;
        ng = ng / nw;
        generate_inputs<Ty, operation>(x, t, m, ns, nw, ng);
    }

    static test_status chk(Ty *x, Ty *y, Ty *mx, Ty *my, cl_int *m,
                           const WorkGroupParams &test_params)
    {
        int ii, i, j, k, n;
        int nw = test_params.local_workgroup_size;
        int ns = test_params.subgroup_size;
        int ng = test_params.global_workgroup_size;
        bs128 work_items_mask = test_params.work_items_mask;
        int nj = (nw + ns - 1) / ns;
        ng = ng / nw;
        Ty tr, rr;

        std::string func_name = (test_params.all_work_item_masks.size() > 0
                                     ? "sub_group_non_uniform_reduce"
                                     : "sub_group_reduce");

        for (k = 0; k < ng; ++k)
        {
            // Map to array indexed to array indexed by local ID and sub
            // group
            for (j = 0; j < nw; ++j)
            {
                mx[j] = x[j];
                my[j] = y[j];
            }

            if (!work_items_mask.any())
            {
                work_items_mask.set();
            }

            for (j = 0; j < nj; ++j)
            {
                ii = j * ns;
                n = ii + ns > nw ? nw - ii : ns;
                std::set<int> active_work_items;
                int catch_frist_active = -1;
                for (i = 0; i < n; ++i)
                {
                    if (work_items_mask.test(i))
                    {
                        if (catch_frist_active == -1)
                        {
                            catch_frist_active = i;
                            tr = mx[ii + i];
                            active_work_items.insert(i);
                            continue;
                        }
                        active_work_items.insert(i);
                        tr = calculate<Ty>(tr, mx[ii + i], operation);
                    }
                }

                if (active_work_items.empty())
                {
                    continue;
                }

                for (const int &active_work_item : active_work_items)
                {
                    rr = my[ii + active_work_item];
                    if (!compare_ordered<Ty>(rr, tr))
                    {
                        log_error("ERROR: %s_%s(%s) "
                                  "mismatch for local id %d in sub group %d in "
                                  "group %d %s\n",
                                  func_name.c_str(), operation_names(operation),
                                  TypeManager<Ty>::name(), active_work_item, j,
                                  k, print_expected_obtained(tr, rr).c_str());
                        return TEST_FAIL;
                    }
                }
            }
            x += nw;
            y += nw;
            m += 4 * nw;
        }

        return TEST_PASS;
    }
};

#endif