[SYCL][Graph] Implement dynamic local accessors

clang/include/clang/Basic/Attr.td

@@ -1637,12 +1637,14 @@ def SYCLType: InheritableAttr {
   let Subjects = SubjectList<[CXXRecord, Enum], ErrorDiag>;
   let LangOpts = [SYCLIsDevice, SilentlyIgnoreSYCLIsHost];
   let Args = [EnumArgument<"Type", "SYCLType", /*is_string=*/true,
-                           ["accessor", "local_accessor", "work_group_memory", "dynamic_work_group_memory",
+                           ["accessor", "local_accessor", "dynamic_local_accessor",
+                           "work_group_memory", "dynamic_work_group_memory",
                             "specialization_id", "kernel_handler", "buffer_location",
                             "no_alias", "accessor_property_list", "group",
                             "private_memory", "aspect", "annotated_ptr", "annotated_arg",
                             "stream", "sampler", "host_pipe", "multi_ptr"],
-                           ["accessor", "local_accessor", "work_group_memory", "dynamic_work_group_memory",
+                           ["accessor", "local_accessor", "dynamic_local_accessor",
+                           "work_group_memory", "dynamic_work_group_memory",
                             "specialization_id", "kernel_handler", "buffer_location",
                             "no_alias", "accessor_property_list", "group",
                             "private_memory", "aspect", "annotated_ptr", "annotated_arg",

clang/include/clang/Sema/SemaSYCL.h

@@ -64,7 +64,8 @@ class SYCLIntegrationHeader {
     kind_stream,
     kind_work_group_memory,
     kind_dynamic_work_group_memory,
-    kind_last = kind_dynamic_work_group_memory
+    kind_dynamic_accessor,
+    kind_last = kind_dynamic_accessor
   };
 
 public:

clang/lib/Sema/SemaSYCL.cpp

@@ -94,7 +94,8 @@ bool SemaSYCL::isSyclType(QualType Ty, SYCLTypeAttr::SYCLType TypeName) {
 
 static bool isSyclAccessorType(QualType Ty) {
   return SemaSYCL::isSyclType(Ty, SYCLTypeAttr::accessor) ||
-         SemaSYCL::isSyclType(Ty, SYCLTypeAttr::local_accessor);
+         SemaSYCL::isSyclType(Ty, SYCLTypeAttr::local_accessor) ||
+         SemaSYCL::isSyclType(Ty, SYCLTypeAttr::dynamic_local_accessor);
 }
 
 // FIXME: Accessor property lists should be modified to use compile-time
@@ -1151,7 +1152,8 @@ static QualType GetSYCLKernelObjectType(const FunctionDecl *KernelCaller) {
 /// \return the target of given SYCL accessor type
 static target getAccessTarget(QualType FieldTy,
                               const ClassTemplateSpecializationDecl *AccTy) {
-  if (SemaSYCL::isSyclType(FieldTy, SYCLTypeAttr::local_accessor))
+  if (SemaSYCL::isSyclType(FieldTy, SYCLTypeAttr::local_accessor) ||
+      SemaSYCL::isSyclType(FieldTy, SYCLTypeAttr::dynamic_local_accessor))
     return local;
 
   return static_cast<target>(
@@ -4796,7 +4798,13 @@ class SyclKernelIntHeaderCreator : public SyclKernelFieldHandler {
       int Dims = static_cast<int>(
           AccTy->getTemplateArgs()[1].getAsIntegral().getExtValue());
       int Info = getAccessTarget(FieldTy, AccTy) | (Dims << 11);
-      Header.addParamDesc(SYCLIntegrationHeader::kind_accessor, Info,
+
+      SYCLIntegrationHeader::kernel_param_kind_t ParamKind =
+          SemaSYCL::isSyclType(FieldTy, SYCLTypeAttr::dynamic_local_accessor)
+              ? SYCLIntegrationHeader::kind_dynamic_accessor
+              : SYCLIntegrationHeader::kind_accessor;
+
+      Header.addParamDesc(ParamKind, Info,
                           CurOffset +
                               offsetOf(RD, BC.getType()->getAsCXXRecordDecl()));
     } else if (SemaSYCL::isSyclType(FieldTy, SYCLTypeAttr::work_group_memory)) {
@@ -4822,8 +4830,12 @@ class SyclKernelIntHeaderCreator : public SyclKernelFieldHandler {
           AccTy->getTemplateArgs()[1].getAsIntegral().getExtValue());
       int Info = getAccessTarget(FieldTy, AccTy) | (Dims << 11);
 
-      Header.addParamDesc(SYCLIntegrationHeader::kind_accessor, Info,
-                          CurOffset + offsetOf(FD, FieldTy));
+      SYCLIntegrationHeader::kernel_param_kind_t ParamKind =
+          SemaSYCL::isSyclType(FieldTy, SYCLTypeAttr::dynamic_local_accessor)
+              ? SYCLIntegrationHeader::kind_dynamic_accessor
+              : SYCLIntegrationHeader::kind_accessor;
+
+      Header.addParamDesc(ParamKind, Info, CurOffset + offsetOf(FD, FieldTy));
     } else if (SemaSYCL::isSyclType(FieldTy, SYCLTypeAttr::stream)) {
       addParam(FD, FieldTy, SYCLIntegrationHeader::kind_stream);
     } else if (SemaSYCL::isSyclType(FieldTy, SYCLTypeAttr::work_group_memory)) {
@@ -6037,6 +6049,7 @@ static const char *paramKind2Str(KernelParamKind K) {
     CASE(pointer);
     CASE(work_group_memory);
     CASE(dynamic_work_group_memory);
+    CASE(dynamic_accessor);
   }
   return "<ERROR>";
 

clang/test/CodeGenSYCL/Inputs/sycl.hpp

@@ -440,6 +440,9 @@ local_accessor: public accessor<dataT,
 #ifdef __SYCL_DEVICE_ONLY__
   void __init(__attribute__((opencl_local)) dataT *Ptr, range<dimensions> AccessRange,
               range<dimensions> MemRange, id<dimensions> Offset) {}
+
+template <typename, int>
+  friend class dynamic_local_accessor;
 #endif
 };
 
@@ -693,6 +696,23 @@ __SYCL_TYPE(dynamic_work_group_memory) dynamic_work_group_memory {
   work_group_memory<DataT> LocalMem;
 };
 
+template <typename DataT, int Dimensions>
+class __attribute__((sycl_special_class))
+__SYCL_TYPE(dynamic_local_accessor) dynamic_local_accessor {
+public:
+  dynamic_local_accessor() = default;
+
+void __init(__attribute__((opencl_local)) DataT *Ptr,
+            range<Dimensions> AccessRange, range<Dimensions> range,
+            id<Dimensions> id) {
+  this->LocalMem.__init(Ptr, AccessRange, range, id);
+}
+  local_accessor<DataT, Dimensions> get() const { return LocalMem; }
+
+private:
+  local_accessor<DataT, Dimensions> LocalMem;
+};
+
 template <typename T, int dimensions = 1,
           typename AllocatorT = int /*fake type as AllocatorT is not used*/>
 class buffer {

clang/test/CodeGenSYCL/dynamic_local_accessor.cpp

@@ -0,0 +1,40 @@
+// RUN: %clang_cc1 -fsycl-is-device -triple spir64-unknown-unknown -disable-llvm-passes -emit-llvm %s -o %t.ll
+// RUN: FileCheck < %t.ll %s --check-prefix CHECK-IR
+// RUN: %clang_cc1 -fsycl-is-device -triple spir64-unknown-unknown -fsycl-int-header=%t.h %s
+// RUN: FileCheck < %t.h %s --check-prefix CHECK-INT-HEADER
+//
+// Tests for dynamic_local_accessor kernel parameter using the dummy implementation in Inputs/sycl.hpp.
+// The first two RUN commands verify that the init call is generated with the correct arguments in LLVM IR
+// and the second two RUN commands verify the contents of the integration header produced by the frontend.
+//
+// CHECK-IR: define dso_local spir_kernel void @
+// CHECK-IR-SAME: ptr addrspace(3) noundef align 4 [[PTR:%[a-zA-Z0-9_]+]]
+//
+// CHECK-IR: [[PTR]].addr = alloca ptr addrspace(3), align 8
+// CHECK-IR: [[PTR]].addr.ascast = addrspacecast ptr [[PTR]].addr to ptr addrspace(4)
+// CHECK-IR: store ptr addrspace(3) [[PTR]], ptr addrspace(4) [[PTR]].addr.ascast, align 8
+// CHECK-IR: [[PTR_LOAD:%[a-zA-Z0-9_]+]] = load ptr addrspace(3), ptr addrspace(4) [[PTR]].addr.ascast, align 8
+//
+// CHECK-IR: call spir_func void @{{.*}}__init{{.*}}(ptr addrspace(4) noundef align 4 dereferenceable_or_null(24) %{{[a-zA-Z0-9_]+}}, ptr addrspace(3) noundef %{{[a-zA-Z0-9_]+}}, ptr noundef byval(%"struct.sycl::_V1::range") align 4 %agg.{{[a-zA-Z0-9_]+}}.ascast.ascast, ptr noundef byval(%"struct.sycl::_V1::range") align 4 %agg.{{[a-zA-Z0-9_]+}}.ascast.ascast, ptr noundef byval(%"struct.sycl::_V1::id") align 4 %agg.{{[a-zA-Z0-9_]+}}.ascast.ascast) #{{[0-9_]+}}
+//
+// CHECK-INT-HEADER: const kernel_param_desc_t kernel_signatures[] = {
+// CHECK-INT-HEADER-NEXT: //--- _ZTSZZ4mainENKUlRN4sycl3_V17handlerEE_clES2_EUlNS0_4itemILi1EEEE_
+// CHECK-INT-HEADER-NEXT: { kernel_param_kind_t::kind_dynamic_accessor, 4064, 0 },
+// CHECK-INT-HEADER-EMPTY:
+// CHECK-INT-HEADER-NEXT: { kernel_param_kind_t::kind_invalid, -987654321, -987654321 },
+// CHECK-INT-HEADER-NEXT: };
+
+#include "Inputs/sycl.hpp"
+
+int main() {
+  sycl::queue Q;
+  sycl::dynamic_local_accessor<int, 1> dynLocalAcc;
+  Q.submit([&](sycl::handler &CGH) {
+    sycl::range<1> ndr;
+    CGH.parallel_for(ndr, [=](sycl::item<1> it) { 
+      auto localAcc = dynLocalAcc.get();
+      auto* ptr = &localAcc;
+});
+  });
+  return 0;
+}

sycl/include/sycl/accessor.hpp

@@ -227,6 +227,10 @@ template <typename DataT, int Dimensions = 1,
           typename PropertyListT = ext::oneapi::accessor_property_list<>>
 class accessor;
 
+namespace ext::oneapi::experimental {
+template <typename, int> class dynamic_local_accessor;
+}
+
 namespace detail {
 
 template <typename... Ts>
@@ -2638,6 +2642,8 @@ class __SYCL_EBO __SYCL_SPECIAL_CLASS __SYCL_TYPE(local_accessor) local_accessor
 
 private:
   friend class sycl::ext::intel::esimd::detail::AccessorPrivateProxy;
+  template <typename, int>
+  friend class ext::oneapi::experimental::dynamic_local_accessor;
 };
 
 template <typename DataT, int Dimensions = 1,

sycl/include/sycl/detail/kernel_desc.hpp

@@ -60,6 +60,7 @@ enum class kernel_param_kind_t {
   kind_stream = 5,
   kind_work_group_memory = 6,
   kind_dynamic_work_group_memory = 7,
+  kind_dynamic_accessor = 8,
   kind_invalid = 0xf, // not a valid kernel kind
 };