diff --git a/clang/test/SemaSYCL/reqd-sub-group-size-ast.cpp b/clang/test/SemaSYCL/reqd-sub-group-size-ast.cpp new file mode 100644 index 0000000000000..67ae007429213 --- /dev/null +++ b/clang/test/SemaSYCL/reqd-sub-group-size-ast.cpp @@ -0,0 +1,110 @@ +// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -sycl-std=2017 -Wno-sycl-2017-compat -ast-dump %s | FileCheck %s + +// The test checks AST of [[intel::reqd_sub_group_size()]] attribute. + +#include "sycl.hpp" + +using namespace cl::sycl; +queue q; + +[[intel::reqd_sub_group_size(4)]] void foo() {} + +class Functor16 { +public: + [[intel::reqd_sub_group_size(16)]] void operator()() const {} +}; + +class Functor { +public: + void operator()() const { + foo(); + } +}; + +// Test that checks template parameter support on member function of class template. +template +class KernelFunctor { +public: + [[intel::reqd_sub_group_size(SIZE)]] void operator()() const {} +}; + +// Test that checks template parameter support on function. +// CHECK: FunctionTemplateDecl {{.*}} func +// CHECK: FunctionDecl {{.*}} func 'void ()' +// CHECK-NEXT: CompoundStmt +// CHECK_NEXT: IntelReqdSubGroupSizeAttr {{.*}} reqd_sub_group_size +// CHECK_NEXT: DeclRefExpr {{.*}} 'int' NonTypeTemplateParm {{.*}} 'N' 'int' +// CHECK: FunctionDecl {{.*}} func 'void ()' +// CHECK-NEXT: TemplateArgument integral 12 +// CHECK-NEXT: CompoundStmt +// CHECK-NEXT: IntelReqdSubGroupSizeAttr {{.*}} reqd_sub_group_size +// CHECK-NEXT: ConstantExpr{{.*}}'int' +// CHECK-NEXT: value: Int 12 +// CHECK-NEXT: SubstNonTypeTemplateParmExpr +// CHECK-NEXT: NonTypeTemplateParmDecl +// CHECK-NEXT: IntegerLiteral {{.*}} 'int' 12 +template +[[intel::reqd_sub_group_size(N)]] void func() {} + +int main() { + q.submit([&](handler &h) { + // CHECK: FunctionDecl {{.*}} {{.*}}kernel_name1 + // CHECK: IntelReqdSubGroupSizeAttr {{.*}} reqd_sub_group_size + // CHECK-NEXT: ConstantExpr {{.*}} 'int' + // CHECK-NEXT: value: Int 16 + // CHECK-NEXT: IntegerLiteral{{.*}}16{{$}} + Functor16 f16; + h.single_task(f16); + + // CHECK: FunctionDecl {{.*}} {{.*}}kernel_name2 + // CHECK: IntelReqdSubGroupSizeAttr {{.*}} reqd_sub_group_size + // CHECK-NEXT: ConstantExpr {{.*}} 'int' + // CHECK-NEXT: value: Int 4 + // CHECK-NEXT: IntegerLiteral{{.*}}4{{$}} + Functor f; + h.single_task(f); + + // CHECK: FunctionDecl {{.*}} {{.*}}kernel_name3 + // CHECK: IntelReqdSubGroupSizeAttr {{.*}} reqd_sub_group_size + // CHECK-NEXT: ConstantExpr {{.*}} 'int' + // CHECK-NEXT: value: Int 2 + // CHECK-NEXT: IntegerLiteral{{.*}}2{{$}} + h.single_task([]() [[intel::reqd_sub_group_size(2)]] {}); + + // CHECK: FunctionDecl {{.*}} {{.*}}kernel_name4 + // CHECK: IntelReqdSubGroupSizeAttr {{.*}} reqd_sub_group_size + // CHECK-NEXT: ConstantExpr {{.*}} 'int' + // CHECK-NEXT: value: Int 4 + // CHECK-NEXT: IntegerLiteral{{.*}}4{{$}} + h.single_task([]() [[intel::reqd_sub_group_size(4)]] { foo(); }); + // CHECK: FunctionDecl {{.*}} {{.*}}kernel_name5 + // CHECK: IntelReqdSubGroupSizeAttr {{.*}} reqd_sub_group_size + // CHECK-NEXT: ConstantExpr {{.*}} 'int' + // CHECK-NEXT: value: Int 6 + // CHECK-NEXT: IntegerLiteral{{.*}}6{{$}} + h.single_task([]() [[intel::reqd_sub_group_size(6)]] {}); + + // CHECK: FunctionDecl {{.*}}kernel_name_6 + // CHECK: IntelReqdSubGroupSizeAttr {{.*}} reqd_sub_group_size + // CHECK-NEXT: ConstantExpr{{.*}}'int' + // CHECK-NEXT: value: Int 10 + // CHECK-NEXT: SubstNonTypeTemplateParmExpr + // CHECK-NEXT: NonTypeTemplateParmDecl + // CHECK-NEXT: IntegerLiteral {{.*}} 'int' 10 + KernelFunctor<10> f2; + h.single_task(f2); + + // Ignore duplicate attribute. + h.single_task( + // CHECK: FunctionDecl {{.*}}kernel_name_7 + // CHECK: IntelReqdSubGroupSizeAttr {{.*}} reqd_sub_group_size + // CHECK-NEXT: ConstantExpr {{.*}} 'int' + // CHECK-NEXT: value: Int 8 + // CHECK-NEXT: IntegerLiteral{{.*}}8{{$}} + // CHECK-NOT: IntelReqdSubGroupSizeAttr + []() [[intel::reqd_sub_group_size(8), + intel::reqd_sub_group_size(8)]] {}); + }); + func<12>(); + return 0; +} diff --git a/clang/test/SemaSYCL/reqd-sub-group-size-device.cpp b/clang/test/SemaSYCL/reqd-sub-group-size-device.cpp deleted file mode 100644 index 4569ca5284e83..0000000000000 --- a/clang/test/SemaSYCL/reqd-sub-group-size-device.cpp +++ /dev/null @@ -1,125 +0,0 @@ -// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -fsyntax-only -sycl-std=2017 -Wno-sycl-2017-compat -verify -DTRIGGER_ERROR %s -// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -sycl-std=2017 -Wno-sycl-2017-compat -ast-dump %s | FileCheck %s - -#include "sycl.hpp" - -using namespace cl::sycl; -queue q; - -[[intel::reqd_sub_group_size(4)]] void foo() {} // expected-note {{conflicting attribute is here}} -// expected-note@-1 {{conflicting attribute is here}} -[[intel::reqd_sub_group_size(32)]] void baz() {} // expected-note {{conflicting attribute is here}} - -// No diagnostic is emitted because the arguments match. -[[intel::reqd_sub_group_size(12)]] void bar(); -[[intel::reqd_sub_group_size(12)]] void bar() {} // OK - -// No diagnostic because the attributes are synonyms with identical behavior. -[[sycl::reqd_sub_group_size(12)]] void bar(); // OK - -// Diagnostic is emitted because the arguments mismatch. -[[intel::reqd_sub_group_size(12)]] void quux(); // expected-note {{previous attribute is here}} -[[intel::reqd_sub_group_size(100)]] void quux(); // expected-warning {{attribute 'reqd_sub_group_size' is already applied with different arguments}} expected-note {{previous attribute is here}} -[[sycl::reqd_sub_group_size(200)]] void quux(); // expected-warning {{attribute 'reqd_sub_group_size' is already applied with different arguments}} - -#ifdef TRIGGER_ERROR -// Make sure there's at least one argument passed. -[[sycl::reqd_sub_group_size]] void quibble(); // expected-error {{'reqd_sub_group_size' attribute takes one argument}} -#endif // TRIGGER_ERROR - -class Functor16 { -public: - [[intel::reqd_sub_group_size(16)]] void operator()() const {} -}; - -class Functor8 { // expected-error {{conflicting attributes applied to a SYCL kernel}} -public: - [[intel::reqd_sub_group_size(8)]] void operator()() const { // expected-note {{conflicting attribute is here}} - foo(); - } -}; - -class Functor4 { -public: - [[intel::reqd_sub_group_size(12)]] void operator()() const {} -}; - -class Functor { -public: - void operator()() const { - foo(); - } -}; - -int main() { - q.submit([&](handler &h) { - Functor16 f16; - h.single_task(f16); - - Functor f; - h.single_task(f); - -#ifdef TRIGGER_ERROR - Functor8 f8; - h.single_task(f8); - - h.single_task([]() { // expected-error {{conflicting attributes applied to a SYCL kernel}} - foo(); - baz(); - }); -#endif - - h.single_task([]() [[intel::reqd_sub_group_size(2)]]{}); - h.single_task([]() [[intel::reqd_sub_group_size(4)]] { foo(); }); - h.single_task([]() [[intel::reqd_sub_group_size(6)]]{}); - - Functor4 f4; - h.single_task(f4); - }); - return 0; -} - -[[intel::reqd_sub_group_size(16)]] SYCL_EXTERNAL void B(); -[[intel::reqd_sub_group_size(16)]] void A() { -} - -[[intel::reqd_sub_group_size(16)]] SYCL_EXTERNAL void B() { - A(); -} - -#ifdef TRIGGER_ERROR -// expected-note@+1 {{conflicting attribute is here}} -[[intel::reqd_sub_group_size(2)]] void sg_size2() {} - -// expected-note@+2 {{conflicting attribute is here}} -// expected-error@+1 {{conflicting attributes applied to a SYCL kernel}} -[[intel::reqd_sub_group_size(4)]] __attribute__((sycl_device)) void sg_size4() { - sg_size2(); -} -#endif - -// CHECK: FunctionDecl {{.*}} {{.*}}kernel_name1 -// CHECK: IntelReqdSubGroupSizeAttr {{.*}} -// CHECK-NEXT: ConstantExpr {{.*}} 'int' -// CHECK-NEXT: value: Int 16 -// CHECK-NEXT: IntegerLiteral{{.*}}16{{$}} -// CHECK: FunctionDecl {{.*}} {{.*}}kernel_name2 -// CHECK: IntelReqdSubGroupSizeAttr {{.*}} -// CHECK-NEXT: ConstantExpr {{.*}} 'int' -// CHECK-NEXT: value: Int 4 -// CHECK-NEXT: IntegerLiteral{{.*}}4{{$}} -// CHECK: FunctionDecl {{.*}} {{.*}}kernel_name5 -// CHECK: IntelReqdSubGroupSizeAttr {{.*}} -// CHECK-NEXT: ConstantExpr {{.*}} 'int' -// CHECK-NEXT: value: Int 2 -// CHECK-NEXT: IntegerLiteral{{.*}}2{{$}} -// CHECK: FunctionDecl {{.*}} {{.*}}kernel_name7 -// CHECK: IntelReqdSubGroupSizeAttr {{.*}} -// CHECK-NEXT: ConstantExpr {{.*}} 'int' -// CHECK-NEXT: value: Int 6 -// CHECK-NEXT: IntegerLiteral{{.*}}6{{$}} -// CHECK: FunctionDecl {{.*}} {{.*}}kernel_name8 -// CHECK: IntelReqdSubGroupSizeAttr {{.*}} -// CHECK-NEXT: ConstantExpr {{.*}} 'int' -// CHECK-NEXT: value: Int 12 -// CHECK-NEXT: IntegerLiteral{{.*}}12{{$}} diff --git a/clang/test/SemaSYCL/reqd-sub-group-size.cpp b/clang/test/SemaSYCL/reqd-sub-group-size.cpp new file mode 100644 index 0000000000000..6b6a51fcbc360 --- /dev/null +++ b/clang/test/SemaSYCL/reqd-sub-group-size.cpp @@ -0,0 +1,134 @@ +// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -fsyntax-only -sycl-std=2017 -Wno-sycl-2017-compat -verify -pedantic %s + +// The test checks functionality of [[intel::reqd_sub_group_size()]] attribute on SYCL kernel. + +#include "sycl.hpp" + +using namespace cl::sycl; +queue q; + +[[intel::reqd_sub_group_size(4)]] void foo() {} // expected-note {{conflicting attribute is here}} +// expected-note@-1 {{conflicting attribute is here}} +[[intel::reqd_sub_group_size(32)]] void baz() {} // expected-note {{conflicting attribute is here}} + +class Functor8 { // expected-error {{conflicting attributes applied to a SYCL kernel}} +public: + [[intel::reqd_sub_group_size(8)]] void operator()() const { // expected-note {{conflicting attribute is here}} + foo(); + } +}; + +int main() { + q.submit([&](handler &h) { + Functor8 f8; + h.single_task(f8); + + h.single_task([]() { // expected-error {{conflicting attributes applied to a SYCL kernel}} + foo(); + baz(); + }); + }); + return 0; +} +[[intel::reqd_sub_group_size(16)]] SYCL_EXTERNAL void B(); +[[intel::reqd_sub_group_size(16)]] void A() { +} + +[[intel::reqd_sub_group_size(16)]] SYCL_EXTERNAL void B() { + A(); +} +// expected-note@+1 {{conflicting attribute is here}} +[[intel::reqd_sub_group_size(2)]] void sg_size2() {} + +// expected-note@+2 {{conflicting attribute is here}} +// expected-error@+1 {{conflicting attributes applied to a SYCL kernel}} +[[intel::reqd_sub_group_size(4)]] __attribute__((sycl_device)) void sg_size4() { + sg_size2(); +} + +// Test that checks support and functionality of reqd_sub_group_size attribute support on function. + +// Tests for incorrect argument values for Intel reqd_sub_group_size attribute. +[[intel::reqd_sub_group_size]] void one() {} // expected-error {{'reqd_sub_group_size' attribute takes one argument}} +[[intel::reqd_sub_group_size(5)]] int a; // expected-error{{'reqd_sub_group_size' attribute only applies to functions}} +[[intel::reqd_sub_group_size("foo")]] void func() {} // expected-error{{integral constant expression must have integral or unscoped enumeration type, not 'const char[4]'}} +[[intel::reqd_sub_group_size(-1)]] void func1() {} // expected-error{{'reqd_sub_group_size' attribute requires a positive integral compile time constant expression}} +[[intel::reqd_sub_group_size(0, 1)]] void arg() {} // expected-error{{'reqd_sub_group_size' attribute takes one argument}} + +// Diagnostic is emitted because the arguments mismatch. +[[intel::reqd_sub_group_size(12)]] void quux(); // expected-note {{previous attribute is here}} +[[intel::reqd_sub_group_size(100)]] void quux(); // expected-warning {{attribute 'reqd_sub_group_size' is already applied with different arguments}} expected-note {{previous attribute is here}} +[[sycl::reqd_sub_group_size(200)]] void quux(); // expected-warning {{attribute 'reqd_sub_group_size' is already applied with different arguments}} + +// Make sure there's at least one argument passed. +[[sycl::reqd_sub_group_size]] void quibble(); // expected-error {{'reqd_sub_group_size' attribute takes one argument}} + +// No diagnostic is emitted because the arguments match. +[[intel::reqd_sub_group_size(12)]] void same(); +[[intel::reqd_sub_group_size(12)]] void same() {} // OK + +// No diagnostic because the attributes are synonyms with identical behavior. +[[sycl::reqd_sub_group_size(12)]] void same(); // OK + +// Test that checks wrong function template instantiation and ensures that the type +// is checked properly when instantiating from the template definition. +template +// expected-error@+3{{'reqd_sub_group_size' attribute requires a positive integral compile time constant expression}} +// expected-error@+2 {{integral constant expression must have integral or unscoped enumeration type, not 'S'}} +// expected-error@+1 {{integral constant expression must have integral or unscoped enumeration type, not 'float'}} +[[intel::reqd_sub_group_size(Ty{})]] void func() {} + +struct S {}; +void test() { + // expected-note@+1{{in instantiation of function template specialization 'func' requested here}} + func(); + // expected-note@+1{{in instantiation of function template specialization 'func' requested here}} + func(); + // expected-note@+1{{in instantiation of function template specialization 'func' requested here}} + func(); +} + +// Test that checks expression is not a constant expression. +// expected-note@+1{{declared here}} +int foo1(); +// expected-error@+2{{expression is not an integral constant expression}} +// expected-note@+1{{non-constexpr function 'foo1' cannot be used in a constant expression}} +[[intel::reqd_sub_group_size(foo1() + 12)]] void func1(); + +// Test that checks expression is a constant expression. +constexpr int bar1() { return 0; } +[[intel::reqd_sub_group_size(bar1() + 12)]] void func2(); // OK + +// Test that checks template parameter support on member function of class template. +template +class KernelFunctor { +public: + // expected-error@+1{{'reqd_sub_group_size' attribute requires a positive integral compile time constant expression}} + [[intel::reqd_sub_group_size(SIZE)]] void operator()() {} +}; + +int check() { + // expected-note@+1{{in instantiation of template class 'KernelFunctor<-1>' requested here}} + KernelFunctor<-1>(); + return 0; +} + +// Test that checks template parameter support on function. +template +// expected-error@+1{{'reqd_sub_group_size' attribute requires a positive integral compile time constant expression}} +[[intel::reqd_sub_group_size(N)]] void func3() {} + +template +[[intel::reqd_sub_group_size(4)]] void func4(); // expected-note {{previous attribute is here}} + +template +[[intel::reqd_sub_group_size(N)]] void func4() {} // expected-warning {{attribute 'reqd_sub_group_size' is already applied with different arguments}} + +int check1() { + // no error expected + func3<12>(); + // expected-note@+1{{in instantiation of function template specialization 'func3<-1>' requested here}} + func3<-1>(); + func4<6>(); // expected-note {{in instantiation of function template specialization 'func4<6>' requested here}} + return 0; +} diff --git a/clang/test/SemaSYCL/sycl-device-reqd-sub-group-size-template.cpp b/clang/test/SemaSYCL/sycl-device-reqd-sub-group-size-template.cpp deleted file mode 100644 index 90908b01ce2ad..0000000000000 --- a/clang/test/SemaSYCL/sycl-device-reqd-sub-group-size-template.cpp +++ /dev/null @@ -1,94 +0,0 @@ -// RUN: %clang_cc1 -fsycl-is-device -fsyntax-only -ast-dump -verify -pedantic %s | FileCheck %s - -// Test that checkes template parameter support for 'reqd_sub_group_size' attribute on sycl device. - -// Test that checks wrong function template instantiation and ensures that the type -// is checked properly when instantiating from the template definition. -template -// expected-error@+3{{'reqd_sub_group_size' attribute requires a positive integral compile time constant expression}} -// expected-error@+2 {{integral constant expression must have integral or unscoped enumeration type, not 'S'}} -// expected-error@+1 {{integral constant expression must have integral or unscoped enumeration type, not 'float'}} -[[intel::reqd_sub_group_size(Ty{})]] void func() {} - -struct S {}; -void test() { - //expected-note@+1{{in instantiation of function template specialization 'func' requested here}} - func(); - //expected-note@+1{{in instantiation of function template specialization 'func' requested here}} - func(); - //expected-note@+1{{in instantiation of function template specialization 'func' requested here}} - func(); -} - -// Test that checks expression is not a constant expression. -// expected-note@+1{{declared here}} -int foo(); -// expected-error@+2{{expression is not an integral constant expression}} -// expected-note@+1{{non-constexpr function 'foo' cannot be used in a constant expression}} -[[intel::reqd_sub_group_size(foo() + 12)]] void func1(); - -// Test that checks expression is a constant expression. -constexpr int bar() { return 0; } -[[intel::reqd_sub_group_size(bar() + 12)]] void func2(); // OK - -// Test that checks template parameter support on member function of class template. -template -class KernelFunctor { -public: - // expected-error@+1{{'reqd_sub_group_size' attribute requires a positive integral compile time constant expression}} - [[intel::reqd_sub_group_size(SIZE)]] void operator()() {} -}; - -int main() { - //expected-note@+1{{in instantiation of template class 'KernelFunctor<-1>' requested here}} - KernelFunctor<-1>(); - // no error expected - KernelFunctor<10>(); - return 0; -} - -// CHECK: ClassTemplateDecl {{.*}} {{.*}} KernelFunctor -// CHECK: ClassTemplateSpecializationDecl {{.*}} {{.*}} class KernelFunctor definition -// CHECK: CXXRecordDecl {{.*}} {{.*}} implicit class KernelFunctor -// CHECK: IntelReqdSubGroupSizeAttr {{.*}} -// CHECK: SubstNonTypeTemplateParmExpr {{.*}} -// CHECK-NEXT: NonTypeTemplateParmDecl {{.*}} -// CHECK-NEXT: IntegerLiteral{{.*}}10{{$}} - -// Test that checks template parameter support on function. -template -// expected-error@+1{{'reqd_sub_group_size' attribute requires a positive integral compile time constant expression}} -[[intel::reqd_sub_group_size(N)]] void func3() {} - -template -[[intel::reqd_sub_group_size(4)]] void func4(); // expected-note {{previous attribute is here}} - -template -[[intel::reqd_sub_group_size(N)]] void func4() {} // expected-warning {{attribute 'reqd_sub_group_size' is already applied with different arguments}} - -int check() { - // no error expected - func3<12>(); - //expected-note@+1{{in instantiation of function template specialization 'func3<-1>' requested here}} - func3<-1>(); - func4<6>(); //expected-note {{in instantiation of function template specialization 'func4<6>' requested here}} - return 0; -} - -// No diagnostic is emitted because the arguments match. Duplicate attribute is silently ignored. -[[intel::reqd_sub_group_size(8)]] -[[intel::reqd_sub_group_size(8)]] void func5() {} - -// CHECK: FunctionTemplateDecl {{.*}} {{.*}} func3 -// CHECK: NonTypeTemplateParmDecl {{.*}} {{.*}} referenced 'int' depth 0 index 0 N -// CHECK: FunctionDecl {{.*}} {{.*}} func3 'void ()' -// CHECK: IntelReqdSubGroupSizeAttr {{.*}} -// CHECK: SubstNonTypeTemplateParmExpr {{.*}} -// CHECK-NEXT: NonTypeTemplateParmDecl {{.*}} -// CHECK-NEXT: IntegerLiteral{{.*}}12{{$}} - -// CHECK: FunctionDecl {{.*}} {{.*}} func5 'void ()' -// CHECK: IntelReqdSubGroupSizeAttr {{.*}} -// CHECK-NEXT: ConstantExpr {{.*}} 'int' -// CHECK-NEXT: value: Int 8 -// CHECK-NEXT: IntegerLiteral{{.*}}8{{$}}