Skip to content

Commit

Permalink
Address comments
Browse files Browse the repository at this point in the history
  • Loading branch information
victor-eds committed Feb 2, 2024
1 parent 96f687e commit 4a510b6
Show file tree
Hide file tree
Showing 2 changed files with 14 additions and 15 deletions.
10 changes: 5 additions & 5 deletions sycl/source/detail/scheduler/graph_builder.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -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()) {
Expand Down Expand Up @@ -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());
}
}
Expand Down
19 changes: 9 additions & 10 deletions sycl/test-e2e/KernelFusion/non-kernel-cg.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -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;
Expand All @@ -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();
Expand All @@ -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();
Expand All @@ -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();
Expand All @@ -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;
Expand All @@ -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) {
Expand All @@ -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();
Expand Down

0 comments on commit 4a510b6

Please sign in to comment.