diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index e6a8b976de883..c1424fa65705e 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -111,6 +111,11 @@ class Util { /// \param Tmpl whether the class is template instantiation or simple record static bool isSyclType(QualType Ty, StringRef Name, bool Tmpl = false); + /// Checks whether given clang type is a standard SYCL API accessor class, + /// the check assumes the type is templated. + /// \param Ty the clang type being checked + static bool isSyclAccessorType(QualType Ty); + /// Checks whether given clang type is a full specialization of the SYCL /// specialization constant class. static bool isSyclSpecConstantType(QualType Ty); @@ -1021,7 +1026,11 @@ static ParamDesc makeParamDesc(ASTContext &Ctx, StringRef Name, QualType Ty) { } /// \return the target of given SYCL accessor type -static target getAccessTarget(const ClassTemplateSpecializationDecl *AccTy) { +static target getAccessTarget(QualType FieldTy, + const ClassTemplateSpecializationDecl *AccTy) { + if (Util::isSyclType(FieldTy, "local_accessor", true /*Tmpl*/)) + return local; + return static_cast( AccTy->getTemplateArgs()[3].getAsIntegral().getExtValue()); } @@ -1615,7 +1624,7 @@ class SyclKernelFieldChecker : public SyclKernelFieldHandler { assert(Util::isSyclSpecialType(Ty) && "Should only be called on sycl special class types."); const RecordDecl *RecD = Ty->getAsRecordDecl(); - if (IsSIMD && !Util::isSyclType(Ty, "accessor", true /*Tmp*/)) + if (IsSIMD && !Util::isSyclAccessorType(Ty)) return SemaRef.Diag(Loc.getBegin(), diag::err_sycl_esimd_not_supported_for_type) << RecD; @@ -1927,19 +1936,24 @@ class SyclKernelDeclCreator : public SyclKernelFieldHandler { } // Additional processing is required for accessor type. - void handleAccessorType(const CXXRecordDecl *RecordDecl, SourceLocation Loc) { + void handleAccessorType(QualType FieldTy, const CXXRecordDecl *RecordDecl, + SourceLocation Loc) { handleAccessorPropertyList(Params.back(), RecordDecl, Loc); - // Get access mode of accessor. - const auto *AccessorSpecializationDecl = - cast(RecordDecl); - const TemplateArgument &AccessModeArg = - AccessorSpecializationDecl->getTemplateArgs().get(2); + + // If "accessor" type check if read only + if (Util::isSyclType(FieldTy, "accessor", true /*Tmpl*/)) { + // Get access mode of accessor. + const auto *AccessorSpecializationDecl = + cast(RecordDecl); + const TemplateArgument &AccessModeArg = + AccessorSpecializationDecl->getTemplateArgs().get(2); + if (isReadOnlyAccessor(AccessModeArg)) + Params.back()->addAttr( + SYCLAccessorReadonlyAttr::CreateImplicit(SemaRef.getASTContext())); + } // Add implicit attribute to parameter decl when it is a read only // SYCL accessor. - if (isReadOnlyAccessor(AccessModeArg)) - Params.back()->addAttr( - SYCLAccessorReadonlyAttr::CreateImplicit(SemaRef.getASTContext())); Params.back()->addAttr( SYCLAccessorPtrAttr::CreateImplicit(SemaRef.getASTContext())); } @@ -1952,8 +1966,7 @@ class SyclKernelDeclCreator : public SyclKernelFieldHandler { const auto *RecordDecl = FieldTy->getAsCXXRecordDecl(); assert(RecordDecl && "The type must be a RecordDecl"); llvm::StringLiteral MethodName = - KernelDecl->hasAttr() && - Util::isSyclType(FieldTy, "accessor", true /*Tmp*/) + KernelDecl->hasAttr() && Util::isSyclAccessorType(FieldTy) ? InitESIMDMethodName : InitMethodName; CXXMethodDecl *InitMethod = getMethodByName(RecordDecl, MethodName); @@ -1978,8 +1991,8 @@ class SyclKernelDeclCreator : public SyclKernelFieldHandler { // added, this code needs to be refactored to call // handleAccessorPropertyList for each class which requires it. if (ParamTy.getTypePtr()->isPointerType() && - Util::isSyclType(FieldTy, "accessor", true /*Tmp*/)) - handleAccessorType(RecordDecl, FD->getBeginLoc()); + Util::isSyclAccessorType(FieldTy)) + handleAccessorType(FieldTy, RecordDecl, FD->getBeginLoc()); } LastParamIndex = ParamIndex; return true; @@ -2073,8 +2086,7 @@ class SyclKernelDeclCreator : public SyclKernelFieldHandler { const auto *RecordDecl = FieldTy->getAsCXXRecordDecl(); assert(RecordDecl && "The type must be a RecordDecl"); llvm::StringLiteral MethodName = - KernelDecl->hasAttr() && - Util::isSyclType(FieldTy, "accessor", true /*Tmp*/) + KernelDecl->hasAttr() && Util::isSyclAccessorType(FieldTy) ? InitESIMDMethodName : InitMethodName; CXXMethodDecl *InitMethod = getMethodByName(RecordDecl, MethodName); @@ -2093,8 +2105,8 @@ class SyclKernelDeclCreator : public SyclKernelFieldHandler { // added, this code needs to be refactored to call // handleAccessorPropertyList for each class which requires it. if (ParamTy.getTypePtr()->isPointerType() && - Util::isSyclType(FieldTy, "accessor", true /*Tmp*/)) - handleAccessorType(RecordDecl, BS.getBeginLoc()); + Util::isSyclAccessorType(FieldTy)) + handleAccessorType(FieldTy, RecordDecl, BS.getBeginLoc()); } LastParamIndex = ParamIndex; return true; @@ -2215,9 +2227,8 @@ class SyclKernelArgsSizeChecker : public SyclKernelFieldHandler { const CXXRecordDecl *RecordDecl = FieldTy->getAsCXXRecordDecl(); assert(RecordDecl && "The type must be a RecordDecl"); llvm::StringLiteral MethodName = - (IsSIMD && Util::isSyclType(FieldTy, "accessor", true /*Tmp*/)) - ? InitESIMDMethodName - : InitMethodName; + (IsSIMD && Util::isSyclAccessorType(FieldTy)) ? InitESIMDMethodName + : InitMethodName; CXXMethodDecl *InitMethod = getMethodByName(RecordDecl, MethodName); assert(InitMethod && "The type must have the __init method"); for (const ParmVarDecl *Param : InitMethod->parameters()) @@ -3124,7 +3135,7 @@ class SyclKernelIntHeaderCreator : public SyclKernelFieldHandler { "Incorrect template args for Accessor Type"); int Dims = static_cast( AccTy->getTemplateArgs()[1].getAsIntegral().getExtValue()); - int Info = getAccessTarget(AccTy) | (Dims << 11); + int Info = getAccessTarget(FieldTy, AccTy) | (Dims << 11); Header.addParamDesc(SYCLIntegrationHeader::kind_accessor, Info, CurOffset + offsetOf(RD, BC.getType()->getAsCXXRecordDecl())); @@ -3134,14 +3145,14 @@ class SyclKernelIntHeaderCreator : public SyclKernelFieldHandler { bool handleSyclSpecialType(FieldDecl *FD, QualType FieldTy) final { const auto *ClassTy = FieldTy->getAsCXXRecordDecl(); assert(ClassTy && "Type must be a C++ record type"); - if (Util::isSyclType(FieldTy, "accessor", true /*Tmp*/)) { + if (Util::isSyclAccessorType(FieldTy)) { const auto *AccTy = cast(FieldTy->getAsRecordDecl()); assert(AccTy->getTemplateArgs().size() >= 2 && "Incorrect template args for Accessor Type"); int Dims = static_cast( AccTy->getTemplateArgs()[1].getAsIntegral().getExtValue()); - int Info = getAccessTarget(AccTy) | (Dims << 11); + int Info = getAccessTarget(FieldTy, AccTy) | (Dims << 11); Header.addParamDesc(SYCLIntegrationHeader::kind_accessor, Info, CurOffset + offsetOf(FD, FieldTy)); @@ -5195,6 +5206,11 @@ bool Util::isSyclType(QualType Ty, StringRef Name, bool Tmpl) { return matchQualifiedTypeName(Ty, Scopes); } +bool Util::isSyclAccessorType(QualType Ty) { + return isSyclType(Ty, "accessor", true /* Tmpl */) || + isSyclType(Ty, "local_accessor", true /* Tmpl */); +} + bool Util::isAccessorPropertyListType(QualType Ty) { std::array Scopes = { Util::MakeDeclContextDesc(Decl::Kind::Namespace, "sycl"), diff --git a/clang/test/CodeGenSYCL/Inputs/sycl.hpp b/clang/test/CodeGenSYCL/Inputs/sycl.hpp index f9ca128301f39..a7424df625128 100644 --- a/clang/test/CodeGenSYCL/Inputs/sycl.hpp +++ b/clang/test/CodeGenSYCL/Inputs/sycl.hpp @@ -291,6 +291,26 @@ class accessor impl; }; +template +class __attribute__((sycl_special_class)) +local_accessor: public accessor { +public: + void use(void) const {} + template + void use(T... args) {} + template + void use(T... args) const {} + _ImplT impl; + +private: +#ifdef __SYCL_DEVICE_ONLY__ + void __init(__attribute__((opencl_local)) dataT *Ptr, range AccessRange, + range MemRange, id Offset) {} +#endif +}; + // TODO: Add support for image_array accessor. // template //class accessor diff --git a/clang/test/CodeGenSYCL/kernel-arg-accessor-pointer.cpp b/clang/test/CodeGenSYCL/kernel-arg-accessor-pointer.cpp index 749b04705e11a..979bbb2b2cf3a 100644 --- a/clang/test/CodeGenSYCL/kernel-arg-accessor-pointer.cpp +++ b/clang/test/CodeGenSYCL/kernel-arg-accessor-pointer.cpp @@ -22,6 +22,8 @@ int main() { access::placeholder::true_t> acc3; + local_accessor acc4; + // kernel_A parameters : int*, sycl::range<1>, sycl::range<1>, sycl::id<1>, // int*, sycl::range<1>, sycl::range<1>,sycl::id<1>. q.submit([&](handler &h) { @@ -67,11 +69,19 @@ int main() { // Using local accessor as a kernel parameter. // kernel_arg_runtime_aligned is generated for pointers from local accessors. q.submit([&](handler &h) { - h.single_task([=]() { + h.single_task([=]() { acc3.use(); }); }); + // Using local_accessor as a kernel parameter. + // kernel_arg_runtime_aligned is generated for pointers from local accessors. + q.submit([&](handler &h) { + h.single_task([=]() { + acc4.use(); + }); + }); + // kernel_acc_raw_ptr parameters : int*, sycl::range<1>, sycl::range<1>, sycl::id<1>, int*. int *rawPtr; q.submit([&](handler &h) { @@ -130,7 +140,7 @@ int main() { // CHECK-NOT: kernel_arg_runtime_aligned // CHECK-NOT: kernel_arg_exclusive_ptr -// CHECK: define {{.*}}spir_kernel void @{{.*}}localAccessor +// CHECK: define {{.*}}spir_kernel void @{{.*}}localAccessorDep // CHECK-SAME: ptr addrspace(1) noundef align 4 [[MEM_ARG1:%[a-zA-Z0-9_]+]], // CHECK-SAME: ptr noundef byval{{.*}}align 4 [[ACC_RANGE1:%[a-zA-Z0-9_]+1]], // CHECK-SAME: ptr noundef byval{{.*}}align 4 [[MEM_RANGE1:%[a-zA-Z0-9_]+2]], @@ -138,6 +148,14 @@ int main() { // CHECK-SAME: !kernel_arg_runtime_aligned ![[#ACCESSORMD2]] // CHECK-SAME: !kernel_arg_exclusive_ptr ![[#ACCESSORMD2]] +// CHECK: define {{.*}}spir_kernel void @{{.*}}localAccessor +// CHECK-SAME: ptr addrspace(3) noundef align 4 [[MEM_ARG1:%[a-zA-Z0-9_]+]], +// CHECK-SAME: ptr noundef byval{{.*}}align 4 [[ACC_RANGE1:%[a-zA-Z0-9_]+1]], +// CHECK-SAME: ptr noundef byval{{.*}}align 4 [[MEM_RANGE1:%[a-zA-Z0-9_]+2]], +// CHECK-SAME: ptr noundef byval{{.*}}align 4 [[OFFSET1:%[a-zA-Z0-9_]+3]] +// CHECK-SAME: !kernel_arg_runtime_aligned ![[#ACCESSORMD2]] +// CHECK-SAME: !kernel_arg_exclusive_ptr ![[#ACCESSORMD2]] + // Check kernel_acc_raw_ptr parameters // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_acc_raw_ptr // CHECK-SAME: ptr addrspace(1) noundef readonly align 4 [[MEM_ARG1:%[a-zA-Z0-9_]+]], diff --git a/clang/test/CodeGenSYCL/no_opaque_kernel-arg-accessor-pointer.cpp b/clang/test/CodeGenSYCL/no_opaque_kernel-arg-accessor-pointer.cpp index 74959c2eddfcc..95d6b687a0d58 100644 --- a/clang/test/CodeGenSYCL/no_opaque_kernel-arg-accessor-pointer.cpp +++ b/clang/test/CodeGenSYCL/no_opaque_kernel-arg-accessor-pointer.cpp @@ -22,6 +22,8 @@ int main() { access::placeholder::true_t> acc3; + local_accessor acc4; + // kernel_A parameters : int*, sycl::range<1>, sycl::range<1>, sycl::id<1>, // int*, sycl::range<1>, sycl::range<1>,sycl::id<1>. q.submit([&](handler &h) { @@ -67,11 +69,19 @@ int main() { // Using local accessor as a kernel parameter. // kernel_arg_runtime_aligned is generated for pointers from local accessors. q.submit([&](handler &h) { - h.single_task([=]() { + h.single_task([=]() { acc3.use(); }); }); + // Using local accessor as a kernel parameter. + // kernel_arg_runtime_aligned is generated for pointers from local accessors. + q.submit([&](handler &h) { + h.single_task([=]() { + acc4.use(); + }); + }); + // kernel_acc_raw_ptr parameters : int*, sycl::range<1>, sycl::range<1>, sycl::id<1>, int*. int *rawPtr; q.submit([&](handler &h) { @@ -125,13 +135,20 @@ int main() { // CHECK-SAME: float addrspace(1)* noundef align 4 [[MEM_ARG1:%[a-zA-Z0-9_]+]] // CHECK-NOT: kernel_arg_runtime_aligned -// CHECK: define {{.*}}spir_kernel void @{{.*}}localAccessor +// CHECK: define {{.*}}spir_kernel void @{{.*}}localAccessorDep // CHECK-SAME: float addrspace(1)* noundef align 4 [[MEM_ARG1:%[a-zA-Z0-9_]+]], // CHECK-SAME: %"struct.sycl::_V1::range.5"* noundef byval{{.*}}align 4 [[ACC_RANGE1:%[a-zA-Z0-9_]+1]], // CHECK-SAME: %"struct.sycl::_V1::range.5"* noundef byval{{.*}}align 4 [[MEM_RANGE1:%[a-zA-Z0-9_]+2]], // CHECK-SAME: %"struct.sycl::_V1::id.6"* noundef byval{{.*}}align 4 [[OFFSET1:%[a-zA-Z0-9_]+3]] // CHECK-SAME: !kernel_arg_runtime_aligned ![[#RTALIGNED2]] +// CHECK: define {{.*}}spir_kernel void @{{.*}}localAccessor +// CHECK-SAME: float addrspace(3)* noundef align 4 [[MEM_ARG1:%[a-zA-Z0-9_]+]], +// CHECK-SAME: %"struct.sycl::_V1::range.5"* noundef byval{{.*}}align 4 [[ACC_RANGE1:%[a-zA-Z0-9_]+1]], +// CHECK-SAME: %"struct.sycl::_V1::range.5"* noundef byval{{.*}}align 4 [[MEM_RANGE1:%[a-zA-Z0-9_]+2]], +// CHECK-SAME: %"struct.sycl::_V1::id.6"* noundef byval{{.*}}align 4 [[OFFSET1:%[a-zA-Z0-9_]+3]] +// CHECK-SAME: !kernel_arg_runtime_aligned ![[#RTALIGNED2]] + // Check kernel_acc_raw_ptr parameters // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_acc_raw_ptr // CHECK-SAME: i32 addrspace(1)* noundef readonly align 4 [[MEM_ARG1:%[a-zA-Z0-9_]+]], diff --git a/clang/test/SemaSYCL/Inputs/sycl.hpp b/clang/test/SemaSYCL/Inputs/sycl.hpp index e1b3c19e5f70c..19adf15143e85 100644 --- a/clang/test/SemaSYCL/Inputs/sycl.hpp +++ b/clang/test/SemaSYCL/Inputs/sycl.hpp @@ -207,6 +207,26 @@ class __attribute__((sycl_special_class)) accessor +class __attribute__((sycl_special_class)) +local_accessor: public accessor { +public: + void use(void) const {} + template + void use(T... args) {} + template + void use(T... args) const {} + _ImplT impl; + +private: +#ifdef __SYCL_DEVICE_ONLY__ + void __init(__attribute__((opencl_local)) dataT *Ptr, range AccessRange, + range MemRange, id Offset) {} +#endif +}; + struct sampler_impl { #ifdef __SYCL_DEVICE_ONLY__ __ocl_sampler_t m_Sampler; diff --git a/clang/test/SemaSYCL/accessors-targets.cpp b/clang/test/SemaSYCL/accessors-targets.cpp index 901e148b0785a..6a06fc86687cf 100644 --- a/clang/test/SemaSYCL/accessors-targets.cpp +++ b/clang/test/SemaSYCL/accessors-targets.cpp @@ -11,7 +11,9 @@ int main() { // Access work-group local memory with read and write access. sycl::accessor - local_acc; + local_acc_dep; + // Access work-group local memory with read and write access. + sycl::local_accessor local_acc; // Access buffer via global memory with read and write access. sycl::accessor @@ -21,6 +23,13 @@ int main() { sycl::access::target::constant_buffer> constant_acc; + q.submit([&](sycl::handler &h) { + h.single_task( + [=] { + local_acc_dep.use(); + }); + }); + q.submit([&](sycl::handler &h) { h.single_task( [=] { @@ -42,6 +51,7 @@ int main() { }); }); } +// CHECK: {{.*}}use_local_dep{{.*}} 'void (__local int *, sycl::range<1>, sycl::range<1>, sycl::id<1>)' // CHECK: {{.*}}use_local{{.*}} 'void (__local int *, sycl::range<1>, sycl::range<1>, sycl::id<1>)' // CHECK: {{.*}}use_global{{.*}} 'void (__global int *, sycl::range<1>, sycl::range<1>, sycl::id<1>)' // CHECK: {{.*}}use_constant{{.*}} 'void (__constant int *, sycl::range<1>, sycl::range<1>, sycl::id<1>)' diff --git a/sycl/include/sycl/access/access.hpp b/sycl/include/sycl/access/access.hpp index b83d7184d87fa..4f75823b7eac7 100644 --- a/sycl/include/sycl/access/access.hpp +++ b/sycl/include/sycl/access/access.hpp @@ -17,7 +17,7 @@ namespace access { enum class target { global_buffer __SYCL2020_DEPRECATED("use 'target::device' instead") = 2014, constant_buffer = 2015, - local = 2016, + local __SYCL2020_DEPRECATED("use `local_accessor` instead") = 2016, image = 2017, host_buffer = 2018, host_image = 2019, diff --git a/sycl/include/sycl/accessor.hpp b/sycl/include/sycl/accessor.hpp index 0fd532bebcd57..3031cf6a395ac 100644 --- a/sycl/include/sycl/accessor.hpp +++ b/sycl/include/sycl/accessor.hpp @@ -300,12 +300,13 @@ class accessor_common { using ConstRefType = const DataT &; using PtrType = detail::const_if_const_AS *; - using AccType = accessor; - // The class which allows to access value of N dimensional accessor using N // subscript operators, e.g. accessor[2][2][3] - template class AccessorSubscript { + template > + class AccessorSubscript { static constexpr int Dims = Dimensions; mutable id MIDs; @@ -2016,8 +2017,7 @@ accessor(buffer, handler, Type1, Type2, Type3, /// \ingroup sycl_api_acc template -class __SYCL_SPECIAL_CLASS accessor : +class __SYCL_SPECIAL_CLASS local_accessor_base : #ifndef __SYCL_DEVICE_ONLY__ public detail::LocalAccessorBaseHost, #endif @@ -2034,7 +2034,9 @@ class __SYCL_SPECIAL_CLASS accessor using AccessorSubscript = - typename AccessorCommonT::template AccessorSubscript; + typename AccessorCommonT::template AccessorSubscript< + Dims, + local_accessor_base>; using ConcreteASPtrType = typename detail::DecoratedType::type *; @@ -2057,7 +2059,7 @@ class __SYCL_SPECIAL_CLASS accessor::template get<0>()) {} protected: @@ -2091,8 +2093,8 @@ class __SYCL_SPECIAL_CLASS accessor> - accessor(handler &, const detail::code_location CodeLoc = - detail::code_location::current()) + local_accessor_base(handler &, const detail::code_location CodeLoc = + detail::code_location::current()) #ifdef __SYCL_DEVICE_ONLY__ : impl(range{1}){} #else @@ -2104,9 +2106,9 @@ class __SYCL_SPECIAL_CLASS accessor> - accessor(handler &, const property_list &propList, - const detail::code_location CodeLoc = - detail::code_location::current()) + local_accessor_base(handler &, const property_list &propList, + const detail::code_location CodeLoc = + detail::code_location::current()) #ifdef __SYCL_DEVICE_ONLY__ : impl(range{1}) { (void)propList; @@ -2120,7 +2122,7 @@ class __SYCL_SPECIAL_CLASS accessor 0)>> - accessor( + local_accessor_base( range AllocationSize, handler &, const detail::code_location CodeLoc = detail::code_location::current()) #ifdef __SYCL_DEVICE_ONLY__ @@ -2135,10 +2137,10 @@ class __SYCL_SPECIAL_CLASS accessor 0)>> - accessor(range AllocationSize, handler &, - const property_list &propList, - const detail::code_location CodeLoc = - detail::code_location::current()) + local_accessor_base(range AllocationSize, handler &, + const property_list &propList, + const detail::code_location CodeLoc = + detail::code_location::current()) #ifdef __SYCL_DEVICE_ONLY__ : impl(AllocationSize) { (void)propList; @@ -2206,7 +2208,9 @@ class __SYCL_SPECIAL_CLASS accessor 1)>> - typename AccessorCommonT::template AccessorSubscript + typename AccessorCommonT::template AccessorSubscript< + Dims - 1, + local_accessor_base> operator[](size_t Index) const { return AccessorSubscript(*this, Index); } @@ -2215,8 +2219,79 @@ class __SYCL_SPECIAL_CLASS accessor(getQualifiedPtr()); } - bool operator==(const accessor &Rhs) const { return impl == Rhs.impl; } - bool operator!=(const accessor &Rhs) const { return !(*this == Rhs); } + bool operator==(const local_accessor_base &Rhs) const { + return impl == Rhs.impl; + } + bool operator!=(const local_accessor_base &Rhs) const { + return !(*this == Rhs); + } +}; + +// TODO: Remove deprecated specialization once no longer needed +template +class __SYCL_SPECIAL_CLASS accessor + : public local_accessor_base { + + using local_acc = + local_accessor_base; + + // Use base classes constructors + using local_acc::local_acc; + +#ifdef __SYCL_DEVICE_ONLY__ + + // __init needs to be defined within the class not through inheritance. + // Map this function to inherited func. + void __init(typename local_acc::ConcreteASPtrType Ptr, + range AccessRange, + range range, + id id) { + local_acc::__init(Ptr, AccessRange, range, id); + } + +public: + // Default constructor for objects later initialized with __init member. + accessor() { + local_acc::impl = detail::InitializedVal::template get<0>(); + } + +#endif +}; + +template +class __SYCL_SPECIAL_CLASS local_accessor + : public local_accessor_base { + + using local_acc = + local_accessor_base; + + // Use base classes constructors + using local_acc::local_acc; + +#ifdef __SYCL_DEVICE_ONLY__ + + // __init needs to be defined within the class not through inheritance. + // Map this function to inherited func. + void __init(typename local_acc::ConcreteASPtrType Ptr, + range AccessRange, + range range, + id id) { + local_acc::__init(Ptr, AccessRange, range, id); + } + +public: + // Default constructor for objects later initialized with __init member. + local_accessor() { + local_acc::impl = detail::InitializedVal::template get<0>(); + } + +#endif }; /// Image accessors. diff --git a/sycl/include/sycl/handler.hpp b/sycl/include/sycl/handler.hpp index 68f4fe815a952..2144955addbc5 100644 --- a/sycl/include/sycl/handler.hpp +++ b/sycl/include/sycl/handler.hpp @@ -1417,6 +1417,11 @@ class __SYCL_EXPORT handler { setArgHelper(ArgIndex, std::move(Arg)); } + template + void set_arg(int ArgIndex, local_accessor Arg) { + setArgHelper(ArgIndex, std::move(Arg)); + } + /// Sets arguments for OpenCL interoperability kernels. /// /// Registers pack of arguments(Args) with indexes starting from 0. diff --git a/sycl/include/sycl/multi_ptr.hpp b/sycl/include/sycl/multi_ptr.hpp index 0dc93c0643efd..53d45d0e27207 100644 --- a/sycl/include/sycl/multi_ptr.hpp +++ b/sycl/include/sycl/multi_ptr.hpp @@ -16,11 +16,12 @@ namespace sycl { __SYCL_INLINE_VER_NAMESPACE(_V1) { -// Forward declaration +// Forward declarations template class accessor; +template class local_accessor; /// Provides constructors for address space qualified and non address space /// qualified pointers to allow interoperability between plain C++ and OpenCL C. @@ -151,6 +152,11 @@ template class multi_ptr { Accessor) : multi_ptr(Accessor.get_pointer()) {} + // Only if Space == local_space || generic_space + template + multi_ptr(local_accessor Accessor) + : multi_ptr(Accessor.get_pointer()) {} + // Only if Space == constant_space template < int dimensions, access::mode Mode, access::placeholder isPlaceholder, @@ -204,6 +210,19 @@ template class multi_ptr { Accessor) : multi_ptr(Accessor.get_pointer()) {} + // Only if Space == local_space || generic_space and element type is const + template < + int dimensions, access::address_space _Space = Space, + typename ET = ElementType, + typename = typename detail::enable_if_t< + _Space == Space && + (Space == access::address_space::generic_space || + Space == access::address_space::local_space) && + std::is_const::value && std::is_same::value>> + multi_ptr( + local_accessor, dimensions> Accessor) + : multi_ptr(Accessor.get_pointer()) {} + // Only if Space == constant_space and element type is const template < int dimensions, access::mode Mode, access::placeholder isPlaceholder, @@ -422,6 +441,16 @@ template class multi_ptr { Accessor) : multi_ptr(Accessor.get_pointer()) {} + // Only if Space == local_space || generic_space + template < + typename ElementType, int dimensions, + access::address_space _Space = Space, + typename = typename detail::enable_if_t< + _Space == Space && (Space == access::address_space::generic_space || + Space == access::address_space::local_space)>> + multi_ptr(local_accessor Accessor) + : multi_ptr(Accessor.get_pointer()) {} + // Only if Space == constant_space template < typename ElementType, int dimensions, access::mode Mode, @@ -546,6 +575,16 @@ template class multi_ptr { Accessor) : multi_ptr(Accessor.get_pointer()) {} + // Only if Space == local_space || generic_space + template < + typename ElementType, int dimensions, + access::address_space _Space = Space, + typename = typename detail::enable_if_t< + _Space == Space && (Space == access::address_space::generic_space || + Space == access::address_space::local_space)>> + multi_ptr(local_accessor Accessor) + : multi_ptr(Accessor.get_pointer()) {} + // Only if Space == constant_space template < typename ElementType, int dimensions, access::mode Mode, @@ -597,6 +636,9 @@ template ) -> multi_ptr; +template +multi_ptr(local_accessor) + -> multi_ptr; #endif template diff --git a/sycl/test/abi/layout_accessors_device.cpp b/sycl/test/abi/layout_accessors_device.cpp index 64ae1a188e463..e068063e69ec5 100644 --- a/sycl/test/abi/layout_accessors_device.cpp +++ b/sycl/test/abi/layout_accessors_device.cpp @@ -39,19 +39,42 @@ SYCL_EXTERNAL void hostAcc(accessor Acc) { (void)Acc.get_size(); } + // CHECK: 0 | class sycl::accessor -// CHECK-NEXT: 0 | class sycl::detail::accessor_common (base) (empty) -// CHECK-NEXT: 0 | class sycl::detail::LocalAccessorBaseDevice<1> impl -// CHECK-NEXT: 0 | class sycl::range<1> AccessRange -// CHECK-NEXT: 0 | class sycl::detail::array<1> (base) -// CHECK-NEXT: 0 | size_t[1] common_array -// CHECK-NEXT: 8 | class sycl::range<1> MemRange -// CHECK-NEXT: 8 | class sycl::detail::array<1> (base) -// CHECK-NEXT: 8 | size_t[1] common_array -// CHECK-NEXT: 16 | class sycl::id<1> Offset -// CHECK-NEXT: 16 | class sycl::detail::array<1> (base) -// CHECK-NEXT: 16 | size_t[1] common_array -// CHECK-NEXT: 24 | ConcreteASPtrType MData +// CHECK-NEXT: 0 | class sycl::local_accessor_base (base) +// CHECK-NEXT: 0 | class sycl::detail::accessor_common (base) (empty) +// CHECK-NEXT: 0 | class sycl::detail::LocalAccessorBaseDevice<1> impl +// CHECK-NEXT: 0 | class sycl::range<1> AccessRange +// CHECK-NEXT: 0 | class sycl::detail::array<1> (base) +// CHECK-NEXT: 0 | size_t[1] common_array +// CHECK-NEXT: 8 | class sycl::range<1> MemRange +// CHECK-NEXT: 8 | class sycl::detail::array<1> (base) +// CHECK-NEXT: 8 | size_t[1] common_array +// CHECK-NEXT: 16 | class sycl::id<1> Offset +// CHECK-NEXT: 16 | class sycl::detail::array<1> (base) +// CHECK-NEXT: 16 | size_t[1] common_array +// CHECK-NEXT: 24 | ConcreteASPtrType MData +// CHECK-NEXT: | [sizeof=32, dsize=32, align=8, +// CHECK-NEXT: | nvsize=32, nvalign=8] + +SYCL_EXTERNAL void hostAcc(local_accessor Acc) { + (void)Acc.get_size(); +} + +// CHECK: 0 | class sycl::local_accessor +// CHECK-NEXT: 0 | class sycl::local_accessor_base (base) +// CHECK-NEXT: 0 | class sycl::detail::accessor_common (base) (empty) +// CHECK-NEXT: 0 | class sycl::detail::LocalAccessorBaseDevice<1> impl +// CHECK-NEXT: 0 | class sycl::range<1> AccessRange +// CHECK-NEXT: 0 | class sycl::detail::array<1> (base) +// CHECK-NEXT: 0 | size_t[1] common_array +// CHECK-NEXT: 8 | class sycl::range<1> MemRange +// CHECK-NEXT: 8 | class sycl::detail::array<1> (base) +// CHECK-NEXT: 8 | size_t[1] common_array +// CHECK-NEXT: 16 | class sycl::id<1> Offset +// CHECK-NEXT: 16 | class sycl::detail::array<1> (base) +// CHECK-NEXT: 16 | size_t[1] common_array +// CHECK-NEXT: 24 | ConcreteASPtrType MData // CHECK-NEXT: | [sizeof=32, dsize=32, align=8, // CHECK-NEXT: | nvsize=32, nvalign=8] diff --git a/sycl/test/abi/layout_accessors_host.cpp b/sycl/test/abi/layout_accessors_host.cpp index eff5d7fde2516..78287c06a6925 100644 --- a/sycl/test/abi/layout_accessors_host.cpp +++ b/sycl/test/abi/layout_accessors_host.cpp @@ -98,16 +98,35 @@ void hostAcc(accessor A (void)Acc.get_size(); } -// CHECK: 0 | class sycl::accessor -// CHECK-NEXT: 0 | class sycl::detail::LocalAccessorBaseHost (base) -// CHECK-NEXT: 0 | class std::shared_ptr impl -// CHECK-NEXT: 0 | class std::__shared_ptr (base) -// CHECK-NEXT: 0 | class std::__shared_ptr_access (base) (empty) -// CHECK-NEXT: 0 | element_type * _M_ptr -// CHECK-NEXT: 8 | class std::__shared_count<__gnu_cxx::_S_atomic> _M_refcount -// CHECK-NEXT: 8 | _Sp_counted_base<(_Lock_policy)2U> * _M_pi -// CHECK-NEXT: 0 | class sycl::detail::accessor_common (base) (empty) -// CHECK-NEXT: 16 | char[16] padding +// CHECK: 0 | class sycl::accessor +// CHECK-NEXT: 0 | class sycl::local_accessor_base (base) +// CHECK-NEXT: 0 | class sycl::detail::LocalAccessorBaseHost (base) +// CHECK-NEXT: 0 | class std::shared_ptr impl +// CHECK-NEXT: 0 | class std::__shared_ptr (base) +// CHECK-NEXT: 0 | class std::__shared_ptr_access (base) (empty) +// CHECK-NEXT: 0 | element_type * _M_ptr +// CHECK-NEXT: 8 | class std::__shared_count<__gnu_cxx::_S_atomic> _M_refcount +// CHECK-NEXT: 8 | _Sp_counted_base<(_Lock_policy)2U> * _M_pi +// CHECK-NEXT: 0 | class sycl::detail::accessor_common (base) (empty) +// CHECK-NEXT: 16 | char[16] padding +// CHECK-NEXT: [sizeof=32, dsize=32, align=8, +// CHECK-NEXT: nvsize=32, nvalign=8] + +void hostAcc(local_accessor Acc) { + (void)Acc.get_size(); +} + +// CHECK: 0 | class sycl::local_accessor +// CHECK-NEXT: 0 | class sycl::local_accessor_base (base) +// CHECK-NEXT: 0 | class sycl::detail::LocalAccessorBaseHost (base) +// CHECK-NEXT: 0 | class std::shared_ptr impl +// CHECK-NEXT: 0 | class std::__shared_ptr (base) +// CHECK-NEXT: 0 | class std::__shared_ptr_access (base) (empty) +// CHECK-NEXT: 0 | element_type * _M_ptr +// CHECK-NEXT: 8 | class std::__shared_count<__gnu_cxx::_S_atomic> _M_refcount +// CHECK-NEXT: 8 | _Sp_counted_base<(_Lock_policy)2U> * _M_pi +// CHECK-NEXT: 0 | class sycl::detail::accessor_common (base) (empty) +// CHECK-NEXT: 16 | char[16] padding // CHECK-NEXT: [sizeof=32, dsize=32, align=8, // CHECK-NEXT: nvsize=32, nvalign=8] diff --git a/sycl/test/abi/user_mangling.cpp b/sycl/test/abi/user_mangling.cpp index bd2fbf014ae4a..2efb0330e8409 100644 --- a/sycl/test/abi/user_mangling.cpp +++ b/sycl/test/abi/user_mangling.cpp @@ -16,6 +16,9 @@ SYCL_EXTERNAL void acc(sycl::accessor) {} +// CHK_DEVICE: define dso_local void @_Z3accN4sycl3_V114local_accessorIiLi1EEE({{.*}}) +SYCL_EXTERNAL void acc(sycl::local_accessor) {} + // CHK-DEVICE: define dso_local spir_func void @_Z3accN4sycl3_V18accessorINS0_3vecIiLi4EEELi1ELNS0_6access4modeE1024ELNS4_6targetE2017ELNS4_11placeholderE0ENS0_3ext6oneapi22accessor_property_listIJEEEEE({{.*}}) SYCL_EXTERNAL void acc(sycl::accessor) {} @@ -46,6 +49,9 @@ void acc(sycl::accessor) {} +// CHK-HOST: define dso_local void @_Z3accN4sycl3_V114local_accessorIiLi1EEE({{.*}}) +void acc(sycl::local_accessor) {} + // CHK-HOST: define dso_local void @_Z3accN4sycl3_V18accessorINS0_3vecIiLi4EEELi1ELNS0_6access4modeE1024ELNS4_6targetE2019ELNS4_11placeholderE0ENS0_3ext6oneapi22accessor_property_listIJEEEEE({{.*}}) void acc(sycl::accessor) {} diff --git a/sycl/test/basic_tests/accessor/addrspace_exposure.cpp b/sycl/test/basic_tests/accessor/addrspace_exposure.cpp index 073af07c4619b..b80ac608905f7 100644 --- a/sycl/test/basic_tests/accessor/addrspace_exposure.cpp +++ b/sycl/test/basic_tests/accessor/addrspace_exposure.cpp @@ -31,7 +31,8 @@ int main() { GlobalBuf.get_access(Cgh); auto ConstantAcc = ConstantBuf.get_access(Cgh); - accessor LocalAcc(Range, Cgh); + local_accessor LocalAcc(Range, Cgh); + accessor LocalAccDep(Range, Cgh); Cgh.single_task([=]() { static_assert(std::is_same::value, @@ -46,6 +47,8 @@ int main() { "Incorrect type from constant accessor"); static_assert(std::is_same::value, "Incorrect type from local accessor"); + static_assert(std::is_same::value, + "Incorrect type from access target::local"); }); }); } diff --git a/sycl/test/basic_tests/set_arg_error.cpp b/sycl/test/basic_tests/set_arg_error.cpp index 90bda90bf0b6d..5febdede1533b 100644 --- a/sycl/test/basic_tests/set_arg_error.cpp +++ b/sycl/test/basic_tests/set_arg_error.cpp @@ -1,3 +1,4 @@ +// RUN: %clangxx %fsycl-host-only -DUSE_DEPRECATED_LOCAL_ACC -fsyntax-only -Xclang -verify -Xclang -verify-ignore-unexpected=note,warning %s // RUN: %clangxx %fsycl-host-only -fsyntax-only -Xclang -verify -Xclang -verify-ignore-unexpected=note,warning %s #include @@ -30,9 +31,14 @@ int main() { sycl::sampler samp(sycl::coordinate_normalization_mode::normalized, sycl::addressing_mode::clamp, sycl::filtering_mode::nearest); +#ifdef USE_DEPRECATED_LOCAL_ACC sycl::accessor local_acc({size}, h); +#else + sycl::local_accessor local_acc({size}, h); +#endif + TriviallyCopyable tc{1, 2}; NonTriviallyCopyable ntc; h.set_arg(0, local_acc); diff --git a/sycl/test/check_device_code/kernel_arguments_as.cpp b/sycl/test/check_device_code/kernel_arguments_as.cpp index 54a1aa8e6476f..5b1bbc6a139fd 100644 --- a/sycl/test/check_device_code/kernel_arguments_as.cpp +++ b/sycl/test/check_device_code/kernel_arguments_as.cpp @@ -1,7 +1,14 @@ +// RUN: %clangxx -DUSE_DEPRECATED_LOCAL_ACC -fsycl-device-only -Xclang -opaque-pointers -emit-llvm %s -S -o %t.ll -I %sycl_include -Wno-sycl-strict -Xclang -verify-ignore-unexpected=note,warning -Xclang -disable-llvm-passes +// RUN: FileCheck %s --input-file %t.ll --check-prefixes=CHECK,CHECK-DISABLE,CHECK-DEP +// // RUN: %clangxx -fsycl-device-only -Xclang -opaque-pointers -emit-llvm %s -S -o %t.ll -I %sycl_include -Wno-sycl-strict -Xclang -verify-ignore-unexpected=note,warning -Xclang -disable-llvm-passes -// RUN: FileCheck %s --input-file %t.ll --check-prefixes=CHECK,CHECK-DISABLE +// RUN: FileCheck %s --input-file %t.ll --check-prefixes=CHECK,CHECK-DISABLE,CHECK-SYCL2020 +// +// RUN: %clangxx -DUSE_DEPRECATED_LOCAL_ACC -fsycl-device-only -Xclang -opaque-pointers -emit-llvm %s -S -o %t.ll -I %sycl_include -Wno-sycl-strict -Xclang -verify-ignore-unexpected=note,warning -Xclang -disable-llvm-passes -D__ENABLE_USM_ADDR_SPACE__ +// RUN: FileCheck %s --input-file %t.ll --check-prefixes=CHECK,CHECK-ENABLE,CHECK-DEP +// // RUN: %clangxx -fsycl-device-only -Xclang -opaque-pointers -emit-llvm %s -S -o %t.ll -I %sycl_include -Wno-sycl-strict -Xclang -verify-ignore-unexpected=note,warning -Xclang -disable-llvm-passes -D__ENABLE_USM_ADDR_SPACE__ -// RUN: FileCheck %s --input-file %t.ll --check-prefixes=CHECK,CHECK-ENABLE +// RUN: FileCheck %s --input-file %t.ll --check-prefixes=CHECK,CHECK-ENABLE,CHECK-SYCL2020 // // Check the address space of the pointer in accessor class. // @@ -9,8 +16,12 @@ // CHECK: %"class.sycl::_V1::accessor[[NUMBER_SUFFIX]]" = type { %"class{{.*}}AccessorImplDevice", %[[UNION:.*]] } // CHECK-DISABLE: %[[UNION]] = type { ptr addrspace(1) } // CHECK-ENABLE: %[[UNION]] = type { ptr addrspace(5) } -// CHECK: %struct.AccWrapper.{{[0-9]+}} = type { %"class.sycl::_V1::accessor.[[NUM:[0-9]+]]" } -// CHECK-NEXT: %"class.sycl::_V1::accessor.[[NUM]]" = type { %"class{{.*}}LocalAccessorBaseDevice", ptr addrspace(3) } +// CHECK-DEP: %struct.AccWrapper.{{[0-9]+}} = type { %"class.sycl::_V1::accessor.[[NUM:[0-9]+]]" } +// CHECK-DEP-NEXT: %"class.sycl::_V1::accessor.[[NUM]]" = type { %"class{{.*}}local_accessor_base" } +// CHECK-DEP-NEXT: %"class.sycl::_V1::local_accessor_base" = type { %"class{{.*}}LocalAccessorBaseDevice", ptr addrspace(3) } +// CHECK-SYCL2020: %struct.AccWrapper.{{[0-9]+}} = type { %"class.sycl::_V1::local_accessor" } +// CHECK-SYCL2020-NEXT: %"class.sycl::_V1::local_accessor" = type { %"class{{.*}}local_accessor_base" } +// CHECK-SYCL2020-NEXT: %"class.sycl::_V1::local_accessor_base" = type { %"class{{.*}}LocalAccessorBaseDevice", ptr addrspace(3) } // // Check that kernel arguments doesn't have generic address space. // @@ -31,9 +42,13 @@ int main() { {sycl::property::buffer::use_host_ptr()}); queue.submit([&](sycl::handler &cgh) { auto acc = buf.get_access(cgh); +#ifdef USE_DEPRECATED_LOCAL_ACC sycl::accessor local_acc(sycl::range<1>(10), cgh); +#else + sycl::local_accessor local_acc(sycl::range<1>(10), cgh); +#endif // USE_DEPRECATED_LOCAL_ACC auto acc_wrapped = AccWrapper{acc}; auto local_acc_wrapped = AccWrapper{local_acc}; cgh.parallel_for( diff --git a/sycl/test/esimd/simd_copy_to_copy_from.cpp b/sycl/test/esimd/simd_copy_to_copy_from.cpp index d79b825c2fa6f..186fe7f099bd7 100644 --- a/sycl/test/esimd/simd_copy_to_copy_from.cpp +++ b/sycl/test/esimd/simd_copy_to_copy_from.cpp @@ -37,15 +37,13 @@ SYCL_EXTERNAL void kernel2(int *ptr) SYCL_ESIMD_FUNCTION { // --- Negative tests. // Incompatible target. -SYCL_EXTERNAL void -kernel3(accessor &buf) - SYCL_ESIMD_FUNCTION { +SYCL_EXTERNAL void kernel3(local_accessor &buf) SYCL_ESIMD_FUNCTION { simd v1(0, 1); simd v0; - // CHECK: simd_copy_to_copy_from.cpp:46{{.*}}error: no matching member function for call to 'copy_from' + // CHECK: simd_copy_to_copy_from.cpp:44{{.*}}error: no matching member function for call to 'copy_from' v0.copy_from(buf, 0); v0 = v0 + v1; - // CHECK: simd_copy_to_copy_from.cpp:49{{.*}}error: no matching member function for call to 'copy_to' + // CHECK: simd_copy_to_copy_from.cpp:47{{.*}}error: no matching member function for call to 'copy_to' v0.copy_to(buf, 0); } @@ -54,7 +52,7 @@ SYCL_EXTERNAL void kernel4(accessor &buf) SYCL_ESIMD_FUNCTION { simd v; - // CHECK: simd_copy_to_copy_from.cpp:58{{.*}}error: no matching member function for call to 'copy_from' + // CHECK: simd_copy_to_copy_from.cpp:56{{.*}}error: no matching member function for call to 'copy_from' v.copy_from(buf, 0); } @@ -63,6 +61,6 @@ SYCL_EXTERNAL void kernel5(accessor &buf) SYCL_ESIMD_FUNCTION { simd v(0, 1); - // CHECK: simd_copy_to_copy_from.cpp:67{{.*}}error: no matching member function for call to 'copy_to' + // CHECK: simd_copy_to_copy_from.cpp:65{{.*}}error: no matching member function for call to 'copy_to' v.copy_to(buf, 0); } diff --git a/sycl/test/extensions/sub_group_as.cpp b/sycl/test/extensions/sub_group_as.cpp index c2fc6dc01bc71..e303ddd91d581 100644 --- a/sycl/test/extensions/sub_group_as.cpp +++ b/sycl/test/extensions/sub_group_as.cpp @@ -26,9 +26,7 @@ int main(int argc, char *argv[]) { queue.submit([&](sycl::handler &cgh) { auto global = buf.get_access(cgh); - sycl::accessor - local(N, cgh); + sycl::local_accessor local(N, cgh); cgh.parallel_for( sycl::nd_range<1>(N, 32), [=](sycl::nd_item<1> it) { diff --git a/sycl/test/multi_ptr/ctad.cpp b/sycl/test/multi_ptr/ctad.cpp index 819cac9a836f3..e8416cae8b4f8 100644 --- a/sycl/test/multi_ptr/ctad.cpp +++ b/sycl/test/multi_ptr/ctad.cpp @@ -16,11 +16,13 @@ int main() { using deviceAcc = sycl::accessor; using globlAcc = sycl::accessor; using constAcc = sycl::accessor; - using localAcc = sycl::accessor; + using localAcc = sycl::local_accessor; + using localAccDep = sycl::accessor; using deviceCTAD = decltype(sycl::multi_ptr(std::declval())); using globlCTAD = decltype(sycl::multi_ptr(std::declval())); using constCTAD = decltype(sycl::multi_ptr(std::declval())); using localCTAD = decltype(sycl::multi_ptr(std::declval())); + using localCTADDep = decltype(sycl::multi_ptr(std::declval())); using deviceMPtr = sycl::multi_ptr; using globlMPtr = sycl::multi_ptr; using constMPtr = sycl::multi_ptr; @@ -30,4 +32,5 @@ int main() { static_assert(std::is_same::value); static_assert(std::is_same::value); static_assert(std::is_same::value); + static_assert(std::is_same::value); } diff --git a/sycl/test/warnings/sycl_2020_deprecations.cpp b/sycl/test/warnings/sycl_2020_deprecations.cpp index bbd3ea323e8e0..674637c6b1fb7 100644 --- a/sycl/test/warnings/sycl_2020_deprecations.cpp +++ b/sycl/test/warnings/sycl_2020_deprecations.cpp @@ -197,5 +197,11 @@ int main() { // expected-warning@+1{{'get_linear_id' is deprecated: use sycl::group::get_group_linear_id() instead}} group.get_linear_id(); + // expected-warning@+2{{'local' is deprecated: use `local_accessor` instead}} + Queue.submit([&](sycl::handler &CGH) { + sycl::accessor + LocalAcc(sycl::range<1>(1), CGH); + }); + return 0; }