[SYCL] Reduce kernel submission overhead (#4524)

This patch slightly reduces the number of instructions
for each submission:

    handler.hpp executes isKernelLambdaCallableWithKernelHandler
    at compile time.
    queue_impl.hpp stores the result of
    has_property<property::queue::in_order>() in a member variable
    and stores getType() in a local variable.
    scheduler.cpp stores the result of getType() in a local variable and
    uses it in all the following if-conditions.

Signed-off-by: Alexander Flegontov [email protected]

Author: alexanderfle

sycl/include/CL/sycl/handler.hpp

@@ -560,9 +560,10 @@ class __SYCL_EXPORT handler {
   template <typename KernelName, typename KernelType, int Dims,
             typename LambdaArgType>
   void StoreLambda(KernelType KernelFunc) {
-    if (detail::isKernelLambdaCallableWithKernelHandler<KernelType,
-                                                        LambdaArgType>() &&
-        MIsHost) {
+    constexpr bool IsCallableWithKernelHandler =
+        detail::isKernelLambdaCallableWithKernelHandler<KernelType,
+                                                        LambdaArgType>();
+    if (IsCallableWithKernelHandler && MIsHost) {
       throw cl::sycl::feature_not_supported(
           "kernel_handler is not yet supported by host device.",
           PI_INVALID_OPERATION);
@@ -590,8 +591,7 @@ class __SYCL_EXPORT handler {
 
     // If the kernel lambda is callable with a kernel_handler argument, manifest
     // the associated kernel handler.
-    if (detail::isKernelLambdaCallableWithKernelHandler<KernelType,
-                                                        LambdaArgType>()) {
+    if constexpr (IsCallableWithKernelHandler) {
       getOrInsertHandlerKernelBundle(/*Insert=*/true);
     }
   }

sycl/source/detail/queue_impl.hpp

@@ -91,7 +91,8 @@ class queue_impl {
              const async_handler &AsyncHandler, const property_list &PropList)
       : MDevice(Device), MContext(Context), MAsyncHandler(AsyncHandler),
         MPropList(PropList), MHostQueue(MDevice->is_host()),
-        MAssertHappenedBuffer(range<1>{1}) {
+        MAssertHappenedBuffer(range<1>{1}),
+        MIsInorder(has_property<property::queue::in_order>()) {
     if (!Context->hasDevice(Device))
       throw cl::sycl::invalid_parameter_error(
           "Queue cannot be constructed with the given context and device "
@@ -114,8 +115,9 @@ class queue_impl {
   /// \param AsyncHandler is a SYCL asynchronous exception handler.
   queue_impl(RT::PiQueue PiQueue, const ContextImplPtr &Context,
              const async_handler &AsyncHandler)
-      : MContext(Context), MAsyncHandler(AsyncHandler), MHostQueue(false),
-        MAssertHappenedBuffer(range<1>{1}) {
+      : MContext(Context), MAsyncHandler(AsyncHandler), MPropList(),
+        MHostQueue(false), MAssertHappenedBuffer(range<1>{1}),
+        MIsInorder(has_property<property::queue::in_order>()) {
 
     MQueues.push_back(pi::cast<RT::PiQueue>(PiQueue));
 
@@ -434,15 +436,15 @@ class queue_impl {
     // Scheduler will later omit events, that are not required to execute tasks.
     // Host and interop tasks, however, are not submitted to low-level runtimes
     // and require separate dependency management.
-    if (has_property<property::queue::in_order>() &&
-        (Handler.getType() == CG::CGTYPE::CodeplayHostTask ||
-         Handler.getType() == CG::CGTYPE::CodeplayInteropTask))
+    const CG::CGTYPE Type = Handler.getType();
+    if (MIsInorder && (Type == CG::CGTYPE::CodeplayHostTask ||
+                       Type == CG::CGTYPE::CodeplayInteropTask))
       Handler.depends_on(MLastEvent);
 
     event Event;
 
     if (PostProcess) {
-      bool IsKernel = Handler.getType() == CG::Kernel;
+      bool IsKernel = Type == CG::Kernel;
       bool KernelUsesAssert = false;
       if (IsKernel)
         KernelUsesAssert =
@@ -456,7 +458,7 @@ class queue_impl {
     } else
       Event = Handler.finalize();
 
-    if (has_property<property::queue::in_order>())
+    if (MIsInorder)
       MLastEvent = Event;
 
     addEvent(Event);
@@ -520,6 +522,7 @@ class queue_impl {
   buffer<AssertHappened, 1> MAssertHappenedBuffer;
 
   event MLastEvent;
+  const bool MIsInorder;
 };
 
 } // namespace detail

sycl/source/detail/scheduler/scheduler.cpp

@@ -71,12 +71,11 @@ void Scheduler::waitForRecordToFinish(MemObjRecord *Record,
 EventImplPtr Scheduler::addCG(std::unique_ptr<detail::CG> CommandGroup,
                               QueueImplPtr Queue) {
   EventImplPtr NewEvent = nullptr;
-  const bool IsKernel = CommandGroup->getType() == CG::Kernel;
+  const CG::CGTYPE Type = CommandGroup->getType();
   std::vector<Command *> AuxiliaryCmds;
-  const bool IsHostKernel = CommandGroup->getType() == CG::RunOnHostIntel;
   std::vector<StreamImplPtr> Streams;
 
-  if (IsKernel) {
+  if (Type == CG::Kernel) {
     Streams = ((CGExecKernel *)CommandGroup.get())->getStreams();
     // Stream's flush buffer memory is mainly initialized in stream's __init
     // method. However, this method is not available on host device.
@@ -146,7 +145,7 @@ EventImplPtr Scheduler::addCG(std::unique_ptr<detail::CG> CommandGroup,
     acquireWriteLock(Lock);
 
     Command *NewCmd = nullptr;
-    switch (CommandGroup->getType()) {
+    switch (Type) {
     case CG::UpdateHost:
       NewCmd = MGraphBuilder.addCGUpdateHost(std::move(CommandGroup),
                                              DefaultHostQueue, AuxiliaryCmds);
@@ -172,7 +171,7 @@ EventImplPtr Scheduler::addCG(std::unique_ptr<detail::CG> CommandGroup,
 
     auto CleanUp = [&]() {
       if (NewCmd && (NewCmd->MDeps.size() == 0 && NewCmd->MUsers.size() == 0)) {
-        if (IsHostKernel)
+        if (Type == CG::RunOnHostIntel)
           static_cast<ExecCGCommand *>(NewCmd)->releaseCG();
 
         NewEvent->setCommand(nullptr);