ahmed/OpenCL-CTS
1
0
Fork 0
mirror of https://github.com/KhronosGroup/OpenCL-CTS.git synced 2026-03-26 08:49:02 +00:00
Code Packages Projects Releases Wiki Activity

Refactor logging of subgroup test start/pass messages (#1361)

Browse Source 
Note that this also corrects the start messages logged for the
sub_group_ballot_bit_count/find_msb/find_lsb tests.

Signed-off-by: Stuart Brady <stuart.brady@arm.com>
This commit is contained in:
Stuart Brady
2021-12-14 17:52:44 +00:00
committed by GitHub
parent 73d71b6a76
commit 1161d788dd
8 changed files with 164 additions and 98 deletions
Download Patch File Download Diff File Expand all files Collapse all files

101
test_conformance/subgroups/subgroup_common_templates.h
View File

@@ -63,6 +63,13 @@ static cl_uint4 generate_bit_mask(cl_uint subgroup_local_id,
// only 4 work_items from subgroup enter the code (are active) // only 4 work_items from subgroup enter the code (are active)
template <typename Ty, SubgroupsBroadcastOp operation> struct BC 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) static void gen(Ty *x, Ty *t, cl_int *m, const WorkGroupParams &test_params)
{ {
int i, ii, j, k, n; int i, ii, j, k, n;
@@ -76,8 +83,6 @@ template <typename Ty, SubgroupsBroadcastOp operation> struct BC
int last_subgroup_size = 0; int last_subgroup_size = 0;
ii = 0; ii = 0;
log_info(" sub_group_%s(%s)...\n", operation_names(operation),
TypeManager<Ty>::name());
if (non_uniform_size) if (non_uniform_size)
{ {
ng++; ng++;
@@ -286,8 +291,6 @@ template <typename Ty, SubgroupsBroadcastOp operation> struct BC
y += nw; y += nw;
m += 4 * nw; m += 4 * nw;
} }
log_info(" sub_group_%s(%s)... passed\n", operation_names(operation),
TypeManager<Ty>::name());
return TEST_PASS; return TEST_PASS;
} }
}; };
@@ -437,6 +440,13 @@ void genrand(Ty *x, Ty *t, cl_int *m, int ns, int nw, int ng)
template <typename Ty, ShuffleOp operation> struct SHF 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) static void gen(Ty *x, Ty *t, cl_int *m, const WorkGroupParams &test_params)
{ {
int i, ii, j, k, l, n, delta; int i, ii, j, k, l, n, delta;
@@ -447,8 +457,6 @@ template <typename Ty, ShuffleOp operation> struct SHF
int d = ns > 100 ? 100 : ns; int d = ns > 100 ? 100 : ns;
ii = 0; ii = 0;
ng = ng / nw; ng = ng / nw;
log_info(" sub_group_%s(%s)...\n", operation_names(operation),
TypeManager<Ty>::name());
for (k = 0; k < ng; ++k) for (k = 0; k < ng; ++k)
{ // for each work_group { // for each work_group
for (j = 0; j < nj; ++j) for (j = 0; j < nj; ++j)
@@ -560,26 +568,29 @@ template <typename Ty, ShuffleOp operation> struct SHF
y += nw; y += nw;
m += 4 * nw; m += 4 * nw;
} }
log_info(" sub_group_%s(%s)... passed\n", operation_names(operation),
TypeManager<Ty>::name());
return TEST_PASS; return TEST_PASS;
} }
}; };
template <typename Ty, ArithmeticOp operation> struct SCEX_NU 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) static void gen(Ty *x, Ty *t, cl_int *m, const WorkGroupParams &test_params)
{ {
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;
int ng = test_params.global_workgroup_size; int ng = test_params.global_workgroup_size;
ng = ng / nw; ng = ng / nw;
std::string func_name;
test_params.work_items_mask.any()
? func_name = "sub_group_non_uniform_scan_exclusive"
: func_name = "sub_group_scan_exclusive";
log_info(" %s_%s(%s)...\n", func_name.c_str(),
operation_names(operation), TypeManager<Ty>::name());
genrand<Ty, operation>(x, t, m, ns, nw, ng); genrand<Ty, operation>(x, t, m, ns, nw, ng);
} }
@@ -595,11 +606,9 @@ template <typename Ty, ArithmeticOp operation> struct SCEX_NU
Ty tr, rr; Ty tr, rr;
ng = ng / nw; ng = ng / nw;
std::string func_name; std::string func_name = (test_params.all_work_item_masks.size() > 0
test_params.work_items_mask.any() ? "sub_group_non_uniform_scan_exclusive"
? func_name = "sub_group_non_uniform_scan_exclusive" : "sub_group_scan_exclusive");
: func_name = "sub_group_scan_exclusive";
// for uniform case take into consideration all workitems // for uniform case take into consideration all workitems
if (!work_items_mask.any()) if (!work_items_mask.any())
@@ -656,8 +665,6 @@ template <typename Ty, ArithmeticOp operation> struct SCEX_NU
m += 4 * nw; m += 4 * nw;
} }
log_info(" %s_%s(%s)... passed\n", func_name.c_str(),
operation_names(operation), TypeManager<Ty>::name());
return TEST_PASS; return TEST_PASS;
} }
}; };
@@ -665,20 +672,24 @@ template <typename Ty, ArithmeticOp operation> struct SCEX_NU
// Test for scan inclusive non uniform functions // Test for scan inclusive non uniform functions
template <typename Ty, ArithmeticOp operation> struct SCIN_NU 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) static void gen(Ty *x, Ty *t, cl_int *m, const WorkGroupParams &test_params)
{ {
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;
int ng = test_params.global_workgroup_size; int ng = test_params.global_workgroup_size;
ng = ng / nw; ng = ng / nw;
std::string func_name;
test_params.work_items_mask.any()
? func_name = "sub_group_non_uniform_scan_inclusive"
: func_name = "sub_group_scan_inclusive";
genrand<Ty, operation>(x, t, m, ns, nw, ng); genrand<Ty, operation>(x, t, m, ns, nw, ng);
log_info(" %s_%s(%s)...\n", func_name.c_str(),
operation_names(operation), TypeManager<Ty>::name());
} }
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,
@@ -694,10 +705,9 @@ template <typename Ty, ArithmeticOp operation> struct SCIN_NU
Ty tr, rr; Ty tr, rr;
ng = ng / nw; ng = ng / nw;
std::string func_name; std::string func_name = (test_params.all_work_item_masks.size() > 0
work_items_mask.any() ? "sub_group_non_uniform_scan_inclusive"
? func_name = "sub_group_non_uniform_scan_inclusive" : "sub_group_scan_inclusive");
: func_name = "sub_group_scan_inclusive";
// for uniform case take into consideration all workitems // for uniform case take into consideration all workitems
if (!work_items_mask.any()) if (!work_items_mask.any())
@@ -771,8 +781,6 @@ template <typename Ty, ArithmeticOp operation> struct SCIN_NU
m += 4 * nw; m += 4 * nw;
} }
log_info(" %s_%s(%s)... passed\n", func_name.c_str(),
operation_names(operation), TypeManager<Ty>::name());
return TEST_PASS; return TEST_PASS;
} }
}; };
@@ -780,6 +788,16 @@ template <typename Ty, ArithmeticOp operation> struct SCIN_NU
// Test for reduce non uniform functions // Test for reduce non uniform functions
template <typename Ty, ArithmeticOp operation> struct RED_NU 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) static void gen(Ty *x, Ty *t, cl_int *m, const WorkGroupParams &test_params)
{ {
@@ -787,13 +805,6 @@ template <typename Ty, ArithmeticOp operation> struct RED_NU
int ns = test_params.subgroup_size; int ns = test_params.subgroup_size;
int ng = test_params.global_workgroup_size; int ng = test_params.global_workgroup_size;
ng = ng / nw; ng = ng / nw;
std::string func_name;
test_params.work_items_mask.any()
? func_name = "sub_group_non_uniform_reduce"
: func_name = "sub_group_reduce";
log_info(" %s_%s(%s)...\n", func_name.c_str(),
operation_names(operation), TypeManager<Ty>::name());
genrand<Ty, operation>(x, t, m, ns, nw, ng); genrand<Ty, operation>(x, t, m, ns, nw, ng);
} }
@@ -809,9 +820,9 @@ template <typename Ty, ArithmeticOp operation> struct RED_NU
ng = ng / nw; ng = ng / nw;
Ty tr, rr; Ty tr, rr;
std::string func_name; std::string func_name = (test_params.all_work_item_masks.size() > 0
work_items_mask.any() ? func_name = "sub_group_non_uniform_reduce" ? "sub_group_non_uniform_reduce"
: func_name = "sub_group_reduce"; : "sub_group_reduce");
for (k = 0; k < ng; ++k) for (k = 0; k < ng; ++k)
{ {
@@ -875,8 +886,6 @@ template <typename Ty, ArithmeticOp operation> struct RED_NU
m += 4 * nw; m += 4 * nw;
} }
log_info(" %s_%s(%s)... passed\n", func_name.c_str(),
operation_names(operation), TypeManager<Ty>::name());
return TEST_PASS; return TEST_PASS;
} }
}; };

32
test_conformance/subgroups/subhelpers.h
View File

@@ -1380,20 +1380,42 @@ template <typename Ty, typename Fns, size_t TSIZE = 0> struct test
const char *kname, const char *src, const char *kname, const char *src,
WorkGroupParams test_params) WorkGroupParams test_params)
{ {
Fns::log_test(test_params, "");
test_status combined_error = TEST_SKIPPED_ITSELF; test_status combined_error = TEST_SKIPPED_ITSELF;
for (auto &mask : test_params.all_work_item_masks) for (auto &mask : test_params.all_work_item_masks)
{ {
test_params.work_items_mask = mask; test_params.work_items_mask = mask;
test_status error = run(device, context, queue, num_elements, kname, test_status error = do_run(device, context, queue, num_elements,
src, test_params); kname, src, test_params);
if (error == TEST_FAIL if (error == TEST_FAIL
|| (error == TEST_PASS && combined_error != TEST_FAIL)) || (error == TEST_PASS && combined_error != TEST_FAIL))
combined_error = error; combined_error = error;
} }
if (combined_error == TEST_PASS)
{
Fns::log_test(test_params, " passed");
}
return combined_error; return combined_error;
}; };
static test_status run(cl_device_id device, cl_context context, 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)
{
Fns::log_test(test_params, "");
int error = do_run(device, context, queue, num_elements, kname, src,
test_params);
if (error == TEST_PASS)
{
Fns::log_test(test_params, " passed");
}
return error;
};
static test_status do_run(cl_device_id device, cl_context context,
cl_command_queue queue, int num_elements, cl_command_queue queue, int num_elements,
const char *kname, const char *src, const char *kname, const char *src,
WorkGroupParams test_params) WorkGroupParams test_params)
@@ -1442,8 +1464,7 @@ template <typename Ty, typename Fns, size_t TSIZE = 0> struct test
log_info("Data type not supported : %s\n", TypeManager<Ty>::name()); log_info("Data type not supported : %s\n", TypeManager<Ty>::name());
return TEST_SKIPPED_ITSELF; return TEST_SKIPPED_ITSELF;
} }
else
{
if (strstr(TypeManager<Ty>::name(), "double")) if (strstr(TypeManager<Ty>::name(), "double"))
{ {
kernel_sstr << "#pragma OPENCL EXTENSION cl_khr_fp64: enable\n"; kernel_sstr << "#pragma OPENCL EXTENSION cl_khr_fp64: enable\n";
@@ -1452,7 +1473,6 @@ template <typename Ty, typename Fns, size_t TSIZE = 0> struct test
{ {
kernel_sstr << "#pragma OPENCL EXTENSION cl_khr_fp16: enable\n"; kernel_sstr << "#pragma OPENCL EXTENSION cl_khr_fp16: enable\n";
} }
}
error = clGetDeviceInfo(device, CL_DEVICE_PLATFORM, sizeof(platform), error = clGetDeviceInfo(device, CL_DEVICE_PLATFORM, sizeof(platform),
(void *)&platform, NULL); (void *)&platform, NULL);

16
test_conformance/subgroups/test_barrier.cpp
View File

@@ -59,6 +59,17 @@ static const char *gbar_source =
// barrier test functions // barrier test functions
template <int Which> struct BAR template <int Which> struct BAR
{ {
static void log_test(const WorkGroupParams &test_params,
const char *extra_text)
{
if (Which == 0)
log_info(" sub_group_barrier(CLK_LOCAL_MEM_FENCE)...%s\n",
extra_text);
else
log_info(" sub_group_barrier(CLK_GLOBAL_MEM_FENCE)...%s\n",
extra_text);
}
static void gen(cl_int *x, cl_int *t, cl_int *m, static void gen(cl_int *x, cl_int *t, cl_int *m,
const WorkGroupParams &test_params) const WorkGroupParams &test_params)
{ {
@@ -103,11 +114,6 @@ template <int Which> struct BAR
ng = ng / nw; ng = ng / nw;
cl_int tr, rr; cl_int tr, rr;
if (Which == 0)
log_info(" sub_group_barrier(CLK_LOCAL_MEM_FENCE)...\n");
else
log_info(" sub_group_barrier(CLK_GLOBAL_MEM_FENCE)...\n");
for (k = 0; k < ng; ++k) for (k = 0; k < ng; ++k)
{ {
// Map to array indexed to array indexed by local ID and sub group // Map to array indexed to array indexed by local ID and sub group

8
test_conformance/subgroups/test_ifp.cpp
View File

@@ -225,6 +225,12 @@ void run_insts(cl_int *x, cl_int *p, int n)
struct IFP struct IFP
{ {
static void log_test(const WorkGroupParams &test_params,
const char *extra_text)
{
log_info(" independent forward progress...%s\n", extra_text);
}
static void gen(cl_int *x, cl_int *t, cl_int *, static void gen(cl_int *x, cl_int *t, cl_int *,
const WorkGroupParams &test_params) const WorkGroupParams &test_params)
{ {
@@ -258,8 +264,6 @@ struct IFP
// We need at least 2 sub groups per group for this test // We need at least 2 sub groups per group for this test
if (nj == 1) return TEST_SKIPPED_ITSELF; if (nj == 1) return TEST_SKIPPED_ITSELF;
log_info(" independent forward progress...\n");
for (k = 0; k < ng; ++k) for (k = 0; k < ng; ++k)
{ {
run_insts(x, t, nj); run_insts(x, t, nj);

9
test_conformance/subgroups/test_subgroup.cpp
View File

@@ -24,6 +24,13 @@ namespace {
// Any/All test functions // Any/All test functions
template <NonUniformVoteOp operation> struct AA template <NonUniformVoteOp operation> struct AA
{ {
static void log_test(const WorkGroupParams &test_params,
const char *extra_text)
{
log_info(" sub_group_%s...%s\n", operation_names(operation),
extra_text);
}
static void gen(cl_int *x, cl_int *t, cl_int *m, static void gen(cl_int *x, cl_int *t, cl_int *m,
const WorkGroupParams &test_params) const WorkGroupParams &test_params)
{ {
@@ -35,7 +42,6 @@ template <NonUniformVoteOp operation> struct AA
int e; int e;
ng = ng / nw; ng = ng / nw;
ii = 0; ii = 0;
log_info(" sub_group_%s...\n", operation_names(operation));
for (k = 0; k < ng; ++k) for (k = 0; k < ng; ++k)
{ {
for (j = 0; j < nj; ++j) for (j = 0; j < nj; ++j)
@@ -124,7 +130,6 @@ template <NonUniformVoteOp operation> struct AA
y += nw; y += nw;
m += 4 * nw; m += 4 * nw;
} }
log_info(" sub_group_%s... passed\n", operation_names(operation));
return TEST_PASS; return TEST_PASS;
} }
}; };

48
test_conformance/subgroups/test_subgroup_ballot.cpp
View File

@@ -23,6 +23,12 @@ namespace {
// Test for ballot functions // Test for ballot functions
template <typename Ty> struct BALLOT template <typename Ty> struct BALLOT
{ {
static void log_test(const WorkGroupParams &test_params,
const char *extra_text)
{
log_info(" sub_group_ballot...%s\n", extra_text);
}
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)
{ {
// no work here // no work here
@@ -30,7 +36,6 @@ template <typename Ty> struct BALLOT
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 non_uniform_size = gws % lws; int non_uniform_size = gws % lws;
log_info(" sub_group_ballot...\n");
} }
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,
@@ -92,7 +97,6 @@ template <typename Ty> struct BALLOT
y += lws; y += lws;
m += 4 * lws; m += 4 * lws;
} }
log_info(" sub_group_ballot... passed\n");
return TEST_PASS; return TEST_PASS;
} }
}; };
@@ -100,6 +104,13 @@ template <typename Ty> struct BALLOT
// Test for bit extract ballot functions // Test for bit extract ballot functions
template <typename Ty, BallotOp operation> struct BALLOT_BIT_EXTRACT template <typename Ty, BallotOp operation> struct BALLOT_BIT_EXTRACT
{ {
static void log_test(const WorkGroupParams &test_params,
const char *extra_text)
{
log_info(" sub_group_ballot_%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) 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, l;
@@ -110,8 +121,6 @@ template <typename Ty, BallotOp operation> struct BALLOT_BIT_EXTRACT
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; int non_uniform_size = gws % lws;
log_info(" sub_group_%s(%s)...\n", operation_names(operation),
TypeManager<Ty>::name());
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
@@ -251,21 +260,24 @@ template <typename Ty, BallotOp operation> struct BALLOT_BIT_EXTRACT
y += lws; y += lws;
m += 4 * lws; m += 4 * lws;
} }
log_info(" sub_group_%s(%s)... passed\n", operation_names(operation),
TypeManager<Ty>::name());
return TEST_PASS; return TEST_PASS;
} }
}; };
template <typename Ty, BallotOp operation> struct BALLOT_INVERSE template <typename Ty, BallotOp operation> struct BALLOT_INVERSE
{ {
static void log_test(const WorkGroupParams &test_params,
const char *extra_text)
{
log_info(" sub_group_inverse_ballot...%s\n", extra_text);
}
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 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 non_uniform_size = gws % lws; int non_uniform_size = gws % lws;
log_info(" sub_group_inverse_ballot...\n");
// no work here // no work here
} }
@@ -341,7 +353,6 @@ template <typename Ty, BallotOp operation> struct BALLOT_INVERSE
m += 4 * lws; m += 4 * lws;
} }
log_info(" sub_group_inverse_ballot... passed\n");
return TEST_PASS; return TEST_PASS;
} }
}; };
@@ -350,6 +361,13 @@ template <typename Ty, BallotOp operation> struct BALLOT_INVERSE
// Test for bit count/inclusive and exclusive scan/ find lsb msb ballot function // Test for bit count/inclusive and exclusive scan/ find lsb msb ballot function
template <typename Ty, BallotOp operation> struct BALLOT_COUNT_SCAN_FIND template <typename Ty, BallotOp operation> struct BALLOT_COUNT_SCAN_FIND
{ {
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) static void gen(Ty *x, Ty *t, cl_int *m, const WorkGroupParams &test_params)
{ {
int wi_id, wg_id, sb_id; int wi_id, wg_id, sb_id;
@@ -362,8 +380,6 @@ template <typename Ty, BallotOp operation> struct BALLOT_COUNT_SCAN_FIND
int last_subgroup_size = 0; int last_subgroup_size = 0;
int current_sbs = 0; int current_sbs = 0;
log_info(" sub_group_%s(%s)...\n", operation_names(operation),
TypeManager<Ty>::name());
if (non_uniform_size) if (non_uniform_size)
{ {
wg_number++; wg_number++;
@@ -562,8 +578,6 @@ template <typename Ty, BallotOp operation> struct BALLOT_COUNT_SCAN_FIND
y += lws; y += lws;
m += 4 * lws; m += 4 * lws;
} }
log_info(" sub_group_ballot_%s(%s)... passed\n",
operation_names(operation), TypeManager<Ty>::name());
return TEST_PASS; return TEST_PASS;
} }
}; };
@@ -571,6 +585,13 @@ template <typename Ty, BallotOp operation> struct BALLOT_COUNT_SCAN_FIND
// test mask functions // test mask functions
template <typename Ty, BallotOp operation> struct SMASK template <typename Ty, BallotOp operation> struct SMASK
{ {
static void log_test(const WorkGroupParams &test_params,
const char *extra_text)
{
log_info(" get_sub_group_%s_mask...%s\n", operation_names(operation),
extra_text);
}
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, l, sb_id;
@@ -579,7 +600,6 @@ template <typename Ty, BallotOp operation> struct SMASK
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;
log_info(" get_sub_group_%s_mask...\n", operation_names(operation));
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
for (sb_id = 0; sb_id < sb_number; ++sb_id) for (sb_id = 0; sb_id < sb_number; ++sb_id)
@@ -655,8 +675,6 @@ template <typename Ty, BallotOp operation> struct SMASK
y += lws; y += lws;
m += 4 * lws; m += 4 * lws;
} }
log_info(" get_sub_group_%s_mask... passed\n",
operation_names(operation));
return TEST_PASS; return TEST_PASS;
} }
}; };

14
test_conformance/subgroups/test_subgroup_clustered_reduce.cpp
View File

@@ -38,15 +38,20 @@ __kernel void test_%s(const __global Type *in, __global int4 *xy, __global Type
// Test for reduce cluster functions // Test for reduce cluster functions
template <typename Ty, ArithmeticOp operation> struct RED_CLU template <typename Ty, ArithmeticOp operation> struct RED_CLU
{ {
static void log_test(const WorkGroupParams &test_params,
const char *extra_text)
{
log_info(" sub_group_clustered_reduce_%s(%s, %d bytes) ...%s\n",
operation_names(operation), TypeManager<Ty>::name(),
sizeof(Ty), extra_text);
}
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 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;
int ng = test_params.global_workgroup_size; int ng = test_params.global_workgroup_size;
ng = ng / nw; ng = ng / nw;
log_info(" sub_group_clustered_reduce_%s(%s, %d bytes) ...\n",
operation_names(operation), TypeManager<Ty>::name(),
sizeof(Ty));
genrand<Ty, operation>(x, t, m, ns, nw, ng); genrand<Ty, operation>(x, t, m, ns, nw, ng);
} }
@@ -124,9 +129,6 @@ template <typename Ty, ArithmeticOp operation> struct RED_CLU
y += nw; y += nw;
m += 4 * nw; m += 4 * nw;
} }
log_info(" sub_group_clustered_reduce_%s(%s, %d bytes) ... passed\n",
operation_names(operation), TypeManager<Ty>::name(),
sizeof(Ty));
return TEST_PASS; return TEST_PASS;
} }
}; };

16
test_conformance/subgroups/test_subgroup_non_uniform_vote.cpp
View File

@@ -22,6 +22,15 @@ namespace {
template <typename T, NonUniformVoteOp operation> struct VOTE template <typename T, NonUniformVoteOp operation> struct VOTE
{ {
static void log_test(const WorkGroupParams &test_params,
const char *extra_text)
{
log_info(" sub_group_%s%s(%s)...%s\n",
(operation == NonUniformVoteOp::elect) ? "" : "non_uniform_",
operation_names(operation), TypeManager<T>::name(),
extra_text);
}
static void gen(T *x, T *t, cl_int *m, const WorkGroupParams &test_params) static void gen(T *x, T *t, cl_int *m, const WorkGroupParams &test_params)
{ {
int i, ii, j, k, n; int i, ii, j, k, n;
@@ -34,10 +43,6 @@ template <typename T, NonUniformVoteOp operation> struct VOTE
int last_subgroup_size = 0; int last_subgroup_size = 0;
ii = 0; ii = 0;
log_info(" sub_group_%s%s(%s)... \n",
(operation == NonUniformVoteOp::elect) ? "" : "non_uniform_",
operation_names(operation), TypeManager<T>::name());
if (operation == NonUniformVoteOp::elect) return; if (operation == NonUniformVoteOp::elect) return;
for (k = 0; k < ng; ++k) for (k = 0; k < ng; ++k)
@@ -192,9 +197,6 @@ template <typename T, NonUniformVoteOp operation> struct VOTE
m += 4 * nw; m += 4 * nw;
} }
log_info(" sub_group_%s%s(%s)... passed\n",
(operation == NonUniformVoteOp::elect) ? "" : "non_uniform_",
operation_names(operation), TypeManager<T>::name());
return TEST_PASS; return TEST_PASS;
} }
}; };