[SYCL][Graph] Implement dynamic_work_group_memory for SYCL-Graphs

clang/include/clang/Basic/Attr.td

@@ -1606,12 +1606,12 @@ 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",
+                           ["accessor", "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",
+                           ["accessor", "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

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

clang/lib/Sema/SemaSYCL.cpp

@@ -2090,7 +2090,9 @@ class SyclKernelFieldChecker : public SyclKernelFieldHandler {
   }
 
   bool handleSyclSpecialType(ParmVarDecl *PD, QualType ParamTy) final {
-    if (!SemaSYCL::isSyclType(ParamTy, SYCLTypeAttr::work_group_memory)) {
+    if (!SemaSYCL::isSyclType(ParamTy, SYCLTypeAttr::work_group_memory) &&
+        !SemaSYCL::isSyclType(ParamTy,
+                              SYCLTypeAttr::dynamic_work_group_memory)) {
       Diag.Report(PD->getLocation(), diag::err_bad_kernel_param_type)
           << ParamTy;
       IsInvalid = true;
@@ -2246,7 +2248,8 @@ class SyclKernelUnionChecker : public SyclKernelFieldHandler {
   }
 
   bool handleSyclSpecialType(ParmVarDecl *PD, QualType ParamTy) final {
-    if (!SemaSYCL::isSyclType(ParamTy, SYCLTypeAttr::work_group_memory))
+    if (!SemaSYCL::isSyclType(ParamTy, SYCLTypeAttr::work_group_memory) &&
+        !SemaSYCL::isSyclType(ParamTy, SYCLTypeAttr::dynamic_work_group_memory))
       unsupportedFreeFunctionParamType(); // TODO
     return true;
   }
@@ -3032,7 +3035,9 @@ class SyclKernelDeclCreator : public SyclKernelFieldHandler {
   }
 
   bool handleSyclSpecialType(ParmVarDecl *PD, QualType ParamTy) final {
-    if (SemaSYCL::isSyclType(ParamTy, SYCLTypeAttr::work_group_memory)) {
+    if (SemaSYCL::isSyclType(ParamTy, SYCLTypeAttr::work_group_memory) ||
+        SemaSYCL::isSyclType(ParamTy,
+                             SYCLTypeAttr::dynamic_work_group_memory)) {
       const auto *RecordDecl = ParamTy->getAsCXXRecordDecl();
       assert(RecordDecl && "The type must be a RecordDecl");
       CXXMethodDecl *InitMethod = getMethodByName(RecordDecl, InitMethodName);
@@ -4544,7 +4549,9 @@ class FreeFunctionKernelBodyCreator : public SyclKernelFieldHandler {
   // TODO: Revisit this approach once https://github.com/intel/llvm/issues/16061
   // is closed.
   bool handleSyclSpecialType(ParmVarDecl *PD, QualType ParamTy) final {
-    if (SemaSYCL::isSyclType(ParamTy, SYCLTypeAttr::work_group_memory)) {
+    if (SemaSYCL::isSyclType(ParamTy, SYCLTypeAttr::work_group_memory) ||
+        SemaSYCL::isSyclType(ParamTy,
+                             SYCLTypeAttr::dynamic_work_group_memory)) {
       const auto *RecordDecl = ParamTy->getAsCXXRecordDecl();
       AccessSpecifier DefaultConstructorAccess;
       auto DefaultConstructor =
@@ -4823,6 +4830,10 @@ class SyclKernelIntHeaderCreator : public SyclKernelFieldHandler {
     } else if (SemaSYCL::isSyclType(FieldTy, SYCLTypeAttr::work_group_memory)) {
       addParam(FieldTy, SYCLIntegrationHeader::kind_work_group_memory,
                offsetOf(RD, BC.getType()->getAsCXXRecordDecl()));
+    } else if (SemaSYCL::isSyclType(FieldTy,
+                                    SYCLTypeAttr::dynamic_work_group_memory)) {
+      addParam(FieldTy, SYCLIntegrationHeader::kind_dynamic_work_group_memory,
+               offsetOf(RD, BC.getType()->getAsCXXRecordDecl()));
     }
     return true;
   }
@@ -4846,6 +4857,10 @@ class SyclKernelIntHeaderCreator : public SyclKernelFieldHandler {
     } else if (SemaSYCL::isSyclType(FieldTy, SYCLTypeAttr::work_group_memory)) {
       addParam(FieldTy, SYCLIntegrationHeader::kind_work_group_memory,
                offsetOf(FD, FieldTy));
+    } else if (SemaSYCL::isSyclType(FieldTy,
+                                    SYCLTypeAttr::dynamic_work_group_memory)) {
+      addParam(FieldTy, SYCLIntegrationHeader::kind_dynamic_work_group_memory,
+               offsetOf(FD, FieldTy));
     } else if (SemaSYCL::isSyclType(FieldTy, SYCLTypeAttr::sampler) ||
                SemaSYCL::isSyclType(FieldTy, SYCLTypeAttr::annotated_ptr) ||
                SemaSYCL::isSyclType(FieldTy, SYCLTypeAttr::annotated_arg)) {
@@ -4870,6 +4885,10 @@ class SyclKernelIntHeaderCreator : public SyclKernelFieldHandler {
   bool handleSyclSpecialType(ParmVarDecl *PD, QualType ParamTy) final {
     if (SemaSYCL::isSyclType(ParamTy, SYCLTypeAttr::work_group_memory))
       addParam(PD, ParamTy, SYCLIntegrationHeader::kind_work_group_memory);
+    else if (SemaSYCL::isSyclType(ParamTy,
+                                  SYCLTypeAttr::dynamic_work_group_memory))
+      addParam(PD, ParamTy,
+               SYCLIntegrationHeader::kind_dynamic_work_group_memory);
     else
       unsupportedFreeFunctionParamType(); // TODO
     return true;
@@ -5993,6 +6012,7 @@ static const char *paramKind2Str(KernelParamKind K) {
     CASE(specialization_constants_buffer);
     CASE(pointer);
     CASE(work_group_memory);
+    CASE(dynamic_work_group_memory);
   }
   return "<ERROR>";
 

clang/test/CodeGenSYCL/Inputs/sycl.hpp

@@ -667,6 +667,19 @@ __SYCL_TYPE(work_group_memory) work_group_memory {
   __attribute((opencl_local)) DataT *Ptr;
 };
 
+template <typename DataT>
+class __attribute__((sycl_special_class))
+__SYCL_TYPE(dynamic_work_group_memory) dynamic_work_group_memory {
+public:
+  dynamic_work_group_memory() = default;
+
+  void __init(__attribute((opencl_local)) DataT *Ptr) { this->LocalMem.__init(Ptr); }
+  work_group_memory<DataT> get() const { return LocalMem; }
+
+private:
+  work_group_memory<DataT> LocalMem;
+};
+
 template <typename T, int dimensions = 1,
           typename AllocatorT = int /*fake type as AllocatorT is not used*/>
 class buffer {

clang/test/CodeGenSYCL/dynamic_work_group_memory.cpp

@@ -0,0 +1,39 @@
+// 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_work_group_memory 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 8 dereferenceable_or_null(8) %{{[a-zA-Z0-9_]+}}, ptr addrspace(3) noundef [[PTR_LOAD]])
+//
+// 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_work_group_memory, {{[4,8]}}, 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_work_group_memory<int> dynMem;
+  Q.submit([&](sycl::handler &CGH) {
+    sycl::range<1> ndr;
+    CGH.parallel_for(ndr, [=](sycl::item<1> it) { 
+      auto localMem = dynMem.get();
+      int *ptr = &localMem; });
+  });
+  return 0;
+}

clang/test/CodeGenSYCL/free_function_int_header.cpp

@@ -2,7 +2,8 @@
 // RUN: FileCheck -input-file=%t.h %s
 // 
 // This test checks integration header contents for free functions with scalar,
-// pointer, non-decomposed struct parameters and work group memory parameters.
+// pointer, non-decomposed struct parameters, work group memory parameters and
+// dynamic work group memory parameters.
 
 #include "mock_properties.hpp"
 #include "sycl.hpp"
@@ -101,6 +102,11 @@ __attribute__((sycl_device))
 void ff_8(sycl::work_group_memory<int>) {
 }
 
+__attribute__((sycl_device))
+[[__sycl_detail__::add_ir_attributes_function("sycl-nd-range-kernel", 0)]] 
+void ff_9(sycl::dynamic_work_group_memory<int>) {
+}
+
 
 // CHECK:      const char* const kernel_names[] = {
 // CHECK-NEXT:   {{.*}}__sycl_kernel_ff_2Piii
@@ -112,6 +118,7 @@ void ff_8(sycl::work_group_memory<int>) {
 // CHECK-NEXT:   {{.*}}__sycl_kernel_ff_6I3Agg7DerivedEvT_T0_i
 // CHECK-NEXT:   {{.*}}__sycl_kernel_ff_7ILi3EEv16KArgWithPtrArrayIXT_EE
 // CHECK-NEXT:   {{.*}}__sycl_kernel_ff_8N4sycl3_V117work_group_memoryIiEE
+// CHECK-NEXT:   {{.*}}__sycl_kernel_ff_9N4sycl3_V125dynamic_work_group_memoryIiEE
 // CHECK-NEXT:   ""
 // CHECK-NEXT: };
 
@@ -158,6 +165,9 @@ void ff_8(sycl::work_group_memory<int>) {
 // CHECK:  //--- _Z18__sycl_kernel_ff_8N4sycl3_V117work_group_memoryIiEE
 // CHECK-NEXT:  { kernel_param_kind_t::kind_work_group_memory, 8, 0 },
 
+// CHECK:  //--- _Z18__sycl_kernel_ff_9N4sycl3_V125dynamic_work_group_memoryIiEE
+// CHECK-NEXT:  { kernel_param_kind_t::kind_dynamic_work_group_memory, 8, 0 },
+
 // CHECK:        { kernel_param_kind_t::kind_invalid, -987654321, -987654321 },
 // CHECK-NEXT: };
 
@@ -324,6 +334,26 @@ void ff_8(sycl::work_group_memory<int>) {
 // CHECK-NEXT: };
 // CHECK-NEXT: }
 
+// CHECK:  // Definition of _Z18__sycl_kernel_ff_9N4sycl3_V125dynamic_work_group_memoryIiEE as a free function kernel 
+//
+// CHECK: Forward declarations of kernel and its argument types:
+// CHECK: template <typename DataT> class dynamic_work_group_memory;
+
+// CHECK: void ff_9(sycl::dynamic_work_group_memory<int>);
+// CHECK-NEXT: static constexpr auto __sycl_shim10() {
+// CHECK-NEXT: return (void (*)(class sycl::dynamic_work_group_memory<int>))ff_9;
+// CHECK-NEXT: }
+// CHECK-NEXT: namespace sycl {
+// CHECK-NEXT: template <>
+// CHECK-NEXT: struct ext::oneapi::experimental::is_kernel<__sycl_shim10()> {
+// CHECK-NEXT: static constexpr bool value = true;
+// CHECK-NEXT: };
+// CHECK-NEXT: template <>
+// CHECK-NEXT: struct ext::oneapi::experimental::is_single_task_kernel<__sycl_shim10()> {
+// CHECK-NEXT: static constexpr bool value = true;
+// CHECK-NEXT: };
+// CHECK-NEXT: }
+
 // CHECK: #include <sycl/kernel_bundle.hpp>
 
 // CHECK: Definition of kernel_id of _Z18__sycl_kernel_ff_2Piii
@@ -397,3 +427,11 @@ void ff_8(sycl::work_group_memory<int>) {
 // CHECK-NEXT: return sycl::detail::get_kernel_id_impl(std::string_view{"_Z18__sycl_kernel_ff_8N4sycl3_V117work_group_memoryIiEE"});
 // CHECK-NEXT: }
 // CHECK-NEXT: }
+//
+// CHECK: // Definition of kernel_id of _Z18__sycl_kernel_ff_9N4sycl3_V125dynamic_work_group_memoryIiEE
+// CHECK-NEXT: namespace sycl {
+// CHECK-NEXT: template <>
+// CHECK-NEXT: kernel_id ext::oneapi::experimental::get_kernel_id<__sycl_shim10()>() {
+// CHECK-NEXT:   return sycl::detail::get_kernel_id_impl(std::string_view{"_Z18__sycl_kernel_ff_9N4sycl3_V125dynamic_work_group_memoryIiEE"});
+// CHECK-NEXT: }
+// CHECK-NEXT: }

clang/test/CodeGenSYCL/free_function_kernel_params.cpp

@@ -1,7 +1,7 @@
 // RUN: %clang_cc1 -internal-isystem %S/Inputs -fsycl-is-device -triple spir64 \
 // RUN: -emit-llvm %s -o - | FileCheck %s
 // This test checks parameter IR generation for free functions with parameters
-// of non-decomposed struct type and work group memory type.
+// of non-decomposed struct type, work group memory type and dynamic work group memory type.
 
 #include "sycl.hpp"
 
@@ -71,3 +71,16 @@ void ff_7(sycl::work_group_memory<int> mem) {
 // CHECK-NEXT:  [[REGISTER:%[a-zA-Z0-9_]+]] = load ptr addrspace(3), ptr addrspace(4) %__arg_Ptr.addr.ascast, align 8
 // CHECK-NEXT:  call spir_func void @{{.*}}work_group_memory{{.*}}__init{{.*}}(ptr addrspace(4) noundef align 8 dereferenceable_or_null(8) %mem.ascast, ptr addrspace(3) noundef [[REGISTER]])
 
+__attribute__((sycl_device))
+[[__sycl_detail__::add_ir_attributes_function("sycl-nd-range-kernel", 0)]]
+void ff_7(sycl::dynamic_work_group_memory<int> DynMem) {
+}
+
+// CHECK:  define dso_local spir_kernel void @{{.*}}__sycl_kernel_ff_7{{.*}}(ptr addrspace(3) noundef align 4 %__arg_Ptr)
+// CHECK:  %__arg_Ptr.addr = alloca ptr addrspace(3), align 8
+// CHECK-NEXT: %DynMem = alloca %"class.sycl::_V1::dynamic_work_group_memory", align 8 
+// CHECK:  %__arg_Ptr.addr.ascast = addrspacecast ptr %__arg_Ptr.addr to ptr addrspace(4)
+// CHECK-NEXT:  %DynMem.ascast = addrspacecast ptr %DynMem to ptr addrspace(4)
+// CHECK:  store ptr addrspace(3) %__arg_Ptr, ptr addrspace(4) %__arg_Ptr.addr.ascast, align 8
+// CHECK-NEXT:  [[REGISTER:%[a-zA-Z0-9_]+]] = load ptr addrspace(3), ptr addrspace(4) %__arg_Ptr.addr.ascast, align 8
+// CHECK-NEXT:  call spir_func void @{{.*}}dynamic_work_group_memory{{.*}}__init{{.*}}(ptr addrspace(4) noundef align 8 dereferenceable_or_null(8) %DynMem.ascast, ptr addrspace(3) noundef [[REGISTER]]) 

clang/test/SemaSYCL/Inputs/sycl.hpp

@@ -452,11 +452,9 @@ class __SYCL_TYPE(multi_ptr) multi_ptr<T, AS, access::decorated::legacy> {
 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:
+  // Default constructor for objects later initialized with __init member.
+  work_group_memory() = default;
   work_group_memory(handler &CGH) {}
 
   void __init(__attribute((opencl_local)) DataT *Ptr) { this->Ptr = Ptr; }
@@ -465,6 +463,19 @@ __SYCL_TYPE(work_group_memory) work_group_memory {
   __attribute((opencl_local)) DataT *Ptr;
 };
 
+template <typename DataT>
+class __attribute__((sycl_special_class))
+__SYCL_TYPE(dynamic_work_group_memory) dynamic_work_group_memory {
+public:
+  dynamic_work_group_memory() = default;
+
+  void __init(__attribute((opencl_local)) DataT *Ptr) { this->LocalMem.__init(Ptr); }
+  work_group_memory<DataT> get() const { return LocalMem; }
+
+private:
+  work_group_memory<DataT> LocalMem;
+};
+
 namespace ext {
 namespace oneapi {
 namespace experimental {

clang/test/SemaSYCL/Inputs/sycl/detail/kernel_desc.hpp

@@ -19,6 +19,7 @@ namespace detail {
     kind_specialization_constants_buffer = 4,
     kind_stream = 5,
     kind_work_group_memory = 6,
+    kind_dynamic_work_group_memory = 7,
     kind_invalid = 0xf, // not a valid kernel kind
   };
 

clang/test/SemaSYCL/free_function_kernel_params.cpp

@@ -1,7 +1,8 @@
 // RUN: %clang_cc1 -internal-isystem %S/Inputs -fsycl-is-device -ast-dump \
 // RUN: %s -o - | FileCheck %s
 // This test checks parameter rewriting for free functions with parameters
-// of type scalar, pointer, non-decomposed struct and work group memory.
+// of type scalar, pointer, non-decomposed struct, work group memory and 
+// dynamic work group memory.
 
 #include "sycl.hpp"
 
@@ -191,3 +192,23 @@ void ff_7(sycl::work_group_memory<int> mem) {
 // CHECK-NEXT: ImplicitCastExpr {{.*}} 'void (*)(sycl::work_group_memory<int>)' <FunctionToPointerDecay>
 // CHECK-NEXT: DeclRefExpr {{.*}} 'void (sycl::work_group_memory<int>)' lvalue Function {{.*}} 'ff_7' 'void (sycl::work_group_memory<int>)'
 // CHECK-NEXT: DeclRefExpr {{.*}} 'sycl::work_group_memory<int>' Var {{.*}} 'mem' 'sycl::work_group_memory<int>'
+
+__attribute__((sycl_device))
+[[__sycl_detail__::add_ir_attributes_function("sycl-nd-range-kernel", 0)]]
+void ff_8(sycl::dynamic_work_group_memory<int> DynMem) {
+}
+// CHECK: FunctionDecl {{.*}}__sycl_kernel{{.*}}'void (__local int *)'
+// CHECK-NEXT: ParmVarDecl {{.*}} used __arg_Ptr '__local int *'
+// CHECK-NEXT: CompoundStmt
+// CHECK-NEXT: DeclStmt
+// CHECK-NEXT: VarDecl {{.*}} used DynMem 'sycl::dynamic_work_group_memory<int>' callinit
+// CHECK-NEXT: CXXConstructExpr {{.*}} 'sycl::dynamic_work_group_memory<int>' 'void () noexcept'
+// CHECK-NEXT: CXXMemberCallExpr {{.*}} 'void'
+// CHECK-NEXT: MemberExpr {{.*}} 'void (__local int *)' lvalue .__init
+// CHECK-NEXT: DeclRefExpr {{.*}} 'sycl::dynamic_work_group_memory<int>' Var {{.*}} 'DynMem' 'sycl::dynamic_work_group_memory<int>'
+// CHECK-NEXT: ImplicitCastExpr {{.*}} '__local int *' <LValueToRValue>
+// CHECK-NEXT: DeclRefExpr {{.*}} '__local int *' lvalue ParmVar {{.*}} '__arg_Ptr' '__local int *'
+// CHECK-NEXT: CallExpr {{.*}} 'void'
+// CHECK-NEXT: ImplicitCastExpr {{.*}} 'void (*)(sycl::dynamic_work_group_memory<int>)' <FunctionToPointerDecay>
+// CHECK-NEXT: DeclRefExpr {{.*}} 'void (sycl::dynamic_work_group_memory<int>)' lvalue Function {{.*}} 'ff_8' 'void (sycl::dynamic_work_group_memory<int>)'
+// CHECK-NEXT: DeclRefExpr {{.*}} 'sycl::dynamic_work_group_memory<int>' Var {{.*}} 'DynMem' 'sycl::dynamic_work_group_memory<int>'

sycl-jit/common/include/Kernel.h

@@ -60,6 +60,7 @@ enum class ParameterKind : uint32_t {
   SpecConstBuffer = 4,
   Stream = 5,
   WorkGroupMemory = 6,
+  DynamicWorkGroupMemory = 7,
   Invalid = 0xF,
 };
 

sycl/include/sycl/detail/kernel_desc.hpp

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