diff --git a/clang/include/clang/Basic/Attr.td b/clang/include/clang/Basic/Attr.td index 53a95c1c8c365..b5b97dfea0c0c 100644 --- a/clang/include/clang/Basic/Attr.td +++ b/clang/include/clang/Basic/Attr.td @@ -1303,7 +1303,7 @@ def SYCLIntelMaxGlobalWorkDim : InheritableAttr { def SYCLIntelNoGlobalWorkOffset : InheritableAttr { let Spellings = [CXX11<"intelfpga","no_global_work_offset">, CXX11<"intel","no_global_work_offset">]; - let Args = [BoolArgument<"Enabled", 1>]; + let Args = [ExprArgument<"Value", /*optional*/1>]; let LangOpts = [SYCLIsDevice, SYCLIsHost]; let Subjects = SubjectList<[Function], ErrorDiag>; let Documentation = [SYCLIntelNoGlobalWorkOffsetAttrDocs]; diff --git a/clang/include/clang/Basic/DiagnosticGroups.td b/clang/include/clang/Basic/DiagnosticGroups.td index e2894e764dba9..4b42a37fd33aa 100644 --- a/clang/include/clang/Basic/DiagnosticGroups.td +++ b/clang/include/clang/Basic/DiagnosticGroups.td @@ -667,10 +667,8 @@ def NSReturnsMismatch : DiagGroup<"nsreturns-mismatch">; def IndependentClassAttribute : DiagGroup<"IndependentClass-attribute">; def UnknownAttributes : DiagGroup<"unknown-attributes">; def IgnoredAttributes : DiagGroup<"ignored-attributes">; -def AdjustedAttributes : DiagGroup<"adjusted-attributes">; def Attributes : DiagGroup<"attributes", [UnknownAttributes, - IgnoredAttributes, - AdjustedAttributes]>; + IgnoredAttributes]>; def UnknownSanitizers : DiagGroup<"unknown-sanitizers">; def UnnamedTypeTemplateArgs : DiagGroup<"unnamed-type-template-args", [CXX98CompatUnnamedTypeTemplateArgs]>; diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td index a0eaf8d7bdb3c..f4511886eb212 100644 --- a/clang/include/clang/Basic/DiagnosticSemaKinds.td +++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td @@ -11134,9 +11134,6 @@ def err_sycl_function_attribute_mismatch : Error< "SYCL kernel without %0 attribute can't call a function with this attribute">; def err_sycl_x_y_z_arguments_must_be_one : Error< "%0 X-, Y- and Z- sizes must be 1 when %1 attribute is used with value 0">; -def warn_boolean_attribute_argument_is_not_valid: Warning< - "The value of %0 attribute should be 0 or 1. Adjusted to 1">, - InGroup; def err_sycl_attibute_cannot_be_applied_here : Error<"%0 attribute cannot be applied to a " "static function or function in an anonymous namespace">; diff --git a/clang/lib/CodeGen/CodeGenFunction.cpp b/clang/lib/CodeGen/CodeGenFunction.cpp index f94a6a6973b6a..5c947cc66b42d 100644 --- a/clang/lib/CodeGen/CodeGenFunction.cpp +++ b/clang/lib/CodeGen/CodeGenFunction.cpp @@ -694,7 +694,12 @@ void CodeGenFunction::EmitOpenCLKernelMetadata(const FunctionDecl *FD, if (const SYCLIntelNoGlobalWorkOffsetAttr *A = FD->getAttr()) { - if (A->getEnabled()) + const Expr *Arg = A->getValue(); + assert(Arg && "Got an unexpected null argument"); + Optional ArgVal = + Arg->getIntegerConstantExpr(FD->getASTContext()); + assert(ArgVal.hasValue() && "Not an integer constant expression"); + if (ArgVal->getBoolValue()) Fn->setMetadata("no_global_work_offset", llvm::MDNode::get(Context, {})); } diff --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp index d1e5855ff2e0b..cc219422c963b 100644 --- a/clang/lib/Sema/SemaDeclAttr.cpp +++ b/clang/lib/Sema/SemaDeclAttr.cpp @@ -5294,24 +5294,18 @@ static void handleNoGlobalWorkOffsetAttr(Sema &S, Decl *D, checkForDuplicateAttribute(S, D, Attr); - uint32_t Enabled = 1; - if (Attr.getNumArgs()) { - const Expr *E = Attr.getArgAsExpr(0); - if (!checkUInt32Argument(S, Attr, E, Enabled, 0, - /*StrictlyUnsigned=*/true)) - return; - } - if (Enabled > 1) - S.Diag(Attr.getLoc(), diag::warn_boolean_attribute_argument_is_not_valid) - << Attr; - if (Attr.getKind() == ParsedAttr::AT_SYCLIntelNoGlobalWorkOffset && checkDeprecatedSYCLAttributeSpelling(S, Attr)) S.Diag(Attr.getLoc(), diag::note_spelling_suggestion) << "'intel::no_global_work_offset'"; - D->addAttr(::new (S.Context) - SYCLIntelNoGlobalWorkOffsetAttr(S.Context, Attr, Enabled)); + // If no attribute argument is specified, set to default value '1'. + Expr *E = Attr.isArgExpr(0) + ? Attr.getArgAsExpr(0) + : IntegerLiteral::Create(S.Context, llvm::APInt(32, 1), + S.Context.IntTy, Attr.getLoc()); + S.addIntelSYCLSingleArgFunctionAttr(D, Attr, + E); } /// Handle the [[intelfpga::doublepump]] and [[intelfpga::singlepump]] attributes. diff --git a/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp b/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp index 5130bad1f7b5c..85fd57c459bf3 100644 --- a/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp +++ b/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp @@ -775,6 +775,12 @@ void Sema::InstantiateAttrs(const MultiLevelTemplateArgumentList &TemplateArgs, *this, TemplateArgs, SYCLIntelMaxGlobalWorkDim, New); continue; } + if (const auto *SYCLIntelNoGlobalWorkOffset = + dyn_cast(TmplAttr)) { + instantiateIntelSYCLFunctionAttr( + *this, TemplateArgs, SYCLIntelNoGlobalWorkOffset, New); + continue; + } // Existing DLL attribute on the instantiation takes precedence. if (TmplAttr->getKind() == attr::DLLExport || TmplAttr->getKind() == attr::DLLImport) { 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 a4d93478134a1..a2d33c2ac2932 100644 --- a/clang/test/CodeGenSYCL/intel-fpga-no-global-work-offset.cpp +++ b/clang/test/CodeGenSYCL/intel-fpga-no-global-work-offset.cpp @@ -1,28 +1,49 @@ -// RUN: %clang_cc1 -fsycl -fsycl-is-device -triple spir64-unknown-unknown-sycldevice -disable-llvm-passes -emit-llvm -o - %s | FileCheck %s +// RUN: %clang_cc1 -fsycl -fsycl-is-device -internal-isystem %S/Inputs -triple spir64-unknown-unknown-sycldevice -disable-llvm-passes -emit-llvm -o - %s | FileCheck %s + +#include "sycl.hpp" + +using namespace cl::sycl; +queue q; class Foo { public: [[intel::no_global_work_offset(1)]] void operator()() const {} }; -template -__attribute__((sycl_kernel)) void kernel(const Func &kernelFunc) { - kernelFunc(); -} +template +class Functor { +public: + [[intel::no_global_work_offset(SIZE)]] void operator()() const {} +}; + +template +[[intel::no_global_work_offset(N)]] void func() {} + +int main() { + q.submit([&](handler &h) { + Foo boo; + h.single_task(boo); + + h.single_task( + []() [[intel::no_global_work_offset]]{}); -void bar() { - Foo boo; - kernel(boo); + h.single_task( + []() [[intel::no_global_work_offset(0)]]{}); - kernel( - []() [[intel::no_global_work_offset]]{}); + Functor<1> f; + h.single_task(f); - kernel( - []() [[intel::no_global_work_offset(0)]]{}); + h.single_task([]() { + func<1>(); + }); + }); + return 0; } -// CHECK: define spir_kernel void @{{.*}}kernel_name1() {{.*}} !no_global_work_offset ![[NUM5:[0-9]+]] -// CHECK: define spir_kernel void @{{.*}}kernel_name2() {{.*}} !no_global_work_offset ![[NUM5]] -// CHECK: define spir_kernel void @{{.*}}kernel_name3() {{.*}} ![[NUM4:[0-9]+]] +// CHECK: define spir_kernel void @{{.*}}kernel_name1"() #0 {{.*}} !no_global_work_offset ![[NUM5:[0-9]+]] +// 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-NOT: ![[NUM4]] = !{i32 0} // CHECK: ![[NUM5]] = !{} diff --git a/clang/test/SemaSYCL/check-notdirect-attribute-propagation.cpp b/clang/test/SemaSYCL/check-notdirect-attribute-propagation.cpp index 777ef128123eb..d2ade080f86cc 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 {{.*}} Enabled + // CHECK: `-SYCLIntelNoGlobalWorkOffsetAttr {{.*}} parallel_for([]() {}); #else parallel_for([]() {}); // expected-error 2 {{conflicting attributes applied to a SYCL kernel or SYCL_EXTERNAL function}} diff --git a/clang/test/SemaSYCL/intel-fpga-no-global-work-offset.cpp b/clang/test/SemaSYCL/intel-fpga-no-global-work-offset.cpp index ff816237d6fb1..c8c54a6913587 100644 --- a/clang/test/SemaSYCL/intel-fpga-no-global-work-offset.cpp +++ b/clang/test/SemaSYCL/intel-fpga-no-global-work-offset.cpp @@ -1,51 +1,50 @@ -// RUN: %clang_cc1 -fsycl -fsycl-is-device -Wno-return-type -fcxx-exceptions -fsyntax-only -ast-dump -verify -pedantic %s | FileCheck %s +// RUN: %clang_cc1 -fsycl -fsycl-is-device -internal-isystem %S/Inputs -Wno-return-type -Wno-sycl-2017-compat -fcxx-exceptions -fsyntax-only -ast-dump -verify -pedantic %s | FileCheck %s + +#include "sycl.hpp" + +using namespace cl::sycl; +queue q; struct FuncObj { //expected-warning@+2 {{attribute 'intelfpga::no_global_work_offset' is deprecated}} //expected-note@+1 {{did you mean to use 'intel::no_global_work_offset' instead?}} - [[intelfpga::no_global_work_offset]] void operator()() {} + [[intelfpga::no_global_work_offset]] void operator()() const {} }; -template -void kernel(Func kernelFunc) { - kernelFunc(); -} - int main() { - // CHECK: SYCLIntelNoGlobalWorkOffsetAttr{{.*}}Enabled - kernel([]() { - FuncObj(); - }); - - // CHECK: SYCLIntelNoGlobalWorkOffsetAttr - // CHECK-NOT: Enabled - kernel( - []() [[intel::no_global_work_offset(0)]]{}); - - // CHECK: SYCLIntelNoGlobalWorkOffsetAttr{{.*}}Enabled - // expected-warning@+2{{'no_global_work_offset' attribute should be 0 or 1. Adjusted to 1}} - kernel( - []() [[intel::no_global_work_offset(42)]]{}); - - // expected-error@+2{{'no_global_work_offset' attribute requires a non-negative integral compile time constant expression}} - kernel( - []() [[intel::no_global_work_offset(-1)]]{}); - - // expected-error@+2{{'no_global_work_offset' attribute requires parameter 0 to be an integer constant}} - kernel( - []() [[intel::no_global_work_offset("foo")]]{}); - - kernel([]() { - // expected-error@+1{{'no_global_work_offset' attribute only applies to functions}} - [[intel::no_global_work_offset(1)]] int a; + q.submit([&](handler &h) { + // CHECK: SYCLIntelNoGlobalWorkOffsetAttr {{.*}} + // CHECK-NEXT: IntegerLiteral{{.*}}1{{$}} + h.single_task(FuncObj()); + + // CHECK: SYCLIntelNoGlobalWorkOffsetAttr {{.*}} + // CHECK-NEXT: IntegerLiteral{{.*}}0{{$}} + h.single_task( + []() [[intel::no_global_work_offset(0)]]{}); + + // CHECK: SYCLIntelNoGlobalWorkOffsetAttr{{.*}} + // CHECK-NEXT: IntegerLiteral{{.*}}42{{$}} + h.single_task( + []() [[intel::no_global_work_offset(42)]]{}); + + // CHECK: SYCLIntelNoGlobalWorkOffsetAttr{{.*}} + // CHECK-NEXT: UnaryOperator{{.*}} 'int' prefix '-' + // CHECK-NEXT-NEXT: IntegerLiteral{{.*}}1{{$}} + h.single_task( + []() [[intel::no_global_work_offset(-1)]]{}); + + // expected-error@+2{{'no_global_work_offset' attribute requires an integer constant}} + h.single_task( + []() [[intel::no_global_work_offset("foo")]]{}); + + h.single_task([]() { + // expected-error@+1{{'no_global_work_offset' attribute only applies to functions}} + [[intel::no_global_work_offset(1)]] int a; + }); + + // expected-warning@+2{{attribute 'no_global_work_offset' is already applied}} + h.single_task( + []() [[intel::no_global_work_offset(0), intel::no_global_work_offset(1)]]{}); }); - - // CHECK: SYCLIntelNoGlobalWorkOffsetAttr{{.*}} - // CHECK-NOT: Enabled - // CHECK: SYCLIntelNoGlobalWorkOffsetAttr{{.*}}Enabled - // expected-warning@+2{{attribute 'no_global_work_offset' is already applied}} - kernel( - []() [[intel::no_global_work_offset(0), intel::no_global_work_offset(1)]]{}); - return 0; } diff --git a/clang/test/SemaSYCL/redeclaration-attribute-propagation.cpp b/clang/test/SemaSYCL/redeclaration-attribute-propagation.cpp index 4b2777c1bb9dd..436f9b3186de2 100644 --- a/clang/test/SemaSYCL/redeclaration-attribute-propagation.cpp +++ b/clang/test/SemaSYCL/redeclaration-attribute-propagation.cpp @@ -1,8 +1,11 @@ -// RUN: %clang_cc1 %s -fsyntax-only -fsycl -fsycl-is-device -triple spir64 -Wno-sycl-2017-compat -verify -// RUN: %clang_cc1 %s -fsyntax-only -fsycl -fsycl-is-device -triple spir64 -DTRIGGER_ERROR -Wno-sycl-2017-compat -verify -// RUN: %clang_cc1 %s -fsyntax-only -ast-dump -fsycl -fsycl-is-device -triple spir64 -Wno-sycl-2017-compat | FileCheck %s +// RUN: %clang_cc1 %s -fsyntax-only -fsycl -fsycl-is-device -internal-isystem %S/Inputs -triple spir64 -Wno-sycl-2017-compat -verify +// RUN: %clang_cc1 %s -fsyntax-only -fsycl -fsycl-is-device -internal-isystem %S/Inputs -triple spir64 -DTRIGGER_ERROR -Wno-sycl-2017-compat -verify +// RUN: %clang_cc1 %s -fsyntax-only -ast-dump -fsycl -fsycl-is-device -internal-isystem %S/Inputs -triple spir64 -Wno-sycl-2017-compat | FileCheck %s -#include "Inputs/sycl.hpp" +#include "sycl.hpp" + +using namespace cl::sycl; +queue q; #ifndef TRIGGER_ERROR //first case - good case @@ -46,23 +49,26 @@ func4() {} // expected-error {{'max_work_group_size' attribute conflicts with '' #endif int main() { + q.submit([&](handler &h) { #ifndef TRIGGER_ERROR - // CHECK-LABEL: FunctionDecl {{.*}} main 'int ()' - // CHECK: `-FunctionDecl {{.*}}test_kernel1 'void ()' - // CHECK: -SYCLIntelMaxWorkGroupSizeAttr {{.*}} Inherited 4 4 4 - // CHECK: -SYCLIntelNoGlobalWorkOffsetAttr {{.*}} Inherited Enabled - // CHECK: `-ReqdWorkGroupSizeAttr {{.*}} 2 2 2 - cl::sycl::kernel_single_task( - []() { func1(); }); + // CHECK-LABEL: FunctionDecl {{.*}} main 'int ()' + // CHECK: `-FunctionDecl {{.*}}test_kernel1 'void ()' + // CHECK: -SYCLIntelMaxWorkGroupSizeAttr {{.*}} Inherited 4 4 4 + // CHECK: -SYCLIntelNoGlobalWorkOffsetAttr {{.*}} + // CHECK: `-ReqdWorkGroupSizeAttr {{.*}} 2 2 2 + h.single_task( + []() { func1(); }); #else - cl::sycl::kernel_single_task( - []() { func2(); }); // expected-error {{conflicting attributes applied to a SYCL kernel or SYCL_EXTERNAL function}} + h.single_task( + []() { func2(); }); // expected-error {{conflicting attributes applied to a SYCL kernel or SYCL_EXTERNAL function}} - cl::sycl::kernel_single_task( - []() { func3(); }); + h.single_task( + []() { func3(); }); - cl::sycl::kernel_single_task( - []() { func4(); }); + h.single_task( + []() { func4(); }); #endif + }); + return 0; } diff --git a/clang/test/SemaSYCL/sycl-device-intel-fpga-no-global-work-offset.cpp b/clang/test/SemaSYCL/sycl-device-intel-fpga-no-global-work-offset.cpp new file mode 100644 index 0000000000000..c6e2bb0475d28 --- /dev/null +++ b/clang/test/SemaSYCL/sycl-device-intel-fpga-no-global-work-offset.cpp @@ -0,0 +1,60 @@ +// RUN: %clang_cc1 -fsycl -fsycl-is-device -fsyntax-only -ast-dump -verify -pedantic %s | FileCheck %s + +// Test that checks template parameter support for 'no_global_work_offset' 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@+1{{'no_global_work_offset' attribute requires an integer constant}} +[[intel::no_global_work_offset(Ty{})]] void func() {} + +struct S {}; +void var() { + //expected-note@+1{{in instantiation of function template specialization 'func' requested here}} + func(); +} + +// Test that checks expression is not a constant expression. +int foo(); +// expected-error@+1{{'no_global_work_offset' attribute requires an integer constant}} +[[intel::no_global_work_offset(foo() + 12)]] void func1(); + +// Test that checks expression is a constant expression. +constexpr int bar() { return 0; } +[[intel::no_global_work_offset(bar() + 12)]] void func2(); // OK + +// Test that checks template parameter suppport on member function of class template. +template +class KernelFunctor { +public: + [[intel::no_global_work_offset(SIZE)]] void operator()() {} +}; + +int main() { + KernelFunctor<1>(); +} + +// CHECK: ClassTemplateDecl {{.*}} {{.*}} KernelFunctor +// CHECK: ClassTemplateSpecializationDecl {{.*}} {{.*}} class KernelFunctor definition +// CHECK: CXXRecordDecl {{.*}} {{.*}} implicit class KernelFunctor +// CHECK: SYCLIntelNoGlobalWorkOffsetAttr {{.*}} +// CHECK: SubstNonTypeTemplateParmExpr {{.*}} +// CHECK-NEXT: NonTypeTemplateParmDecl {{.*}} +// CHECK-NEXT: IntegerLiteral{{.*}}1{{$}} + +// Test that checks template parameter suppport on function. +template +[[intel::no_global_work_offset(N)]] void func3() {} + +int check() { + func3<1>(); + return 0; +} + +// CHECK: FunctionTemplateDecl {{.*}} {{.*}} func3 +// CHECK: NonTypeTemplateParmDecl {{.*}} {{.*}} referenced 'int' depth 0 index 0 N +// CHECK: FunctionDecl {{.*}} {{.*}} func3 'void ()' +// CHECK: SYCLIntelNoGlobalWorkOffsetAttr {{.*}} +// CHECK: SubstNonTypeTemplateParmExpr {{.*}} +// CHECK-NEXT: NonTypeTemplateParmDecl {{.*}} +// CHECK-NEXT: IntegerLiteral{{.*}}1{{$}}