From 903f1bf65dfe15956295eb9379f5706568d858a9 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?K=C3=A9vin=20Petit?= Date: Thu, 30 Sep 2021 13:33:18 +0100 Subject: [PATCH] Report unsupported extended subgroup tests as skipped rather than passed (#1301) MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit * Report unsupported extended subgroup tests as skipped rather than passed Also don't check the presence of extensions for each sub-test. Signed-off-by: Kévin Petit * address review comments --- test_conformance/subgroups/subhelpers.h | 17 +---------------- .../subgroups/test_subgroup_ballot.cpp | 10 +++++++--- .../test_subgroup_clustered_reduce.cpp | 12 +++++++----- .../subgroups/test_subgroup_extended_types.cpp | 12 +++++++----- .../test_subgroup_non_uniform_arithmetic.cpp | 15 +++++++++------ .../test_subgroup_non_uniform_vote.cpp | 13 +++++++------ .../subgroups/test_subgroup_shuffle.cpp | 10 +++++++--- .../test_subgroup_shuffle_relative.cpp | 12 +++++++----- 8 files changed, 52 insertions(+), 49 deletions(-) diff --git a/test_conformance/subgroups/subhelpers.h b/test_conformance/subgroups/subhelpers.h index 9232cded..0d497fb3 100644 --- a/test_conformance/subgroups/subhelpers.h +++ b/test_conformance/subgroups/subhelpers.h @@ -33,10 +33,9 @@ extern cl_half_rounding_mode g_rounding_mode; struct WorkGroupParams { WorkGroupParams(size_t gws, size_t lws, - const std::vector &req_ext = {}, const std::vector &all_wim = {}) : global_workgroup_size(gws), local_workgroup_size(lws), - required_extensions(req_ext), all_work_item_masks(all_wim) + all_work_item_masks(all_wim) { subgroup_size = 0; work_items_mask = 0; @@ -49,7 +48,6 @@ struct WorkGroupParams uint32_t work_items_mask; int dynsc; bool use_core_subgroups; - std::vector required_extensions; std::vector all_work_item_masks; }; @@ -1297,19 +1295,6 @@ template struct test } } - for (std::string extension : test_params.required_extensions) - { - if (!is_extension_available(device, extension.c_str())) - { - log_info("The extension %s not supported on this device. SKIP " - "testing - kernel %s data type %s\n", - extension.c_str(), kname, TypeManager::name()); - return TEST_PASS; - } - kernel_sstr << "#pragma OPENCL EXTENSION " + extension - + ": enable\n"; - } - error = clGetDeviceInfo(device, CL_DEVICE_PLATFORM, sizeof(platform), (void *)&platform, NULL); test_error(error, "clGetDeviceInfo failed for CL_DEVICE_PLATFORM"); diff --git a/test_conformance/subgroups/test_subgroup_ballot.cpp b/test_conformance/subgroups/test_subgroup_ballot.cpp index f2e4060b..9a2da5d9 100644 --- a/test_conformance/subgroups/test_subgroup_ballot.cpp +++ b/test_conformance/subgroups/test_subgroup_ballot.cpp @@ -926,11 +926,15 @@ template int run_non_uniform_broadcast_for_type(RunTestForType rft) int test_subgroup_functions_ballot(cl_device_id device, cl_context context, cl_command_queue queue, int num_elements) { - std::vector required_extensions = { "cl_khr_subgroup_ballot" }; + if (!is_extension_available(device, "cl_khr_subgroup_ballot")) + { + log_info("cl_khr_subgroup_ballot is not supported on this device, " + "skipping test.\n"); + return TEST_SKIPPED_ITSELF; + } constexpr size_t global_work_size = 170; constexpr size_t local_work_size = 64; - WorkGroupParams test_params(global_work_size, local_work_size, - required_extensions); + WorkGroupParams test_params(global_work_size, local_work_size); RunTestForType rft(device, context, queue, num_elements, test_params); // non uniform broadcast functions diff --git a/test_conformance/subgroups/test_subgroup_clustered_reduce.cpp b/test_conformance/subgroups/test_subgroup_clustered_reduce.cpp index 588e9cee..87507e37 100644 --- a/test_conformance/subgroups/test_subgroup_clustered_reduce.cpp +++ b/test_conformance/subgroups/test_subgroup_clustered_reduce.cpp @@ -305,13 +305,15 @@ int test_subgroup_functions_clustered_reduce(cl_device_id device, cl_command_queue queue, int num_elements) { - std::vector required_extensions = { - "cl_khr_subgroup_clustered_reduce" - }; + if (!is_extension_available(device, "cl_khr_subgroup_clustered_reduce")) + { + log_info("cl_khr_subgroup_clustered_reduce is not supported on this " + "device, skipping test.\n"); + return TEST_SKIPPED_ITSELF; + } constexpr size_t global_work_size = 2000; constexpr size_t local_work_size = 200; - WorkGroupParams test_params(global_work_size, local_work_size, - required_extensions); + WorkGroupParams test_params(global_work_size, local_work_size); RunTestForType rft(device, context, queue, num_elements, test_params); int error = run_cluster_red_add_max_min_mul_for_type(rft); diff --git a/test_conformance/subgroups/test_subgroup_extended_types.cpp b/test_conformance/subgroups/test_subgroup_extended_types.cpp index 98401b8e..b281f618 100644 --- a/test_conformance/subgroups/test_subgroup_extended_types.cpp +++ b/test_conformance/subgroups/test_subgroup_extended_types.cpp @@ -59,13 +59,15 @@ int test_subgroup_functions_extended_types(cl_device_id device, cl_command_queue queue, int num_elements) { - std::vector required_extensions = { - "cl_khr_subgroup_extended_types" - }; + if (!is_extension_available(device, "cl_khr_subgroup_extended_types")) + { + log_info("cl_khr_subgroup_extended_types is not supported on this " + "device, skipping test.\n"); + return TEST_SKIPPED_ITSELF; + } constexpr size_t global_work_size = 2000; constexpr size_t local_work_size = 200; - WorkGroupParams test_params(global_work_size, local_work_size, - required_extensions); + WorkGroupParams test_params(global_work_size, local_work_size); RunTestForType rft(device, context, queue, num_elements, test_params); int error = run_broadcast_for_extended_type(rft); diff --git a/test_conformance/subgroups/test_subgroup_non_uniform_arithmetic.cpp b/test_conformance/subgroups/test_subgroup_non_uniform_arithmetic.cpp index eb46ff09..6c44249e 100644 --- a/test_conformance/subgroups/test_subgroup_non_uniform_arithmetic.cpp +++ b/test_conformance/subgroups/test_subgroup_non_uniform_arithmetic.cpp @@ -434,17 +434,20 @@ int test_subgroup_functions_non_uniform_arithmetic(cl_device_id device, cl_command_queue queue, int num_elements) { - std::vector required_extensions = { - "cl_khr_subgroup_non_uniform_arithmetic" - }; + if (!is_extension_available(device, + "cl_khr_subgroup_non_uniform_arithmetic")) + { + log_info("cl_khr_subgroup_non_uniform_arithmetic is not supported on " + "this device, skipping test.\n"); + return TEST_SKIPPED_ITSELF; + } std::vector masks{ 0xffffffff, 0x55aaaa55, 0x5555aaaa, 0xaaaa5555, 0x0f0ff0f0, 0x0f0f0f0f, 0xff0000ff, 0xff00ff00, 0x00ffff00, 0x80000000, 0xaaaaaaaa }; constexpr size_t global_work_size = 2000; constexpr size_t local_work_size = 200; - WorkGroupParams test_params(global_work_size, local_work_size, - required_extensions, masks); + WorkGroupParams test_params(global_work_size, local_work_size, masks); RunTestForType rft(device, context, queue, num_elements, test_params); int error = run_functions_add_mul_max_min_for_type(rft); @@ -470,4 +473,4 @@ int test_subgroup_functions_non_uniform_arithmetic(cl_device_id device, error |= run_functions_logical_and_or_xor_for_type(rft); return error; -} \ No newline at end of file +} diff --git a/test_conformance/subgroups/test_subgroup_non_uniform_vote.cpp b/test_conformance/subgroups/test_subgroup_non_uniform_vote.cpp index 2b00b4dd..484e9b6b 100644 --- a/test_conformance/subgroups/test_subgroup_non_uniform_vote.cpp +++ b/test_conformance/subgroups/test_subgroup_non_uniform_vote.cpp @@ -272,17 +272,18 @@ int test_subgroup_functions_non_uniform_vote(cl_device_id device, cl_command_queue queue, int num_elements) { - std::vector required_extensions = { - "cl_khr_subgroup_non_uniform_vote" - }; - + if (!is_extension_available(device, "cl_khr_subgroup_non_uniform_vote")) + { + log_info("cl_khr_subgroup_non_uniform_vote is not supported on this " + "device, skipping test.\n"); + return TEST_SKIPPED_ITSELF; + } std::vector masks{ 0xffffffff, 0x55aaaa55, 0x5555aaaa, 0xaaaa5555, 0x0f0ff0f0, 0x0f0f0f0f, 0xff0000ff, 0xff00ff00, 0x00ffff00, 0x80000000 }; constexpr size_t global_work_size = 170; constexpr size_t local_work_size = 64; - WorkGroupParams test_params(global_work_size, local_work_size, - required_extensions, masks); + WorkGroupParams test_params(global_work_size, local_work_size, masks); RunTestForType rft(device, context, queue, num_elements, test_params); int error = run_vote_all_equal_for_type(rft); diff --git a/test_conformance/subgroups/test_subgroup_shuffle.cpp b/test_conformance/subgroups/test_subgroup_shuffle.cpp index 049f0982..37b27ced 100644 --- a/test_conformance/subgroups/test_subgroup_shuffle.cpp +++ b/test_conformance/subgroups/test_subgroup_shuffle.cpp @@ -55,11 +55,15 @@ template int run_shuffle_for_type(RunTestForType rft) int test_subgroup_functions_shuffle(cl_device_id device, cl_context context, cl_command_queue queue, int num_elements) { - std::vector required_extensions{ "cl_khr_subgroup_shuffle" }; + if (!is_extension_available(device, "cl_khr_subgroup_shuffle")) + { + log_info("cl_khr_subgroup_shuffle is not supported on this device, " + "skipping test.\n"); + return TEST_SKIPPED_ITSELF; + } constexpr size_t global_work_size = 2000; constexpr size_t local_work_size = 200; - WorkGroupParams test_params(global_work_size, local_work_size, - required_extensions); + WorkGroupParams test_params(global_work_size, local_work_size); RunTestForType rft(device, context, queue, num_elements, test_params); int error = run_shuffle_for_type(rft); diff --git a/test_conformance/subgroups/test_subgroup_shuffle_relative.cpp b/test_conformance/subgroups/test_subgroup_shuffle_relative.cpp index 6000c970..11401e80 100644 --- a/test_conformance/subgroups/test_subgroup_shuffle_relative.cpp +++ b/test_conformance/subgroups/test_subgroup_shuffle_relative.cpp @@ -56,13 +56,15 @@ int test_subgroup_functions_shuffle_relative(cl_device_id device, cl_command_queue queue, int num_elements) { - std::vector required_extensions = { - "cl_khr_subgroup_shuffle_relative" - }; + if (!is_extension_available(device, "cl_khr_subgroup_shuffle_relative")) + { + log_info("cl_khr_subgroup_shuffle_relative is not supported on this " + "device, skipping test.\n"); + return TEST_SKIPPED_ITSELF; + } constexpr size_t global_work_size = 2000; constexpr size_t local_work_size = 200; - WorkGroupParams test_params(global_work_size, local_work_size, - required_extensions); + WorkGroupParams test_params(global_work_size, local_work_size); RunTestForType rft(device, context, queue, num_elements, test_params); int error = run_shuffle_relative_for_type(rft);