[SYCL] Add clang implementation for accessor property no_alias

clang/lib/CodeGen/CodeGenModule.cpp

@@ -1580,6 +1580,9 @@ void CodeGenModule::GenOpenCLArgMetadata(llvm::Function *Fn,
   // MDNode for the intel_buffer_location attribute.
   SmallVector<llvm::Metadata *, 8> argSYCLBufferLocationAttr;
 
+  // MDNode for accessor no_alias property.
+  SmallVector<llvm::Metadata *, 8> argAccessorNoAliasPropertyAttr;
+
   // MDNode for listing ESIMD kernel pointer arguments originating from
   // accessors
   SmallVector<llvm::Metadata *, 8> argESIMDAccPtrs;
@@ -1687,17 +1690,25 @@ void CodeGenModule::GenOpenCLArgMetadata(llvm::Function *Fn,
                     SYCLBufferLocationAttr->getLocationID()))
               : llvm::ConstantAsMetadata::get(CGF->Builder.getInt32(-1)));
 
+      if (parm->hasAttr<RestrictAttr>())
+        argAccessorNoAliasPropertyAttr.push_back(
+            llvm::ConstantAsMetadata::get(CGF->Builder.getInt1(true)));
+
       if (FD->hasAttr<SYCLSimdAttr>())
         argESIMDAccPtrs.push_back(llvm::ConstantAsMetadata::get(
             CGF->Builder.getInt1(parm->hasAttr<SYCLSimdAccessorPtrAttr>())));
     }
 
   bool IsEsimdFunction = FD && FD->hasAttr<SYCLSimdAttr>();
 
-  if (LangOpts.SYCLIsDevice && !IsEsimdFunction)
+  if (LangOpts.SYCLIsDevice && !IsEsimdFunction) {
+
+    Fn->setMetadata("noalias", llvm::MDNode::get(
+                                   VMContext, argAccessorNoAliasPropertyAttr));
+
     Fn->setMetadata("kernel_arg_buffer_location",
                     llvm::MDNode::get(VMContext, argSYCLBufferLocationAttr));
-  else {
+  } else {
     Fn->setMetadata("kernel_arg_addr_space",
                     llvm::MDNode::get(VMContext, addressQuals));
     Fn->setMetadata("kernel_arg_access_qual",

clang/lib/Sema/SemaSYCL.cpp

@@ -90,6 +90,10 @@ class Util {
   /// accessor_property_list class.
   static bool isAccessorPropertyListType(const 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(const QualType &Ty);
@@ -1742,11 +1746,21 @@ 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) {
+    if (PropTy.isRestrictQualified()) {
+      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,
@@ -4313,6 +4327,18 @@ bool Util::isSyclKernelHandlerType(const QualType &Ty) {
   return matchQualifiedTypeName(Ty, Scopes);
 }
 
+bool Util::isSyclAccessorNoAliasPropertyType(QualType Ty) {
+  std::array<DeclContextDesc, 6> 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(const QualType &Ty) {
   const StringRef &PropertyName = "buffer_location";
   const StringRef &InstanceName = "instance";

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 <bool> class instance {};
+};
+} // namespace property
+} // namespace ONEAPI
+
 namespace ONEAPI {
 template <typename... properties>
 class accessor_property_list {};

clang/test/CodeGenSYCL/accessor_no_alias_property.cpp

@@ -0,0 +1,19 @@
+// RUN: %clang_cc1 -fsycl-is-device -triple spir64-unknown-unknown-sycldevice -emit-llvm %s -o - | FileCheck %s
+
+// CHECK: define {{.*}}spir_kernel void @_ZTSZ4mainE15kernel_function{{.*}} !noalias
+
+#include "Inputs/sycl.hpp"
+
+int main() {
+  cl::sycl::accessor<int, 1, cl::sycl::access::mode::read_write,
+                     cl::sycl::access::target::global_buffer,
+                     cl::sycl::access::placeholder::false_t,
+                     cl::sycl::ONEAPI::accessor_property_list<
+                         cl::sycl::ONEAPI::property::no_alias::instance<true>>>
+      accessorA;
+  cl::sycl::kernel_single_task<class kernel_function>(
+      [=]() {
+        accessorA.use();
+      });
+  return 0;
+}

clang/test/CodeGenSYCL/disable_loop_pipelining.cpp

@@ -29,8 +29,7 @@ int main() {
   return 0;
 }
 
-// CHECK: define dso_local spir_kernel void @"{{.*}}test_kernel1"() #0 !kernel_arg_buffer_location ![[NUM4:[0-9]+]] !disable_loop_pipelining ![[NUM5:[0-9]+]]
-// CHECK: define dso_local spir_kernel void @"{{.*}}test_kernel2"() #0 !kernel_arg_buffer_location ![[NUM4]]
-// CHECK: define dso_local spir_kernel void @"{{.*}}test_kernel3"() #0 !kernel_arg_buffer_location ![[NUM4]] !disable_loop_pipelining ![[NUM5]]
-// CHECK: ![[NUM4]] = !{}
+// CHECK: define dso_local spir_kernel void @"{{.*}}test_kernel1"() #0 {{.*}} !disable_loop_pipelining ![[NUM5:[0-9]+]]
+// CHECK: define dso_local spir_kernel void @"{{.*}}test_kernel2"() #0 {{.*}}
+// CHECK: define dso_local spir_kernel void @"{{.*}}test_kernel3"() #0 {{.*}} !disable_loop_pipelining ![[NUM5]]
 // CHECK: ![[NUM5]] = !{i32 1}

clang/test/CodeGenSYCL/initiation_interval.cpp

@@ -39,11 +39,10 @@ int main() {
   return 0;
 }
 
-// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name1"() #0 !kernel_arg_buffer_location ![[NUM0:[0-9]+]] !initiation_interval ![[NUM1:[0-9]+]]
-// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name2"() #0 !kernel_arg_buffer_location ![[NUM0]] !initiation_interval ![[NUM42:[0-9]+]]
-// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name3"() #0 !kernel_arg_buffer_location ![[NUM0]] !initiation_interval ![[NUM2:[0-9]+]]
-// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name4"() #0 !kernel_arg_buffer_location ![[NUM0]]
-// CHECK: ![[NUM0]] = !{}
+// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name1"() #0 {{.*}} !initiation_interval ![[NUM1:[0-9]+]]
+// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name2"() #0 {{.*}} !initiation_interval ![[NUM42:[0-9]+]]
+// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name3"() #0 {{.*}} !initiation_interval ![[NUM2:[0-9]+]]
+// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name4"() #0 {{.*}}
 // CHECK: ![[NUM1]] = !{i32 1}
 // CHECK: ![[NUM42]] = !{i32 42}
 // CHECK: ![[NUM2]] = !{i32 2}