diff --git a/clang/include/clang/Basic/AttrDocs.td b/clang/include/clang/Basic/AttrDocs.td index 73cd65eb7aa4f..d1f740e3bfef9 100644 --- a/clang/include/clang/Basic/AttrDocs.td +++ b/clang/include/clang/Basic/AttrDocs.td @@ -2418,7 +2418,8 @@ attribute was applied. This effect is equivalent to annotating restrict on **all** kernel pointer arguments in an OpenCL or SPIR-V kernel. If ``intel::kernel_args_restrict`` is applied to a function called from a device -kernel, the attribute is not ignored and it is propagated to the kernel. +kernel, the attribute is ignored and it is not propagated to the kernel to +match with new SYCL 2020 spec. The attribute forms an unchecked assertion, in that implementations do not need to check/confirm the pre-condition in any way. If a user applies @@ -2711,6 +2712,10 @@ device operation, guiding the FPGA backend to insert the appropriate number of registers to break-up the combinational logic circuit, and thereby controlling the length of the longest combinational path. +If ``intel::scheduler_target_fmax_mhz`` is applied to a function called from a +device kernel, the attribute is ignored and it is not propagated to the kernel +to match with new SYCL 2020 spec. + .. code-block:: c++ [[intel::scheduler_target_fmax_mhz(4)]] void foo() {} @@ -2741,6 +2746,10 @@ function object). If 1, compiler doesn't use the global work offset values for the device function. Valid values are 0 and 1. If used without argument, value of 1 is set implicitly. +If ``intel::no_global_work_offset`` is applied to a function called from a +device kernel, the attribute is ignored and it is not propagated to the kernel +to match with new SYCL 2020 spec. + .. code-block:: c++ [[intel::no_global_work_offset]] diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index c64da6ba38f7f..ff8235a63d643 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -346,11 +346,9 @@ static void collectSYCLAttributes(Sema &S, FunctionDecl *FD, llvm::copy_if(FD->getAttrs(), std::back_inserter(Attrs), [](Attr *A) { // FIXME: Make this list self-adapt as new SYCL attributes are added. return isa(A); + SYCLSimdAttr>(A); }); // Allow the kernel attribute "use_stall_enable_clusters" only on lambda @@ -372,6 +370,9 @@ static void collectSYCLAttributes(Sema &S, FunctionDecl *FD, llvm::copy_if(FD->getAttrs(), std::back_inserter(Attrs), [](Attr *A) { return isa(A); }); } diff --git a/clang/test/CodeGenSYCL/intel-fpga-no-global-work-offset.cpp b/clang/test/CodeGenSYCL/intel-fpga-no-global-work-offset.cpp index f8edf692f66e2..4b0f86efc2b32 100644 --- a/clang/test/CodeGenSYCL/intel-fpga-no-global-work-offset.cpp +++ b/clang/test/CodeGenSYCL/intel-fpga-no-global-work-offset.cpp @@ -1,4 +1,4 @@ -// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -triple spir64-unknown-unknown-sycldevice -disable-llvm-passes -emit-llvm -o - %s | FileCheck %s +// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -triple spir64-unknown-unknown-sycldevice -disable-llvm-passes -sycl-std=2020 -emit-llvm -o - %s | FileCheck %s #include "sycl.hpp" @@ -16,8 +16,7 @@ class Functor { [[intel::no_global_work_offset(SIZE)]] void operator()() const {} }; -template -[[intel::no_global_work_offset(N)]] void func() {} +[[intel::no_global_work_offset(1)]] void func() {} int main() { q.submit([&](handler &h) { @@ -33,9 +32,9 @@ int main() { Functor<1> f; h.single_task(f); - h.single_task([]() { - func<1>(); - }); + // Test attribute is not propagated. + h.single_task( + []() { func(); }); }); return 0; } @@ -44,6 +43,6 @@ int main() { // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name2"() #0 {{.*}} !no_global_work_offset ![[NUM5]] // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name3"() #0 {{.*}} ![[NUM4:[0-9]+]] // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name4"() #0 {{.*}} !no_global_work_offset ![[NUM5]] -// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name5"() #0 {{.*}} !no_global_work_offset ![[NUM5]] +// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name5"() #0 {{.*}} ![[NUM5]] // CHECK-NOT: ![[NUM4]] = !{i32 0} // CHECK: ![[NUM5]] = !{} diff --git a/clang/test/CodeGenSYCL/scheduler-target-fmax-mhz.cpp b/clang/test/CodeGenSYCL/scheduler-target-fmax-mhz.cpp index 21c981220528e..839a9a69354ed 100644 --- a/clang/test/CodeGenSYCL/scheduler-target-fmax-mhz.cpp +++ b/clang/test/CodeGenSYCL/scheduler-target-fmax-mhz.cpp @@ -1,25 +1,49 @@ -// RUN: %clang_cc1 -fsycl-is-device -disable-llvm-passes -triple spir64-unknown-unknown-sycldevice -emit-llvm -o - %s | FileCheck %s +// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -triple spir64-unknown-unknown-sycldevice -disable-llvm-passes -sycl-std=2020 -emit-llvm -o - %s | FileCheck %s -#include "Inputs/sycl.hpp" -[[intel::scheduler_target_fmax_mhz(5)]] void -func() {} +#include "sycl.hpp" + +using namespace cl::sycl; +queue q; + +class Foo { +public: + [[intel::scheduler_target_fmax_mhz(5)]] void operator()() const {} +}; template -[[intel::scheduler_target_fmax_mhz(N)]] void zoo() {} +class Functor { +public: + [[intel::scheduler_target_fmax_mhz(N)]] void operator()() const {} +}; + +[[intel::scheduler_target_fmax_mhz(2)]] void zoo() {} int main() { - cl::sycl::kernel_single_task( - []() [[intel::scheduler_target_fmax_mhz(2)]]{}); + q.submit([&](handler &h) { + // Test attribute argument size. + Foo boo; + h.single_task(boo); - cl::sycl::kernel_single_task( - []() { func(); }); + // Test attribute is applied on lambda. + h.single_task( + []() [[intel::scheduler_target_fmax_mhz(42)]]{}); - cl::sycl::kernel_single_task( - []() { zoo<75>(); }); + // Test template argument. + Functor<75> f; + h.single_task(f); + + // Test attribute is not propagated. + h.single_task( + []() { zoo(); }); + }); + return 0; } -// CHECK: define {{.*}}spir_kernel void @{{.*}}test_kernel1() {{.*}} !scheduler_target_fmax_mhz ![[PARAM1:[0-9]+]] -// CHECK: define {{.*}}spir_kernel void @{{.*}}test_kernel2() {{.*}} !scheduler_target_fmax_mhz ![[PARAM2:[0-9]+]] -// CHECK: define {{.*}}spir_kernel void @{{.*}}test_kernel3() {{.*}} !scheduler_target_fmax_mhz ![[PARAM3:[0-9]+]] -// CHECK: ![[PARAM1]] = !{i32 2} -// CHECK: ![[PARAM2]] = !{i32 5} -// CHECK: ![[PARAM3]] = !{i32 75} + +// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name1"() #0 {{.*}} !scheduler_target_fmax_mhz ![[NUM5:[0-9]+]] +// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name2"() #0 {{.*}} !scheduler_target_fmax_mhz ![[NUM42:[0-9]+]] +// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name3"() #0 {{.*}} !scheduler_target_fmax_mhz ![[NUM75:[0-9]+]] +// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name4"() #0 {{.*}} ![[NUM0:[0-9]+]] +// CHECK: ![[NUM0]] = !{} +// CHECK: ![[NUM5]] = !{i32 5} +// CHECK: ![[NUM42]] = !{i32 42} +// CHECK: ![[NUM75]] = !{i32 75} diff --git a/clang/test/SemaSYCL/check-notdirect-attribute-propagation.cpp b/clang/test/SemaSYCL/check-notdirect-attribute-propagation.cpp index 2a0fe658114f9..4abd5234c6301 100644 --- a/clang/test/SemaSYCL/check-notdirect-attribute-propagation.cpp +++ b/clang/test/SemaSYCL/check-notdirect-attribute-propagation.cpp @@ -46,7 +46,7 @@ void invoke_foo2() { // CHECK-LABEL: FunctionDecl {{.*}} invoke_foo2 'void ()' // CHECK: `-FunctionDecl {{.*}}KernelName 'void ()' // CHECK: -IntelReqdSubGroupSizeAttr {{.*}} - // CHECK: `-SYCLIntelNoGlobalWorkOffsetAttr {{.*}} + // CHECK-NOT: `-SYCLIntelNoGlobalWorkOffsetAttr {{.*}} parallel_for([]() {}); #else parallel_for([]() {}); // expected-error 3 {{conflicting attributes applied to a SYCL kernel or SYCL_EXTERNAL function}} diff --git a/clang/test/SemaSYCL/intel-restrict.cpp b/clang/test/SemaSYCL/intel-restrict.cpp index bd51eb6b5c055..1a8f817895561 100644 --- a/clang/test/SemaSYCL/intel-restrict.cpp +++ b/clang/test/SemaSYCL/intel-restrict.cpp @@ -27,7 +27,7 @@ int main() { []() [[intel::kernel_args_restrict]] {}); // CHECK-LABEL: FunctionDecl {{.*}}test_kernel3 - // CHECK: SYCLIntelKernelArgsRestrictAttr + // CHECK-NOT: SYCLIntelKernelArgsRestrictAttr kernel( []() { func_do_not_ignore(); }); } diff --git a/clang/test/SemaSYCL/redeclaration-attribute-propagation.cpp b/clang/test/SemaSYCL/redeclaration-attribute-propagation.cpp index ccacb1f4beee4..4768964c0736f 100644 --- a/clang/test/SemaSYCL/redeclaration-attribute-propagation.cpp +++ b/clang/test/SemaSYCL/redeclaration-attribute-propagation.cpp @@ -53,10 +53,7 @@ int main() { #ifndef TRIGGER_ERROR // CHECK-LABEL: FunctionDecl {{.*}} main 'int ()' // CHECK: `-FunctionDecl {{.*}}test_kernel1 'void ()' - // CHECK: -SYCLIntelNoGlobalWorkOffsetAttr {{.*}} - // CHECK-NEXT: ConstantExpr {{.*}} 'int' - // CHECK-NEXT: value: Int 1 - // CHECK-NEXT: IntegerLiteral{{.*}}1{{$}} + // CHECK-NOT: SYCLIntelNoGlobalWorkOffsetAttr {{.*}} // CHECK: -SYCLIntelMaxWorkGroupSizeAttr {{.*}} Inherited // CHECK-NEXT: ConstantExpr {{.*}} 'int' // CHECK-NEXT: value: Int 4 diff --git a/clang/test/SemaSYCL/scheduler_target_fmax_mhz.cpp b/clang/test/SemaSYCL/scheduler_target_fmax_mhz.cpp index 984ceacd8b67e..636d36422b699 100644 --- a/clang/test/SemaSYCL/scheduler_target_fmax_mhz.cpp +++ b/clang/test/SemaSYCL/scheduler_target_fmax_mhz.cpp @@ -1,60 +1,77 @@ -// RUN: %clang_cc1 %s -fsyntax-only -ast-dump -fsycl-is-device -triple spir64 -Wno-sycl-2017-compat -verify | FileCheck %s +// RUN: %clang_cc1 -fsycl-is-device -verify %s -#include "Inputs/sycl.hpp" +// Test that checks scheduler_target_fmax_mhz attribute support on Function. + +// Test for deprecated spelling of scheduler_target_fmax_mhz attribute. // expected-warning@+2 {{attribute 'intelfpga::scheduler_target_fmax_mhz' is deprecated}} // expected-note@+1 {{did you mean to use 'intel::scheduler_target_fmax_mhz' instead?}} -[[intelfpga::scheduler_target_fmax_mhz(2)]] void -func() {} +[[intelfpga::scheduler_target_fmax_mhz(2)]] void deprecate() {} + +// Tests for incorrect argument values for Intel FPGA scheduler_target_fmax_mhz function attribute. +[[intel::scheduler_target_fmax_mhz(0)]] int Var = 0; // expected-error{{'scheduler_target_fmax_mhz' attribute only applies to functions}} + +[[intel::scheduler_target_fmax_mhz(1048577)]] void correct() {} // OK + +[[intel::scheduler_target_fmax_mhz("foo")]] void func() {} // expected-error{{integral constant expression must have integral or unscoped enumeration type, not 'const char [4]'}} + +[[intel::scheduler_target_fmax_mhz(-1)]] void func1() {} // expected-error{{'scheduler_target_fmax_mhz' attribute requires a non-negative integral compile time constant expression}} + +[[intel::scheduler_target_fmax_mhz(0, 1)]] void func2() {} // expected-error{{'scheduler_target_fmax_mhz' attribute takes one argument}} + +// Tests for Intel FPGA scheduler_target_fmax_mhz function attribute duplication. +// No diagnostic is emitted because the arguments match. Duplicate attribute is silently ignored. +[[intel::scheduler_target_fmax_mhz(2)]] +[[intel::scheduler_target_fmax_mhz(2)]] void func3() {} // No diagnostic is emitted because the arguments match. -[[intel::scheduler_target_fmax_mhz(12)]] void bar(); -[[intel::scheduler_target_fmax_mhz(12)]] void bar() {} // OK +[[intel::scheduler_target_fmax_mhz(4)]] void func4(); +[[intel::scheduler_target_fmax_mhz(4)]] void func4(); // OK // Diagnostic is emitted because the arguments mismatch. -[[intel::scheduler_target_fmax_mhz(12)]] void baz(); // expected-note {{previous attribute is here}} -[[intel::scheduler_target_fmax_mhz(100)]] void baz(); // expected-warning {{attribute 'scheduler_target_fmax_mhz' is already applied with different arguments}} +[[intel::scheduler_target_fmax_mhz(2)]] // expected-note {{previous attribute is here}} +[[intel::scheduler_target_fmax_mhz(4)]] void func5() {} // expected-warning {{attribute 'scheduler_target_fmax_mhz' is already applied with different arguments}} + +[[intel::scheduler_target_fmax_mhz(1)]] void func6(); // expected-note {{previous attribute is here}} +[[intel::scheduler_target_fmax_mhz(3)]] void func6(); // expected-warning {{attribute 'scheduler_target_fmax_mhz' is already applied with different arguments}} +// Tests that check template parameter support for Intel FPGA scheduler_target_fmax_mhz function attributes. template -[[intel::scheduler_target_fmax_mhz(N)]] void zoo() {} - -int main() { - // CHECK-LABEL: FunctionDecl {{.*}}test_kernel1 'void ()' - // CHECK: SYCLIntelSchedulerTargetFmaxMhzAttr {{.*}} - // CHECK-NEXT: ConstantExpr {{.*}} 'int' - // CHECK-NEXT: value: Int 5 - // CHECK-NEXT: IntegerLiteral {{.*}} 'int' 5 - // expected-warning@+3 {{attribute 'intelfpga::scheduler_target_fmax_mhz' is deprecated}} - // expected-note@+2 {{did you mean to use 'intel::scheduler_target_fmax_mhz' instead?}} - cl::sycl::kernel_single_task( - []() [[intelfpga::scheduler_target_fmax_mhz(5)]]{}); - - // CHECK-LABEL: FunctionDecl {{.*}}test_kernel2 'void ()' - // CHECK: SYCLIntelSchedulerTargetFmaxMhzAttr {{.*}} - // CHECK-NEXT: ConstantExpr {{.*}} 'int' - // CHECK-NEXT: value: Int 2 - // CHECK-NEXT: IntegerLiteral {{.*}} 'int' 2 - cl::sycl::kernel_single_task( - []() { func(); }); - - // CHECK-LABEL: FunctionDecl {{.*}}test_kernel3 'void ()' - // CHECK: SYCLIntelSchedulerTargetFmaxMhzAttr {{.*}} - // CHECK-NEXT: ConstantExpr {{.*}} 'int' - // CHECK-NEXT: value: Int 75 - // CHECK-NEXT: SubstNonTypeTemplateParmExpr {{.*}} 'int' - // CHECK-NEXT: NonTypeTemplateParmDecl {{.*}} referenced 'int' depth 0 index 0 N - // CHECK-NEXT: IntegerLiteral {{.*}} 'int' 75 - cl::sycl::kernel_single_task( - []() { zoo<75>(); }); - - [[intel::scheduler_target_fmax_mhz(0)]] int Var = 0; // expected-error{{'scheduler_target_fmax_mhz' attribute only applies to functions}} - - cl::sycl::kernel_single_task( - []() [[intel::scheduler_target_fmax_mhz(1048577)]]{}); // OK - - cl::sycl::kernel_single_task( - []() [[intel::scheduler_target_fmax_mhz(-4)]]{}); // expected-error{{'scheduler_target_fmax_mhz' attribute requires a non-negative integral compile time constant expression}} - - cl::sycl::kernel_single_task( - []() [[intel::scheduler_target_fmax_mhz(1), // expected-note {{previous attribute is here}} - intel::scheduler_target_fmax_mhz(2)]]{}); // expected-warning{{attribute 'scheduler_target_fmax_mhz' is already applied with different arguments}} +[[intel::scheduler_target_fmax_mhz(N)]] void func7(); // expected-error {{'scheduler_target_fmax_mhz' attribute requires a non-negative integral compile time constant expression}} + +template +[[intel::scheduler_target_fmax_mhz(10)]] void func8(); // expected-note {{previous attribute is here}} +template +[[intel::scheduler_target_fmax_mhz(size)]] void func8() {} // expected-warning {{attribute 'scheduler_target_fmax_mhz' is already applied with different arguments}} + +void checkTemplates() { + func7<4>(); // OK + func7<-1>(); // expected-note {{in instantiation of function template specialization 'func7<-1>' requested here}} + func7<0>(); // OK + func8<20>(); // expected-note {{in instantiation of function template specialization 'func8<20>' requested here}} +} + +// Test that checks expression is not a constant expression. +// expected-note@+1{{declared here}} +int baz(); +// expected-error@+2{{expression is not an integral constant expression}} +// expected-note@+1{{non-constexpr function 'baz' cannot be used in a constant expression}} +[[intel::scheduler_target_fmax_mhz(baz() + 1)]] void func9(); + +// Test that checks expression is a constant expression. +constexpr int bar() { return 0; } +[[intel::scheduler_target_fmax_mhz(bar() + 2)]] void func10(); // 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@+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::scheduler_target_fmax_mhz(Ty{})]] void func11() {} + +struct S {}; +void test() { + //expected-note@+1{{in instantiation of function template specialization 'func11' requested here}} + func11(); + //expected-note@+1{{in instantiation of function template specialization 'func11' requested here}} + func11(); } diff --git a/clang/test/SemaSYCL/scheduler_target_fmax_mhz_ast.cpp b/clang/test/SemaSYCL/scheduler_target_fmax_mhz_ast.cpp new file mode 100644 index 0000000000000..2c9ca6e380a9d --- /dev/null +++ b/clang/test/SemaSYCL/scheduler_target_fmax_mhz_ast.cpp @@ -0,0 +1,97 @@ +// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -sycl-std=2020 -ast-dump %s | FileCheck %s + +// Tests for AST of Intel FPGA scheduler_target_fmax_mhz function attribute. +#include "sycl.hpp" + +sycl::queue deviceQueue; + +// CHECK: FunctionDecl {{.*}} func1 'void ()' +// CHECK-NEXT: CompoundStmt +// CHECK-NEXT: SYCLIntelSchedulerTargetFmaxMhzAttr {{.*}} +// CHECK-NEXT: ConstantExpr{{.*}}'int' +// CHECK-NEXT: value: Int 4 +// CHECK-NEXT: IntegerLiteral {{.*}} 'int' 4 +[[intel::scheduler_target_fmax_mhz(4)]] void func1() {} + +// Test that checks template parameter support on function. +// CHECK: FunctionTemplateDecl {{.*}} func2 +// CHECK: FunctionDecl {{.*}} func2 'void ()' +// CHECK-NEXT: CompoundStmt +// CHECK_NEXT: SYCLIntelSchedulerTargetFmaxMhzAttr {{.*}} +// CHECK_NEXT: DeclRefExpr {{.*}} 'int' NonTypeTemplateParm {{.*}} 'N' 'int' +// CHECK: FunctionDecl {{.*}} func2 'void ()' +// CHECK-NEXT: TemplateArgument integral 6 +// CHECK-NEXT: CompoundStmt +// CHECK-NEXT: SYCLIntelSchedulerTargetFmaxMhzAttr {{.*}} +// CHECK-NEXT: ConstantExpr{{.*}}'int' +// CHECK-NEXT: value: Int 6 +// CHECK-NEXT: SubstNonTypeTemplateParmExpr +// CHECK-NEXT: NonTypeTemplateParmDecl +// CHECK-NEXT: IntegerLiteral {{.*}} 'int' 6 +template +[[intel::scheduler_target_fmax_mhz(N)]] void func2() {} + +// No diagnostic is emitted because the arguments match. Duplicate attribute is silently ignored. +// CHECK: FunctionDecl {{.*}} {{.*}} func3 'void ()' +// CHECK: SYCLIntelSchedulerTargetFmaxMhzAttr {{.*}} +// CHECK-NEXT: ConstantExpr {{.*}} 'int' +// CHECK-NEXT: value: Int 10 +// CHECK-NEXT: IntegerLiteral{{.*}}10{{$}} +[[intel::scheduler_target_fmax_mhz(10)]] +[[intel::scheduler_target_fmax_mhz(10)]] void func3() {} + +class KernelFunctor { +public: + void operator()() const { + func1(); + } +}; + +// Test that checks template parameter support on class member function. +template +class KernelFunctor2 { +public: + [[intel::scheduler_target_fmax_mhz(N)]] void operator()() const { + } +}; + +int main() { + deviceQueue.submit([&](sycl::handler &h) { + // CHECK-LABEL: FunctionDecl {{.*}}kernel_name_1 + // CHECK-NOT: SYCLIntelSchedulerTargetFmaxMhzAttr + KernelFunctor f1; + h.single_task(f1); + + // CHECK-LABEL: FunctionDecl {{.*}}kernel_name_2 + // CHECK: SYCLIntelSchedulerTargetFmaxMhzAttr {{.*}} + // CHECK-NEXT: ConstantExpr{{.*}}'int' + // CHECK-NEXT: value: Int 3 + // CHECK-NEXT: SubstNonTypeTemplateParmExpr + // CHECK-NEXT: NonTypeTemplateParmDecl + // CHECK-NEXT: IntegerLiteral {{.*}} 'int' 3 + KernelFunctor2<3> f2; + h.single_task(f2); + + // CHECK-LABEL: FunctionDecl {{.*}}kernel_name_3 + // CHECK: SYCLIntelSchedulerTargetFmaxMhzAttr {{.*}} + // CHECK-NEXT: ConstantExpr{{.*}}'int' + // CHECK-NEXT: value: Int 4 + // CHECK-NEXT: IntegerLiteral {{.*}} 'int' 4 + h.single_task( + []() [[intel::scheduler_target_fmax_mhz(4)]]{}); + + // Ignore duplicate attribute. + h.single_task( + // CHECK-LABEL: FunctionDecl {{.*}}kernel_name_4 + // CHECK: SYCLIntelSchedulerTargetFmaxMhzAttr {{.*}} + // CHECK-NEXT: ConstantExpr {{.*}} 'int' + // CHECK-NEXT: value: Int 6 + // CHECK-NEXT: IntegerLiteral{{.*}}6{{$}} + []() [[intel::scheduler_target_fmax_mhz(6), + intel::scheduler_target_fmax_mhz(6)]]{}); + }); + + func2<6>(); + + return 0; +} diff --git a/clang/test/SemaSYCL/sycl-device-scheduler_target_fmax_mhz-template.cpp b/clang/test/SemaSYCL/sycl-device-scheduler_target_fmax_mhz-template.cpp deleted file mode 100644 index 407488f9bb0ae..0000000000000 --- a/clang/test/SemaSYCL/sycl-device-scheduler_target_fmax_mhz-template.cpp +++ /dev/null @@ -1,95 +0,0 @@ -// RUN: %clang_cc1 -fsycl-is-device -fsyntax-only -ast-dump -verify -pedantic %s | FileCheck %s - -// Test that checkes template parameter support for 'scheduler_target_fmax_mhz' 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@+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::scheduler_target_fmax_mhz(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(); -} - -// 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::scheduler_target_fmax_mhz(foo() + 1)]] void func1(); - -// Test that checks expression is a constant expression. -constexpr int bar() { return 0; } -[[intel::scheduler_target_fmax_mhz(bar() + 2)]] void func2(); // OK - -// Test that checks template parameter support on member function of class template. -template -class KernelFunctor { -public: - // expected-error@+1{{'scheduler_target_fmax_mhz' attribute requires a non-negative integral compile time constant expression}} - [[intel::scheduler_target_fmax_mhz(SIZE)]] void operator()() {} -}; - -int main() { - //expected-note@+1{{in instantiation of template class 'KernelFunctor<-1>' requested here}} - KernelFunctor<-1>(); - // no error expected - KernelFunctor<2>(); -} - -// CHECK: ClassTemplateDecl {{.*}} {{.*}} KernelFunctor -// CHECK: ClassTemplateSpecializationDecl {{.*}} {{.*}} class KernelFunctor definition -// CHECK: TemplateArgument integral 2 -// CHECK: CXXRecordDecl {{.*}} {{.*}} implicit class KernelFunctor -// CHECK: SYCLIntelSchedulerTargetFmaxMhzAttr {{.*}} -// CHECK-NEXT: ConstantExpr {{.*}} 'int' -// CHECK-NEXT: value: Int 2 -// CHECK-NEXT: SubstNonTypeTemplateParmExpr {{.*}} -// CHECK-NEXT: NonTypeTemplateParmDecl {{.*}} -// CHECK-NEXT: IntegerLiteral{{.*}}2{{$}} - -// Test that checks template parameter support on function. -template -// expected-error@+1{{'scheduler_target_fmax_mhz' attribute requires a non-negative integral compile time constant expression}} -[[intel::scheduler_target_fmax_mhz(N)]] void func3() {} - -template -[[intel::scheduler_target_fmax_mhz(4)]] void func4(); // expected-note {{previous attribute is here}} - -template -[[intel::scheduler_target_fmax_mhz(N)]] void func4() {} // expected-warning {{attribute 'scheduler_target_fmax_mhz' is already applied with different arguments}} - -int check() { - // no error expected - func3<3>(); - //expected-note@+1{{in instantiation of function template specialization 'func3<-1>' requested here}} - func3<-1>(); - //expected-note@+1 {{in instantiation of function template specialization 'func4<6>' requested here}} - func4<6>(); - return 0; -} - -// No diagnostic is emitted because the arguments match. Duplicate attribute is silently ignored. -[[intel::scheduler_target_fmax_mhz(8)]] -[[intel::scheduler_target_fmax_mhz(8)]] void func5() {} - -// CHECK: FunctionDecl {{.*}} {{.*}} func3 'void ()' -// CHECK: TemplateArgument integral 3 -// CHECK: SYCLIntelSchedulerTargetFmaxMhzAttr {{.*}} -// CHECK-NEXT: ConstantExpr {{.*}} 'int' -// CHECK-NEXT: value: Int 3 -// CHECK-NEXT: SubstNonTypeTemplateParmExpr {{.*}} -// CHECK-NEXT: NonTypeTemplateParmDecl {{.*}} -// CHECK-NEXT: IntegerLiteral{{.*}}3{{$}} - -// CHECK: FunctionDecl {{.*}} {{.*}} func5 'void ()' -// CHECK: SYCLIntelSchedulerTargetFmaxMhzAttr {{.*}} -// CHECK-NEXT: ConstantExpr {{.*}} 'int' -// CHECK-NEXT: value: Int 8 -// CHECK-NEXT: IntegerLiteral{{.*}}8{{$}}