From ecb23100015cc7a3fde497bd11cda5b75db094dc Mon Sep 17 00:00:00 2001 From: Alexey Sachkov Date: Fri, 27 Oct 2023 06:31:36 -0700 Subject: [PATCH 1/2] Refactor `vector_swizzles` test This is a partial refactoring of the test, which aims to: - reduce kernel size, which allows to speed up JIT compilation and therefore test execution - improve error reporting from the test, by returning more than just a single boolean value for all of the checks performed by kernels Main changes: - introduced new template for swizzle test - "basic" `.swizzle<>()` checks now report results back to host where their verification happens instead of doing so on device side --- tests/common/common_python_vec.py | 117 ++++++++++++++++++-------- tests/common/vector_swizzles.template | 8 ++ 2 files changed, 89 insertions(+), 36 deletions(-) diff --git a/tests/common/common_python_vec.py b/tests/common/common_python_vec.py index c47a74a41..6e2980b6c 100644 --- a/tests/common/common_python_vec.py +++ b/tests/common/common_python_vec.py @@ -180,6 +180,42 @@ class ReverseData: CHECK(resArray[0]); """) +# Unlike in the template above, some of the checks which are done by tests for +# swizzles are outlined to host code. +# This is done to reduce kernel size and have more information about a problem +# in case of a test failure. +swizzle_kernel_template = Template(""" + { + auto vecBuffer = sycl::buffer, 1>( + sycl::range<1>(total_per_element_swizzle_test_cases)); + bool resArray[1] = {true}; + { + sycl::buffer boolBuffer(resArray, sycl::range<1>(1)); + testQueue.submit([&](sycl::handler &cgh) { + sycl::accessor resAcc(boolBuffer, cgh, sycl::write_only); + sycl::accessor vecAcc(vecBuffer, cgh, sycl::write_only); + + cgh.single_task([=]() { + ${test} + }); + }); + } + INFO("Checking ${testName}"); + CHECK(resArray[0]); + + ${type} in_order_vals[] = {${in_order_vals}}; + ${type} reversed_vals[] = {${reversed_vals}}; + ${type} in_order_reversed_pair_vals[] = {${in_order_pair_vals}}; + ${type} reverse_order_reversed_pair_vals[] = {${reverse_order_pair_vals}}; + + auto vecAcc = vecBuffer.get_host_access(); + CHECK(check_vector_values<${type}, ${size}>(vecAcc[in_order], in_order_vals)); + CHECK(check_vector_values<${type}, ${size}>(vecAcc[reverse_order], reversed_vals)); + CHECK(check_vector_values<${type}, ${size}>(vecAcc[in_order_reversed_pair], in_order_reversed_pair_vals)); + CHECK(check_vector_values<${type}, ${size}>(vecAcc[reverse_order_reversed_pair], reverse_order_reversed_pair_vals)); + } +""") + test_func_template = Template(""" void ${func_name}(util::logger &log) { @@ -224,6 +260,33 @@ def wrap_with_kernel(type_str, kernel_name, test_name, test_string): testName=test_name, test=test_string)) +def wrap_with_swizzle_kernel(type_str, vec_size, in_order_vals, reversed_vals, + in_order_pair_vals, reverse_order_pair_vals, kernel_name, test_name, + test_string): + """ + Wraps |test_string| inside a kernel with |kernel_name|. + + Wraps kernels with checks for fp16 and fp64 when appropriate. + The necessity for extension checks is determined based on |type_str| + + Unlike |wrap_with_kernel| above, this function accepts several extra + arguments, because kernels for swizzle tests do some of the validation + checks on host to reduce kernel size and make error messages more detailed. + """ + + return wrap_with_extension_checks(type_str, + swizzle_kernel_template.substitute( + kernelName=remove_namespaces_whitespaces(kernel_name), + testName=test_name, + test=test_string, + type=type_str, + size=vec_size, + in_order_vals=in_order_vals, + reversed_vals=reversed_vals, + in_order_pair_vals=in_order_pair_vals, + reverse_order_pair_vals=reverse_order_pair_vals)) + + def wrap_with_test_func(test_name, type_str, test, additional=''): """ @@ -419,37 +482,23 @@ class SwizzleData: """) swizzle_elem_template = Template( - """ sycl::vec<${type}, ${size}> inOrderSwizzleFunctionVec {swizzledVec.template swizzle<${in_order_swiz_indexes}>()}; - if (!check_vector_values<${type}, ${size}>(inOrderSwizzleFunctionVec, in_order_vals)) { - resAcc[0] = false; - } - sycl::vec<${type}, ${size}> reverseOrderSwizzleFunctionVec {swizzledVec.template swizzle<${reverse_order_swiz_indexes}>()}; - if (!check_vector_values<${type}, ${size}>(reverseOrderSwizzleFunctionVec, reversed_vals)) { - resAcc[0] = false; - } - sycl::vec<${type}, ${size}> inOrderReversedPairSwizzleFunctionVec {swizzledVec.template swizzle<${in_order_reversed_pair_swiz_indexes}>()}; - if (!check_vector_values<${type}, ${size}>(inOrderReversedPairSwizzleFunctionVec, in_order_reversed_pair_vals)) { - resAcc[0] = false; - } - sycl::vec<${type}, ${size}> reverseOrderReversedPairSwizzleFunctionVec {swizzledVec.template swizzle<${reverse_order_reversed_pair_swiz_indexes}>()}; - if (!check_vector_values<${type}, ${size}>(reverseOrderReversedPairSwizzleFunctionVec, reverse_order_reversed_pair_vals)) { - resAcc[0] = false; - } + """ + vecAcc[in_order] = swizzledVec.template swizzle<${in_order_swiz_indexes}>(); + vecAcc[reverse_order] = swizzledVec.template swizzle<${reverse_order_swiz_indexes}>(); + vecAcc[in_order_reversed_pair] = swizzledVec.template swizzle<${in_order_reversed_pair_swiz_indexes}>(); + vecAcc[reverse_order_reversed_pair] = swizzledVec.template swizzle<${reverse_order_reversed_pair_swiz_indexes}>(); """) swizzle_full_test_template = Template( """ sycl::vec<${type}, ${size}> ${name}DimTestVec = sycl::vec<${type}, ${size}>(${testVecValues}); - ${type} in_order_vals[] = {${in_order_vals}}; sycl::vec<${type}, ${size}> inOrderSwizzleFunctionVec {${name}DimTestVec.template swizzle<${in_order_swiz_indexes}>()}; + vecAcc[in_order] = inOrderSwizzleFunctionVec; if (!check_equal_type_bool>(inOrderSwizzleFunctionVec)) { resAcc[0] = false; } if (!check_vector_size<${type}, ${size}>(inOrderSwizzleFunctionVec)) { resAcc[0] = false; } - if (!check_vector_values<${type}, ${size}>(inOrderSwizzleFunctionVec, in_order_vals)) { - resAcc[0] = false; - } if (!check_vector_size_byte_size<${type}, ${size}>(inOrderSwizzleFunctionVec)) { resAcc[0] = false; } @@ -459,17 +508,14 @@ class SwizzleData: } #endif // SYCL_CTS_ENABLE_FULL_CONFORMANCE - ${type} reversed_vals[] = {${reversed_vals}}; sycl::vec<${type}, ${size}> reverseOrderSwizzleFunctionVec {${name}DimTestVec.template swizzle<${reverse_order_swiz_indexes}>()}; + vecAcc[reverse_order] = reverseOrderSwizzleFunctionVec; if (!check_equal_type_bool>(reverseOrderSwizzleFunctionVec)) { resAcc[0] = false; } if (!check_vector_size<${type}, ${size}>(reverseOrderSwizzleFunctionVec)) { resAcc[0] = false; } - if (!check_vector_values<${type}, ${size}>(reverseOrderSwizzleFunctionVec, reversed_vals)) { - resAcc[0] = false; - } if (!check_vector_size_byte_size<${type}, ${size}>(reverseOrderSwizzleFunctionVec)) { resAcc[0] = false; } @@ -479,17 +525,14 @@ class SwizzleData: } #endif // SYCL_CTS_ENABLE_FULL_CONFORMANCE - ${type} in_order_reversed_pair_vals[] = {${in_order_pair_vals}}; sycl::vec<${type}, ${size}> inOrderReversedPairSwizzleFunctionVec {${name}DimTestVec.template swizzle<${in_order_reversed_pair_swiz_indexes}>()}; + vecAcc[in_order_reversed_pair] = inOrderReversedPairSwizzleFunctionVec; if (!check_equal_type_bool>(inOrderReversedPairSwizzleFunctionVec)) { resAcc[0] = false; } if (!check_vector_size<${type}, ${size}>(inOrderReversedPairSwizzleFunctionVec)) { resAcc[0] = false; } - if (!check_vector_values<${type}, ${size}>(inOrderReversedPairSwizzleFunctionVec, in_order_reversed_pair_vals)) { - resAcc[0] = false; - } if (!check_vector_size_byte_size<${type}, ${size}>(inOrderReversedPairSwizzleFunctionVec)) { resAcc[0] = false; } @@ -499,17 +542,14 @@ class SwizzleData: } #endif // SYCL_CTS_ENABLE_FULL_CONFORMANCE - ${type} reverse_order_reversed_pair_vals[] = {${reverse_order_pair_vals}}; sycl::vec<${type}, ${size}> reverseOrderReversedPairSwizzleFunctionVec {${name}DimTestVec.template swizzle<${reverse_order_reversed_pair_swiz_indexes}>()}; + vecAcc[reverse_order_reversed_pair] = reverseOrderReversedPairSwizzleFunctionVec; if (!check_equal_type_bool>(reverseOrderReversedPairSwizzleFunctionVec)) { resAcc[0] = false; } if (!check_vector_size<${type}, ${size}>(reverseOrderReversedPairSwizzleFunctionVec)) { resAcc[0] = false; } - if (!check_vector_values<${type}, ${size}>(reverseOrderReversedPairSwizzleFunctionVec, reverse_order_reversed_pair_vals)) { - resAcc[0] = false; - } if (!check_vector_size_byte_size<${type}, ${size}>(reverseOrderReversedPairSwizzleFunctionVec)) { resAcc[0] = false; } @@ -559,8 +599,9 @@ def substitute_swizzles_templates(type_str, size, index_subset, value_subset, co swap_pairs(Data.swizzle_elem_list_dict[size])), reverse_order_reversed_pair_swiz_indexes=', '.join( swap_pairs(Data.swizzle_elem_list_dict[size][::-1]))) - string += wrap_with_kernel( - type_str, + string += wrap_with_swizzle_kernel( + type_str, str(size), ', '.join(val_list), ', '.join(val_list[::-1]), + ', '.join(swap_pairs(val_list)), ', '.join(swap_pairs(val_list[::-1])), 'KERNEL_' + type_str + str(size) + index_string, 'vec<' + type_str + ', ' + str(size) + '>.' + index_string, @@ -590,8 +631,12 @@ def gen_swizzle_test(type_str, convert_type_str, as_type_str, size): swap_pairs(Data.vals_list_dict[size])), reverse_order_pair_vals=', '.join( swap_pairs(Data.vals_list_dict[size][::-1]))) - string += wrap_with_kernel( - type_str, 'ELEM_KERNEL_' + type_str + str(size) + + string += wrap_with_swizzle_kernel( + type_str, str(size), ', '.join(Data.vals_list_dict[size]), + ', '.join(Data.vals_list_dict[size][::-1]), + ', '.join(swap_pairs(Data.vals_list_dict[size])), + ', '.join(swap_pairs(Data.vals_list_dict[size][::-1])), + 'ELEM_KERNEL_' + type_str + str(size) + ''.join(Data.swizzle_elem_list_dict[size][:size]).replace( 'sycl::elem::', ''), 'vec<' + type_str + ', ' + str(size) + '> .swizzle<' + diff --git a/tests/common/vector_swizzles.template b/tests/common/vector_swizzles.template index 6f9a97e0b..abd8dbb37 100644 --- a/tests/common/vector_swizzles.template +++ b/tests/common/vector_swizzles.template @@ -34,6 +34,14 @@ using namespace sycl_cts; * and used like a normal vector */ class TEST_NAME : public util::test_base { + enum per_element_swizzle_test_case : int { + in_order = 0, + reverse_order, + in_order_reversed_pair, + reverse_order_reversed_pair, + total_per_element_swizzle_test_cases + }; + public: /** return information about this test */ From 35fbe6511cd9621700083bdbff1fd558e33ea478 Mon Sep 17 00:00:00 2001 From: Alexey Sachkov Date: Thu, 9 Nov 2023 07:43:15 -0800 Subject: [PATCH 2/2] [NFC] Record a bunch of FIXMEs in the code --- tests/common/common_python_vec.py | 15 +++++++++++++++ 1 file changed, 15 insertions(+) diff --git a/tests/common/common_python_vec.py b/tests/common/common_python_vec.py index 6e2980b6c..98f0b0fc0 100644 --- a/tests/common/common_python_vec.py +++ b/tests/common/common_python_vec.py @@ -448,6 +448,17 @@ def get_types(): types.append(Data.fixed_width_type_dict[(sign, base_type)]) return types +# FIXME: We shouldn't accumulate results of *all* checks into a single boolean, +# because it makes debugging failures really hard as there is no indication +# about what exactly went wrong. +# FIXME: vector swizzles should be tested not only on device, but also on host. +# Therefore, the test should ideally perform the same calculations on both host +# and device and then compare to ensure that results are equivalend and match +# reference. +# FIXME: consider reducing amount of calls to 'check_convert_as_all_types' +# swizzle_template invokes 'check_convert_as_all_types' and this template +# is "instantiated" for every simple swizzle there is (such as xyzz, zyxw, etc.) +# which seems overly exessive. class SwizzleData: swizzle_template = Template( """ sycl::vec<${type}, ${size}> ${name}DimTestVec = sycl::vec<${type}, ${size}>(${testVecValues}); @@ -709,6 +720,10 @@ def get_reverse_type(type_str): reverse_type_str = type_str return reverse_type_str +# FIXME: Move this and other functions to generate_vector_swizzles.py +# Reason for the TODO above is that this function and several more it calls are +# not really common and only used to generate vector_swizzles test. +# FIXME: The test (main template and others) should be updated to use Catch2 def make_swizzles_tests(type_str, input_file, output_file): if type_str == 'bool': Data.vals_list_dict = cast_to_bool(Data.vals_list_dict)