[SYCL][Docs] Rename restrict property to unaliased

clang/lib/CodeGen/CGCall.cpp

@@ -3030,7 +3030,7 @@ static bool hasSYCLRestrictPropertyIRAttr(const VarDecl *Arg,
   return std::any_of(
       NameValuePairs.begin(), NameValuePairs.end(),
       [](const std::pair<std::string, std::string> &NameValuePair) {
-        return NameValuePair.first == "sycl-restrict";
+        return NameValuePair.first == "sycl-unaliased";
       });
 }
 

clang/test/CodeGenSYCL/sycl_unaliased_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-unaliased", 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_nounaliased>([a, b, c]() { c[0] = a[0] + b[0]; });
+    // CHECK-DAG: define {{.*}}spir_kernel {{.*}}kernel_nounaliased(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_unaliased1>([a, b, c]() { c[0] = a[0] + b[0]; });
+    // CHECK-DAG: define {{.*}}spir_kernel {{.*}}kernel_unaliased1(ptr addrspace(1) noalias noundef align 4 "sycl-unaliased" %{{.*}}, ptr addrspace(1) noundef align 4 %{{.*}}, ptr addrspace(1) noundef align 4 %{{.*}})
+  }
+  {
+    int *a;
+    AnnotatedIntPtr b;
+    int *c;
+    kernel<class kernel_unaliased2>([a, b, c]() { c[0] = a[0] + b[0]; });
+    // CHECK-DAG: define {{.*}}spir_kernel {{.*}}kernel_unaliased2(ptr addrspace(1) noundef align 4 %{{.*}}, ptr addrspace(1) noalias noundef align 4 "sycl-unaliased" %{{.*}}, ptr addrspace(1) noundef align 4 %{{.*}})
+  }
+  {
+    int *a;
+    int *b;
+    AnnotatedIntPtr c;
+    kernel<class kernel_unaliased3>([a, b, c]() { c[0] = a[0] + b[0]; });
+    // CHECK-DAG: define {{.*}}spir_kernel {{.*}}kernel_unaliased3(ptr addrspace(1) noundef align 4 %{{.*}}, ptr addrspace(1) noundef align 4 %{{.*}}, ptr addrspace(1) noalias noundef align 4 "sycl-unaliased" %{{.*}})
+  }
+  {
+    AnnotatedIntPtr a;
+    AnnotatedIntPtr b;
+    int *c;
+    kernel<class kernel_unaliased4>([a, b, c]() { c[0] = a[0] + b[0]; });
+    // CHECK-DAG: define {{.*}}spir_kernel {{.*}}kernel_unaliased4(ptr addrspace(1) noalias noundef align 4 "sycl-unaliased" %{{.*}}, ptr addrspace(1) noalias noundef align 4 "sycl-unaliased" %{{.*}}, ptr addrspace(1) noundef align 4 %{{.*}})
+  }
+  {
+    AnnotatedIntPtr a;
+    int *b;
+    AnnotatedIntPtr c;
+    kernel<class kernel_unaliased5>([a, b, c]() { c[0] = a[0] + b[0]; });
+    // CHECK-DAG: define {{.*}}spir_kernel {{.*}}kernel_unaliased5(ptr addrspace(1) noalias noundef align 4 "sycl-unaliased" %{{.*}}, ptr addrspace(1) noundef align 4 %{{.*}}, ptr addrspace(1) noalias noundef align 4 "sycl-unaliased" %{{.*}})
+  }
+  {
+    int *a;
+    AnnotatedIntPtr b;
+    AnnotatedIntPtr c;
+    kernel<class kernel_unaliased6>([a, b, c]() { c[0] = a[0] + b[0]; });
+    // CHECK-DAG: define {{.*}}spir_kernel {{.*}}kernel_unaliased6(ptr addrspace(1) noundef align 4 %{{.*}}, ptr addrspace(1) noalias noundef align 4 "sycl-unaliased" %{{.*}}, ptr addrspace(1) noalias noundef align 4 "sycl-unaliased" %{{.*}})
+  }
+  {
+    AnnotatedIntPtr a;
+    AnnotatedIntPtr b;
+    AnnotatedIntPtr c;
+    kernel<class kernel_unaliased7>([a, b, c]() { c[0] = a[0] + b[0]; });
+    // CHECK-DAG: define {{.*}}spir_kernel {{.*}}kernel_unaliased7(ptr addrspace(1) noalias noundef align 4 "sycl-unaliased" %{{.*}}, ptr addrspace(1) noalias noundef align 4 "sycl-unaliased" %{{.*}}, ptr addrspace(1) noalias noundef align 4 "sycl-unaliased" %{{.*}})
+  }
+}

sycl/doc/extensions/experimental/sycl_ext_oneapi_kernel_arg_properties.asciidoc

@@ -88,29 +88,29 @@ implementation supports.
 |Initial version of this extension.
 |===
 
-=== `restrict` property
+=== `unaliased` property
 
-The `restrict` property defined here is only meaningful on the kernel arguments
+The `unaliased` property defined here is only meaningful on the kernel arguments
 when the kernel argument is a pointer type. It is ignored for other types.
 
 This property is not meaningful within the kernel body.
 
 
 ```c++
 namespace sycl::ext::oneapi::experimental {
-struct restrict_key {
-  using value_t = property_value<restrict_key>;
+struct unaliased_key {
+  using value_t = property_value<unaliased_key>;
 };
 
-inline constexpr restrict_key::value_t restrict;
+inline constexpr unaliased_key::value_t unaliased;
 
 template<typename T, typename PropertyListT>
 struct is_property_key_of<
-  restrict_key, annotated_ptr<T, PropertyListT>> : std::true_type {};
+  unaliased_key, annotated_ptr<T, PropertyListT>> : std::true_type {};
 
 template<typename T, typename PropertyListT>
 struct is_property_key_of<
-  restrict_key, annotated_arg<T, PropertyListT>> : std::true_type {};
+  unaliased_key, annotated_arg<T, PropertyListT>> : std::true_type {};
 } // namespace sycl::ext::oneapi::experimental
 ```
 === `alignment` property
@@ -152,7 +152,7 @@ struct is_property_key_of<
 a|
 [source,c++]
 ----
-restrict
+unaliased
 ----
 a|
 This is an assertion by the application that the pointer kernel arguments marked
@@ -195,8 +195,8 @@ using sycl::ext::oneapi::experimental;
   int* ptr_b = ...;
 
   // Add properties
-  auto arg_a = annotated_ptr(ptr_a, properties{restrict, alignment<32>});
-  auto arg_n = annotated_arg(ptr_b, properties{restrict});
+  auto arg_a = annotated_ptr(ptr_a, properties{unaliased, alignment<32>});
+  auto arg_n = annotated_arg(ptr_b, properties{unaliased});
   ...
 
   q.single_task([=] {

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

@@ -66,12 +66,12 @@ 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>;
+struct unaliased_key
+    : detail::compile_time_property_key<detail::PropKind::Unaliased> {
+  using value_t = property_value<unaliased_key>;
 };
 
-inline constexpr restrict_key::value_t restrict;
+inline constexpr unaliased_key::value_t unaliased;
 
 struct alignment_key
     : detail::compile_time_property_key<detail::PropKind::Alignment> {
@@ -82,15 +82,15 @@ 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>
+struct is_valid_property<T, unaliased_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>>
+struct is_property_key_of<unaliased_key, annotated_ptr<T, PropertyListT>>
     : std::true_type {};
 
 template <typename T, typename PropertyListT>
@@ -102,7 +102,7 @@ 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>>
+struct is_property_key_of<unaliased_key, annotated_arg<T, PropertyListT>>
     : std::true_type {};
 
 template <> struct propagateToPtrAnnotation<alignment_key> : std::true_type {};
@@ -113,8 +113,8 @@ 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";
+template <> struct PropertyMetaInfo<unaliased_key::value_t> {
+  static constexpr const char *name = "sycl-unaliased";
   static constexpr std::nullptr_t value = nullptr;
 };
 

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

@@ -222,7 +222,7 @@ enum PropKind : uint32_t {
   Deterministic = 77,
   InitializeToIdentity = 78,
   WorkGroupScratchSize = 79,
-  Restrict = 80,
+  Unaliased = 80,
   EventMode = 81,
   NativeLocalBlockIO = 82,
   // PropKindSize must always be the last value.

sycl/test-e2e/Annotated_arg_ptr/annotated_arg_restrict.cpp → sycl/test-e2e/Annotated_arg_ptr/annotated_arg_unaliased.cpp

@@ -2,7 +2,7 @@
 // RUN: %{run} %t.out
 // REQUIRES: aspect-usm_shared_allocations
 
-// Checks that restrict annotated_arg works in device code.
+// Checks that unaliased annotated_arg works in device code.
 
 #include <sycl/detail/core.hpp>
 #include <sycl/ext/oneapi/experimental/annotated_arg/annotated_arg.hpp>
@@ -15,7 +15,7 @@ int main() {
 
   int *Ptr = sycl::malloc_shared<int>(1, Q);
   syclexp::annotated_arg<int *,
-                         decltype(syclexp::properties(syclexp::restrict))>
+                         decltype(syclexp::properties(syclexp::unaliased))>
       AnnotArg{Ptr};
   Q.submit([&](sycl::handler &CGH) {
      CGH.single_task([=]() { *AnnotArg = 42; });

sycl/test-e2e/Annotated_arg_ptr/annotated_ptr_restrict.cpp → sycl/test-e2e/Annotated_arg_ptr/annotated_ptr_unaliased.cpp

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

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

@@ -9,7 +9,7 @@ int main() {
 
   auto Ptr = sycl::malloc_shared<int>(1, Q);
   syclexp::annotated_arg<int *,
-                         decltype(syclexp::properties(syclexp::restrict))>
+                         decltype(syclexp::properties(syclexp::unaliased))>
       AnnotArg{Ptr};
   Q.submit([&](sycl::handler &CGH) {
      CGH.single_task([=]() { *AnnotArg = 42; });
@@ -19,4 +19,4 @@ int main() {
   return 0;
 }
 
-// CHECK-IR: spir_kernel void @_ZTSZZ4mainENKUlRN4sycl3_V17handlerEE_clES2_EUlvE_(ptr addrspace(1) noalias noundef align 4 "sycl-restrict"
+// CHECK-IR: spir_kernel void @_ZTSZZ4mainENKUlRN4sycl3_V17handlerEE_clES2_EUlvE_(ptr addrspace(1) noalias noundef align 4 "sycl-unaliased"

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

@@ -8,7 +8,7 @@ int main() {
   sycl::queue Q;
 
   auto Ptr = sycl::malloc_shared<int>(1, Q);
-  syclexp::annotated_ptr<int, decltype(syclexp::properties(syclexp::restrict))>
+  syclexp::annotated_ptr<int, decltype(syclexp::properties(syclexp::unaliased))>
       AnnotPtr{Ptr};
   Q.submit([&](sycl::handler &CGH) {
      CGH.single_task([=]() { *AnnotPtr = 42; });
@@ -18,4 +18,4 @@ int main() {
   return 0;
 }
 
-// CHECK-IR: spir_kernel void @_ZTSZZ4mainENKUlRN4sycl3_V17handlerEE_clES2_EUlvE_(ptr addrspace(1) noalias noundef align 4 "sycl-restrict"
+// CHECK-IR: spir_kernel void @_ZTSZZ4mainENKUlRN4sycl3_V17handlerEE_clES2_EUlvE_(ptr addrspace(1) noalias noundef align 4 "sycl-unaliased"