diff --git a/clang/lib/CodeGen/CGCall.cpp b/clang/lib/CodeGen/CGCall.cpp index 75b185c9b7230..49562adffdd99 100644 --- a/clang/lib/CodeGen/CGCall.cpp +++ b/clang/lib/CodeGen/CGCall.cpp @@ -2771,7 +2771,8 @@ void CodeGenFunction::EmitFunctionProlog(const CGFunctionInfo &FI, if (Arg->getType().isRestrictQualified() || (CurCodeDecl && CurCodeDecl->hasAttr() && - Arg->getType()->isPointerType())) + Arg->getType()->isPointerType()) || + (Arg->hasAttr() && Arg->getType()->isPointerType())) AI->addAttr(llvm::Attribute::NoAlias); } diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index e32779f942ab6..fa903b92cfa0e 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -101,6 +101,10 @@ class Util { /// accessor_property_list class. static bool isAccessorPropertyListType(QualType Ty); + /// Checks whether given clang type is a full specialization of the SYCL + /// no_alias class. + static bool isSyclAccessorNoAliasPropertyType(QualType Ty); + /// Checks whether given clang type is a full specialization of the SYCL /// buffer_location class. static bool isSyclBufferLocationType(QualType Ty); @@ -1758,11 +1762,19 @@ class SyclKernelDeclCreator : public SyclKernelFieldHandler { for (TemplateArgument::pack_iterator Prop = TemplArg.pack_begin(); Prop != TemplArg.pack_end(); ++Prop) { QualType PropTy = Prop->getAsType(); + if (Util::isSyclAccessorNoAliasPropertyType(PropTy)) + handleNoAliasProperty(Param, PropTy, Loc); if (Util::isSyclBufferLocationType(PropTy)) handleBufferLocationProperty(Param, PropTy, Loc); } } + void handleNoAliasProperty(ParmVarDecl *Param, QualType PropTy, + SourceLocation Loc) { + ASTContext &Ctx = SemaRef.getASTContext(); + Param->addAttr(RestrictAttr::CreateImplicit(Ctx, Loc)); + } + // Obtain an integer value stored in a template parameter of buffer_location // property to pass it to buffer_location kernel attribute void handleBufferLocationProperty(ParmVarDecl *Param, QualType PropTy, @@ -4415,6 +4427,18 @@ bool Util::isSyclKernelHandlerType(QualType Ty) { return matchQualifiedTypeName(Ty, Scopes); } +bool Util::isSyclAccessorNoAliasPropertyType(QualType Ty) { + std::array Scopes = { + Util::DeclContextDesc{Decl::Kind::Namespace, "cl"}, + Util::DeclContextDesc{Decl::Kind::Namespace, "sycl"}, + Util::DeclContextDesc{Decl::Kind::Namespace, "ONEAPI"}, + Util::DeclContextDesc{Decl::Kind::Namespace, "property"}, + Util::DeclContextDesc{Decl::Kind::CXXRecord, "no_alias"}, + Util::DeclContextDesc{Decl::Kind::ClassTemplateSpecialization, + "instance"}}; + return matchQualifiedTypeName(Ty, Scopes); +} + bool Util::isSyclBufferLocationType(QualType Ty) { std::array Scopes = { Util::MakeDeclContextDesc(Decl::Kind::Namespace, "cl"), diff --git a/clang/test/CodeGenSYCL/Inputs/sycl.hpp b/clang/test/CodeGenSYCL/Inputs/sycl.hpp index f041bd230a9ff..6f0b1281b86da 100644 --- a/clang/test/CodeGenSYCL/Inputs/sycl.hpp +++ b/clang/test/CodeGenSYCL/Inputs/sycl.hpp @@ -103,6 +103,15 @@ struct buffer_location { } // namespace property } // namespace INTEL +namespace ONEAPI { +namespace property { +// Compile time known accessor property +struct no_alias { + template class instance {}; +}; +} // namespace property +} // namespace ONEAPI + namespace ONEAPI { template class accessor_property_list {}; diff --git a/clang/test/CodeGenSYCL/accessor_no_alias_property.cpp b/clang/test/CodeGenSYCL/accessor_no_alias_property.cpp new file mode 100644 index 0000000000000..682480cf9569f --- /dev/null +++ b/clang/test/CodeGenSYCL/accessor_no_alias_property.cpp @@ -0,0 +1,36 @@ +// RUN: %clang_cc1 -fsycl-is-device -triple spir64-unknown-unknown-sycldevice -emit-llvm %s -o - | FileCheck %s +// check that noalias parameter attribute is emitted when no_alias accessor property is used +// CHECK: define {{.*}}spir_kernel void @_ZTSZ4mainE16kernel_function1({{.*}} noalias {{.*}} %_arg_, {{.*}}) + +// check that noalias parameter attribute is NOT emitted when it is not used +// CHECK: define {{.*}}spir_kernel void @_ZTSZ4mainE16kernel_function2{{.*}} !kernel_arg_buffer_location +// CHECK-NOT: define {{.*}}spir_kernel void @_ZTSZ4mainE16kernel_function2({{.*}} noalias {{.*}} + +#include "Inputs/sycl.hpp" + +int main() { + cl::sycl::accessor>> + accessorA; + + cl::sycl::accessor>> + accessorB; + + cl::sycl::kernel_single_task( + [=]() { + accessorA.use(); + }); + + cl::sycl::kernel_single_task( + [=]() { + accessorB.use(); + }); + return 0; +}