diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index e45b038273d77..7d53638c8eff3 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -3950,13 +3950,26 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler { } // Default inits the type, then calls the init-method in the body. + // A type may not have a public default constructor as per its spec so + // typically if this is the case the default constructor will be private and + // in such cases we must manually override the access specifier from private + // to public just for the duration of this default initialization. + // TODO: Revisit this approach once https://github.com/intel/llvm/issues/16061 + // is closed. bool handleSpecialType(FieldDecl *FD, QualType Ty) { + const auto *RecordDecl = Ty->getAsCXXRecordDecl(); + AccessSpecifier DefaultConstructorAccess; + auto DefaultConstructor = + std::find_if(RecordDecl->ctor_begin(), RecordDecl->ctor_end(), + [](auto it) { return it->isDefaultConstructor(); }); + DefaultConstructorAccess = DefaultConstructor->getAccess(); + DefaultConstructor->setAccess(AS_public); + addFieldInit(FD, Ty, std::nullopt, InitializationKind::CreateDefault(KernelCallerSrcLoc)); - + DefaultConstructor->setAccess(DefaultConstructorAccess); addFieldMemberExpr(FD, Ty); - const auto *RecordDecl = Ty->getAsCXXRecordDecl(); createSpecialMethodCall(RecordDecl, getInitMethodName(), BodyStmts); CXXMethodDecl *FinalizeMethod = getMethodByName(RecordDecl, FinalizeMethodName); @@ -3970,9 +3983,17 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler { } bool handleSpecialType(const CXXBaseSpecifier &BS, QualType Ty) { - const auto *RecordDecl = Ty->getAsCXXRecordDecl(); + const auto *BaseRecordDecl = BS.getType()->getAsCXXRecordDecl(); + AccessSpecifier DefaultConstructorAccess; + auto DefaultConstructor = + std::find_if(BaseRecordDecl->ctor_begin(), BaseRecordDecl->ctor_end(), + [](auto it) { return it->isDefaultConstructor(); }); + DefaultConstructorAccess = DefaultConstructor->getAccess(); + DefaultConstructor->setAccess(AS_public); + addBaseInit(BS, Ty, InitializationKind::CreateDefault(KernelCallerSrcLoc)); - createSpecialMethodCall(RecordDecl, getInitMethodName(), BodyStmts); + DefaultConstructor->setAccess(DefaultConstructorAccess); + createSpecialMethodCall(BaseRecordDecl, getInitMethodName(), BodyStmts); return true; } @@ -4669,16 +4690,21 @@ class SyclKernelIntHeaderCreator : public SyclKernelFieldHandler { bool handleSyclSpecialType(const CXXRecordDecl *RD, const CXXBaseSpecifier &BC, QualType FieldTy) final { - 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(FieldTy, AccTy) | (Dims << 11); - Header.addParamDesc(SYCLIntegrationHeader::kind_accessor, Info, - CurOffset + - offsetOf(RD, BC.getType()->getAsCXXRecordDecl())); + if (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(FieldTy, AccTy) | (Dims << 11); + Header.addParamDesc(SYCLIntegrationHeader::kind_accessor, Info, + CurOffset + + offsetOf(RD, BC.getType()->getAsCXXRecordDecl())); + } else if (SemaSYCL::isSyclType(FieldTy, SYCLTypeAttr::work_group_memory)) { + addParam(FieldTy, SYCLIntegrationHeader::kind_work_group_memory, + offsetOf(RD, BC.getType()->getAsCXXRecordDecl())); + } return true; } diff --git a/clang/test/SemaSYCL/Inputs/sycl.hpp b/clang/test/SemaSYCL/Inputs/sycl.hpp index 337320764de2f..5df1550ed2dcb 100644 --- a/clang/test/SemaSYCL/Inputs/sycl.hpp +++ b/clang/test/SemaSYCL/Inputs/sycl.hpp @@ -448,6 +448,23 @@ class __SYCL_TYPE(multi_ptr) multi_ptr { pointer_t m_Pointer; }; +// Dummy implementation of work_group_memory for use in SemaSYCL tests. +template +class __attribute__((sycl_special_class)) +__SYCL_TYPE(work_group_memory) work_group_memory { + +// Default constructor for objects later initialized with __init member. + work_group_memory() = default; + +public: + work_group_memory(handler &CGH) {} + + void __init(__attribute((opencl_local)) DataT *Ptr) { this->Ptr = Ptr; } + void use() const {} +private: + __attribute((opencl_local)) DataT *Ptr; +}; + namespace ext { namespace oneapi { namespace experimental { diff --git a/clang/test/SemaSYCL/work_group_memory_inheritance.cpp b/clang/test/SemaSYCL/work_group_memory_inheritance.cpp new file mode 100644 index 0000000000000..6f1f6badbdc59 --- /dev/null +++ b/clang/test/SemaSYCL/work_group_memory_inheritance.cpp @@ -0,0 +1,43 @@ +// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -ast-dump -sycl-std=2020 %s | FileCheck %s + +// Check that AST is correctly generated for kernel arguments that inherit from work group memory. + +#include "sycl.hpp" + +sycl::queue myQueue; + +struct WorkGroupMemoryDerived : + sycl::work_group_memory { +}; + +int main() { + myQueue.submit([&](sycl::handler &h) { + WorkGroupMemoryDerived DerivedObject{ h }; + h.parallel_for([=] { + DerivedObject.use(); + }); + }); + return 0; +} + +// CHECK: FunctionDecl {{.*}}kernel{{.*}} 'void (__local int *)' +// CHECK-NEXT: ParmVarDecl {{.*}}used _arg__base '__local int *' +// CHECK-NEXT: CompoundStmt {{.*}} +// CHECK-NEXT: DeclStmt {{.*}} +// CHECK-NEXT: VarDecl {{.*}} used __SYCLKernel {{.*}} cinit +// CHECK-NEXT: InitListExpr {{.*}} +// CHECK-NEXT: InitListExpr {{.*}} 'WorkGroupMemoryDerived' +// CHECK-NEXT: CXXConstructExpr {{.*}} 'sycl::work_group_memory' 'void () noexcept' +// CHECK-NEXT: CXXMemberCallExpr {{.*}} 'void' +// CHECK-NEXT: MemberExpr {{.*}} 'void (__local int *)' lvalue .__init {{.*}} +// CHECK-NEXT: MemberExpr {{.*}} 'WorkGroupMemoryDerived' lvalue .DerivedObject +// CHECK-NEXT: DeclRefExpr {{.*}} lvalue Var {{.*}} '__SYCLKernel' +// CHECK-NEXT: ImplicitCastExpr {{.*}} '__local int *' +// CHECK-NEXT: DeclRefExpr {{.*}} '__local int *' lvalue ParmVar {{.*}} '_arg__base' '__local int *' +// CHECK-NEXT: CompoundStmt {{.*}} +// CHECK-NEXT: CXXOperatorCallExpr {{.*}} 'void' '()' +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'auto (*)() const -> void' +// CHECK-NEXT: DeclRefExpr {{.*}}'auto () const -> void' lvalue CXXMethod {{.*}} 'operator()' 'auto () const -> void' +// CHECK-NEXT: ImplicitCastExpr {{.*}} +// CHECK-NEXT: DeclRefExpr {{.*}}lvalue Var {{.*}} '__SYCLKernel' + diff --git a/sycl/include/sycl/ext/oneapi/experimental/work_group_memory.hpp b/sycl/include/sycl/ext/oneapi/experimental/work_group_memory.hpp index 7870ebd3ca73e..d7c9138a2a23d 100644 --- a/sycl/include/sycl/ext/oneapi/experimental/work_group_memory.hpp +++ b/sycl/include/sycl/ext/oneapi/experimental/work_group_memory.hpp @@ -36,6 +36,9 @@ class work_group_memory_impl { } // namespace detail namespace ext::oneapi::experimental { +struct indeterminate_t {}; +inline constexpr indeterminate_t indeterminate; + template class __SYCL_SPECIAL_CLASS __SYCL_TYPE(work_group_memory) work_group_memory : sycl::detail::work_group_memory_impl { @@ -46,8 +49,20 @@ class __SYCL_SPECIAL_CLASS __SYCL_TYPE(work_group_memory) work_group_memory using decoratedPtr = typename sycl::detail::DecoratedType< value_type, access::address_space::local_space>::type *; -public: + // Frontend requires special types to have a default constructor in order to + // have a uniform way of initializing an object of special type to then call + // the __init method on it. This is purely an implementation detail and not + // part of the spec. + // TODO: Revisit this once https://github.com/intel/llvm/issues/16061 is + // closed. work_group_memory() = default; + +#ifdef __SYCL_DEVICE_ONLY__ + void __init(decoratedPtr ptr) { this->ptr = ptr; } +#endif + +public: + work_group_memory(const indeterminate_t &) {}; work_group_memory(const work_group_memory &rhs) = default; work_group_memory &operator=(const work_group_memory &rhs) = default; template ptr = ptr; } -#endif + private: decoratedPtr ptr; }; diff --git a/sycl/test-e2e/WorkGroupMemory/swap_test.cpp b/sycl/test-e2e/WorkGroupMemory/swap_test.cpp index 13fbde212a47d..7552774edcbbf 100644 --- a/sycl/test-e2e/WorkGroupMemory/swap_test.cpp +++ b/sycl/test-e2e/WorkGroupMemory/swap_test.cpp @@ -49,7 +49,7 @@ template void swap_scalar(T &a, T &b) { syclexp::work_group_memory temp{cgh}; sycl::nd_range<1> ndr{size, wgsize}; cgh.parallel_for(ndr, [=](sycl::nd_item<1> it) { - syclexp::work_group_memory temp2; + syclexp::work_group_memory temp2{syclexp::indeterminate}; temp2 = temp; // temp and temp2 have the same underlying data temp = acc_a[0]; acc_a[0] = acc_b[0]; @@ -264,7 +264,7 @@ void swap_array_2d(T (&a)[N][N], T (&b)[N][N], size_t batch_size) { const auto j = it.get_global_id()[1]; temp[i][j] = acc_a[i][j]; acc_a[i][j] = acc_b[i][j]; - syclexp::work_group_memory temp2; + syclexp::work_group_memory temp2{syclexp::indeterminate}; temp2 = temp; acc_b[i][j] = temp2[i][j]; });