Skip to content

[SYCL] Support or diagnose use of namespace std types as kernel type names #1963

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

Merged
merged 21 commits into from
Sep 14, 2020
Merged

[SYCL] Support or diagnose use of namespace std types as kernel type names #1963

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
21 commits
Select commit Hold shift + click to select a range
46c54c6
[SYCL] Support or diagnose use of namespace std types as kernel type …
premanandrao Apr 16, 2020
5e8004d
Emit diagnostic when kernel names use nullptr_t
jtmott-intel May 26, 2020
26b38f9
Merge from branch sycl
jtmott-intel Jun 23, 2020
c21b9df
Disallow std names in kernel names
jtmott-intel Jul 2, 2020
1d345ec
Clang format patch
jtmott-intel Jul 2, 2020
d2d55b0
Fixing tests, specifically tests with std kernels
jtmott-intel Jul 10, 2020
e678857
Merge branch 'sycl'
jtmott-intel Jul 14, 2020
e6afe60
Fixing tests, specifically tests with std kernels
jtmott-intel Jul 16, 2020
07e8899
Merge branch sycl
jtmott-intel Jul 31, 2020
e52236f
Clang format
jtmott-intel Jul 31, 2020
0340297
Merge branch sycl
jtmott-intel Jul 31, 2020
c7c155f
Merge branch sycl
jtmott-intel Aug 12, 2020
5752ac2
Turn on unnamed lambda by default
jtmott-intel Aug 15, 2020
be84c38
Revert "Turn on unnamed lambda by default"
jtmott-intel Aug 18, 2020
fce4460
Fixing tests, specifically tests with std kernels
jtmott-intel Aug 18, 2020
8e3ea41
Merge branch sycl
jtmott-intel Aug 18, 2020
bdee3e1
Merge branch sycl
jtmott-intel Aug 21, 2020
04bd440
Clang format and update expected error line numbers
jtmott-intel Aug 21, 2020
736118f
Merge branch sycl
jtmott-intel Sep 1, 2020
a96b556
Fix bad merge resolution
jtmott-intel Sep 5, 2020
8d0087e
Revert "[SYCL] Support or diagnose use of namespace std types as kern…
jtmott-intel Sep 11, 2020
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
1 change: 1 addition & 0 deletions clang/include/clang/Basic/DiagnosticSemaKinds.td
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -10977,6 +10977,7 @@ def err_sycl_kernel_incorrectly_named : Error<
"kernel %select{name is missing"
"|needs to have a globally-visible name"
"|name is invalid. Unscoped enum requires fixed underlying type"
"|name cannot be a type in the \"std\" namespace"
"}0">;
def err_sycl_kernel_not_function_object
: Error<"kernel parameter must be a lambda or function object">;
Expand Down
14 changes: 13 additions & 1 deletion clang/lib/Sema/SemaSYCL.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -2755,6 +2755,13 @@ void SYCLIntegrationHeader::emitFwdDecl(raw_ostream &O, const Decl *D,
}
break;
}

if (NS->isStdNamespace()) {
Copy link
Contributor

Choose a reason for hiding this comment

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

This whole thing is a non-issue if unnamed lambda compiler option is enabled even if the user provides a kernel name. Because then we don't template on the lambda name type but on the string representation of it. Therefore we don't need to emit an error in that case.

Copy link
Contributor

Choose a reason for hiding this comment

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

It seems like we could always do that: Template KernelInfo on the string representation of the name rather than the name type. Then we would never need to forward declare the name type.

Copy link
Contributor Author

@jtmott-intel jtmott-intel Aug 15, 2020 •
edited
Loading

Choose a reason for hiding this comment

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

I started experimenting with templating KernelInfo on the string representation of the name, and probably went further down that rabbit hole than I should have before I realized I was reimplementing what the existing unnamed lambda logic does already. And so in this latest commit -- meant to prompt discussion -- I flipped a 0 to a 1 (don't laugh :-P) and turned on the unnamed lambda CLI option by default. This seems to deliver exactly what we've been wanting. KernelInfo is templated on the string representation of the name, nullptr_t and other std names work, and the existing code will already skip forward decls when unnamed lambda is on. I assume, though, that the idea of turning on unnamed lambdas by default was already discussed in some meeting somewhere. Is there some reason it can't be this simple?

@rolandschulz @erichkeane @premanandrao @elizabethandrews

Copy link
Contributor

Choose a reason for hiding this comment

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

Right, we shouldn't emit this diagnostic if UnnamedLambdaSupport

Diag.Report(KernelLocation, diag::err_sycl_kernel_incorrectly_named)
<< /* name cannot be a type in the std namespace */ 3;
return;
}

++NamespaceCnt;
const StringRef NSInlinePrefix = NS->isInline() ? "inline " : "";
NSStr.insert(
Expand Down Expand Up @@ -2837,8 +2844,13 @@ void SYCLIntegrationHeader::emitForwardClassDecls(
;
const CXXRecordDecl *RD = T->getAsCXXRecordDecl();

if (!RD)
if (!RD) {
if (T->isNullPtrType())
Diag.Report(KernelLocation, diag::err_sycl_kernel_incorrectly_named)
<< /* name cannot be a type in the std namespace */ 3;

return;
}

// see if this is a template specialization ...
if (const auto *TSD = dyn_cast<ClassTemplateSpecializationDecl>(RD)) {
Expand Down
63 changes: 63 additions & 0 deletions clang/test/CodeGenSYCL/stdtypes_kernel_type.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -0,0 +1,63 @@
// RUN: %clang_cc1 -fsycl -fsycl-is-device -fsycl-int-header=%t.h -DCHECK_ERROR -verify %s
// RUN: %clang_cc1 -fsycl -fsycl-is-device -triple spir64-unknown-unknown-sycldevice -fsycl-int-header=%t.h %s
// RUN: FileCheck -input-file=%t.h %s
//
// CHECK: #include <CL/sycl/detail/defines.hpp>
// CHECK-NEXT: #include <CL/sycl/detail/kernel_desc.hpp>
//
// CHECK: static constexpr
// CHECK-NEXT: const char* const kernel_names[] = {
// CHECK-NEXT: "_ZTSm",
// CHECK-NEXT: "_ZTSl"
// CHECK-NEXT: };
//
// CHECK: static constexpr
// CHECK-NEXT: const kernel_param_desc_t kernel_signatures[] = {
// CHECK-NEXT: //--- _ZTSm
// CHECK-EMPTY:
// CHECK-NEXT: //--- _ZTSl
// CHECK-EMPTY:
// CHECK-NEXT: };
//
// CHECK: static constexpr
// CHECK-NEXT: const unsigned kernel_signature_start[] = {
// CHECK-NEXT: 0, // _ZTSm
// CHECK-NEXT: 1 // _ZTSl
// CHECK-NEXT: };

// CHECK: template <> struct KernelInfo<unsigned long> {
// CHECK: template <> struct KernelInfo<long> {

void usage() {
}

namespace std {
typedef long unsigned int size_t;
typedef long int ptrdiff_t;
typedef decltype(nullptr) nullptr_t;
class T;
class U;
} // namespace std

template <typename T>
struct Templated_kernel_name;

template <typename name, typename Func>
__attribute__((sycl_kernel)) void kernel_single_task(Func kernelFunc) {
kernelFunc();
}

int main() {
#ifdef CHECK_ERROR
kernel_single_task<std::nullptr_t>([=]() {}); // expected-error {{kernel name cannot be a type in the "std" namespace}}
kernel_single_task<std::T>([=]() {}); // expected-error {{kernel name cannot be a type in the "std" namespace}}
kernel_single_task<Templated_kernel_name<std::nullptr_t>>([=]() {}); // expected-error {{kernel name cannot be a type in the "std" namespace}}
Copy link
Contributor

Choose a reason for hiding this comment

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

I think this should be allowed. The KernelName here is Templated_kernel_name<std::nullptr_t> This is not in the 'std' namespace. I do not think there is any spec restriction which says we cannot template a kernel with types from std namespace.

I think this is why the tests are failing as well. Changing this will probably fix the tests without the need for 'SpecializationKernelName'. I think this change might be breaking valid code

Copy link
Contributor

Choose a reason for hiding this comment

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

The problem here is that the justification here is the same. Whether you do kernel_single_taskstd::string or kernel_single_task<some_templatestd::string>, we cannot actually put it in the int-header, because we cannot forward declare std::string.

nullptr_t and std::size_t/ptrdiff_t are less troublesome, since we know they are well defined typedefs (and can just spell them that way), but other types we cannot.

Copy link
Contributor

@premanandrao premanandrao Jul 16, 2020 •
edited
Loading

Choose a reason for hiding this comment

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

the standard doesn't allow us to forward declare anything in std.

@jtmott-intel, if a Kernel Name is not a std-type itself, but is templated on a std-type instead, like: foo<std::some_type>, what is the case where we need to forward-declare the std-type? We only need to forward-declare foo, right?

Copy link
Contributor

@erichkeane erichkeane Jul 16, 2020 •
edited
Loading

Choose a reason for hiding this comment

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

@premanadrao: How would we get away with that? We cannot spell foo<std::string> without spelling std::string.

Copy link
Contributor

Choose a reason for hiding this comment

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

In current code, I think std::some_type is also forward declared. There is no check on template parameters IIRC. However, if std header is included prior to integration header, I see no reason why we need to forward declare std::sometype

Copy link
Contributor

Choose a reason for hiding this comment

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

I still think we have to prohibit everything in namespace 'std', with some allowance for the types where desugaring/canonicalizing results in a non-std type (basically, anything that was typedef'ed to a builtin type).

Copy link
Contributor Author

Choose a reason for hiding this comment

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

@erichkeane I think that's what we have now in this PR. The if (NS->isStdNamespace()) { prohibits everything in namespace std. Some types such as std::size_t desugar to a builtin type such as unsigned long and are still allowed to work. Only nullptr_t is singled out, since std::nullptr_t also desugars to a builtin type but one that we explicitly check for and disallow. Is there a specific case you had in mind that should be changed?

Copy link
Contributor

Choose a reason for hiding this comment

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

No just being clear. I wouldn't mind allowing std::nullptr_t as well, since we can just canonicalize it as decltype(nullptr_t) though.

Copy link
Contributor

Choose a reason for hiding this comment

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

I don't know what to do here. Things need to be forward declarable for us to do the right thing in the int-header. std types aren't. So we don't really have a choice. But we have codes using std-types in the name. Therefore if we go ahead with this solution we'll break codes. And the workaround shown above (with SpecializationKernelName_XYZ) doesn't work for someone who needs to write a kernel which works with any type. They can't provide a custom name for each without asking their users to provide a name. I hope @mkinsner has an idea.

Copy link
Contributor

Choose a reason for hiding this comment

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

Just talked with @mkinsner . The intention is that that names and their template arguments are forward declarable. This excludes std types. We think it is OK for us to not support names with std names as template arguments. For those use cases users should use unnamed lambdas.

kernel_single_task<Templated_kernel_name<std::U>>([=]() {}); // expected-error {{kernel name cannot be a type in the "std" namespace}}
#endif

// Although in the std namespace, these resolve to builtins such as `int` that are allowed in kernel names
kernel_single_task<std::size_t>([=]() {});
Copy link
Contributor

Choose a reason for hiding this comment

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

I think that any exceptions we have to allow std types like size_t and ptrdiff_t, we need to just canonicalize them. So this should be kernel_single_task. Otherwise we're still violating the same rule as with std::string.

Copy link
Contributor Author

Choose a reason for hiding this comment

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

The complication with std::size_t and a couple others like it is by the time we get to Sema::ConstructOpenCLKernel, the template argument type already evaluates to the builtin type unsigned long, and not convertible to a Decl. At this stage of the compilation, it doesn't seem like I can distinguish between the user spelling "std::size_t" or "unsigned long". Maybe I need an AST visitor to check how the user spelled the kernel name type in source?

Copy link
Contributor

Choose a reason for hiding this comment

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

I think you misunderstood my comment (and apparently github ate the 'so this should be kernel_single_task<unsigned long>> line) ... I think we need to prohibit any of the classes in the std namespace, but size_t/ptrdiff_t/etc are both typedefs, so just canonicalizing them (aka, desugaring the typedef) should be sufficient.

Copy link
Contributor Author

Choose a reason for hiding this comment

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

OK. If I understand what you mean, then size_t/ptrdiff_t are already desurgared. A kernel_single_task<std::size_t> is written to the integration header as kernel_single_task<unsigned long>. That's why those names still work as kernel names.

kernel_single_task<std::ptrdiff_t>([=]() {});

return 0;
}
20 changes: 16 additions & 4 deletions clang/test/SemaSYCL/unnamed-kernel.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -11,6 +11,11 @@ template <typename T>
class KernelName;
}

namespace std {
typedef struct {
} max_align_t;
} // namespace std

struct MyWrapper {
private:
class InvalidKernelName0 {};
Expand Down Expand Up @@ -41,15 +46,15 @@ struct MyWrapper {

#ifndef __SYCL_UNNAMED_LAMBDA__
// expected-error@+4 {{kernel needs to have a globally-visible name}}
// expected-note@16 {{InvalidKernelName0 declared here}}
// expected-note@21 {{InvalidKernelName0 declared here}}
#endif
q.submit([&](cl::sycl::handler &h) {
h.single_task<InvalidKernelName0>([] {});
});

#ifndef __SYCL_UNNAMED_LAMBDA__
// expected-error@+4 {{kernel needs to have a globally-visible name}}
// expected-note@17 {{InvalidKernelName3 declared here}}
// expected-note@22 {{InvalidKernelName3 declared here}}
#endif
q.submit([&](cl::sycl::handler &h) {
h.single_task<namespace1::KernelName<InvalidKernelName3>>([] {});
Expand All @@ -60,10 +65,17 @@ struct MyWrapper {
h.single_task<ValidAlias>([] {});
});

#ifndef __SYCL_UNNAMED_LAMBDA__
// expected-error@+3 {{kernel name cannot be a type in the "std" namespace}}
#endif
q.submit([&](cl::sycl::handler &h) {
h.single_task<std::max_align_t>([] {});
});

using InvalidAlias = InvalidKernelName4;
#ifndef __SYCL_UNNAMED_LAMBDA__
// expected-error@+4 {{kernel needs to have a globally-visible name}}
// expected-note@18 {{InvalidKernelName4 declared here}}
// expected-note@23 {{InvalidKernelName4 declared here}}
#endif
q.submit([&](cl::sycl::handler &h) {
h.single_task<InvalidAlias>([] {});
Expand All @@ -72,7 +84,7 @@ struct MyWrapper {
using InvalidAlias1 = InvalidKernelName5;
#ifndef __SYCL_UNNAMED_LAMBDA__
// expected-error@+4 {{kernel needs to have a globally-visible name}}
// expected-note@19 {{InvalidKernelName5 declared here}}
// expected-note@24 {{InvalidKernelName5 declared here}}
#endif
q.submit([&](cl::sycl::handler &h) {
h.single_task<namespace1::KernelName<InvalidAlias1>>([] {});
Expand Down
15 changes: 6 additions & 9 deletions sycl/test/group-algorithm/broadcast.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -15,14 +15,11 @@
using namespace sycl;
using namespace sycl::ONEAPI;

template <typename InputContainer, typename OutputContainer>
class broadcast_kernel;

template <typename InputContainer, typename OutputContainer>
template <typename kernel_name, typename InputContainer,
typename OutputContainer>
void test(queue q, InputContainer input, OutputContainer output) {
typedef typename InputContainer::value_type InputT;
typedef typename OutputContainer::value_type OutputT;
typedef class broadcast_kernel<InputContainer, OutputContainer> kernel_name;
size_t N = input.size();
size_t G = 4;
range<2> R(G, G);
Expand Down Expand Up @@ -63,7 +60,7 @@ int main() {
std::array<int, 3> output;
std::iota(input.begin(), input.end(), 1);
std::fill(output.begin(), output.end(), false);
test(q, input, output);
test<class KernelName_EFL>(q, input, output);
}

// Test pointer type
Expand All @@ -74,7 +71,7 @@ int main() {
input[i] = static_cast<int *>(0x0) + i;
}
std::fill(output.begin(), output.end(), static_cast<int *>(0x0));
test(q, input, output);
test<class KernelName_NrqELzFQToOSPsRNMi>(q, input, output);
}

// Test user-defined type
Expand All @@ -88,7 +85,7 @@ int main() {
std::complex<float>(0, 1) + (float)i * std::complex<float>(2, 2);
}
std::fill(output.begin(), output.end(), std::complex<float>(0, 0));
test(q, input, output);
test<class KernelName_rCblcml>(q, input, output);
}
{
std::array<std::complex<double>, N> input;
Expand All @@ -98,7 +95,7 @@ int main() {
std::complex<double>(0, 1) + (double)i * std::complex<double>(2, 2);
}
std::fill(output.begin(), output.end(), std::complex<float>(0, 0));
test(q, input, output);
test<class KernelName_NCWhjnQ>(q, input, output);
}
std::cout << "Test passed." << std::endl;
}
39 changes: 22 additions & 17 deletions sycl/test/group-algorithm/exclusive_scan.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -24,7 +24,7 @@
using namespace sycl;
using namespace sycl::ONEAPI;

template <class BinaryOperation, int TestNumber>
template <class SpecializationKernelName, int TestNumber>
class exclusive_scan_kernel;

// std::exclusive_scan isn't implemented yet, so use serial implementation
Expand All @@ -44,17 +44,17 @@ OutputIterator exclusive_scan(InputIterator first, InputIterator last,
}
} // namespace emu

template <typename InputContainer, typename OutputContainer,
class BinaryOperation>
template <typename SpecializationKernelName, typename InputContainer,
typename OutputContainer, class BinaryOperation>
void test(queue q, InputContainer input, OutputContainer output,
BinaryOperation binary_op,
typename OutputContainer::value_type identity) {
typedef typename InputContainer::value_type InputT;
typedef typename OutputContainer::value_type OutputT;
typedef class exclusive_scan_kernel<BinaryOperation, 0> kernel_name0;
typedef class exclusive_scan_kernel<BinaryOperation, 1> kernel_name1;
typedef class exclusive_scan_kernel<BinaryOperation, 2> kernel_name2;
typedef class exclusive_scan_kernel<BinaryOperation, 3> kernel_name3;
typedef class exclusive_scan_kernel<SpecializationKernelName, 0> kernel_name0;
typedef class exclusive_scan_kernel<SpecializationKernelName, 1> kernel_name1;
typedef class exclusive_scan_kernel<SpecializationKernelName, 2> kernel_name2;
typedef class exclusive_scan_kernel<SpecializationKernelName, 3> kernel_name3;
OutputT init = 42;
size_t N = input.size();
size_t G = 16;
Expand Down Expand Up @@ -159,19 +159,24 @@ int main() {
std::iota(input.begin(), input.end(), 0);
std::fill(output.begin(), output.end(), 0);

test(q, input, output, plus<>(), 0);
test(q, input, output, minimum<>(), std::numeric_limits<int>::max());
test(q, input, output, maximum<>(), std::numeric_limits<int>::lowest());
test<class KernelNamePlusV>(q, input, output, plus<>(), 0);
test<class KernelNameMinimumV>(q, input, output, minimum<>(),
std::numeric_limits<int>::max());
test<class KernelNameMaximumV>(q, input, output, maximum<>(),
std::numeric_limits<int>::lowest());

test(q, input, output, plus<int>(), 0);
test(q, input, output, minimum<int>(), std::numeric_limits<int>::max());
test(q, input, output, maximum<int>(), std::numeric_limits<int>::lowest());
test<class KernelNamePlusI>(q, input, output, plus<int>(), 0);
test<class KernelNameMinimumI>(q, input, output, minimum<int>(),
std::numeric_limits<int>::max());
test<class KernelNameMaximumI>(q, input, output, maximum<int>(),
std::numeric_limits<int>::lowest());

#ifdef SPIRV_1_3
test(q, input, output, multiplies<int>(), 1);
test(q, input, output, bit_or<int>(), 0);
test(q, input, output, bit_xor<int>(), 0);
test(q, input, output, bit_and<int>(), ~0);
test<class KernelName_VzAPutpBRRJrQPB>(q, input, output, multiplies<int>(),
1);
test<class KernelName_UXdGbr>(q, input, output, bit_or<int>(), 0);
test<class KernelName_saYaodNyJknrPW>(q, input, output, bit_xor<int>(), 0);
test<class KernelName_GPcuAlvAOjrDyP>(q, input, output, bit_and<int>(), ~0);
#endif // SPIRV_1_3

std::cout << "Test passed." << std::endl;
Expand Down
40 changes: 23 additions & 17 deletions sycl/test/group-algorithm/inclusive_scan.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -24,7 +24,7 @@
using namespace sycl;
using namespace sycl::ONEAPI;

template <class BinaryOperation, int TestNumber>
template <class SpecializationKernelName, int TestNumber>
class inclusive_scan_kernel;

// std::inclusive_scan isn't implemented yet, so use serial implementation
Expand All @@ -44,17 +44,17 @@ OutputIterator inclusive_scan(InputIterator first, InputIterator last,
}
} // namespace emu

template <typename InputContainer, typename OutputContainer,
class BinaryOperation>
template <typename SpecializationKernelName, typename InputContainer,
typename OutputContainer, class BinaryOperation>
void test(queue q, InputContainer input, OutputContainer output,
BinaryOperation binary_op,
typename OutputContainer::value_type identity) {
typedef typename InputContainer::value_type InputT;
typedef typename OutputContainer::value_type OutputT;
typedef class inclusive_scan_kernel<BinaryOperation, 0> kernel_name0;
typedef class inclusive_scan_kernel<BinaryOperation, 1> kernel_name1;
typedef class inclusive_scan_kernel<BinaryOperation, 2> kernel_name2;
typedef class inclusive_scan_kernel<BinaryOperation, 3> kernel_name3;
typedef class inclusive_scan_kernel<SpecializationKernelName, 0> kernel_name0;
typedef class inclusive_scan_kernel<SpecializationKernelName, 1> kernel_name1;
typedef class inclusive_scan_kernel<SpecializationKernelName, 2> kernel_name2;
typedef class inclusive_scan_kernel<SpecializationKernelName, 3> kernel_name3;
OutputT init = 42;
size_t N = input.size();
size_t G = 16;
Expand Down Expand Up @@ -159,19 +159,25 @@ int main() {
std::iota(input.begin(), input.end(), 0);
std::fill(output.begin(), output.end(), 0);

test(q, input, output, plus<>(), 0);
test(q, input, output, minimum<>(), std::numeric_limits<int>::max());
test(q, input, output, maximum<>(), std::numeric_limits<int>::lowest());
test<class KernelNamePlusV>(q, input, output, plus<>(), 0);
test<class KernelNameMinimumV>(q, input, output, minimum<>(),
std::numeric_limits<int>::max());
test<class KernelNameMaximumV>(q, input, output, maximum<>(),
std::numeric_limits<int>::lowest());

test(q, input, output, plus<int>(), 0);
test(q, input, output, minimum<int>(), std::numeric_limits<int>::max());
test(q, input, output, maximum<int>(), std::numeric_limits<int>::lowest());
test<class KernelNamePlusI>(q, input, output, plus<int>(), 0);
test<class KernelNameMinimumI>(q, input, output, minimum<int>(),
std::numeric_limits<int>::max());
test<class KernelNameMaximumI>(q, input, output, maximum<int>(),
std::numeric_limits<int>::lowest());

#ifdef SPIRV_1_3
test(q, input, output, multiplies<int>(), 1);
test(q, input, output, bit_or<int>(), 0);
test(q, input, output, bit_xor<int>(), 0);
test(q, input, output, bit_and<int>(), ~0);
test<class KernelName_zMyjxUrBgeUGoxmDwhvJ>(q, input, output,
multiplies<int>(), 1);
test<class KernelName_SljjtroxNRaAXoVnT>(q, input, output, bit_or<int>(), 0);
test<class KernelName_yXIZfjwjxQGiPeQAnc>(q, input, output, bit_xor<int>(),
0);
test<class KernelName_xGnAnMYHvqekCk>(q, input, output, bit_and<int>(), ~0);
#endif // SPIRV_1_3

std::cout << "Test passed." << std::endl;
Expand Down
Loading