|
| 1 | +//==----------- KernelFusion.cpp - Kernel Fusion scheduler unit tests ------==// |
| 2 | +// |
| 3 | +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. |
| 4 | +// See https://llvm.org/LICENSE.txt for license information. |
| 5 | +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception |
| 6 | +// |
| 7 | +//===----------------------------------------------------------------------===// |
| 8 | + |
| 9 | +#include "SchedulerTest.hpp" |
| 10 | +#include "SchedulerTestUtils.hpp" |
| 11 | + |
| 12 | +#include <helpers/PiMock.hpp> |
| 13 | +#include <helpers/ScopedEnvVar.hpp> |
| 14 | +#include <helpers/TestKernel.hpp> |
| 15 | + |
| 16 | +#include <vector> |
| 17 | + |
| 18 | +using namespace sycl; |
| 19 | +using EventImplPtr = std::shared_ptr<detail::event_impl>; |
| 20 | + |
| 21 | +template <typename T, int Dim> |
| 22 | +detail::Command *CreateTaskCommand(MockScheduler &MS, |
| 23 | + detail::QueueImplPtr DevQueue, |
| 24 | + buffer<T, Dim> &buf) { |
| 25 | + MockHandlerCustomFinalize MockCGH(DevQueue, false); |
| 26 | + |
| 27 | + auto acc = buf.get_access(static_cast<sycl::handler &>(MockCGH)); |
| 28 | + |
| 29 | + kernel_bundle KernelBundle = |
| 30 | + sycl::get_kernel_bundle<sycl::bundle_state::input>( |
| 31 | + DevQueue->get_context()); |
| 32 | + auto ExecBundle = sycl::build(KernelBundle); |
| 33 | + MockCGH.use_kernel_bundle(ExecBundle); |
| 34 | + MockCGH.single_task<TestKernel<>>([] {}); |
| 35 | + |
| 36 | + auto CmdGrp = MockCGH.finalize(); |
| 37 | + |
| 38 | + std::vector<detail::Command *> ToEnqueue; |
| 39 | + detail::Command *NewCmd = MS.addCG(std::move(CmdGrp), DevQueue, ToEnqueue); |
| 40 | + EXPECT_EQ(ToEnqueue.size(), 0u); |
| 41 | + return NewCmd; |
| 42 | +} |
| 43 | + |
| 44 | +bool CheckTestExecRequirements(const platform &plt) { |
| 45 | + if (plt.is_host()) { |
| 46 | + std::cout << "Not run due to host-only environment\n"; |
| 47 | + return false; |
| 48 | + } |
| 49 | + // This test only contains device image for SPIR-V capable devices. |
| 50 | + if (plt.get_backend() != sycl::backend::opencl && |
| 51 | + plt.get_backend() != sycl::backend::ext_oneapi_level_zero) { |
| 52 | + std::cout << "Only OpenCL and Level Zero are supported for this test\n"; |
| 53 | + return false; |
| 54 | + } |
| 55 | + return true; |
| 56 | +} |
| 57 | + |
| 58 | +bool containsCommand(detail::Command *Cmd, |
| 59 | + std::vector<detail::Command *> &List) { |
| 60 | + return std::find(List.begin(), List.end(), Cmd) != List.end(); |
| 61 | +} |
| 62 | + |
| 63 | +bool dependsOnViaDep(detail::Command *Dependent, detail::Command *Dependee) { |
| 64 | + return std::find_if(Dependent->MDeps.begin(), Dependent->MDeps.end(), |
| 65 | + [=](detail::DepDesc &Desc) { |
| 66 | + return Desc.MDepCommand == Dependee; |
| 67 | + }) != Dependent->MDeps.end(); |
| 68 | +} |
| 69 | + |
| 70 | +bool dependsOnViaEvent(detail::Command *Dependent, detail::Command *Dependee) { |
| 71 | + auto &DepEvents = Dependent->getPreparedDepsEvents(); |
| 72 | + return std::find_if(DepEvents.begin(), DepEvents.end(), |
| 73 | + [=](const EventImplPtr &Ev) { |
| 74 | + return Ev->getCommand() && Ev->getCommand() == Dependee; |
| 75 | + }) != DepEvents.end(); |
| 76 | +} |
| 77 | + |
| 78 | +TEST_F(SchedulerTest, CancelKernelFusion) { |
| 79 | + unittest::PiMock Mock; |
| 80 | + platform Plt = Mock.getPlatform(); |
| 81 | + if (!CheckTestExecRequirements(Plt)) |
| 82 | + return; |
| 83 | + |
| 84 | + queue QueueDev(context(Plt), default_selector_v); |
| 85 | + MockScheduler MS; |
| 86 | + |
| 87 | + detail::QueueImplPtr QueueDevImpl = detail::getSyclObjImpl(QueueDev); |
| 88 | + |
| 89 | + // Test scenario: Create four memory objects (buffers) and one command for |
| 90 | + // each memory object before starting fusion. Then start fusion, again adding |
| 91 | + // one command with a requirement for each of the memory objects. Then cancel |
| 92 | + // fusion and check for correct dependencies. |
| 93 | + |
| 94 | + buffer<int, 1> b1{range<1>{4}}; |
| 95 | + buffer<int, 1> b2{range<1>{4}}; |
| 96 | + buffer<int, 1> b3{range<1>{4}}; |
| 97 | + buffer<int, 1> b4{range<1>{4}}; |
| 98 | + |
| 99 | + auto *nonFusionCmd1 = CreateTaskCommand(MS, QueueDevImpl, b1); |
| 100 | + auto *nonFusionCmd2 = CreateTaskCommand(MS, QueueDevImpl, b2); |
| 101 | + auto *nonFusionCmd3 = CreateTaskCommand(MS, QueueDevImpl, b3); |
| 102 | + auto *nonFusionCmd4 = CreateTaskCommand(MS, QueueDevImpl, b4); |
| 103 | + |
| 104 | + MS.startFusion(QueueDevImpl); |
| 105 | + |
| 106 | + auto *fusionCmd1 = CreateTaskCommand(MS, QueueDevImpl, b1); |
| 107 | + auto *fusionCmd2 = CreateTaskCommand(MS, QueueDevImpl, b2); |
| 108 | + auto *fusionCmd3 = CreateTaskCommand(MS, QueueDevImpl, b3); |
| 109 | + auto *fusionCmd4 = CreateTaskCommand(MS, QueueDevImpl, b4); |
| 110 | + |
| 111 | + std::vector<detail::Command *> ToEnqueue; |
| 112 | + MS.cancelFusion(QueueDevImpl, ToEnqueue); |
| 113 | + |
| 114 | + // The list of commands filled by cancelFusion should contain the four |
| 115 | + // commands submitted while in fusion mode, plus the placeholder command. |
| 116 | + EXPECT_EQ(ToEnqueue.size(), 5u); |
| 117 | + EXPECT_TRUE(containsCommand(fusionCmd1, ToEnqueue)); |
| 118 | + EXPECT_TRUE(containsCommand(fusionCmd2, ToEnqueue)); |
| 119 | + EXPECT_TRUE(containsCommand(fusionCmd3, ToEnqueue)); |
| 120 | + EXPECT_TRUE(containsCommand(fusionCmd4, ToEnqueue)); |
| 121 | + |
| 122 | + // Each of the commands submitted while in fusion mode should have exactly one |
| 123 | + // dependency on the command not participating in fusion, but accessing the |
| 124 | + // same memory object. |
| 125 | + EXPECT_TRUE(dependsOnViaDep(fusionCmd1, nonFusionCmd1)); |
| 126 | + EXPECT_EQ(fusionCmd1->MDeps.size(), 1u); |
| 127 | + EXPECT_TRUE(dependsOnViaDep(fusionCmd2, nonFusionCmd2)); |
| 128 | + EXPECT_EQ(fusionCmd2->MDeps.size(), 1u); |
| 129 | + EXPECT_TRUE(dependsOnViaDep(fusionCmd3, nonFusionCmd3)); |
| 130 | + EXPECT_EQ(fusionCmd3->MDeps.size(), 1u); |
| 131 | + EXPECT_TRUE(dependsOnViaDep(fusionCmd4, nonFusionCmd4)); |
| 132 | + EXPECT_EQ(fusionCmd4->MDeps.size(), 1u); |
| 133 | + |
| 134 | + // There should be one placeholder command in the command list. |
| 135 | + auto FusionCmdIt = std::find_if( |
| 136 | + ToEnqueue.begin(), ToEnqueue.end(), [](detail::Command *Cmd) { |
| 137 | + return Cmd->getType() == sycl::_V1::detail::Command::FUSION; |
| 138 | + }); |
| 139 | + EXPECT_NE(FusionCmdIt, ToEnqueue.end()); |
| 140 | + |
| 141 | + // Check that the placeholder command has an event dependency on each of the |
| 142 | + // commands submitted while in fusion mode. |
| 143 | + auto *placeHolderCmd = |
| 144 | + static_cast<detail::KernelFusionCommand *>(*FusionCmdIt); |
| 145 | + EXPECT_EQ(placeHolderCmd->getPreparedDepsEvents().size(), 4u); |
| 146 | + EXPECT_TRUE(dependsOnViaEvent(placeHolderCmd, fusionCmd2)); |
| 147 | + EXPECT_TRUE(dependsOnViaEvent(placeHolderCmd, fusionCmd3)); |
| 148 | + EXPECT_TRUE(dependsOnViaEvent(placeHolderCmd, fusionCmd4)); |
| 149 | + EXPECT_TRUE(dependsOnViaEvent(placeHolderCmd, fusionCmd1)); |
| 150 | +} |
0 commit comments