From 9f739ecb1e430a13e51593095f86f58f9d4b53e1 Mon Sep 17 00:00:00 2001 From: "Larsen, Steffen" Date: Thu, 14 Nov 2024 07:56:28 -0800 Subject: [PATCH 1/2] [SYCL] Implement restrict kernel argument property 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 --- clang/lib/CodeGen/CGCall.cpp | 22 +++++- .../CodeGenSYCL/sycl_restrict_property.cpp | 79 +++++++++++++++++++ .../properties.hpp | 24 ++++++ .../sycl/ext/oneapi/properties/property.hpp | 3 +- .../annotated_arg_restrict.cpp | 27 +++++++ .../annotated_ptr_restrict.cpp | 28 +++++++ .../extensions/annotated_arg/restrict.cpp | 22 ++++++ .../extensions/annotated_ptr/restrict.cpp | 21 +++++ 8 files changed, 222 insertions(+), 4 deletions(-) create mode 100644 clang/test/CodeGenSYCL/sycl_restrict_property.cpp create mode 100644 sycl/test-e2e/Annotated_arg_ptr/annotated_arg_restrict.cpp create mode 100644 sycl/test-e2e/Annotated_arg_ptr/annotated_ptr_restrict.cpp create mode 100644 sycl/test/check_device_code/extensions/annotated_arg/restrict.cpp create mode 100644 sycl/test/check_device_code/extensions/annotated_ptr/restrict.cpp diff --git a/clang/lib/CodeGen/CGCall.cpp b/clang/lib/CodeGen/CGCall.cpp index a78b74e616b12..85a934a68f633 100644 --- a/clang/lib/CodeGen/CGCall.cpp +++ b/clang/lib/CodeGen/CGCall.cpp @@ -2995,6 +2995,21 @@ namespace { }; } +static bool HasSYCLRestrictPropertyIRAttr(const VarDecl *Arg, + const ASTContext &Context) { + auto *IRAttr = Arg->getAttr(); + if (!IRAttr) + return false; + + SmallVector, 4> NameValuePairs = + IRAttr->getAttributeNameValuePairs(Context); + return std::any_of( + NameValuePairs.begin(), NameValuePairs.end(), + [](const std::pair &NameValuePair) { + return NameValuePair.first == "sycl-restrict"; + }); +} + void CodeGenFunction::EmitFunctionProlog(const CGFunctionInfo &FI, llvm::Function *Fn, const FunctionArgList &Args) { @@ -3219,9 +3234,10 @@ void CodeGenFunction::EmitFunctionProlog(const CGFunctionInfo &FI, // Set 'noalias' if an argument type has the `restrict` qualifier. if (Arg->getType().isRestrictQualified() || - (CurCodeDecl && - CurCodeDecl->hasAttr() && - Arg->getType()->isPointerType()) || + (Arg->getType()->isPointerType() && + ((CurCodeDecl && + CurCodeDecl->hasAttr()) || + HasSYCLRestrictPropertyIRAttr(Arg, getContext()))) || (Arg->hasAttr() && Arg->getType()->isPointerType())) AI->addAttr(llvm::Attribute::NoAlias); } diff --git a/clang/test/CodeGenSYCL/sycl_restrict_property.cpp b/clang/test/CodeGenSYCL/sycl_restrict_property.cpp new file mode 100644 index 0000000000000..b62cac27159a6 --- /dev/null +++ b/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 +__attribute__((sycl_kernel)) void kernel(const Func &kernelFunc) { + kernelFunc(); +} + +int main() { + { + int *a; + int *b; + int *c; + kernel([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([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([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([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([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([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([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([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" %{{.*}}) + } +} diff --git a/sycl/include/sycl/ext/oneapi/experimental/common_annotated_properties/properties.hpp b/sycl/include/sycl/ext/oneapi/experimental/common_annotated_properties/properties.hpp index dd540a14232b3..04258b9280fb1 100644 --- a/sycl/include/sycl/ext/oneapi/experimental/common_annotated_properties/properties.hpp +++ b/sycl/include/sycl/ext/oneapi/experimental/common_annotated_properties/properties.hpp @@ -66,6 +66,13 @@ struct propagateToPtrAnnotation> //===----------------------------------------------------------------------===// // Common properties of annotated_arg/annotated_ptr //===----------------------------------------------------------------------===// +struct restrict_key + : detail::compile_time_property_key { + using value_t = property_value; +}; + +inline constexpr restrict_key::value_t restrict; + struct alignment_key : detail::compile_time_property_key { template @@ -74,10 +81,18 @@ struct alignment_key template inline constexpr alignment_key::value_t alignment; +template +struct is_valid_property + : std::bool_constant::value> {}; + template struct is_valid_property> : std::bool_constant::value> {}; +template +struct is_property_key_of> + : std::true_type {}; + template struct is_property_key_of> : std::true_type {}; @@ -86,6 +101,10 @@ template struct is_property_key_of> : std::true_type {}; +template +struct is_property_key_of> + : std::true_type {}; + template <> struct propagateToPtrAnnotation : std::true_type {}; namespace detail { @@ -94,6 +113,11 @@ template struct PropertyMetaInfo> { static constexpr int value = N; }; +template <> struct PropertyMetaInfo { + static constexpr const char *name = "sycl-restrict"; + static constexpr std::nullptr_t value = nullptr; +}; + } // namespace detail } // namespace experimental diff --git a/sycl/include/sycl/ext/oneapi/properties/property.hpp b/sycl/include/sycl/ext/oneapi/properties/property.hpp index 5b147d93f7e95..f080d3eb78200 100644 --- a/sycl/include/sycl/ext/oneapi/properties/property.hpp +++ b/sycl/include/sycl/ext/oneapi/properties/property.hpp @@ -221,8 +221,9 @@ enum PropKind : uint32_t { Prefetch = 76, Deterministic = 77, InitializeToIdentity = 78, + Restrict = 79, // PropKindSize must always be the last value. - PropKindSize = 79, + PropKindSize = 80, }; template struct PropertyToKind { diff --git a/sycl/test-e2e/Annotated_arg_ptr/annotated_arg_restrict.cpp b/sycl/test-e2e/Annotated_arg_ptr/annotated_arg_restrict.cpp new file mode 100644 index 0000000000000..eefaa25c66286 --- /dev/null +++ b/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 +#include +#include + +namespace syclexp = sycl::ext::oneapi::experimental; + +int main() { + sycl::queue Q; + + int *Ptr = sycl::malloc_shared(1, Q); + syclexp::annotated_arg + AnnotArg{Ptr}; + Q.submit([&](sycl::handler &CGH) { + CGH.single_task([=]() { *AnnotArg = 42; }); + }).wait(); + assert(*Ptr == 42); + free(Ptr, Q); + + return 0; +} diff --git a/sycl/test-e2e/Annotated_arg_ptr/annotated_ptr_restrict.cpp b/sycl/test-e2e/Annotated_arg_ptr/annotated_ptr_restrict.cpp new file mode 100644 index 0000000000000..42343b6986ae6 --- /dev/null +++ b/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 +#include +#include + +namespace syclexp = sycl::ext::oneapi::experimental; + +int main() { + sycl::queue Q; + + auto Ptr = sycl::malloc_shared(1, Q); + syclexp::annotated_ptr + 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) diff --git a/sycl/test/check_device_code/extensions/annotated_arg/restrict.cpp b/sycl/test/check_device_code/extensions/annotated_arg/restrict.cpp new file mode 100644 index 0000000000000..305d681823ea3 --- /dev/null +++ b/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(1, Q); + syclexp::annotated_arg + 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) diff --git a/sycl/test/check_device_code/extensions/annotated_ptr/restrict.cpp b/sycl/test/check_device_code/extensions/annotated_ptr/restrict.cpp new file mode 100644 index 0000000000000..c894954444a7c --- /dev/null +++ b/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(1, Q); + syclexp::annotated_ptr + 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) From bc43053f07055735dd1519fec6e562da7fc576d7 Mon Sep 17 00:00:00 2001 From: "Larsen, Steffen" Date: Fri, 15 Nov 2024 04:07:21 -0800 Subject: [PATCH 2/2] Rename check function Signed-off-by: Larsen, Steffen --- clang/lib/CodeGen/CGCall.cpp | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/clang/lib/CodeGen/CGCall.cpp b/clang/lib/CodeGen/CGCall.cpp index 85a934a68f633..be3e405678dca 100644 --- a/clang/lib/CodeGen/CGCall.cpp +++ b/clang/lib/CodeGen/CGCall.cpp @@ -2995,7 +2995,7 @@ namespace { }; } -static bool HasSYCLRestrictPropertyIRAttr(const VarDecl *Arg, +static bool hasSYCLRestrictPropertyIRAttr(const VarDecl *Arg, const ASTContext &Context) { auto *IRAttr = Arg->getAttr(); if (!IRAttr) @@ -3237,7 +3237,7 @@ void CodeGenFunction::EmitFunctionProlog(const CGFunctionInfo &FI, (Arg->getType()->isPointerType() && ((CurCodeDecl && CurCodeDecl->hasAttr()) || - HasSYCLRestrictPropertyIRAttr(Arg, getContext()))) || + hasSYCLRestrictPropertyIRAttr(Arg, getContext()))) || (Arg->hasAttr() && Arg->getType()->isPointerType())) AI->addAttr(llvm::Attribute::NoAlias); }