[SYCL] Add clang implementation for accessor property no_alias (#3452)

no_alias is a compile-time-constant accessor property.  
It exists in cl::sycl::ONEAPI::property scope. 
When applied the LLVM IR should contain the noalias parameter attribute
on the kernel argument corresponding to that accessor.
This patch uses clang's RestrictAttr to emit `noalias` i.e., llvm::Attribute::NoAlias.

Author: schittir

clang/lib/CodeGen/CGCall.cpp

@@ -2771,7 +2771,8 @@ void CodeGenFunction::EmitFunctionProlog(const CGFunctionInfo &FI,
         if (Arg->getType().isRestrictQualified() ||
             (CurCodeDecl &&
              CurCodeDecl->hasAttr<SYCLIntelKernelArgsRestrictAttr>() &&
-             Arg->getType()->isPointerType()))
+             Arg->getType()->isPointerType()) ||
+            (Arg->hasAttr<RestrictAttr>() && Arg->getType()->isPointerType()))
           AI->addAttr(llvm::Attribute::NoAlias);
       }
 

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<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(QualType Ty) {
   std::array<DeclContextDesc, 6> Scopes = {
       Util::MakeDeclContextDesc(Decl::Kind::Namespace, "cl"),

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,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<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::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::INTEL::property::buffer_location::instance<1>>>
+      accessorB;
+
+  cl::sycl::kernel_single_task<class kernel_function1>(
+      [=]() {
+        accessorA.use();
+      });
+
+  cl::sycl::kernel_single_task<class kernel_function2>(
+      [=]() {
+        accessorB.use();
+      });
+  return 0;
+}