Remove unused variables in subgroup tests (#1460)

Signed-off-by: Stuart Brady <stuart.brady@arm.com>
This commit is contained in:
Stuart Brady
2022-08-04 15:04:14 +01:00
committed by GitHub
parent c12bff46c6
commit 4ee8022230
6 changed files with 4 additions and 33 deletions

View File

@@ -481,7 +481,7 @@ template <typename Ty, ShuffleOp operation> struct SHF
static void gen(Ty *x, Ty *t, cl_int *m, const WorkGroupParams &test_params) static void gen(Ty *x, Ty *t, cl_int *m, const WorkGroupParams &test_params)
{ {
int i, ii, j, k, n, delta; int i, ii, j, k, n;
cl_uint l; cl_uint l;
int nw = test_params.local_workgroup_size; int nw = test_params.local_workgroup_size;
int ns = test_params.subgroup_size; int ns = test_params.subgroup_size;

View File

@@ -1496,7 +1496,6 @@ template <typename Ty, typename Fns, size_t TSIZE = 0> struct test
size_t tmp; size_t tmp;
cl_int error; cl_int error;
int subgroup_size, num_subgroups; int subgroup_size, num_subgroups;
size_t realSize;
size_t global = test_params.global_workgroup_size; size_t global = test_params.global_workgroup_size;
size_t local = test_params.local_workgroup_size; size_t local = test_params.local_workgroup_size;
clProgramWrapper program; clProgramWrapper program;

View File

@@ -79,7 +79,6 @@ template <int Which> struct BAR
int ng = test_params.global_workgroup_size; int ng = test_params.global_workgroup_size;
int nj = (nw + ns - 1) / ns; int nj = (nw + ns - 1) / ns;
ng = ng / nw; ng = ng / nw;
int e;
ii = 0; ii = 0;
for (k = 0; k < ng; ++k) for (k = 0; k < ng; ++k)

View File

@@ -134,23 +134,6 @@ template <NonUniformVoteOp operation> struct AA
} }
}; };
static const char *any_source = "__kernel void test_any(const __global Type "
"*in, __global int4 *xy, __global Type *out)\n"
"{\n"
" int gid = get_global_id(0);\n"
" XY(xy,gid);\n"
" out[gid] = sub_group_any(in[gid]);\n"
"}\n";
static const char *all_source = "__kernel void test_all(const __global Type "
"*in, __global int4 *xy, __global Type *out)\n"
"{\n"
" int gid = get_global_id(0);\n"
" XY(xy,gid);\n"
" out[gid] = sub_group_all(in[gid]);\n"
"}\n";
template <typename T> template <typename T>
int run_broadcast_scan_reduction_for_type(RunTestForType rft) int run_broadcast_scan_reduction_for_type(RunTestForType rft)
{ {

View File

@@ -190,14 +190,13 @@ template <typename Ty, BallotOp operation> struct BALLOT_BIT_EXTRACT
static void gen(Ty *x, Ty *t, cl_int *m, const WorkGroupParams &test_params) static void gen(Ty *x, Ty *t, cl_int *m, const WorkGroupParams &test_params)
{ {
int wi_id, sb_id, wg_id, l; int wi_id, sb_id, wg_id;
int gws = test_params.global_workgroup_size; int gws = test_params.global_workgroup_size;
int lws = test_params.local_workgroup_size; int lws = test_params.local_workgroup_size;
int sbs = test_params.subgroup_size; int sbs = test_params.subgroup_size;
int sb_number = (lws + sbs - 1) / sbs; int sb_number = (lws + sbs - 1) / sbs;
int wg_number = gws / lws; int wg_number = gws / lws;
int limit_sbs = sbs > 100 ? 100 : sbs; int limit_sbs = sbs > 100 ? 100 : sbs;
int non_uniform_size = gws % lws;
for (wg_id = 0; wg_id < wg_number; ++wg_id) for (wg_id = 0; wg_id < wg_number; ++wg_id)
{ // for each work_group { // for each work_group
@@ -235,7 +234,7 @@ template <typename Ty, BallotOp operation> struct BALLOT_BIT_EXTRACT
static test_status chk(Ty *x, Ty *y, Ty *mx, Ty *my, cl_int *m, static test_status chk(Ty *x, Ty *y, Ty *mx, Ty *my, cl_int *m,
const WorkGroupParams &test_params) const WorkGroupParams &test_params)
{ {
int wi_id, wg_id, l, sb_id; int wi_id, wg_id, sb_id;
int gws = test_params.global_workgroup_size; int gws = test_params.global_workgroup_size;
int lws = test_params.local_workgroup_size; int lws = test_params.local_workgroup_size;
int sbs = test_params.subgroup_size; int sbs = test_params.subgroup_size;
@@ -351,10 +350,6 @@ template <typename Ty, BallotOp operation> struct BALLOT_INVERSE
static void gen(Ty *x, Ty *t, cl_int *m, const WorkGroupParams &test_params) static void gen(Ty *x, Ty *t, cl_int *m, const WorkGroupParams &test_params)
{ {
int gws = test_params.global_workgroup_size;
int lws = test_params.local_workgroup_size;
int sbs = test_params.subgroup_size;
int non_uniform_size = gws % lws;
// no work here // no work here
} }
@@ -398,9 +393,6 @@ template <typename Ty, BallotOp operation> struct BALLOT_INVERSE
{ {
current_sbs = wg_offset + sbs > lws ? lws - wg_offset : sbs; current_sbs = wg_offset + sbs > lws ? lws - wg_offset : sbs;
} }
// take index of array where info which work_item will
// be broadcast its value is stored
int midx = 4 * wg_offset + 2;
// take subgroup local id of this work_item // take subgroup local id of this work_item
// Check result // Check result
for (wi_id = 0; wi_id < current_sbs; ++wi_id) for (wi_id = 0; wi_id < current_sbs; ++wi_id)
@@ -461,7 +453,6 @@ template <typename Ty, BallotOp operation> struct BALLOT_COUNT_SCAN_FIND
{ {
wg_number++; wg_number++;
} }
int e;
for (wg_id = 0; wg_id < wg_number; ++wg_id) for (wg_id = 0; wg_id < wg_number; ++wg_id)
{ // for each work_group { // for each work_group
if (non_uniform_size && wg_id == wg_number - 1) if (non_uniform_size && wg_id == wg_number - 1)
@@ -683,7 +674,7 @@ template <typename Ty, BallotOp operation> struct SMASK
static void gen(Ty *x, Ty *t, cl_int *m, const WorkGroupParams &test_params) static void gen(Ty *x, Ty *t, cl_int *m, const WorkGroupParams &test_params)
{ {
int wi_id, wg_id, l, sb_id; int wi_id, wg_id, sb_id;
int gws = test_params.global_workgroup_size; int gws = test_params.global_workgroup_size;
int lws = test_params.local_workgroup_size; int lws = test_params.local_workgroup_size;
int sbs = test_params.subgroup_size; int sbs = test_params.subgroup_size;

View File

@@ -102,7 +102,6 @@ template <typename Ty, ArithmeticOp operation> struct RED_CLU
{ {
int ii = j * ns; int ii = j * ns;
int n = ii + ns > nw ? nw - ii : ns; int n = ii + ns > nw ? nw - ii : ns;
int midx = 4 * ii + 2;
std::vector<Ty> clusters_results; std::vector<Ty> clusters_results;
int clusters_counter = ns / test_params.cluster_size; int clusters_counter = ns / test_params.cluster_size;
clusters_results.resize(clusters_counter); clusters_results.resize(clusters_counter);