Skip to content

Commit 5afb184

Browse files
steffenlarsensteffenlarsen
authored
[SYCL][Docs] Rename restrict property to unaliased (#16814)
This commit renames the `restrict` kernel argument property to `unaliased` to avoid conflicting with common code patterns in C/C++ projects. --------- Signed-off-by: Larsen, Steffen <[email protected]>
1 parent e856440 commit 5afb184

File tree

10 files changed

+109
-109
lines changed

10 files changed

+109
-109
lines changed

‎clang/lib/CodeGen/CGCall.cpp

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -3030,7 +3030,7 @@ static bool hasSYCLRestrictPropertyIRAttr(const VarDecl *Arg,
30303030
return std::any_of(
30313031
NameValuePairs.begin(), NameValuePairs.end(),
30323032
[](const std::pair<std::string, std::string> &NameValuePair) {
3033-
return NameValuePair.first == "sycl-restrict";
3033+
return NameValuePair.first == "sycl-unaliased";
30343034
});
30353035
}
30363036

‎clang/test/CodeGenSYCL/sycl_restrict_property.cpp

Lines changed: 0 additions & 79 deletions
This file was deleted.

‎clang/test/CodeGenSYCL/sycl_unaliased_property.cpp

Lines changed: 79 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,79 @@
1+
// RUN: %clang_cc1 -fsycl-is-device %s -emit-llvm -triple spir64-unknown-unknown -o - | FileCheck %s
2+
3+
struct __attribute__((sycl_special_class))
4+
[[__sycl_detail__::sycl_type(annotated_arg)]]
5+
AnnotatedIntPtr {
6+
void __init([[__sycl_detail__::add_ir_attributes_kernel_parameter(
7+
"sycl-unaliased", nullptr)]]
8+
__attribute__((opencl_global)) int* InPtr) {
9+
Ptr = InPtr;
10+
}
11+
12+
int &operator[](unsigned I) const { return Ptr[I]; }
13+
14+
__attribute__((opencl_global)) int *Ptr;
15+
};
16+
17+
template <typename name, typename Func>
18+
__attribute__((sycl_kernel)) void kernel(const Func &kernelFunc) {
19+
kernelFunc();
20+
}
21+
22+
int main() {
23+
{
24+
int *a;
25+
int *b;
26+
int *c;
27+
kernel<class kernel_nounaliased>([a, b, c]() { c[0] = a[0] + b[0]; });
28+
// 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 %{{.*}})
29+
}
30+
{
31+
AnnotatedIntPtr a;
32+
int *b;
33+
int *c;
34+
kernel<class kernel_unaliased1>([a, b, c]() { c[0] = a[0] + b[0]; });
35+
// 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 %{{.*}})
36+
}
37+
{
38+
int *a;
39+
AnnotatedIntPtr b;
40+
int *c;
41+
kernel<class kernel_unaliased2>([a, b, c]() { c[0] = a[0] + b[0]; });
42+
// 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 %{{.*}})
43+
}
44+
{
45+
int *a;
46+
int *b;
47+
AnnotatedIntPtr c;
48+
kernel<class kernel_unaliased3>([a, b, c]() { c[0] = a[0] + b[0]; });
49+
// 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" %{{.*}})
50+
}
51+
{
52+
AnnotatedIntPtr a;
53+
AnnotatedIntPtr b;
54+
int *c;
55+
kernel<class kernel_unaliased4>([a, b, c]() { c[0] = a[0] + b[0]; });
56+
// 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 %{{.*}})
57+
}
58+
{
59+
AnnotatedIntPtr a;
60+
int *b;
61+
AnnotatedIntPtr c;
62+
kernel<class kernel_unaliased5>([a, b, c]() { c[0] = a[0] + b[0]; });
63+
// 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" %{{.*}})
64+
}
65+
{
66+
int *a;
67+
AnnotatedIntPtr b;
68+
AnnotatedIntPtr c;
69+
kernel<class kernel_unaliased6>([a, b, c]() { c[0] = a[0] + b[0]; });
70+
// 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" %{{.*}})
71+
}
72+
{
73+
AnnotatedIntPtr a;
74+
AnnotatedIntPtr b;
75+
AnnotatedIntPtr c;
76+
kernel<class kernel_unaliased7>([a, b, c]() { c[0] = a[0] + b[0]; });
77+
// 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" %{{.*}})
78+
}
79+
}

‎sycl/doc/extensions/experimental/sycl_ext_oneapi_kernel_arg_properties.asciidoc

Lines changed: 10 additions & 10 deletions
Original file line numberDiff line numberDiff line change
@@ -88,29 +88,29 @@ implementation supports.
8888
|Initial version of this extension.
8989
|===
9090

91-
=== `restrict` property
91+
=== `unaliased` property
9292

93-
The `restrict` property defined here is only meaningful on the kernel arguments
93+
The `unaliased` property defined here is only meaningful on the kernel arguments
9494
when the kernel argument is a pointer type. It is ignored for other types.
9595

9696
This property is not meaningful within the kernel body.
9797

9898

9999
```c++
100100
namespace sycl::ext::oneapi::experimental {
101-
struct restrict_key {
102-
using value_t = property_value<restrict_key>;
101+
struct unaliased_key {
102+
using value_t = property_value<unaliased_key>;
103103
};
104104

105-
inline constexpr restrict_key::value_t restrict;
105+
inline constexpr unaliased_key::value_t unaliased;
106106

107107
template<typename T, typename PropertyListT>
108108
struct is_property_key_of<
109-
restrict_key, annotated_ptr<T, PropertyListT>> : std::true_type {};
109+
unaliased_key, annotated_ptr<T, PropertyListT>> : std::true_type {};
110110

111111
template<typename T, typename PropertyListT>
112112
struct is_property_key_of<
113-
restrict_key, annotated_arg<T, PropertyListT>> : std::true_type {};
113+
unaliased_key, annotated_arg<T, PropertyListT>> : std::true_type {};
114114
} // namespace sycl::ext::oneapi::experimental
115115
```
116116
=== `alignment` property
@@ -152,7 +152,7 @@ struct is_property_key_of<
152152
a|
153153
[source,c++]
154154
----
155-
restrict
155+
unaliased
156156
----
157157
a|
158158
This is an assertion by the application that the pointer kernel arguments marked
@@ -195,8 +195,8 @@ using sycl::ext::oneapi::experimental;
195195
int* ptr_b = ...;
196196

197197
// Add properties
198-
auto arg_a = annotated_ptr(ptr_a, properties{restrict, alignment<32>});
199-
auto arg_n = annotated_arg(ptr_b, properties{restrict});
198+
auto arg_a = annotated_ptr(ptr_a, properties{unaliased, alignment<32>});
199+
auto arg_n = annotated_arg(ptr_b, properties{unaliased});
200200
...
201201

202202
q.single_task([=] {

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

Lines changed: 9 additions & 9 deletions
Original file line numberDiff line numberDiff line change
@@ -66,12 +66,12 @@ struct propagateToPtrAnnotation<property_value<PropKeyT, PropValuesTs...>>
6666
//===----------------------------------------------------------------------===//
6767
// Common properties of annotated_arg/annotated_ptr
6868
//===----------------------------------------------------------------------===//
69-
struct restrict_key
70-
: detail::compile_time_property_key<detail::PropKind::Restrict> {
71-
using value_t = property_value<restrict_key>;
69+
struct unaliased_key
70+
: detail::compile_time_property_key<detail::PropKind::Unaliased> {
71+
using value_t = property_value<unaliased_key>;
7272
};
7373

74-
inline constexpr restrict_key::value_t restrict;
74+
inline constexpr unaliased_key::value_t unaliased;
7575

7676
struct alignment_key
7777
: detail::compile_time_property_key<detail::PropKind::Alignment> {
@@ -82,15 +82,15 @@ struct alignment_key
8282
template <int K> inline constexpr alignment_key::value_t<K> alignment;
8383

8484
template <typename T>
85-
struct is_valid_property<T, restrict_key::value_t>
85+
struct is_valid_property<T, unaliased_key::value_t>
8686
: std::bool_constant<std::is_pointer<T>::value> {};
8787

8888
template <typename T, int W>
8989
struct is_valid_property<T, alignment_key::value_t<W>>
9090
: std::bool_constant<std::is_pointer<T>::value> {};
9191

9292
template <typename T, typename PropertyListT>
93-
struct is_property_key_of<restrict_key, annotated_ptr<T, PropertyListT>>
93+
struct is_property_key_of<unaliased_key, annotated_ptr<T, PropertyListT>>
9494
: std::true_type {};
9595

9696
template <typename T, typename PropertyListT>
@@ -102,7 +102,7 @@ struct is_property_key_of<alignment_key, annotated_arg<T, PropertyListT>>
102102
: std::true_type {};
103103

104104
template <typename T, typename PropertyListT>
105-
struct is_property_key_of<restrict_key, annotated_arg<T, PropertyListT>>
105+
struct is_property_key_of<unaliased_key, annotated_arg<T, PropertyListT>>
106106
: std::true_type {};
107107

108108
template <> struct propagateToPtrAnnotation<alignment_key> : std::true_type {};
@@ -113,8 +113,8 @@ template <int N> struct PropertyMetaInfo<alignment_key::value_t<N>> {
113113
static constexpr int value = N;
114114
};
115115

116-
template <> struct PropertyMetaInfo<restrict_key::value_t> {
117-
static constexpr const char *name = "sycl-restrict";
116+
template <> struct PropertyMetaInfo<unaliased_key::value_t> {
117+
static constexpr const char *name = "sycl-unaliased";
118118
static constexpr std::nullptr_t value = nullptr;
119119
};
120120

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

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -222,7 +222,7 @@ enum PropKind : uint32_t {
222222
Deterministic = 77,
223223
InitializeToIdentity = 78,
224224
WorkGroupScratchSize = 79,
225-
Restrict = 80,
225+
Unaliased = 80,
226226
EventMode = 81,
227227
NativeLocalBlockIO = 82,
228228
// PropKindSize must always be the last value.

‎sycl/test-e2e/Annotated_arg_ptr/annotated_arg_restrict.cpp renamed to ‎sycl/test-e2e/Annotated_arg_ptr/annotated_arg_unaliased.cpp

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -2,7 +2,7 @@
22
// RUN: %{run} %t.out
33
// REQUIRES: aspect-usm_shared_allocations
44

5-
// Checks that restrict annotated_arg works in device code.
5+
// Checks that unaliased annotated_arg works in device code.
66

77
#include <sycl/detail/core.hpp>
88
#include <sycl/ext/oneapi/experimental/annotated_arg/annotated_arg.hpp>
@@ -15,7 +15,7 @@ int main() {
1515

1616
int *Ptr = sycl::malloc_shared<int>(1, Q);
1717
syclexp::annotated_arg<int *,
18-
decltype(syclexp::properties(syclexp::restrict))>
18+
decltype(syclexp::properties(syclexp::unaliased))>
1919
AnnotArg{Ptr};
2020
Q.submit([&](sycl::handler &CGH) {
2121
CGH.single_task([=]() { *AnnotArg = 42; });

‎sycl/test-e2e/Annotated_arg_ptr/annotated_ptr_restrict.cpp renamed to ‎sycl/test-e2e/Annotated_arg_ptr/annotated_ptr_unaliased.cpp

Lines changed: 3 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -2,7 +2,7 @@
22
// RUN: %{run} %t.out
33
// REQUIRES: aspect-usm_shared_allocations
44

5-
// Checks that restrict annotated_ptr works in device code.
5+
// Checks that unaliased annotated_ptr works in device code.
66

77
#include <sycl/detail/core.hpp>
88
#include <sycl/ext/oneapi/experimental/annotated_ptr/annotated_ptr.hpp>
@@ -14,7 +14,7 @@ int main() {
1414
sycl::queue Q;
1515

1616
auto Ptr = sycl::malloc_shared<int>(1, Q);
17-
syclexp::annotated_ptr<int, decltype(syclexp::properties(syclexp::restrict))>
17+
syclexp::annotated_ptr<int, decltype(syclexp::properties(syclexp::unaliased))>
1818
AnnotPtr{Ptr};
1919
Q.submit([&](sycl::handler &CGH) {
2020
CGH.single_task([=]() { *AnnotPtr = 42; });
@@ -25,4 +25,4 @@ int main() {
2525
return 0;
2626
}
2727

28-
// CHECK-IR: spir_kernel void @_ZTSZZ4mainENKUlRN4sycl3_V17handlerEE_clES2_EUlvE_(ptr addrspace(1) noalias noundef align 4 "sycl-restrict" %_arg_AnnotPtr)
28+
// 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 renamed to ‎sycl/test/check_device_code/extensions/annotated_arg/unaliased.cpp

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -9,7 +9,7 @@ int main() {
99

1010
auto Ptr = sycl::malloc_shared<int>(1, Q);
1111
syclexp::annotated_arg<int *,
12-
decltype(syclexp::properties(syclexp::restrict))>
12+
decltype(syclexp::properties(syclexp::unaliased))>
1313
AnnotArg{Ptr};
1414
Q.submit([&](sycl::handler &CGH) {
1515
CGH.single_task([=]() { *AnnotArg = 42; });
@@ -19,4 +19,4 @@ int main() {
1919
return 0;
2020
}
2121

22-
// CHECK-IR: spir_kernel void @_ZTSZZ4mainENKUlRN4sycl3_V17handlerEE_clES2_EUlvE_(ptr addrspace(1) noalias noundef align 4 "sycl-restrict"
22+
// 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 renamed to ‎sycl/test/check_device_code/extensions/annotated_ptr/unaliased.cpp

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -8,7 +8,7 @@ int main() {
88
sycl::queue Q;
99

1010
auto Ptr = sycl::malloc_shared<int>(1, Q);
11-
syclexp::annotated_ptr<int, decltype(syclexp::properties(syclexp::restrict))>
11+
syclexp::annotated_ptr<int, decltype(syclexp::properties(syclexp::unaliased))>
1212
AnnotPtr{Ptr};
1313
Q.submit([&](sycl::handler &CGH) {
1414
CGH.single_task([=]() { *AnnotPtr = 42; });
@@ -18,4 +18,4 @@ int main() {
1818
return 0;
1919
}
2020

21-
// CHECK-IR: spir_kernel void @_ZTSZZ4mainENKUlRN4sycl3_V17handlerEE_clES2_EUlvE_(ptr addrspace(1) noalias noundef align 4 "sycl-restrict"
21+
// CHECK-IR: spir_kernel void @_ZTSZZ4mainENKUlRN4sycl3_V17handlerEE_clES2_EUlvE_(ptr addrspace(1) noalias noundef align 4 "sycl-unaliased"

0 commit comments

Comments
 (0)