Move MKernelIsCooperative to KernelData

Author: vinser52

sycl/source/detail/handler_impl.hpp

@@ -106,7 +106,6 @@ class handler_impl {
 
   ur_kernel_cache_config_t MKernelCacheConfig = UR_KERNEL_CACHE_CONFIG_DEFAULT;
 
-  bool MKernelIsCooperative = false;
   uint32_t MKernelWorkGroupMemorySize = 0;
 
   // Extra information for bindless image copy

sycl/source/detail/kernel_data.hpp

@@ -115,6 +115,12 @@ class KernelData {
 
   bool usesAssert() const { return MDeviceKernelInfoPtr->usesAssert(); }
 
+  bool isCooperative() const { return MKernelIsCooperative; }
+
+  void setCooperative(bool IsCooperative) {
+    MKernelIsCooperative = IsCooperative;
+  }
+
   bool usesClusterLaunch() const { return MKernelUsesClusterLaunch; }
 
   template <int Dims_> void setClusterDimensions(sycl::range<Dims_> N) {
@@ -138,6 +144,7 @@ class KernelData {
   /// The list of arguments for the kernel.
   std::vector<detail::ArgDesc> MArgs;
 
+  bool MKernelIsCooperative = false;
   bool MKernelUsesClusterLaunch = false;
 
   /// Struct that encodes global size, local size, ...

sycl/source/handler.cpp

@@ -664,7 +664,7 @@ event handler::finalize() {
                          MKernel.get(), toKernelNameStrT(MKernelName),
                          *impl->MKernelData.getDeviceKernelInfoPtr(), RawEvents,
                          ResultEvent.get(), nullptr, impl->MKernelCacheConfig,
-                         impl->MKernelIsCooperative,
+                         impl->MKernelData.isCooperative(),
                          impl->MKernelData.usesClusterLaunch(),
                          impl->MKernelWorkGroupMemorySize, BinImage,
                          impl->MKernelData.getKernelFuncPtr(),
@@ -731,7 +731,7 @@ event handler::finalize() {
         toKernelNameStrT(MKernelName),
         *impl->MKernelData.getDeviceKernelInfoPtr(), std::move(MStreamStorage),
         std::move(impl->MAuxiliaryResources), getType(),
-        impl->MKernelCacheConfig, impl->MKernelIsCooperative,
+        impl->MKernelCacheConfig, impl->MKernelData.isCooperative(),
         impl->MKernelData.usesClusterLaunch(), impl->MKernelWorkGroupMemorySize,
         MCodeLoc));
     break;
@@ -2023,7 +2023,7 @@ void handler::setKernelCacheConfig(handler::StableKernelCacheConfig Config) {
 }
 
 void handler::setKernelIsCooperative(bool KernelIsCooperative) {
-  impl->MKernelIsCooperative = KernelIsCooperative;
+  impl->MKernelData.setCooperative(KernelIsCooperative);
 }
 
 #ifndef __INTEL_PREVIEW_BREAKING_CHANGES

sycl/unittests/program_manager/arg_mask/EliminatedArgMask.cpp

@@ -151,9 +151,9 @@ class MockHandler : public sycl::handler {
           std::move(impl->MKernelData).getArgs(), CGH->MKernelName.data(),
           *impl->MKernelData.getDeviceKernelInfoPtr(),
           std::move(CGH->MStreamStorage), std::move(impl->MAuxiliaryResources),
-          impl->MCGType, {}, impl->MKernelIsCooperative,
-          impl->MKernelUsesClusterLaunch, impl->MKernelWorkGroupMemorySize,
-          CGH->MCodeLoc));
+          impl->MCGType, {}, impl->MKernelData.isCooperative(),
+          impl->MKernelData.usesClusterLaunch(),
+          impl->MKernelWorkGroupMemorySize, CGH->MCodeLoc));
       break;
     }
     default:

sycl/unittests/scheduler/SchedulerTestUtils.hpp

@@ -309,7 +309,8 @@ class MockHandlerCustomFinalize : public MockHandler {
           std::move(impl->MKernelBundle), std::move(CGData), getArgs(),
           getKernelName(), *impl->MKernelData.getDeviceKernelInfoPtr(),
           getStreamStorage(), impl->MAuxiliaryResources, getType(), {},
-          impl->MKernelIsCooperative, impl->MKernelUsesClusterLaunch,
+          impl->MKernelData.isCooperative(),
+          impl->MKernelData.usesClusterLaunch(),
           impl->MKernelWorkGroupMemorySize, getCodeLoc()));
       break;
     }

sycl/unittests/scheduler/StreamInitDependencyOnHost.cpp

@@ -36,7 +36,8 @@ class MockHandlerStreamInit : public MockHandler {
           getArgs(), getKernelName(),
           *impl->MKernelData.getDeviceKernelInfoPtr(), getStreamStorage(),
           std::move(impl->MAuxiliaryResources), getType(), {},
-          impl->MKernelIsCooperative, impl->MKernelUsesClusterLaunch,
+          impl->MKernelData.isCooperative(),
+          impl->MKernelData.usesClusterLaunch(),
           impl->MKernelWorkGroupMemorySize, getCodeLoc()));
       break;
     }