Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Jump to bottom

Improve error handling in subgroup tests #1352

Merged
merged 2 commits into from
Nov 16, 2021
Merged

Improve error handling in subgroup tests #1352

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
2 commits
Select commit Hold shift + click to select a range
d15f91a
MPGCOMP-14761 Improve error handling in subgroup tests
StuartDBrady Nov 3, 2021
90031a7
Add missing newline
StuartDBrady Nov 15, 2021
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
1 change: 1 addition & 0 deletions test_common/harness/errorHelpers.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -62,6 +62,7 @@ static int vlog_win32(const char *format, ...);
return TEST_FAIL; \
}
#define test_error(errCode, msg) test_error_ret(errCode, msg, errCode)
jlewis-austin marked this conversation as resolved.
Show resolved Hide resolved
#define test_error_fail(errCode, msg) test_error_ret(errCode, msg, TEST_FAIL)
#define test_error_ret(errCode, msg, retValue) \
{ \
auto errCodeResult = errCode; \
Expand Down
20 changes: 10 additions & 10 deletions test_conformance/subgroups/subgroup_common_templates.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -168,8 +168,8 @@ template <typename Ty, SubgroupsBroadcastOp operation> struct BC
}
}

static int chk(Ty *x, Ty *y, Ty *mx, Ty *my, cl_int *m,
const WorkGroupParams &test_params)
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;
Expand Down Expand Up @@ -499,8 +499,8 @@ template <typename Ty, ShuffleOp operation> struct SHF
}
}

static int chk(Ty *x, Ty *y, Ty *mx, Ty *my, cl_int *m,
const WorkGroupParams &test_params)
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 nw = test_params.local_workgroup_size;
Expand Down Expand Up @@ -583,8 +583,8 @@ template <typename Ty, ArithmeticOp operation> struct SCEX_NU
genrand<Ty, operation>(x, t, m, ns, nw, ng);
}

static int chk(Ty *x, Ty *y, Ty *mx, Ty *my, cl_int *m,
const WorkGroupParams &test_params)
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;
Expand Down Expand Up @@ -689,8 +689,8 @@ template <typename Ty, ArithmeticOp operation> struct SCIN_NU
operation_names(operation), TypeManager<Ty>::name());
}

static int chk(Ty *x, Ty *y, Ty *mx, Ty *my, cl_int *m,
const WorkGroupParams &test_params)
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;
Expand Down Expand Up @@ -805,8 +805,8 @@ template <typename Ty, ArithmeticOp operation> struct RED_NU
genrand<Ty, operation>(x, t, m, ns, nw, ng);
}

static int chk(Ty *x, Ty *y, Ty *mx, Ty *my, cl_int *m,
const WorkGroupParams &test_params)
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;
Expand Down
58 changes: 36 additions & 22 deletions test_conformance/subgroups/subhelpers.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -1375,25 +1375,31 @@ static int run_kernel(cl_context context, cl_command_queue queue,
// Driver for testing a single built in function
template <typename Ty, typename Fns, size_t TSIZE = 0> struct test
{
static int mrun(cl_device_id device, cl_context context,
cl_command_queue queue, int num_elements, const char *kname,
const char *src, WorkGroupParams test_params)
static test_status mrun(cl_device_id device, cl_context context,
cl_command_queue queue, int num_elements,
const char *kname, const char *src,
WorkGroupParams test_params)
{
int error = TEST_PASS;
test_status combined_error = TEST_SKIPPED_ITSELF;
for (auto &mask : test_params.all_work_item_masks)
{
test_params.work_items_mask = mask;
error |= run(device, context, queue, num_elements, kname, src,
test_params);
test_status error = run(device, context, queue, num_elements, kname,
src, test_params);

if (error == TEST_FAIL
|| (error == TEST_PASS && combined_error != TEST_FAIL))
combined_error = error;
}
return error;
return combined_error;
};
static int run(cl_device_id device, cl_context context,
cl_command_queue queue, int num_elements, const char *kname,
const char *src, WorkGroupParams test_params)
static test_status run(cl_device_id device, cl_context context,
cl_command_queue queue, int num_elements,
const char *kname, const char *src,
WorkGroupParams test_params)
{
size_t tmp;
int error;
cl_int error;
int subgroup_size, num_subgroups;
size_t realSize;
size_t global = test_params.global_workgroup_size;
Expand Down Expand Up @@ -1434,7 +1440,7 @@ template <typename Ty, typename Fns, size_t TSIZE = 0> struct test
if (!TypeManager<Ty>::type_supported(device))
{
log_info("Data type not supported : %s\n", TypeManager<Ty>::name());
return 0;
return TEST_SKIPPED_ITSELF;
}
else
{
Expand All @@ -1450,7 +1456,7 @@ template <typename Ty, typename Fns, size_t TSIZE = 0> struct test

error = clGetDeviceInfo(device, CL_DEVICE_PLATFORM, sizeof(platform),
(void *)&platform, NULL);
test_error(error, "clGetDeviceInfo failed for CL_DEVICE_PLATFORM");
test_error_fail(error, "clGetDeviceInfo failed for CL_DEVICE_PLATFORM");
if (test_params.use_core_subgroups)
{
kernel_sstr
Expand All @@ -1465,12 +1471,12 @@ template <typename Ty, typename Fns, size_t TSIZE = 0> struct test

error = create_single_kernel_helper(context, &program, &kernel, 1,
&kernel_src, kname);
if (error != 0) return error;
if (error != CL_SUCCESS) return TEST_FAIL;

// Determine some local dimensions to use for the test.
error = get_max_common_work_group_size(
context, kernel, test_params.global_workgroup_size, &local);
test_error(error, "get_max_common_work_group_size failed");
test_error_fail(error, "get_max_common_work_group_size failed");

// Limit it a bit so we have muliple work groups
// Ideally this will still be large enough to give us multiple
Expand Down Expand Up @@ -1543,7 +1549,7 @@ template <typename Ty, typename Fns, size_t TSIZE = 0> struct test
input_array_size * sizeof(Ty), sgmap.data(),
global * sizeof(cl_int4), odata.data(),
output_array_size * sizeof(Ty), TSIZE * sizeof(Ty));
test_error(error, "Running kernel first time failed");
test_error_fail(error, "Running kernel first time failed");

// Generate the desired input for the kernel

Expand All @@ -1553,13 +1559,18 @@ template <typename Ty, typename Fns, size_t TSIZE = 0> struct test
input_array_size * sizeof(Ty), sgmap.data(),
global * sizeof(cl_int4), odata.data(),
output_array_size * sizeof(Ty), TSIZE * sizeof(Ty));
test_error(error, "Running kernel second time failed");
test_error_fail(error, "Running kernel second time failed");

// Check the result
error = Fns::chk(idata.data(), odata.data(), mapin.data(),
mapout.data(), sgmap.data(), test_params);
test_error(error, "Data verification failed");
return TEST_PASS;
test_status status = Fns::chk(idata.data(), odata.data(), mapin.data(),
mapout.data(), sgmap.data(), test_params);
// Detailed failure and skip messages should be logged by Fns::gen
// and Fns::chk.
if (status == TEST_FAIL)
{
test_fail("Data verification failed\n");
}
return status;
}
};

Expand Down Expand Up @@ -1625,7 +1636,10 @@ struct RunTestForType
test_params_);
}

return error;
// If we return TEST_SKIPPED_ITSELF here, then an entire suite may be
// reported as having been skipped even if some tests within it
// passed, as the status codes are erroneously ORed together:
return error == TEST_FAIL ? TEST_FAIL : TEST_PASS;
}

private:
Expand Down
10 changes: 5 additions & 5 deletions test_conformance/subgroups/test_barrier.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -92,8 +92,8 @@ template <int Which> struct BAR
}
}

static int chk(cl_int *x, cl_int *y, cl_int *mx, cl_int *my, cl_int *m,
const WorkGroupParams &test_params)
static test_status chk(cl_int *x, cl_int *y, cl_int *mx, cl_int *my,
cl_int *m, const WorkGroupParams &test_params)
{
int ii, i, j, k, n;
int nw = test_params.local_workgroup_size;
Expand Down Expand Up @@ -133,7 +133,7 @@ template <int Which> struct BAR
"id %d in sub group %d in group %d expected "
"%d got %d\n",
i, j, k, tr, rr);
return -1;
return TEST_FAIL;
}
}
}
Expand All @@ -143,7 +143,7 @@ template <int Which> struct BAR
m += 2 * nw;
}

return 0;
return TEST_PASS;
}
};

Expand Down Expand Up @@ -187,4 +187,4 @@ int test_barrier_functions_ext(cl_device_id device, cl_context context,
}

return test_barrier_functions(device, context, queue, num_elements, false);
}
}
12 changes: 6 additions & 6 deletions test_conformance/subgroups/test_ifp.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -245,8 +245,8 @@ struct IFP
}
}

static int chk(cl_int *x, cl_int *y, cl_int *t, cl_int *, cl_int *,
const WorkGroupParams &test_params)
static test_status chk(cl_int *x, cl_int *y, cl_int *t, cl_int *, cl_int *,
const WorkGroupParams &test_params)
{
int i, k;
int nw = test_params.local_workgroup_size;
Expand All @@ -255,8 +255,8 @@ struct IFP
int nj = (nw + ns - 1) / ns;
ng = ng / nw;

// We need at least 2 sub groups per group for this tes
if (nj == 1) return 0;
// We need at least 2 sub groups per group for this test
if (nj == 1) return TEST_SKIPPED_ITSELF;

log_info(" independent forward progress...\n");

Expand All @@ -270,14 +270,14 @@ struct IFP
log_error(
"ERROR: mismatch at element %d in work group %d\n", i,
k);
return -1;
return TEST_FAIL;
}
}
x += nj * (NUM_LOC + 1);
y += NUM_LOC;
}

return 0;
return TEST_PASS;
}
};

Expand Down
4 changes: 2 additions & 2 deletions test_conformance/subgroups/test_subgroup.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -68,8 +68,8 @@ template <NonUniformVoteOp operation> struct AA
}
}

static int chk(cl_int *x, cl_int *y, cl_int *mx, cl_int *my, cl_int *m,
const WorkGroupParams &test_params)
static test_status chk(cl_int *x, cl_int *y, cl_int *mx, cl_int *my,
cl_int *m, const WorkGroupParams &test_params)
{
int ii, i, j, k, n;
int ng = test_params.global_workgroup_size;
Expand Down
20 changes: 10 additions & 10 deletions test_conformance/subgroups/test_subgroup_ballot.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -33,8 +33,8 @@ template <typename Ty> struct BALLOT
log_info(" sub_group_ballot...\n");
}

static int chk(Ty *x, Ty *y, Ty *mx, Ty *my, cl_int *m,
const WorkGroupParams &test_params)
static test_status chk(Ty *x, Ty *y, Ty *mx, Ty *my, cl_int *m,
const WorkGroupParams &test_params)
{
int wi_id, wg_id, sb_id;
int gws = test_params.global_workgroup_size;
Expand Down Expand Up @@ -146,8 +146,8 @@ template <typename Ty, BallotOp operation> struct BALLOT_BIT_EXTRACT
}
}

static int chk(Ty *x, Ty *y, Ty *mx, Ty *my, cl_int *m,
const WorkGroupParams &test_params)
static test_status chk(Ty *x, Ty *y, Ty *mx, Ty *my, cl_int *m,
const WorkGroupParams &test_params)
{
int wi_id, wg_id, l, sb_id;
int gws = test_params.global_workgroup_size;
Expand Down Expand Up @@ -269,8 +269,8 @@ template <typename Ty, BallotOp operation> struct BALLOT_INVERSE
// no work here
}

static int chk(Ty *x, Ty *y, Ty *mx, Ty *my, cl_int *m,
const WorkGroupParams &test_params)
static test_status chk(Ty *x, Ty *y, Ty *mx, Ty *my, cl_int *m,
const WorkGroupParams &test_params)
{
int wi_id, wg_id, sb_id;
int gws = test_params.global_workgroup_size;
Expand Down Expand Up @@ -444,8 +444,8 @@ template <typename Ty, BallotOp operation> struct BALLOT_COUNT_SCAN_FIND
return mask;
}

static int chk(Ty *x, Ty *y, Ty *mx, Ty *my, cl_int *m,
const WorkGroupParams &test_params)
static test_status chk(Ty *x, Ty *y, Ty *mx, Ty *my, cl_int *m,
const WorkGroupParams &test_params)
{
int wi_id, wg_id, sb_id;
int gws = test_params.global_workgroup_size;
Expand Down Expand Up @@ -617,8 +617,8 @@ template <typename Ty, BallotOp operation> struct SMASK
}
}

static int chk(Ty *x, Ty *y, Ty *mx, Ty *my, cl_int *m,
const WorkGroupParams &test_params)
static test_status chk(Ty *x, Ty *y, Ty *mx, Ty *my, cl_int *m,
const WorkGroupParams &test_params)
{
int wi_id, wg_id, sb_id;
int gws = test_params.global_workgroup_size;
Expand Down
4 changes: 2 additions & 2 deletions test_conformance/subgroups/test_subgroup_clustered_reduce.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -50,8 +50,8 @@ template <typename Ty, ArithmeticOp operation> struct RED_CLU
genrand<Ty, operation>(x, t, m, ns, nw, ng);
}

static int chk(Ty *x, Ty *y, Ty *mx, Ty *my, cl_int *m,
const WorkGroupParams &test_params)
static test_status chk(Ty *x, Ty *y, Ty *mx, Ty *my, cl_int *m,
const WorkGroupParams &test_params)
{
int nw = test_params.local_workgroup_size;
int ns = test_params.subgroup_size;
Expand Down
4 changes: 2 additions & 2 deletions test_conformance/subgroups/test_subgroup_non_uniform_vote.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -83,8 +83,8 @@ template <typename T, NonUniformVoteOp operation> struct VOTE
}
}

static int chk(T *x, T *y, T *mx, T *my, cl_int *m,
const WorkGroupParams &test_params)
static test_status chk(T *x, T *y, T *mx, T *my, cl_int *m,
const WorkGroupParams &test_params)
{
int ii, i, j, k, n;
int nw = test_params.local_workgroup_size;
Expand Down