Move MKernelWorkGroupMemorySize to KernelData

Author: vinser52

sycl/source/detail/handler_impl.hpp

@@ -106,8 +106,6 @@ class handler_impl {
 
   ur_kernel_cache_config_t MKernelCacheConfig = UR_KERNEL_CACHE_CONFIG_DEFAULT;
 
-  uint32_t MKernelWorkGroupMemorySize = 0;
-
   // Extra information for bindless image copy
   ur_image_desc_t MSrcImageDesc = {};
   ur_image_desc_t MDstImageDesc = {};

sycl/source/detail/kernel_data.hpp

@@ -128,6 +128,14 @@ class KernelData {
     MNDRDesc.setClusterDimensions(N);
   }
 
+  uint32_t getKernelWorkGroupMemorySize() const {
+    return MKernelWorkGroupMemorySize;
+  }
+
+  void setKernelWorkGroupMemorySize(uint32_t Size) {
+    MKernelWorkGroupMemorySize = Size;
+  }
+
   void processArg(void *Ptr, const detail::kernel_param_kind_t &Kind,
                   const int Size, const size_t Index, size_t &IndexShift,
                   bool IsKernelCreatedFromSource);
@@ -146,6 +154,7 @@ class KernelData {
 
   bool MKernelIsCooperative = false;
   bool MKernelUsesClusterLaunch = false;
+  uint32_t MKernelWorkGroupMemorySize = 0;
 
   /// Struct that encodes global size, local size, ...
   detail::NDRDescT MNDRDesc;

sycl/source/handler.cpp

@@ -666,8 +666,8 @@ event handler::finalize() {
                          ResultEvent.get(), nullptr, impl->MKernelCacheConfig,
                          impl->MKernelData.isCooperative(),
                          impl->MKernelData.usesClusterLaunch(),
-                         impl->MKernelWorkGroupMemorySize, BinImage,
-                         impl->MKernelData.getKernelFuncPtr(),
+                         impl->MKernelData.getKernelWorkGroupMemorySize(),
+                         BinImage, impl->MKernelData.getKernelFuncPtr(),
                          impl->MKernelData.getKernelNumArgs(),
                          impl->MKernelData.getKernelParamDescGetter(),
                          impl->MKernelData.hasSpecialCaptures());
@@ -732,8 +732,8 @@ event handler::finalize() {
         *impl->MKernelData.getDeviceKernelInfoPtr(), std::move(MStreamStorage),
         std::move(impl->MAuxiliaryResources), getType(),
         impl->MKernelCacheConfig, impl->MKernelData.isCooperative(),
-        impl->MKernelData.usesClusterLaunch(), impl->MKernelWorkGroupMemorySize,
-        MCodeLoc));
+        impl->MKernelData.usesClusterLaunch(),
+        impl->MKernelData.getKernelWorkGroupMemorySize(), MCodeLoc));
     break;
   }
   case detail::CGType::CopyAccToPtr:
@@ -2068,7 +2068,7 @@ void handler::setKernelClusterLaunch(sycl::range<1> ClusterSize) {
 void handler::setKernelWorkGroupMem(size_t Size) {
   throwIfGraphAssociated<syclex::detail::UnsupportedGraphFeatures::
                              sycl_ext_oneapi_work_group_scratch_memory>();
-  impl->MKernelWorkGroupMemorySize = Size;
+  impl->MKernelData.setKernelWorkGroupMemorySize(Size);
 }
 
 void handler::ext_oneapi_graph(

sycl/unittests/program_manager/arg_mask/EliminatedArgMask.cpp

@@ -153,7 +153,7 @@ class MockHandler : public sycl::handler {
           std::move(CGH->MStreamStorage), std::move(impl->MAuxiliaryResources),
           impl->MCGType, {}, impl->MKernelData.isCooperative(),
           impl->MKernelData.usesClusterLaunch(),
-          impl->MKernelWorkGroupMemorySize, CGH->MCodeLoc));
+          impl->MKernelData.getKernelWorkGroupMemorySize(), CGH->MCodeLoc));
       break;
     }
     default:

sycl/unittests/scheduler/SchedulerTestUtils.hpp

@@ -311,7 +311,7 @@ class MockHandlerCustomFinalize : public MockHandler {
           getStreamStorage(), impl->MAuxiliaryResources, getType(), {},
           impl->MKernelData.isCooperative(),
           impl->MKernelData.usesClusterLaunch(),
-          impl->MKernelWorkGroupMemorySize, getCodeLoc()));
+          impl->MKernelData.getKernelWorkGroupMemorySize(), getCodeLoc()));
       break;
     }
     case sycl::detail::CGType::CodeplayHostTask: {

sycl/unittests/scheduler/StreamInitDependencyOnHost.cpp

@@ -38,7 +38,7 @@ class MockHandlerStreamInit : public MockHandler {
           std::move(impl->MAuxiliaryResources), getType(), {},
           impl->MKernelData.isCooperative(),
           impl->MKernelData.usesClusterLaunch(),
-          impl->MKernelWorkGroupMemorySize, getCodeLoc()));
+          impl->MKernelData.getKernelWorkGroupMemorySize(), getCodeLoc()));
       break;
     }
     default: