mirror of
https://github.com/KhronosGroup/OpenCL-CTS.git
synced 2026-03-19 06:09:01 +00:00
Improve error handling in subgroup tests (#1352)
* MPGCOMP-14761 Improve error handling in subgroup tests Signed-off-by: Stuart Brady <stuart.brady@arm.com> * Add missing newline
This commit is contained in:
@@ -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)
|
||||
#define test_error_fail(errCode, msg) test_error_ret(errCode, msg, TEST_FAIL)
|
||||
#define test_error_ret(errCode, msg, retValue) \
|
||||
{ \
|
||||
auto errCodeResult = errCode; \
|
||||
|
||||
@@ -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;
|
||||
@@ -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;
|
||||
@@ -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;
|
||||
@@ -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;
|
||||
@@ -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;
|
||||
|
||||
@@ -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;
|
||||
@@ -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
|
||||
{
|
||||
@@ -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
|
||||
@@ -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
|
||||
@@ -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
|
||||
|
||||
@@ -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;
|
||||
}
|
||||
};
|
||||
|
||||
@@ -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:
|
||||
|
||||
@@ -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;
|
||||
@@ -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;
|
||||
}
|
||||
}
|
||||
}
|
||||
@@ -143,7 +143,7 @@ template <int Which> struct BAR
|
||||
m += 2 * nw;
|
||||
}
|
||||
|
||||
return 0;
|
||||
return TEST_PASS;
|
||||
}
|
||||
};
|
||||
|
||||
|
||||
@@ -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;
|
||||
@@ -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");
|
||||
|
||||
@@ -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;
|
||||
}
|
||||
};
|
||||
|
||||
|
||||
@@ -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;
|
||||
|
||||
@@ -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;
|
||||
@@ -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;
|
||||
@@ -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;
|
||||
@@ -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;
|
||||
@@ -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;
|
||||
|
||||
@@ -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;
|
||||
|
||||
@@ -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;
|
||||
|
||||
Reference in New Issue
Block a user