From c6334641a48e37f96d14c56ffed8835f4f7c9e9a Mon Sep 17 00:00:00 2001 From: Andrew Savonichev Date: Mon, 14 Oct 2019 15:52:09 +0300 Subject: [PATCH 1/5] [SYCL] Add intel::kernel_args_restrict attribute When applied to a device function that is invoked as a device kernel, the attribute is a hint to the compiler that no pointer argument to the kernel which is defined through an accessor (not USM), will alias any other pointer kernel argument that was defined through an accessor. This effect is equivalent to annotating restrict on all kernel pointer arguments in an OpenCL or SPIR-V kernel. Signed-off-by: Andrew Savonichev --- clang/include/clang/Basic/Attr.td | 6 ++ clang/include/clang/Basic/AttrDocs.td | 37 ++++++++++++ .../include/clang/Basic/AttributeCommonInfo.h | 7 +++ clang/lib/CodeGen/CGCall.cpp | 4 +- clang/lib/Sema/SemaDeclAttr.cpp | 13 ++++- clang/lib/Sema/SemaSYCL.cpp | 26 +++++++-- clang/lib/Sema/SemaType.cpp | 8 ++- clang/test/CodeGenSYCL/intel-restrict.cpp | 57 +++++++++++++++++++ ...a-attribute-supported-attributes-list.test | 1 + clang/test/SemaSYCL/intel-restrict.cpp | 33 +++++++++++ 10 files changed, 185 insertions(+), 7 deletions(-) create mode 100644 clang/test/CodeGenSYCL/intel-restrict.cpp create mode 100644 clang/test/SemaSYCL/intel-restrict.cpp diff --git a/clang/include/clang/Basic/Attr.td b/clang/include/clang/Basic/Attr.td index 55932dc4ec8e7..7f8b4dc1a9980 100644 --- a/clang/include/clang/Basic/Attr.td +++ b/clang/include/clang/Basic/Attr.td @@ -1091,6 +1091,12 @@ def SYCLDeviceIndirectlyCallable : InheritableAttr { let LangOpts = [SYCLIsDevice]; let Documentation = [SYCLDeviceIndirectlyCallableDocs]; } +def SYCLIntelKernelArgsRestrict : InheritableAttr { + let Spellings = [ CXX11<"intel", "kernel_args_restrict"> ]; + let Subjects = SubjectList<[ FunctionLike ], ErrorDiag>; + let LangOpts = [ SYCLIsDevice, SYCLIsHost ]; + let Documentation = [ SYCLIntelKernelArgsRestrictDocs ]; +} def C11NoReturn : InheritableAttr { let Spellings = [Keyword<"_Noreturn">]; diff --git a/clang/include/clang/Basic/AttrDocs.td b/clang/include/clang/Basic/AttrDocs.td index ce78fcc139ea0..75d2e4589be05 100644 --- a/clang/include/clang/Basic/AttrDocs.td +++ b/clang/include/clang/Basic/AttrDocs.td @@ -1835,6 +1835,43 @@ loads). }]; } +def SYCLIntelKernelArgsRestrictDocs : Documentation { + let Category = DocCatVariable; + let Heading = "kernel_args_restrict"; + let Content = [{ +The attribute ``intel::kernel_args_restrict`` is legal on device functions, and +can be ignored on non-device functions. When applied to a function, lambda, or +function call operator (of a functor), the attribute is a hint to the compiler +equivalent to specifying the C99 restrict attribute on all pointer arguments or +the pointer member of any accessors, which are a function argument, lambda +capture, or functor member, of the callable to which the 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, propagation of the attribute to any caller(s), including up to a kernel +boundary, is implementation defined and not guaranteed through this +extension. 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 +``inte::_kernel_args_restrict`` to a kernel, but there is in fact aliasing +between kernel pointer arguments at runtime, the behavior is undefined. + +The attribute-token ``intel::kernel_args_restrict`` shall appear at most once in +each attribute-list and no attribute-argument-clause shall be present. The +attribute may be applied to the function-type in a function declaration. The +first declaration of a function shall specify the +``intel::kernel_args_restrict`` attribute if any declaration of that function +specifies the ``intel::kernel_args_restrict`` attribute. If a function is +declared with the ``intel::kernel_args_restrict`` attribute in one translation +unit and the same function is declared without the +``intel::kernel_args_restrict`` attribute in another translation unit, the +program is ill-formed and no diagnostic is required. + +The ``intel::kernel_args_restrict`` attribute has an effect when applied to a +function, and no effect otherwise. + }]; +} + def SYCLIntelFPGAIVDepAttrDocs : Documentation { let Category = DocCatVariable; let Heading = "ivdep"; diff --git a/clang/include/clang/Basic/AttributeCommonInfo.h b/clang/include/clang/Basic/AttributeCommonInfo.h index 545e7e9a2b47e..88eec79ed4ef2 100644 --- a/clang/include/clang/Basic/AttributeCommonInfo.h +++ b/clang/include/clang/Basic/AttributeCommonInfo.h @@ -148,6 +148,13 @@ class AttributeCommonInfo { return SyntaxUsed == AS_CXX11 || isAlignasAttribute(); } + bool isAllowedOnLambdas() const { + // FIXME: Eventually we want to do a list here populated via tablegen. But + // we want C++ attributes to be permissible on Lambdas, and get propagated + // to the call operator declaration. + return getParsedKind() == AT_SYCLIntelKernelArgsRestrict; + } + bool isC2xAttribute() const { return SyntaxUsed == AS_C2x; } bool isKeywordAttribute() const { diff --git a/clang/lib/CodeGen/CGCall.cpp b/clang/lib/CodeGen/CGCall.cpp index 228eb5e44747a..d33576ba204eb 100644 --- a/clang/lib/CodeGen/CGCall.cpp +++ b/clang/lib/CodeGen/CGCall.cpp @@ -2412,7 +2412,9 @@ void CodeGenFunction::EmitFunctionProlog(const CGFunctionInfo &FI, } } - if (Arg->getType().isRestrictQualified()) + if (Arg->getType().isRestrictQualified() || + (CurCodeDecl && + CurCodeDecl->hasAttr())) AI->addAttr(llvm::Attribute::NoAlias); // LLVM expects swifterror parameters to be used in very restricted diff --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp index 353a8cd4ed91d..8796698712366 100644 --- a/clang/lib/Sema/SemaDeclAttr.cpp +++ b/clang/lib/Sema/SemaDeclAttr.cpp @@ -6752,6 +6752,13 @@ static void handleMSAllocatorAttr(Sema &S, Decl *D, const ParsedAttr &AL) { // Top Level Sema Entry Points //===----------------------------------------------------------------------===// +static bool IsDeclLambdaCallOperator(Decl *D) { + if (const auto *MD = dyn_cast(D)) + return MD->getParent()->isLambda() && + MD->getOverloadedOperator() == OverloadedOperatorKind::OO_Call; + return false; +} + /// ProcessDeclAttribute - Apply the specific attribute to the specified decl if /// the attribute applies to decls. If the attribute is a type attribute, just /// silently ignore it if a GNU attribute. @@ -6763,7 +6770,8 @@ static void ProcessDeclAttribute(Sema &S, Scope *scope, Decl *D, // Ignore C++11 attributes on declarator chunks: they appertain to the type // instead. - if (AL.isCXX11Attribute() && !IncludeCXX11Attributes) + if (AL.isCXX11Attribute() && !IncludeCXX11Attributes && + (!IsDeclLambdaCallOperator(D) || !AL.isAllowedOnLambdas())) return; // Unknown attributes are automatically warned on. Target-specific attributes @@ -7516,6 +7524,9 @@ static void ProcessDeclAttribute(Sema &S, Scope *scope, Decl *D, case ParsedAttr::AT_RenderScriptKernel: handleSimpleAttribute(S, D, AL); break; + case ParsedAttr::AT_SYCLIntelKernelArgsRestrict: + handleSimpleAttribute(S, D, AL); + break; // XRay attributes. case ParsedAttr::AT_XRayInstrument: handleSimpleAttribute(S, D, AL); diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index e440e21449e87..a97ffeb3af574 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -411,12 +411,14 @@ class MarkDeviceFunction : public RecursiveASTVisitor { // Attributes applied to SYCLKernel are also included void CollectPossibleKernelAttributes(FunctionDecl *SYCLKernel, llvm::SmallPtrSet &Attrs) { + typedef std::pair ChildParentPair; llvm::SmallPtrSet Visited; - llvm::SmallVector WorkList; - WorkList.push_back(SYCLKernel); + llvm::SmallVector WorkList; + WorkList.push_back({SYCLKernel, nullptr}); while (!WorkList.empty()) { - FunctionDecl *FD = WorkList.back(); + FunctionDecl *FD = WorkList.back().first; + FunctionDecl *ParentFD = WorkList.back().second; WorkList.pop_back(); if (!Visited.insert(FD).second) continue; // We've already seen this Decl @@ -425,6 +427,18 @@ class MarkDeviceFunction : public RecursiveASTVisitor { Attrs.insert(A); else if (auto *A = FD->getAttr()) Attrs.insert(A); + else if (auto *A = FD->getAttr()) { + // Allow the intel::kernel_args_restrict only on the lambda (functor + // object) function, that is called directly from a kernel (i.e. the one + // passed to the parallel_for function). Emit a warning and ignore all + // other cases. + if (ParentFD == SYCLKernel) { + Attrs.insert(A); + } else { + SemaRef.Diag(A->getLocation(), diag::warn_attribute_ignored) << A; + FD->dropAttr(); + } + } // TODO: vec_len_hint should be handled here @@ -436,7 +450,7 @@ class MarkDeviceFunction : public RecursiveASTVisitor { if (auto *Callee = dyn_cast(CI->getDecl())) { Callee = Callee->getCanonicalDecl(); if (!Visited.count(Callee)) - WorkList.push_back(Callee); + WorkList.push_back({Callee, FD}); } } } @@ -1296,6 +1310,10 @@ void Sema::MarkDevice(void) { } break; } + case attr::Kind::SYCLIntelKernelArgsRestrict: { + SYCLKernel->addAttr(A); + break; + } // TODO: vec_len_hint should be handled here default: // Seeing this means that CollectPossibleKernelAttributes was diff --git a/clang/lib/Sema/SemaType.cpp b/clang/lib/Sema/SemaType.cpp index ee1b869c74366..85753842e0cf3 100644 --- a/clang/lib/Sema/SemaType.cpp +++ b/clang/lib/Sema/SemaType.cpp @@ -208,6 +208,11 @@ namespace { return chunkIndex == declarator.getNumTypeObjects(); } + bool isProcessingLambdaExpr() const { + return declarator.isFunctionDeclarator() && + declarator.getContext() == DeclaratorContext::LambdaExprContext; + } + unsigned getCurrentChunkIndex() const { return chunkIndex; } @@ -7582,7 +7587,8 @@ static void processTypeAttrs(TypeProcessingState &state, QualType &type, switch (attr.getKind()) { default: // A C++11 attribute on a declarator chunk must appertain to a type. - if (attr.isCXX11Attribute() && TAL == TAL_DeclChunk) { + if (attr.isCXX11Attribute() && TAL == TAL_DeclChunk && + (!state.isProcessingLambdaExpr() || !attr.isAllowedOnLambdas())) { state.getSema().Diag(attr.getLoc(), diag::err_attribute_not_type_attr) << attr; attr.setUsedAsTypeAttr(); diff --git a/clang/test/CodeGenSYCL/intel-restrict.cpp b/clang/test/CodeGenSYCL/intel-restrict.cpp new file mode 100644 index 0000000000000..da2b0102e7ccd --- /dev/null +++ b/clang/test/CodeGenSYCL/intel-restrict.cpp @@ -0,0 +1,57 @@ +// RUN: %clang %s -S -emit-llvm --sycl -o - | FileCheck %s + +#include "CL/sycl.hpp" + +constexpr auto sycl_read_write = cl::sycl::access::mode::read_write; +constexpr auto sycl_global_buffer = cl::sycl::access::target::global_buffer; + +template +struct foostr { + Acc1Ty A; + Acc2Ty B; + foostr(Acc1Ty A, Acc2Ty B): A(A), B(B) {} + [[intel::kernel_args_restrict]] + void operator()() { + A[0] = B[0]; + } +}; + +int foo(int X) { + int A[] = { 42 }; + int B[] = { 0 }; + { + cl::sycl::queue Q; + cl::sycl::buffer BufA(A, 1); + cl::sycl::buffer BufB(B, 1); + + // CHECK: define {{.*}} spir_kernel {{.*}}kernel_norestrict{{.*}}(i32 addrspace(1)* %{{.*}} i32 addrspace(1)* %{{.*}} + + Q.submit([&](cl::sycl::handler& cgh) { + auto AccA = BufA.get_access(cgh); + auto AccB = BufB.get_access(cgh); + cgh.single_task( + [=]() { + AccB[0] = AccA[0]; + }); + }); + + // CHECK: define {{.*}} spir_kernel {{.*}}kernel_restrict{{.*}}(i32 addrspace(1)* noalias %{{.*}} i32 addrspace(1)* noalias %{{.*}} + Q.submit([&](cl::sycl::handler& cgh) { + auto AccA = BufA.get_access(cgh); + auto AccB = BufB.get_access(cgh); + cgh.single_task( + [=]() [[intel::kernel_args_restrict]] { + AccB[0] = AccA[0]; + }); + }); + + // CHECK: define {{.*}} spir_kernel {{.*}}kernel_restrict_struct{{.*}}(i32 addrspace(1)* noalias %{{.*}} i32 addrspace(1)* noalias %{{.*}} + Q.submit([&](cl::sycl::handler& cgh) { + auto AccA = BufA.get_access(cgh); + auto AccB = BufB.get_access(cgh); + foostr f(AccA, AccB); + cgh.single_task(f); + }); + } + return B[0]; +} diff --git a/clang/test/Misc/pragma-attribute-supported-attributes-list.test b/clang/test/Misc/pragma-attribute-supported-attributes-list.test index b8a3154eb51ef..ab5609c203356 100644 --- a/clang/test/Misc/pragma-attribute-supported-attributes-list.test +++ b/clang/test/Misc/pragma-attribute-supported-attributes-list.test @@ -131,6 +131,7 @@ // CHECK-NEXT: ReturnsTwice (SubjectMatchRule_function) // CHECK-NEXT: SYCLDevice (SubjectMatchRule_function) // CHECK-NEXT: SYCLDeviceIndirectlyCallable (SubjectMatchRule_function) +// CHECK-NEXT: SYCLIntelKernelArgsRestrict (SubjectMatchRule_hasType_functionType) // CHECK-NEXT: SYCLKernel (SubjectMatchRule_function) // CHECK-NEXT: ScopedLockable (SubjectMatchRule_record) // CHECK-NEXT: Section (SubjectMatchRule_function, SubjectMatchRule_variable_is_global, SubjectMatchRule_objc_method, SubjectMatchRule_objc_property) diff --git a/clang/test/SemaSYCL/intel-restrict.cpp b/clang/test/SemaSYCL/intel-restrict.cpp new file mode 100644 index 0000000000000..5349ee2d0b8e5 --- /dev/null +++ b/clang/test/SemaSYCL/intel-restrict.cpp @@ -0,0 +1,33 @@ +// RUN: %clang %s -fsyntax-only --sycl -DCHECKDIAG -Xclang -verify +// RUN: %clang %s -fsyntax-only -Xclang -ast-dump --sycl | FileCheck %s + +[[intel::kernel_args_restrict]] // expected-warning{{'kernel_args_restrict' attribute ignored}} +void func_ignore() {} + + +struct Functor { + [[intel::kernel_args_restrict]] + void operator()() {} +}; + +template +__attribute__((sycl_kernel)) void kernel(Func kernelFunc) { + kernelFunc(); +} + +int main() { + // CHECK-LABEL: FunctionDecl {{.*}} _ZTSZ4mainE12test_kernel1 + // CHECK: SYCLIntelKernelArgsRestrictAttr + kernel( + Functor()); + + // CHECK-LABEL: FunctionDecl {{.*}} _ZTSZ4mainE12test_kernel2 + // CHECK: SYCLIntelKernelArgsRestrictAttr + kernel( + []() [[intel::kernel_args_restrict]] {}); + + // CHECK-LABEL: FunctionDecl {{.*}} _ZTSZ4mainE12test_kernel3 + // CHECK-NOT: SYCLIntelKernelArgsRestrictAttr + kernel( + []() {func_ignore();}); +} From 487ab078ad43620feb7c7805f82c0e689e34cdef Mon Sep 17 00:00:00 2001 From: Andrew Savonichev Date: Fri, 18 Oct 2019 18:44:08 +0300 Subject: [PATCH 2/5] Do not set noalias attribute for non-pointer arguments Signed-off-by: Andrew Savonichev --- clang/lib/CodeGen/CGCall.cpp | 3 ++- clang/test/CodeGenSYCL/intel-restrict.cpp | 11 +++++++++++ 2 files changed, 13 insertions(+), 1 deletion(-) diff --git a/clang/lib/CodeGen/CGCall.cpp b/clang/lib/CodeGen/CGCall.cpp index d33576ba204eb..8ee08d670e4c5 100644 --- a/clang/lib/CodeGen/CGCall.cpp +++ b/clang/lib/CodeGen/CGCall.cpp @@ -2414,7 +2414,8 @@ void CodeGenFunction::EmitFunctionProlog(const CGFunctionInfo &FI, if (Arg->getType().isRestrictQualified() || (CurCodeDecl && - CurCodeDecl->hasAttr())) + CurCodeDecl->hasAttr() && + Arg->getType()->isPointerType())) AI->addAttr(llvm::Attribute::NoAlias); // LLVM expects swifterror parameters to be used in very restricted diff --git a/clang/test/CodeGenSYCL/intel-restrict.cpp b/clang/test/CodeGenSYCL/intel-restrict.cpp index da2b0102e7ccd..71b5a4833890f 100644 --- a/clang/test/CodeGenSYCL/intel-restrict.cpp +++ b/clang/test/CodeGenSYCL/intel-restrict.cpp @@ -52,6 +52,17 @@ int foo(int X) { foostr f(AccA, AccB); cgh.single_task(f); }); + + // CHECK: define {{.*}} spir_kernel {{.*}}kernel_restrict_other_params{{.*}}(i32 addrspace(1)* noalias %{{.*}} i32 addrspace(1)* noalias %{{.*}}, i32 %_arg_9) + int num = 42; + Q.submit([&](cl::sycl::handler& cgh) { + auto AccA = BufA.get_access(cgh); + auto AccB = BufB.get_access(cgh); + cgh.single_task( + [=]() [[intel::kernel_args_restrict]] { + AccB[0] = AccA[0] = num; + }); + }); } return B[0]; } From 56b53c0c2558999e0b6086838d6956cd4ff1e73a Mon Sep 17 00:00:00 2001 From: Andrew Savonichev Date: Mon, 21 Oct 2019 15:03:59 +0300 Subject: [PATCH 3/5] Change wording in the attribute documentation Signed-off-by: Andrew Savonichev --- clang/include/clang/Basic/AttrDocs.td | 8 ++++---- 1 file changed, 4 insertions(+), 4 deletions(-) diff --git a/clang/include/clang/Basic/AttrDocs.td b/clang/include/clang/Basic/AttrDocs.td index 75d2e4589be05..14ca103987c72 100644 --- a/clang/include/clang/Basic/AttrDocs.td +++ b/clang/include/clang/Basic/AttrDocs.td @@ -1849,11 +1849,11 @@ 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, propagation of the attribute to any caller(s), including up to a kernel -boundary, is implementation defined and not guaranteed through this -extension. The attribute forms an unchecked assertion, in that implementations +kernel, the attribute is ignored and it is not propagated to a kernel. + +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 -``inte::_kernel_args_restrict`` to a kernel, but there is in fact aliasing +``intel::_kernel_args_restrict`` to a kernel, but there is in fact aliasing between kernel pointer arguments at runtime, the behavior is undefined. The attribute-token ``intel::kernel_args_restrict`` shall appear at most once in From 862433c372933bd1f83c99251e9afa4fdd7d0ab1 Mon Sep 17 00:00:00 2001 From: Andrew Savonichev Date: Mon, 21 Oct 2019 15:08:03 +0300 Subject: [PATCH 4/5] Reword functor -> function object Signed-off-by: Andrew Savonichev --- clang/include/clang/Basic/AttrDocs.td | 12 ++++++------ clang/lib/Sema/SemaSYCL.cpp | 2 +- clang/test/SemaSYCL/intel-restrict.cpp | 4 ++-- 3 files changed, 9 insertions(+), 9 deletions(-) diff --git a/clang/include/clang/Basic/AttrDocs.td b/clang/include/clang/Basic/AttrDocs.td index 14ca103987c72..cee65e7797063 100644 --- a/clang/include/clang/Basic/AttrDocs.td +++ b/clang/include/clang/Basic/AttrDocs.td @@ -1841,12 +1841,12 @@ def SYCLIntelKernelArgsRestrictDocs : Documentation { let Content = [{ The attribute ``intel::kernel_args_restrict`` is legal on device functions, and can be ignored on non-device functions. When applied to a function, lambda, or -function call operator (of a functor), the attribute is a hint to the compiler -equivalent to specifying the C99 restrict attribute on all pointer arguments or -the pointer member of any accessors, which are a function argument, lambda -capture, or functor member, of the callable to which the attribute was -applied. This effect is equivalent to annotating restrict on **all** kernel -pointer arguments in an OpenCL or SPIR-V kernel. +function call operator (of a function object), the attribute is a hint to the +compiler equivalent to specifying the C99 restrict attribute on all pointer +arguments or the pointer member of any accessors, which are a function argument, +lambda capture, or function object member, of the callable to which the +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 ignored and it is not propagated to a kernel. diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index a97ffeb3af574..70f2a44a05245 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -428,7 +428,7 @@ class MarkDeviceFunction : public RecursiveASTVisitor { else if (auto *A = FD->getAttr()) Attrs.insert(A); else if (auto *A = FD->getAttr()) { - // Allow the intel::kernel_args_restrict only on the lambda (functor + // Allow the intel::kernel_args_restrict only on the lambda (function // object) function, that is called directly from a kernel (i.e. the one // passed to the parallel_for function). Emit a warning and ignore all // other cases. diff --git a/clang/test/SemaSYCL/intel-restrict.cpp b/clang/test/SemaSYCL/intel-restrict.cpp index 5349ee2d0b8e5..4de7cd2750687 100644 --- a/clang/test/SemaSYCL/intel-restrict.cpp +++ b/clang/test/SemaSYCL/intel-restrict.cpp @@ -5,7 +5,7 @@ void func_ignore() {} -struct Functor { +struct FuncObj { [[intel::kernel_args_restrict]] void operator()() {} }; @@ -19,7 +19,7 @@ int main() { // CHECK-LABEL: FunctionDecl {{.*}} _ZTSZ4mainE12test_kernel1 // CHECK: SYCLIntelKernelArgsRestrictAttr kernel( - Functor()); + FuncObj()); // CHECK-LABEL: FunctionDecl {{.*}} _ZTSZ4mainE12test_kernel2 // CHECK: SYCLIntelKernelArgsRestrictAttr From 5adac8746c0717c26d2e890733e43165d19a375a Mon Sep 17 00:00:00 2001 From: Andrew Savonichev Date: Mon, 21 Oct 2019 15:30:56 +0300 Subject: [PATCH 5/5] Change subject from FunctionLike to Function; add tests for subject FunctionLike also applies to function pointers, which is probably not what we want. Signed-off-by: Andrew Savonichev --- clang/include/clang/Basic/Attr.td | 2 +- .../test/Misc/pragma-attribute-supported-attributes-list.test | 2 +- clang/test/SemaSYCL/intel-restrict.cpp | 4 +++- 3 files changed, 5 insertions(+), 3 deletions(-) diff --git a/clang/include/clang/Basic/Attr.td b/clang/include/clang/Basic/Attr.td index 7f8b4dc1a9980..6e08f7fb4daf5 100644 --- a/clang/include/clang/Basic/Attr.td +++ b/clang/include/clang/Basic/Attr.td @@ -1093,7 +1093,7 @@ def SYCLDeviceIndirectlyCallable : InheritableAttr { } def SYCLIntelKernelArgsRestrict : InheritableAttr { let Spellings = [ CXX11<"intel", "kernel_args_restrict"> ]; - let Subjects = SubjectList<[ FunctionLike ], ErrorDiag>; + let Subjects = SubjectList<[Function], ErrorDiag>; let LangOpts = [ SYCLIsDevice, SYCLIsHost ]; let Documentation = [ SYCLIntelKernelArgsRestrictDocs ]; } diff --git a/clang/test/Misc/pragma-attribute-supported-attributes-list.test b/clang/test/Misc/pragma-attribute-supported-attributes-list.test index ab5609c203356..653c6a246a908 100644 --- a/clang/test/Misc/pragma-attribute-supported-attributes-list.test +++ b/clang/test/Misc/pragma-attribute-supported-attributes-list.test @@ -131,7 +131,7 @@ // CHECK-NEXT: ReturnsTwice (SubjectMatchRule_function) // CHECK-NEXT: SYCLDevice (SubjectMatchRule_function) // CHECK-NEXT: SYCLDeviceIndirectlyCallable (SubjectMatchRule_function) -// CHECK-NEXT: SYCLIntelKernelArgsRestrict (SubjectMatchRule_hasType_functionType) +// CHECK-NEXT: SYCLIntelKernelArgsRestrict (SubjectMatchRule_function) // CHECK-NEXT: SYCLKernel (SubjectMatchRule_function) // CHECK-NEXT: ScopedLockable (SubjectMatchRule_record) // CHECK-NEXT: Section (SubjectMatchRule_function, SubjectMatchRule_variable_is_global, SubjectMatchRule_objc_method, SubjectMatchRule_objc_property) diff --git a/clang/test/SemaSYCL/intel-restrict.cpp b/clang/test/SemaSYCL/intel-restrict.cpp index 4de7cd2750687..20fe10c9596f0 100644 --- a/clang/test/SemaSYCL/intel-restrict.cpp +++ b/clang/test/SemaSYCL/intel-restrict.cpp @@ -4,7 +4,6 @@ [[intel::kernel_args_restrict]] // expected-warning{{'kernel_args_restrict' attribute ignored}} void func_ignore() {} - struct FuncObj { [[intel::kernel_args_restrict]] void operator()() {} @@ -13,6 +12,9 @@ struct FuncObj { template __attribute__((sycl_kernel)) void kernel(Func kernelFunc) { kernelFunc(); +#ifdef CHECKDIAG + [[intel::kernel_args_restrict]] int invalid = 42; // expected-error{{'kernel_args_restrict' attribute only applies to functions}} +#endif } int main() {