Added support for free function kernels

Author: Konrad Kusiak

clang/include/clang/Sema/SemaSYCL.h

@@ -667,7 +667,7 @@ class SemaSYCL : public SemaBase {
   // Used to check whether the function represented by FD is a SYCL
   // free function kernel or not.
   bool isFreeFunction(const FunctionDecl *FD);
-
+  
   StmtResult BuildSYCLKernelCallStmt(FunctionDecl *FD, CompoundStmt *Body);
 };
 

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;
   }
@@ -4847,9 +4858,9 @@ class SyclKernelIntHeaderCreator : public SyclKernelFieldHandler {
       addParam(FieldTy, SYCLIntegrationHeader::kind_work_group_memory,
                offsetOf(FD, FieldTy));
     } else if (SemaSYCL::isSyclType(FieldTy,
-			            SYCLTypeAttr::dynamic_work_group_memory)) {
+                                    SYCLTypeAttr::dynamic_work_group_memory)) {
       addParam(FieldTy, SYCLIntegrationHeader::kind_dynamic_work_group_memory,
-	       offsetOf(FD, FieldTy));
+               offsetOf(FD, FieldTy));
     } else if (SemaSYCL::isSyclType(FieldTy, SYCLTypeAttr::sampler) ||
                SemaSYCL::isSyclType(FieldTy, SYCLTypeAttr::annotated_ptr) ||
                SemaSYCL::isSyclType(FieldTy, SYCLTypeAttr::annotated_arg)) {
@@ -4874,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;

sycl/include/sycl/ext/oneapi/experimental/graph.hpp

@@ -8,7 +8,6 @@
 
 #pragma once
 
-#include <cstddef>
 #include <sycl/accessor.hpp>               // for detail::AccessorBaseHost
 #include <sycl/context.hpp>                // for context
 #include <sycl/detail/export.hpp>          // for __SYCL_EXPORT
@@ -21,7 +20,8 @@
 #include <sycl/device.hpp> // for device
 #include <sycl/ext/oneapi/experimental/detail/properties/graph_properties.hpp> // for graph properties classes
 #include <sycl/ext/oneapi/experimental/work_group_memory.hpp> // for dynamic_work_group_memory
-#include <sycl/nd_range.hpp>                   // for range, nd_range
+#include <sycl/ext/oneapi/properties/properties.hpp> // for empty_properties_t
+#include <sycl/nd_range.hpp>                         // for range, nd_range
 #include <sycl/properties/property_traits.hpp> // for is_property, is_property_of
 #include <sycl/property_list.hpp>              // for property_list
 
@@ -49,6 +49,7 @@ enum class graph_state {
 // Forward declare ext::oneapi::experimental classes
 template <graph_state State> class command_graph;
 class raw_kernel_arg;
+template <typename, typename> class work_group_memory;
 
 namespace detail {
 // List of sycl features and extensions which are not supported by graphs. Used
@@ -503,6 +504,7 @@ class command_graph<graph_state::executable>
 namespace detail {
 class __SYCL_EXPORT dynamic_parameter_base {
 public:
+  dynamic_parameter_base() = default;
   dynamic_parameter_base(
       sycl::ext::oneapi::experimental::command_graph<graph_state::modifiable>
           Graph);
@@ -548,12 +550,13 @@ class dynamic_work_group_memory_base
 public:
   dynamic_work_group_memory_base() = default;
   dynamic_work_group_memory_base(
-      experimental::command_graph<graph_state::modifiable> Graph, size_t Size)
-      :
+      [[maybe_unused]] experimental::command_graph<graph_state::modifiable>
+          Graph,
+      [[maybe_unused]] size_t Size)
 #ifndef __SYCL_DEVICE_ONLY__
-        dynamic_parameter_base(Graph),
+      : dynamic_parameter_base(Graph), BufferSize(Size)
 #endif
-        BufferSize(Size) {
+  {
   }
 
 private:
@@ -565,23 +568,23 @@ class dynamic_work_group_memory_base
 };
 } // namespace detail
 
-template <typename DataT,
-          typename = std::enable_if_t<detail::is_unbounded_array_v<DataT>>>
+template <typename DataT, typename PropertyListT = empty_properties_t>
 class __SYCL_SPECIAL_CLASS
 __SYCL_TYPE(dynamic_work_group_memory) dynamic_work_group_memory
     : public detail::dynamic_work_group_memory_base {
-private:
-  work_group_memory<DataT> WorkGroupMem;
-
-#ifdef __SYCL_DEVICE_ONLY__
-  using value_type = std::remove_all_extents_t<DataT>;
-  using decoratedPtr = typename sycl::detail::DecoratedType<
-      value_type, access::address_space::local_space>::type *;
+public:
+  // Check that DataT is an unbounded array type.
+  static_assert(std::is_array_v<DataT> && std::extent_v<DataT, 0> == 0);
+  static_assert(std::is_same_v<PropertyListT, empty_properties_t>);
 
-  void __init(decoratedPtr Ptr) { this->WorkGroupMem.__init(Ptr); }
-#endif
+  // 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.
+  dynamic_work_group_memory() = default;
 
-public:
   /// Constructs a new dynamic_work_group_memory object.
   /// @param Graph The graph associated with this object.
   /// @param Num Number of elements in the unbounded array DataT.
@@ -590,24 +593,35 @@ __SYCL_TYPE(dynamic_work_group_memory) dynamic_work_group_memory
       : detail::dynamic_work_group_memory_base(
             Graph, Num * sizeof(std::remove_extent_t<DataT>)) {}
 
+  work_group_memory<DataT, PropertyListT> get() const {
+#ifndef __SYCL_DEVICE_ONLY__
+    throw sycl::exception(sycl::make_error_code(errc::invalid),
+                          "Error: dynamic_work_group_memory::get() can be only "
+                          "called on the device!");
+#endif
+    return WorkGroupMem;
+  }
+
   /// Updates on the host this dynamic_work_group_memory and all registered
   /// nodes with a new buffer size.
   /// @param Num The new number of elements in the unbounded array.
-  void update(size_t Num) {
+  void update([[maybe_unused]] size_t Num) {
 #ifndef __SYCL_DEVICE_ONLY__
     detail::dynamic_parameter_base::updateWorkGroupMem(
         Num * sizeof(std::remove_extent_t<DataT>));
 #endif
   }
-  work_group_memory<DataT> get() const { return WorkGroupMem; }
 
-  // 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.
-  dynamic_work_group_memory() = default;
+private:
+  work_group_memory<DataT, PropertyListT> WorkGroupMem;
+
+#ifdef __SYCL_DEVICE_ONLY__
+  using value_type = std::remove_all_extents_t<DataT>;
+  using decoratedPtr = typename sycl::detail::DecoratedType<
+      value_type, access::address_space::local_space>::type *;
+
+  void __init(decoratedPtr Ptr) { this->WorkGroupMem.__init(Ptr); }
+#endif
 };
 
 template <typename ValueT>
@@ -684,4 +698,14 @@ struct hash<sycl::ext::oneapi::experimental::dynamic_parameter<ValueT>> {
     return std::hash<decltype(ID)>()(ID);
   }
 };
+
+template <typename DataT>
+struct hash<sycl::ext::oneapi::experimental::dynamic_work_group_memory<DataT>> {
+  size_t operator()(
+      const sycl::ext::oneapi::experimental::dynamic_work_group_memory<DataT>
+          &DynWorkGroupMem) const {
+    auto ID = sycl::detail::getSyclObjImpl(DynWorkGroupMem)->getID();
+    return std::hash<decltype(ID)>()(ID);
+  }
+};
 } // namespace std

sycl/include/sycl/ext/oneapi/experimental/work_group_memory.hpp

@@ -7,9 +7,9 @@
 
 #pragma once
 
-#include "sycl/ext/oneapi/experimental/graph.hpp"
 #include <sycl/access/access.hpp>
 #include <sycl/detail/defines.hpp>
+#include <sycl/ext/oneapi/experimental/graph.hpp>
 #include <sycl/ext/oneapi/properties/properties.hpp>
 #include <sycl/multi_ptr.hpp>
 

sycl/include/sycl/handler.hpp

@@ -161,8 +161,8 @@ __SYCL_EXPORT void *async_malloc_from_pool(sycl::handler &h, size_t size,
 
 namespace ext::oneapi::experimental::detail {
 class graph_impl;
-class dynamic_work_group_memory_base;
 class dynamic_parameter_base;
+class dynamic_work_group_memory_base;
 } // namespace ext::oneapi::experimental::detail
 namespace detail {
 
@@ -687,7 +687,7 @@ class __SYCL_EXPORT handler {
   void
   setArgHelper(int ArgIndex,
                ext::oneapi::experimental::detail::dynamic_work_group_memory_base
-                   &DynWorkGroupMemParam);
+                   &DynWorkGroupBase);
 
   // setArgHelper for the raw_kernel_arg extension type.
   void setArgHelper(int ArgIndex,
@@ -1889,13 +1889,15 @@ class __SYCL_EXPORT handler {
   }
 
   // set_arg for graph dynamic_work_group_memory
-  template <typename DataT>
-  void set_arg(int argIndex,
-               ext::oneapi::experimental::dynamic_work_group_memory<DataT>
-                   &dynWorkGroupMem) {
+  template <typename DataT, typename PropertyListT =
+                                ext::oneapi::experimental::empty_properties_t>
+  void set_arg(
+      int argIndex,
+      ext::oneapi::experimental::dynamic_work_group_memory<DataT, PropertyListT>
+          &dynWorkGroupMem) {
     ext::oneapi::experimental::detail::dynamic_work_group_memory_base
-        &dynWorkGroupMemImpl = dynWorkGroupMem;
-    setArgHelper(argIndex, dynWorkGroupMemImpl);
+        &dynWorkGroupBase = dynWorkGroupMem;
+    setArgHelper(argIndex, dynWorkGroupBase);
   }
 
   // set_arg for the raw_kernel_arg extension type.

sycl/source/detail/graph_impl.hpp

@@ -1588,7 +1588,7 @@ class dynamic_parameter_impl {
   // Dynamic command-groups which will be updated
   std::vector<DynamicCGInfo> MDynCGs;
 
-  std::shared_ptr<graph_impl> MGraph;
+  std::weak_ptr<graph_impl> MGraph;
   std::vector<std::byte> MValueStorage;
 
 private:

sycl/source/detail/handler_impl.hpp

@@ -199,11 +199,6 @@ class handler_impl {
   std::vector<std::shared_ptr<detail::work_group_memory_impl>>
       MWorkGroupMemoryObjects;
 
-  /// List of dynamic work group memory objects associated with this handler
-  std::vector<std::shared_ptr<
-      ext::oneapi::experimental::detail::dynamic_work_group_memory_base>>
-      MDynWorkGroupMemoryParams;
-
   /// Potential event mode for the result event of the command.
   ext::oneapi::experimental::event_mode_enum MEventMode =
       ext::oneapi::experimental::event_mode_enum::none;

sycl/source/handler.cpp

@@ -7,7 +7,6 @@
 //===----------------------------------------------------------------------===//
 
 #include "sycl/detail/helpers.hpp"
-#include "sycl/ext/oneapi/experimental/graph.hpp"
 #include "ur_api.h"
 #include <algorithm>
 
@@ -34,6 +33,7 @@
 #include <sycl/info/info_desc.hpp>
 #include <sycl/stream.hpp>
 
+#include "sycl/ext/oneapi/experimental/graph.hpp"
 #include <sycl/ext/oneapi/bindless_images_memory.hpp>
 #include <sycl/ext/oneapi/experimental/work_group_memory.hpp>
 #include <sycl/ext/oneapi/memcpy2d.hpp>
@@ -1058,18 +1058,14 @@ void handler::setArgHelper(int ArgIndex, detail::work_group_memory_impl &Arg) {
 void handler::setArgHelper(
     int ArgIndex,
     ext::oneapi::experimental::detail::dynamic_work_group_memory_base
-        &DynWorkGroupMemParam) {
+        &DynWorkGroupBase) {
 
-  impl->MDynWorkGroupMemoryParams.push_back(
-      std::make_shared<
-          ext::oneapi::experimental::detail::dynamic_work_group_memory_base>(
-          DynWorkGroupMemParam));
   addArg(detail::kernel_param_kind_t::kind_dynamic_work_group_memory,
-         impl->MDynWorkGroupMemoryParams.back().get(), 0, ArgIndex);
+         &DynWorkGroupBase, 0, ArgIndex);
 
   // Register the dynamic parameter with the handler for later association
   // with the node being added
-  registerDynamicParameter(DynWorkGroupMemParam, ArgIndex);
+  registerDynamicParameter(DynWorkGroupBase, ArgIndex);
 }
 
 // The argument can take up more space to store additional information about
@@ -2102,7 +2098,7 @@ void handler::registerDynamicParameter(
   }
 
   auto Paraimpl = detail::getSyclObjImpl(DynamicParamBase);
-  if (Paraimpl->MGraph != this->impl->MGraph) {
+  if (Paraimpl->MGraph.lock() != this->impl->MGraph) {
     throw sycl::exception(
         make_error_code(errc::invalid),
         "Cannot use a Dynamic Parameter with a node associated with a graph "