[SYCL] Add indeterminate constructor to work group memory interface

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<ClassTemplateSpecializationDecl>(FieldTy->getAsRecordDecl());
-    assert(AccTy->getTemplateArgs().size() >= 2 &&
-           "Incorrect template args for Accessor Type");
-    int Dims = static_cast<int>(
-        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<ClassTemplateSpecializationDecl>(FieldTy->getAsRecordDecl());
+      assert(AccTy->getTemplateArgs().size() >= 2 &&
+             "Incorrect template args for Accessor Type");
+      int Dims = static_cast<int>(
+          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;
   }
 

clang/test/SemaSYCL/Inputs/sycl.hpp

@@ -448,6 +448,23 @@ class __SYCL_TYPE(multi_ptr) multi_ptr<T, AS, access::decorated::legacy> {
   pointer_t m_Pointer;
 };
 
+// Dummy implementation of work_group_memory for use in SemaSYCL tests.
+template <typename DataT>
+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 {

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> {
+};
+
+int main() {
+  myQueue.submit([&](sycl::handler &h) {
+    WorkGroupMemoryDerived DerivedObject{ h };
+    h.parallel_for<class kernel>([=] {
+          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<int>' '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 *' <LValueToRValue>
+// CHECK-NEXT: DeclRefExpr {{.*}} '__local int *' lvalue ParmVar {{.*}} '_arg__base' '__local int *'
+// CHECK-NEXT: CompoundStmt {{.*}}
+// CHECK-NEXT: CXXOperatorCallExpr {{.*}} 'void' '()'
+// CHECK-NEXT: ImplicitCastExpr {{.*}} 'auto (*)() const -> void' <FunctionToPointerDecay>
+// CHECK-NEXT: DeclRefExpr {{.*}}'auto () const -> void' lvalue CXXMethod {{.*}} 'operator()' 'auto () const -> void'
+// CHECK-NEXT: ImplicitCastExpr {{.*}}
+// CHECK-NEXT: DeclRefExpr {{.*}}lvalue Var {{.*}} '__SYCLKernel'
+

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 <typename DataT, typename PropertyListT = empty_properties_t>
 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 <typename T = DataT,
@@ -73,9 +88,7 @@ class __SYCL_SPECIAL_CLASS __SYCL_TYPE(work_group_memory) work_group_memory
     *ptr = value;
     return *this;
   }
-#ifdef __SYCL_DEVICE_ONLY__
-  void __init(decoratedPtr ptr) { this->ptr = ptr; }
-#endif
+
 private:
   decoratedPtr ptr;
 };

sycl/test-e2e/WorkGroupMemory/swap_test.cpp

@@ -49,7 +49,7 @@ template <typename T> void swap_scalar(T &a, T &b) {
       syclexp::work_group_memory<T> temp{cgh};
       sycl::nd_range<1> ndr{size, wgsize};
       cgh.parallel_for(ndr, [=](sycl::nd_item<1> it) {
-        syclexp::work_group_memory<T> temp2;
+        syclexp::work_group_memory<T> 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<T[N][N]> temp2;
+        syclexp::work_group_memory<T[N][N]> temp2{syclexp::indeterminate};
         temp2 = temp;
         acc_b[i][j] = temp2[i][j];
       });