[SYCL] Implement restrict kernel argument property

clang/lib/CodeGen/CGCall.cpp

@@ -3020,6 +3020,21 @@ namespace {
   };
 }
 
+static bool hasSYCLRestrictPropertyIRAttr(const VarDecl *Arg,
+                                          const ASTContext &Context) {
+  auto *IRAttr = Arg->getAttr<SYCLAddIRAttributesKernelParameterAttr>();
+  if (!IRAttr)
+    return false;
+
+  SmallVector<std::pair<std::string, std::string>, 4> NameValuePairs =
+      IRAttr->getAttributeNameValuePairs(Context);
+  return std::any_of(
+      NameValuePairs.begin(), NameValuePairs.end(),
+      [](const std::pair<std::string, std::string> &NameValuePair) {
+        return NameValuePair.first == "sycl-restrict";
+      });
+}
+
 void CodeGenFunction::EmitFunctionProlog(const CGFunctionInfo &FI,
                                          llvm::Function *Fn,
                                          const FunctionArgList &Args) {
@@ -3244,9 +3259,10 @@ void CodeGenFunction::EmitFunctionProlog(const CGFunctionInfo &FI,
 
         // Set 'noalias' if an argument type has the `restrict` qualifier.
         if (Arg->getType().isRestrictQualified() ||
-            (CurCodeDecl &&
-             CurCodeDecl->hasAttr<SYCLIntelKernelArgsRestrictAttr>() &&
-             Arg->getType()->isPointerType()) ||
+            (Arg->getType()->isPointerType() &&
+             ((CurCodeDecl &&
+               CurCodeDecl->hasAttr<SYCLIntelKernelArgsRestrictAttr>()) ||
+              hasSYCLRestrictPropertyIRAttr(Arg, getContext()))) ||
             (Arg->hasAttr<RestrictAttr>() && Arg->getType()->isPointerType()))
           AI->addAttr(llvm::Attribute::NoAlias);
       }

clang/test/CodeGenSYCL/sycl_restrict_property.cpp

@@ -0,0 +1,79 @@
+// RUN: %clang_cc1 -fsycl-is-device %s -emit-llvm -triple spir64-unknown-unknown -o - | FileCheck %s
+
+struct __attribute__((sycl_special_class))
+      [[__sycl_detail__::sycl_type(annotated_arg)]]
+    AnnotatedIntPtr {
+  void __init([[__sycl_detail__::add_ir_attributes_kernel_parameter(
+                  "sycl-restrict", nullptr)]]
+              __attribute__((opencl_global)) int* InPtr) {
+    Ptr = InPtr;
+  }
+
+  int &operator[](unsigned I) const { return Ptr[I]; }
+
+  __attribute__((opencl_global)) int *Ptr;
+};
+
+template <typename name, typename Func>
+__attribute__((sycl_kernel)) void kernel(const Func &kernelFunc) {
+  kernelFunc();
+}
+
+int main() {
+  {
+    int *a;
+    int *b;
+    int *c;
+    kernel<class kernel_norestrict>([a, b, c]() { c[0] = a[0] + b[0]; });
+    // CHECK-DAG: define {{.*}}spir_kernel {{.*}}kernel_norestrict(ptr addrspace(1) noundef align 4 %{{.*}}, ptr addrspace(1) noundef align 4 %{{.*}}, ptr addrspace(1) noundef align 4 %{{.*}})
+  }
+  {
+    AnnotatedIntPtr a;
+    int *b;
+    int *c;
+    kernel<class kernel_restrict1>([a, b, c]() { c[0] = a[0] + b[0]; });
+    // CHECK-DAG: define {{.*}}spir_kernel {{.*}}kernel_restrict1(ptr addrspace(1) noalias noundef align 4 "sycl-restrict" %{{.*}}, ptr addrspace(1) noundef align 4 %{{.*}}, ptr addrspace(1) noundef align 4 %{{.*}})
+  }
+  {
+    int *a;
+    AnnotatedIntPtr b;
+    int *c;
+    kernel<class kernel_restrict2>([a, b, c]() { c[0] = a[0] + b[0]; });
+    // CHECK-DAG: define {{.*}}spir_kernel {{.*}}kernel_restrict2(ptr addrspace(1) noundef align 4 %{{.*}}, ptr addrspace(1) noalias noundef align 4 "sycl-restrict" %{{.*}}, ptr addrspace(1) noundef align 4 %{{.*}})
+  }
+  {
+    int *a;
+    int *b;
+    AnnotatedIntPtr c;
+    kernel<class kernel_restrict3>([a, b, c]() { c[0] = a[0] + b[0]; });
+    // CHECK-DAG: define {{.*}}spir_kernel {{.*}}kernel_restrict3(ptr addrspace(1) noundef align 4 %{{.*}}, ptr addrspace(1) noundef align 4 %{{.*}}, ptr addrspace(1) noalias noundef align 4 "sycl-restrict" %{{.*}})
+  }
+  {
+    AnnotatedIntPtr a;
+    AnnotatedIntPtr b;
+    int *c;
+    kernel<class kernel_restrict4>([a, b, c]() { c[0] = a[0] + b[0]; });
+    // CHECK-DAG: define {{.*}}spir_kernel {{.*}}kernel_restrict4(ptr addrspace(1) noalias noundef align 4 "sycl-restrict" %{{.*}}, ptr addrspace(1) noalias noundef align 4 "sycl-restrict" %{{.*}}, ptr addrspace(1) noundef align 4 %{{.*}})
+  }
+  {
+    AnnotatedIntPtr a;
+    int *b;
+    AnnotatedIntPtr c;
+    kernel<class kernel_restrict5>([a, b, c]() { c[0] = a[0] + b[0]; });
+    // CHECK-DAG: define {{.*}}spir_kernel {{.*}}kernel_restrict5(ptr addrspace(1) noalias noundef align 4 "sycl-restrict" %{{.*}}, ptr addrspace(1) noundef align 4 %{{.*}}, ptr addrspace(1) noalias noundef align 4 "sycl-restrict" %{{.*}})
+  }
+  {
+    int *a;
+    AnnotatedIntPtr b;
+    AnnotatedIntPtr c;
+    kernel<class kernel_restrict6>([a, b, c]() { c[0] = a[0] + b[0]; });
+    // CHECK-DAG: define {{.*}}spir_kernel {{.*}}kernel_restrict6(ptr addrspace(1) noundef align 4 %{{.*}}, ptr addrspace(1) noalias noundef align 4 "sycl-restrict" %{{.*}}, ptr addrspace(1) noalias noundef align 4 "sycl-restrict" %{{.*}})
+  }
+  {
+    AnnotatedIntPtr a;
+    AnnotatedIntPtr b;
+    AnnotatedIntPtr c;
+    kernel<class kernel_restrict7>([a, b, c]() { c[0] = a[0] + b[0]; });
+    // CHECK-DAG: define {{.*}}spir_kernel {{.*}}kernel_restrict7(ptr addrspace(1) noalias noundef align 4 "sycl-restrict" %{{.*}}, ptr addrspace(1) noalias noundef align 4 "sycl-restrict" %{{.*}}, ptr addrspace(1) noalias noundef align 4 "sycl-restrict" %{{.*}})
+  }
+}

sycl/include/sycl/ext/oneapi/experimental/common_annotated_properties/properties.hpp

@@ -66,6 +66,13 @@ struct propagateToPtrAnnotation<property_value<PropKeyT, PropValuesTs...>>
 //===----------------------------------------------------------------------===//
 //        Common properties of annotated_arg/annotated_ptr
 //===----------------------------------------------------------------------===//
+struct restrict_key
+    : detail::compile_time_property_key<detail::PropKind::Restrict> {
+  using value_t = property_value<restrict_key>;
+};
+
+inline constexpr restrict_key::value_t restrict;
+
 struct alignment_key
     : detail::compile_time_property_key<detail::PropKind::Alignment> {
   template <int K>
@@ -74,10 +81,18 @@ struct alignment_key
 
 template <int K> inline constexpr alignment_key::value_t<K> alignment;
 
+template <typename T>
+struct is_valid_property<T, restrict_key::value_t>
+    : std::bool_constant<std::is_pointer<T>::value> {};
+
 template <typename T, int W>
 struct is_valid_property<T, alignment_key::value_t<W>>
     : std::bool_constant<std::is_pointer<T>::value> {};
 
+template <typename T, typename PropertyListT>
+struct is_property_key_of<restrict_key, annotated_ptr<T, PropertyListT>>
+    : std::true_type {};
+
 template <typename T, typename PropertyListT>
 struct is_property_key_of<alignment_key, annotated_ptr<T, PropertyListT>>
     : std::true_type {};
@@ -86,6 +101,10 @@ template <typename T, typename PropertyListT>
 struct is_property_key_of<alignment_key, annotated_arg<T, PropertyListT>>
     : std::true_type {};
 
+template <typename T, typename PropertyListT>
+struct is_property_key_of<restrict_key, annotated_arg<T, PropertyListT>>
+    : std::true_type {};
+
 template <> struct propagateToPtrAnnotation<alignment_key> : std::true_type {};
 
 namespace detail {
@@ -94,6 +113,11 @@ template <int N> struct PropertyMetaInfo<alignment_key::value_t<N>> {
   static constexpr int value = N;
 };
 
+template <> struct PropertyMetaInfo<restrict_key::value_t> {
+  static constexpr const char *name = "sycl-restrict";
+  static constexpr std::nullptr_t value = nullptr;
+};
+
 } // namespace detail
 
 } // namespace experimental

sycl/include/sycl/ext/oneapi/properties/property.hpp

@@ -222,8 +222,9 @@ enum PropKind : uint32_t {
   Deterministic = 77,
   InitializeToIdentity = 78,
   WorkGroupScratchSize = 79,
+  Restrict = 80,
   // PropKindSize must always be the last value.
-  PropKindSize = 80,
+  PropKindSize = 81,
 };
 
 template <typename PropertyT> struct PropertyToKind {

sycl/test-e2e/Annotated_arg_ptr/annotated_arg_restrict.cpp

@@ -0,0 +1,27 @@
+// RUN: %{build} -o %t.out
+// RUN: %{run} %t.out
+// REQUIRES: aspect-usm_shared_allocations
+
+// Checks that restrict annotated_arg works in device code.
+
+#include <sycl/detail/core.hpp>
+#include <sycl/ext/oneapi/experimental/annotated_arg/annotated_arg.hpp>
+#include <sycl/usm.hpp>
+
+namespace syclexp = sycl::ext::oneapi::experimental;
+
+int main() {
+  sycl::queue Q;
+
+  int *Ptr = sycl::malloc_shared<int>(1, Q);
+  syclexp::annotated_arg<int *,
+                         decltype(syclexp::properties(syclexp::restrict))>
+      AnnotArg{Ptr};
+  Q.submit([&](sycl::handler &CGH) {
+     CGH.single_task([=]() { *AnnotArg = 42; });
+   }).wait();
+  assert(*Ptr == 42);
+  free(Ptr, Q);
+
+  return 0;
+}

sycl/test-e2e/Annotated_arg_ptr/annotated_ptr_restrict.cpp

@@ -0,0 +1,28 @@
+// RUN: %{build} -o %t.out
+// RUN: %{run} %t.out
+// REQUIRES: aspect-usm_shared_allocations
+
+// Checks that restrict annotated_ptr works in device code.
+
+#include <sycl/detail/core.hpp>
+#include <sycl/ext/oneapi/experimental/annotated_ptr/annotated_ptr.hpp>
+#include <sycl/usm.hpp>
+
+namespace syclexp = sycl::ext::oneapi::experimental;
+
+int main() {
+  sycl::queue Q;
+
+  auto Ptr = sycl::malloc_shared<int>(1, Q);
+  syclexp::annotated_ptr<int, decltype(syclexp::properties(syclexp::restrict))>
+      AnnotPtr{Ptr};
+  Q.submit([&](sycl::handler &CGH) {
+     CGH.single_task([=]() { *AnnotPtr = 42; });
+   }).wait();
+  assert(*Ptr == 42);
+  free(Ptr, Q);
+
+  return 0;
+}
+
+// CHECK-IR: spir_kernel void @_ZTSZZ4mainENKUlRN4sycl3_V17handlerEE_clES2_EUlvE_(ptr addrspace(1) noalias noundef align 4 "sycl-restrict" %_arg_AnnotPtr)

sycl/test/check_device_code/extensions/annotated_arg/restrict.cpp

@@ -0,0 +1,22 @@
+// RUN: %clangxx -fsycl-device-only -S -emit-llvm %s -o - | FileCheck %s  --check-prefix CHECK-IR
+
+#include "sycl/sycl.hpp"
+
+namespace syclexp = sycl::ext::oneapi::experimental;
+
+int main() {
+  sycl::queue Q;
+
+  auto Ptr = sycl::malloc_shared<int>(1, Q);
+  syclexp::annotated_arg<int *,
+                         decltype(syclexp::properties(syclexp::restrict))>
+      AnnotArg{Ptr};
+  Q.submit([&](sycl::handler &CGH) {
+     CGH.single_task([=]() { *AnnotArg = 42; });
+   }).wait();
+  free(Ptr, Q);
+
+  return 0;
+}
+
+// CHECK-IR: spir_kernel void @_ZTSZZ4mainENKUlRN4sycl3_V17handlerEE_clES2_EUlvE_(ptr addrspace(1) noalias noundef align 4 "sycl-restrict" %_arg_AnnotArg)

sycl/test/check_device_code/extensions/annotated_ptr/restrict.cpp

@@ -0,0 +1,21 @@
+// RUN: %clangxx -fsycl-device-only -S -emit-llvm %s -o - | FileCheck %s  --check-prefix CHECK-IR
+
+#include "sycl/sycl.hpp"
+
+namespace syclexp = sycl::ext::oneapi::experimental;
+
+int main() {
+  sycl::queue Q;
+
+  auto Ptr = sycl::malloc_shared<int>(1, Q);
+  syclexp::annotated_ptr<int, decltype(syclexp::properties(syclexp::restrict))>
+      AnnotPtr{Ptr};
+  Q.submit([&](sycl::handler &CGH) {
+     CGH.single_task([=]() { *AnnotPtr = 42; });
+   }).wait();
+  free(Ptr, Q);
+
+  return 0;
+}
+
+// CHECK-IR: spir_kernel void @_ZTSZZ4mainENKUlRN4sycl3_V17handlerEE_clES2_EUlvE_(ptr addrspace(1) noalias noundef align 4 "sycl-restrict" %_arg_AnnotPtr)