Skip to content

Commit afa8821

Browse files
steffenlarsensteffenlarsen
authored
[SYCL] Implement restrict kernel argument property (#16090)
This commit implements the restrict property for annotated_arg and annotated_ptr, as specified in sycl_ext_oneapi_kernel_arg_properties. --------- Signed-off-by: Larsen, Steffen <[email protected]>
1 parent 6a9de65 commit afa8821

File tree

8 files changed

+222
-4
lines changed

8 files changed

+222
-4
lines changed

‎clang/lib/CodeGen/CGCall.cpp‎

Lines changed: 19 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -3020,6 +3020,21 @@ namespace {
30203020
};
30213021
}
30223022

3023+
static bool hasSYCLRestrictPropertyIRAttr(const VarDecl *Arg,
3024+
const ASTContext &Context) {
3025+
auto *IRAttr = Arg->getAttr<SYCLAddIRAttributesKernelParameterAttr>();
3026+
if (!IRAttr)
3027+
return false;
3028+
3029+
SmallVector<std::pair<std::string, std::string>, 4> NameValuePairs =
3030+
IRAttr->getAttributeNameValuePairs(Context);
3031+
return std::any_of(
3032+
NameValuePairs.begin(), NameValuePairs.end(),
3033+
[](const std::pair<std::string, std::string> &NameValuePair) {
3034+
return NameValuePair.first == "sycl-restrict";
3035+
});
3036+
}
3037+
30233038
void CodeGenFunction::EmitFunctionProlog(const CGFunctionInfo &FI,
30243039
llvm::Function *Fn,
30253040
const FunctionArgList &Args) {
@@ -3244,9 +3259,10 @@ void CodeGenFunction::EmitFunctionProlog(const CGFunctionInfo &FI,
32443259

32453260
// Set 'noalias' if an argument type has the `restrict` qualifier.
32463261
if (Arg->getType().isRestrictQualified() ||
3247-
(CurCodeDecl &&
3248-
CurCodeDecl->hasAttr<SYCLIntelKernelArgsRestrictAttr>() &&
3249-
Arg->getType()->isPointerType()) ||
3262+
(Arg->getType()->isPointerType() &&
3263+
((CurCodeDecl &&
3264+
CurCodeDecl->hasAttr<SYCLIntelKernelArgsRestrictAttr>()) ||
3265+
hasSYCLRestrictPropertyIRAttr(Arg, getContext()))) ||
32503266
(Arg->hasAttr<RestrictAttr>() && Arg->getType()->isPointerType()))
32513267
AI->addAttr(llvm::Attribute::NoAlias);
32523268
}

‎clang/test/CodeGenSYCL/sycl_restrict_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-restrict", 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_norestrict>([a, b, c]() { c[0] = a[0] + b[0]; });
28+
// 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 %{{.*}})
29+
}
30+
{
31+
AnnotatedIntPtr a;
32+
int *b;
33+
int *c;
34+
kernel<class kernel_restrict1>([a, b, c]() { c[0] = a[0] + b[0]; });
35+
// 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 %{{.*}})
36+
}
37+
{
38+
int *a;
39+
AnnotatedIntPtr b;
40+
int *c;
41+
kernel<class kernel_restrict2>([a, b, c]() { c[0] = a[0] + b[0]; });
42+
// 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 %{{.*}})
43+
}
44+
{
45+
int *a;
46+
int *b;
47+
AnnotatedIntPtr c;
48+
kernel<class kernel_restrict3>([a, b, c]() { c[0] = a[0] + b[0]; });
49+
// 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" %{{.*}})
50+
}
51+
{
52+
AnnotatedIntPtr a;
53+
AnnotatedIntPtr b;
54+
int *c;
55+
kernel<class kernel_restrict4>([a, b, c]() { c[0] = a[0] + b[0]; });
56+
// 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 %{{.*}})
57+
}
58+
{
59+
AnnotatedIntPtr a;
60+
int *b;
61+
AnnotatedIntPtr c;
62+
kernel<class kernel_restrict5>([a, b, c]() { c[0] = a[0] + b[0]; });
63+
// 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" %{{.*}})
64+
}
65+
{
66+
int *a;
67+
AnnotatedIntPtr b;
68+
AnnotatedIntPtr c;
69+
kernel<class kernel_restrict6>([a, b, c]() { c[0] = a[0] + b[0]; });
70+
// 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" %{{.*}})
71+
}
72+
{
73+
AnnotatedIntPtr a;
74+
AnnotatedIntPtr b;
75+
AnnotatedIntPtr c;
76+
kernel<class kernel_restrict7>([a, b, c]() { c[0] = a[0] + b[0]; });
77+
// 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" %{{.*}})
78+
}
79+
}

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

Lines changed: 24 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -66,6 +66,13 @@ 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>;
72+
};
73+
74+
inline constexpr restrict_key::value_t restrict;
75+
6976
struct alignment_key
7077
: detail::compile_time_property_key<detail::PropKind::Alignment> {
7178
template <int K>
@@ -74,10 +81,18 @@ struct alignment_key
7481

7582
template <int K> inline constexpr alignment_key::value_t<K> alignment;
7683

84+
template <typename T>
85+
struct is_valid_property<T, restrict_key::value_t>
86+
: std::bool_constant<std::is_pointer<T>::value> {};
87+
7788
template <typename T, int W>
7889
struct is_valid_property<T, alignment_key::value_t<W>>
7990
: std::bool_constant<std::is_pointer<T>::value> {};
8091

92+
template <typename T, typename PropertyListT>
93+
struct is_property_key_of<restrict_key, annotated_ptr<T, PropertyListT>>
94+
: std::true_type {};
95+
8196
template <typename T, typename PropertyListT>
8297
struct is_property_key_of<alignment_key, annotated_ptr<T, PropertyListT>>
8398
: std::true_type {};
@@ -86,6 +101,10 @@ template <typename T, typename PropertyListT>
86101
struct is_property_key_of<alignment_key, annotated_arg<T, PropertyListT>>
87102
: std::true_type {};
88103

104+
template <typename T, typename PropertyListT>
105+
struct is_property_key_of<restrict_key, annotated_arg<T, PropertyListT>>
106+
: std::true_type {};
107+
89108
template <> struct propagateToPtrAnnotation<alignment_key> : std::true_type {};
90109

91110
namespace detail {
@@ -94,6 +113,11 @@ template <int N> struct PropertyMetaInfo<alignment_key::value_t<N>> {
94113
static constexpr int value = N;
95114
};
96115

116+
template <> struct PropertyMetaInfo<restrict_key::value_t> {
117+
static constexpr const char *name = "sycl-restrict";
118+
static constexpr std::nullptr_t value = nullptr;
119+
};
120+
97121
} // namespace detail
98122

99123
} // namespace experimental

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

Lines changed: 2 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -222,8 +222,9 @@ enum PropKind : uint32_t {
222222
Deterministic = 77,
223223
InitializeToIdentity = 78,
224224
WorkGroupScratchSize = 79,
225+
Restrict = 80,
225226
// PropKindSize must always be the last value.
226-
PropKindSize = 80,
227+
PropKindSize = 81,
227228
};
228229

229230
template <typename PropertyT> struct PropertyToKind {

‎sycl/test-e2e/Annotated_arg_ptr/annotated_arg_restrict.cpp‎

Lines changed: 27 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,27 @@
1+
// RUN: %{build} -o %t.out
2+
// RUN: %{run} %t.out
3+
// REQUIRES: aspect-usm_shared_allocations
4+
5+
// Checks that restrict annotated_arg works in device code.
6+
7+
#include <sycl/detail/core.hpp>
8+
#include <sycl/ext/oneapi/experimental/annotated_arg/annotated_arg.hpp>
9+
#include <sycl/usm.hpp>
10+
11+
namespace syclexp = sycl::ext::oneapi::experimental;
12+
13+
int main() {
14+
sycl::queue Q;
15+
16+
int *Ptr = sycl::malloc_shared<int>(1, Q);
17+
syclexp::annotated_arg<int *,
18+
decltype(syclexp::properties(syclexp::restrict))>
19+
AnnotArg{Ptr};
20+
Q.submit([&](sycl::handler &CGH) {
21+
CGH.single_task([=]() { *AnnotArg = 42; });
22+
}).wait();
23+
assert(*Ptr == 42);
24+
free(Ptr, Q);
25+
26+
return 0;
27+
}

‎sycl/test-e2e/Annotated_arg_ptr/annotated_ptr_restrict.cpp‎

Lines changed: 28 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,28 @@
1+
// RUN: %{build} -o %t.out
2+
// RUN: %{run} %t.out
3+
// REQUIRES: aspect-usm_shared_allocations
4+
5+
// Checks that restrict annotated_ptr works in device code.
6+
7+
#include <sycl/detail/core.hpp>
8+
#include <sycl/ext/oneapi/experimental/annotated_ptr/annotated_ptr.hpp>
9+
#include <sycl/usm.hpp>
10+
11+
namespace syclexp = sycl::ext::oneapi::experimental;
12+
13+
int main() {
14+
sycl::queue Q;
15+
16+
auto Ptr = sycl::malloc_shared<int>(1, Q);
17+
syclexp::annotated_ptr<int, decltype(syclexp::properties(syclexp::restrict))>
18+
AnnotPtr{Ptr};
19+
Q.submit([&](sycl::handler &CGH) {
20+
CGH.single_task([=]() { *AnnotPtr = 42; });
21+
}).wait();
22+
assert(*Ptr == 42);
23+
free(Ptr, Q);
24+
25+
return 0;
26+
}
27+
28+
// 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‎

Lines changed: 22 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,22 @@
1+
// RUN: %clangxx -fsycl-device-only -S -emit-llvm %s -o - | FileCheck %s --check-prefix CHECK-IR
2+
3+
#include "sycl/sycl.hpp"
4+
5+
namespace syclexp = sycl::ext::oneapi::experimental;
6+
7+
int main() {
8+
sycl::queue Q;
9+
10+
auto Ptr = sycl::malloc_shared<int>(1, Q);
11+
syclexp::annotated_arg<int *,
12+
decltype(syclexp::properties(syclexp::restrict))>
13+
AnnotArg{Ptr};
14+
Q.submit([&](sycl::handler &CGH) {
15+
CGH.single_task([=]() { *AnnotArg = 42; });
16+
}).wait();
17+
free(Ptr, Q);
18+
19+
return 0;
20+
}
21+
22+
// 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‎

Lines changed: 21 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,21 @@
1+
// RUN: %clangxx -fsycl-device-only -S -emit-llvm %s -o - | FileCheck %s --check-prefix CHECK-IR
2+
3+
#include "sycl/sycl.hpp"
4+
5+
namespace syclexp = sycl::ext::oneapi::experimental;
6+
7+
int main() {
8+
sycl::queue Q;
9+
10+
auto Ptr = sycl::malloc_shared<int>(1, Q);
11+
syclexp::annotated_ptr<int, decltype(syclexp::properties(syclexp::restrict))>
12+
AnnotPtr{Ptr};
13+
Q.submit([&](sycl::handler &CGH) {
14+
CGH.single_task([=]() { *AnnotPtr = 42; });
15+
}).wait();
16+
free(Ptr, Q);
17+
18+
return 0;
19+
}
20+
21+
// CHECK-IR: spir_kernel void @_ZTSZZ4mainENKUlRN4sycl3_V17handlerEE_clES2_EUlvE_(ptr addrspace(1) noalias noundef align 4 "sycl-restrict" %_arg_AnnotPtr)

0 commit comments

Comments
 (0)