diff --git a/clang/lib/CodeGen/CGCall.cpp b/clang/lib/CodeGen/CGCall.cpp index 388d1d2666804..81f3a5e9094bf 100644 --- a/clang/lib/CodeGen/CGCall.cpp +++ b/clang/lib/CodeGen/CGCall.cpp @@ -3020,6 +3020,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) { @@ -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() && - 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 a473da3b8aa7c..c5cec5d5639a3 100644 --- a/sycl/include/sycl/ext/oneapi/properties/property.hpp +++ b/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 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)