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

Refactor vector_swizzles test #825

Merged
merged 3 commits into from
Nov 9, 2023
Merged

Refactor vector_swizzles test #825

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
3 commits
Select commit Hold shift + click to select a range
ecb2310
Refactor `vector_swizzles` test
AlexeySachkov Oct 27, 2023
78644ec
Merge remote-tracking branch 'origin/SYCL-2020' into private/asachkov…
AlexeySachkov Oct 31, 2023
35fbe65
[NFC] Record a bunch of FIXMEs in the code
AlexeySachkov Nov 9, 2023
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
132 changes: 96 additions & 36 deletions tests/common/common_python_vec.py
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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<sycl::vec<${type}, ${size}>, 1>(
sycl::range<1>(total_per_element_swizzle_test_cases));
bool resArray[1] = {true};
{
sycl::buffer<bool, 1> 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<class ${kernelName}>([=]() {
${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) {

Expand Down Expand Up @@ -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=''):
"""
Expand Down Expand Up @@ -385,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});
Expand Down Expand Up @@ -419,37 +493,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<sycl::vec<${type}, ${size}>>(inOrderSwizzleFunctionVec)) {
resAcc[0] = false;
}
if (!check_vector_size<${type}, ${size}>(inOrderSwizzleFunctionVec)) {
resAcc[0] = false;
}
if (!check_vector_values<${type}, ${size}>(inOrderSwizzleFunctionVec, in_order_vals)) {
Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Note for reviewers: there are two templates for kernel code:

  • one for vectors with 4 or less elements, which is used to exercise all simple swizzles like .xyzw; modified swizzle_elem_template is its part
  • another one for bigger vectors, where only per-element swizzles are tested; swizzle_full_test_template is used for it

Kernels for both categories of vectors do the same checks to ensure that .swizzle method returns elements in the right order, they are just described by two different code templates.

resAcc[0] = false;
}
if (!check_vector_size_byte_size<${type}, ${size}>(inOrderSwizzleFunctionVec)) {
resAcc[0] = false;
}
Expand All @@ -459,17 +519,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<sycl::vec<${type}, ${size}>>(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;
}
Expand All @@ -479,17 +536,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<sycl::vec<${type}, ${size}>>(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;
}
Expand All @@ -499,17 +553,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<sycl::vec<${type}, ${size}>>(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;
}
Expand Down Expand Up @@ -559,8 +610,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,
Expand Down Expand Up @@ -590,8 +642,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<' +
Expand Down Expand Up @@ -664,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)
Expand Down
8 changes: 8 additions & 0 deletions tests/common/vector_swizzles.template
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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
*/
Expand Down