From 4a510b6ae620452b7c494a14391098328e6c1042 Mon Sep 17 00:00:00 2001
From: Victor Perez <victor.perez@codeplay.com>
Date: Wed, 31 Jan 2024 14:44:08 +0000
Subject: [PATCH] Address comments

---
 .../source/detail/scheduler/graph_builder.cpp | 10 +++++-----
 sycl/test-e2e/KernelFusion/non-kernel-cg.cpp  | 19 +++++++++----------
 2 files changed, 14 insertions(+), 15 deletions(-)

diff --git a/sycl/source/detail/scheduler/graph_builder.cpp b/sycl/source/detail/scheduler/graph_builder.cpp
index 66823c088463a..b3b4343b530ad 100644
--- a/sycl/source/detail/scheduler/graph_builder.cpp
+++ b/sycl/source/detail/scheduler/graph_builder.cpp
@@ -950,10 +950,10 @@ Scheduler::GraphBuildResult Scheduler::GraphBuilder::addCG(
   if (!NewCmd)
     throw runtime_error("Out of host memory", PI_ERROR_OUT_OF_HOST_MEMORY);
 
-  // Host tasks cannot participate in fusion. They take the regular route. If
-  // they create any requirement or event dependency on any of the kernels in
-  // the fusion list, this will lead to cancellation of the fusion in the
-  // GraphProcessor.
+  // Only device kernel command groups can participate in fusion. Otherwise,
+  // command groups take the regular route. If they create any requirement or
+  // event dependency on any of the kernels in the fusion list, this will lead
+  // to cancellation of the fusion in the GraphProcessor.
   auto QUniqueID = std::hash<sycl::detail::queue_impl *>()(Queue.get());
   if (isInFusionMode(QUniqueID)) {
     if (NewCmd->isFusable()) {
@@ -1016,7 +1016,7 @@ Scheduler::GraphBuildResult Scheduler::GraphBuilder::addCG(
       std::string s;
       std::stringstream ss(s);
       ss << "Not fusing '" << NewCmd->getTypeString()
-         << "' command group. Can only fuse device kernel command groups";
+         << "' command group. Can only fuse device kernel command groups.";
       printFusionWarning(ss.str());
     }
   }
diff --git a/sycl/test-e2e/KernelFusion/non-kernel-cg.cpp b/sycl/test-e2e/KernelFusion/non-kernel-cg.cpp
index bfdb61d9f7356..cc156c461548b 100644
--- a/sycl/test-e2e/KernelFusion/non-kernel-cg.cpp
+++ b/sycl/test-e2e/KernelFusion/non-kernel-cg.cpp
@@ -9,7 +9,6 @@
 
 // Test non-kernel device command groups are not fused
 
-#include "sycl/detail/pi.h"
 #include <sycl/sycl.hpp>
 
 using namespace sycl;
@@ -26,7 +25,7 @@ int main() {
   auto *src = malloc_device<float>(count, q);
 
   {
-    // CHECK: Not fusing 'copy acc to ptr' command group. Can only fuse device kernel command groups
+    // CHECK: Not fusing 'copy acc to ptr' command group. Can only fuse device kernel command groups.
     buffer<float> src(dataSize);
     std::shared_ptr<float> dst(new float[dataSize]);
     fw.start_fusion();
@@ -38,7 +37,7 @@ int main() {
   }
 
   {
-    // CHECK: Not fusing 'copy ptr to acc' command group. Can only fuse device kernel command groups
+    // CHECK: Not fusing 'copy ptr to acc' command group. Can only fuse device kernel command groups.
     buffer<float> dst(dataSize);
     std::shared_ptr<float> src(new float[dataSize]);
     fw.start_fusion();
@@ -50,7 +49,7 @@ int main() {
   }
 
   {
-    // CHECK: Not fusing 'copy acc to acc' command group. Can only fuse device kernel command groups
+    // CHECK: Not fusing 'copy acc to acc' command group. Can only fuse device kernel command groups.
     buffer<float> dst(dataSize);
     buffer<float> src(dataSize);
     fw.start_fusion();
@@ -63,14 +62,14 @@ int main() {
   }
 
   {
-    // CHECK: Not fusing 'barrier' command group. Can only fuse device kernel command groups
+    // CHECK: Not fusing 'barrier' command group. Can only fuse device kernel command groups.
     fw.start_fusion();
     q.submit([&](handler &cgh) { cgh.ext_oneapi_barrier(); });
     fw.complete_fusion();
   }
 
   {
-    // CHECK: Not fusing 'barrier waitlist' command group. Can only fuse device kernel command groups
+    // CHECK: Not fusing 'barrier waitlist' command group. Can only fuse device kernel command groups.
     buffer<float> dst(dataSize);
     buffer<float> src(dataSize);
     std::vector<event> event_list;
@@ -85,7 +84,7 @@ int main() {
   }
 
   {
-    // CHECK: Not fusing 'fill' command group. Can only fuse device kernel command groups
+    // CHECK: Not fusing 'fill' command group. Can only fuse device kernel command groups.
     buffer<float> dst(dataSize);
     fw.start_fusion();
     q.submit([&](handler &cgh) {
@@ -96,21 +95,21 @@ int main() {
   }
 
   {
-    // CHECK: Not fusing 'copy usm' command group. Can only fuse device kernel command groups
+    // CHECK: Not fusing 'copy usm' command group. Can only fuse device kernel command groups.
     fw.start_fusion();
     q.submit([&](handler &cgh) { cgh.memcpy(dst, src, count); });
     fw.complete_fusion();
   }
 
   {
-    // CHECK: Not fusing 'fill usm' command group. Can only fuse device kernel command groups
+    // CHECK: Not fusing 'fill usm' command group. Can only fuse device kernel command groups.
     fw.start_fusion();
     q.submit([&](handler &cgh) { cgh.fill(dst, Pattern, count); });
     fw.complete_fusion();
   }
 
   {
-    // CHECK: Not fusing 'prefetch usm' command group. Can only fuse device kernel command groups
+    // CHECK: Not fusing 'prefetch usm' command group. Can only fuse device kernel command groups.
     fw.start_fusion();
     q.submit([&](handler &cgh) { cgh.prefetch(dst, count); });
     fw.complete_fusion();