Skip to content

Commit

Permalink
[SYCL] Fix an issue when spec constants are set in a CG w/o a kernel (#…
Browse files Browse the repository at this point in the history 
…6826)

Fix an issue introduced by #6595.
Removing the version check also removed the implicit check for whether
we have a command group at all, which caused a CTS regression.
  • Loading branch information
@sergey-semenov
sergey-semenov authored Sep 27, 2022
1 parent 593ba5f commit b333cee
Show file tree
Hide file tree
Showing 2 changed files with 123 additions and 104 deletions.
212 changes: 108 additions & 104 deletions sycl/source/handler.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -93,122 +93,126 @@ event handler::finalize() {
return MLastEvent;
MIsFinalized = true;

std::shared_ptr<detail::kernel_bundle_impl> KernelBundleImpPtr = nullptr;
// If there were uses of set_specialization_constant build the kernel_bundle
KernelBundleImpPtr = getOrInsertHandlerKernelBundle(/*Insert=*/false);
if (KernelBundleImpPtr) {
// Make sure implicit non-interop kernel bundles have the kernel
if (!KernelBundleImpPtr->isInterop() &&
!MImpl->isStateExplicitKernelBundle()) {
kernel_id KernelID =
detail::ProgramManager::getInstance().getSYCLKernelID(MKernelName);
bool KernelInserted =
KernelBundleImpPtr->add_kernel(KernelID, MQueue->get_device());
// If kernel was not inserted and the bundle is in input mode we try
// building it and trying to find the kernel in executable mode
if (!KernelInserted &&
KernelBundleImpPtr->get_bundle_state() == bundle_state::input) {
auto KernelBundle =
const auto &type = getType();
if (type == detail::CG::Kernel) {
// If there were uses of set_specialization_constant build the kernel_bundle
std::shared_ptr<detail::kernel_bundle_impl> KernelBundleImpPtr =
getOrInsertHandlerKernelBundle(/*Insert=*/false);
if (KernelBundleImpPtr) {
// Make sure implicit non-interop kernel bundles have the kernel
if (!KernelBundleImpPtr->isInterop() &&
!MImpl->isStateExplicitKernelBundle()) {
kernel_id KernelID =
detail::ProgramManager::getInstance().getSYCLKernelID(MKernelName);
bool KernelInserted =
KernelBundleImpPtr->add_kernel(KernelID, MQueue->get_device());
// If kernel was not inserted and the bundle is in input mode we try
// building it and trying to find the kernel in executable mode
if (!KernelInserted &&
KernelBundleImpPtr->get_bundle_state() == bundle_state::input) {
auto KernelBundle =
detail::createSyclObjFromImpl<kernel_bundle<bundle_state::input>>(
KernelBundleImpPtr);
kernel_bundle<bundle_state::executable> ExecKernelBundle =
build(KernelBundle);
KernelBundleImpPtr = detail::getSyclObjImpl(ExecKernelBundle);
setHandlerKernelBundle(KernelBundleImpPtr);
KernelInserted =
KernelBundleImpPtr->add_kernel(KernelID, MQueue->get_device());
}
// If the kernel was not found in executable mode we throw an exception
if (!KernelInserted)
throw sycl::exception(make_error_code(errc::runtime),
"Failed to add kernel to kernel bundle.");
}

switch (KernelBundleImpPtr->get_bundle_state()) {
case bundle_state::input: {
// Underlying level expects kernel_bundle to be in executable state
kernel_bundle<bundle_state::executable> ExecBundle = build(
detail::createSyclObjFromImpl<kernel_bundle<bundle_state::input>>(
KernelBundleImpPtr);
kernel_bundle<bundle_state::executable> ExecKernelBundle =
build(KernelBundle);
KernelBundleImpPtr = detail::getSyclObjImpl(ExecKernelBundle);
KernelBundleImpPtr));
KernelBundleImpPtr = detail::getSyclObjImpl(ExecBundle);
setHandlerKernelBundle(KernelBundleImpPtr);
KernelInserted =
KernelBundleImpPtr->add_kernel(KernelID, MQueue->get_device());
break;
}
case bundle_state::executable:
// Nothing to do
break;
case bundle_state::object:
assert(0 && "Expected that the bundle is either in input or executable "
"states.");
break;
}
// If the kernel was not found in executable mode we throw an exception
if (!KernelInserted)
throw sycl::exception(make_error_code(errc::runtime),
"Failed to add kernel to kernel bundle.");
}

switch (KernelBundleImpPtr->get_bundle_state()) {
case bundle_state::input: {
// Underlying level expects kernel_bundle to be in executable state
kernel_bundle<bundle_state::executable> ExecBundle = build(
detail::createSyclObjFromImpl<kernel_bundle<bundle_state::input>>(
KernelBundleImpPtr));
KernelBundleImpPtr = detail::getSyclObjImpl(ExecBundle);
setHandlerKernelBundle(KernelBundleImpPtr);
break;
}
case bundle_state::executable:
// Nothing to do
break;
case bundle_state::object:
assert(0 && "Expected that the bundle is either in input or executable "
"states.");
break;
}
}
if (MRequirements.size() + MEvents.size() + MStreamStorage.size() == 0) {
// if user does not add a new dependency to the dependency graph, i.e.
// the graph is not changed, then this faster path is used to submit
// kernel bypassing scheduler and avoiding CommandGroup, Command objects
// creation.

const auto &type = getType();
if (type == detail::CG::Kernel &&
MRequirements.size() + MEvents.size() + MStreamStorage.size() == 0) {
// if user does not add a new dependency to the dependency graph, i.e.
// the graph is not changed, then this faster path is used to submit kernel
// bypassing scheduler and avoiding CommandGroup, Command objects creation.

std::vector<RT::PiEvent> RawEvents;
detail::EventImplPtr NewEvent;
RT::PiEvent *OutEvent = nullptr;

auto EnqueueKernel = [&]() {
// 'Result' for single point of return
pi_int32 Result = PI_ERROR_INVALID_VALUE;

if (MQueue->is_host()) {
MHostKernel->call(
MNDRDesc, (NewEvent) ? NewEvent->getHostProfilingInfo() : nullptr);
Result = PI_SUCCESS;
} else {
if (MQueue->getPlugin().getBackend() ==
backend::ext_intel_esimd_emulator) {
MQueue->getPlugin().call<detail::PiApiKind::piEnqueueKernelLaunch>(
nullptr, reinterpret_cast<pi_kernel>(MHostKernel->getPtr()),
MNDRDesc.Dims, &MNDRDesc.GlobalOffset[0], &MNDRDesc.GlobalSize[0],
&MNDRDesc.LocalSize[0], 0, nullptr, nullptr);
std::vector<RT::PiEvent> RawEvents;
detail::EventImplPtr NewEvent;
RT::PiEvent *OutEvent = nullptr;

auto EnqueueKernel = [&]() {
// 'Result' for single point of return
pi_int32 Result = PI_ERROR_INVALID_VALUE;

if (MQueue->is_host()) {
MHostKernel->call(MNDRDesc, (NewEvent)
? NewEvent->getHostProfilingInfo()
: nullptr);
Result = PI_SUCCESS;
} else {
Result = enqueueImpKernel(MQueue, MNDRDesc, MArgs, KernelBundleImpPtr,
MKernel, MKernelName, MOSModuleHandle,
RawEvents, OutEvent, nullptr);
if (MQueue->getPlugin().getBackend() ==
backend::ext_intel_esimd_emulator) {
MQueue->getPlugin().call<detail::PiApiKind::piEnqueueKernelLaunch>(
nullptr, reinterpret_cast<pi_kernel>(MHostKernel->getPtr()),
MNDRDesc.Dims, &MNDRDesc.GlobalOffset[0],
&MNDRDesc.GlobalSize[0], &MNDRDesc.LocalSize[0], 0, nullptr,
nullptr);
Result = PI_SUCCESS;
} else {
Result = enqueueImpKernel(
MQueue, MNDRDesc, MArgs, KernelBundleImpPtr, MKernel,
MKernelName, MOSModuleHandle, RawEvents, OutEvent, nullptr);
}
}
return Result;
};

bool DiscardEvent = false;
if (MQueue->has_discard_events_support()) {
// Kernel only uses assert if it's non interop one
bool KernelUsesAssert =
!(MKernel && MKernel->isInterop()) &&
detail::ProgramManager::getInstance().kernelUsesAssert(
MOSModuleHandle, MKernelName);
DiscardEvent = !KernelUsesAssert;
}
return Result;
};

bool DiscardEvent = false;
if (MQueue->has_discard_events_support()) {
// Kernel only uses assert if it's non interop one
bool KernelUsesAssert =
!(MKernel && MKernel->isInterop()) &&
detail::ProgramManager::getInstance().kernelUsesAssert(
MOSModuleHandle, MKernelName);
DiscardEvent = !KernelUsesAssert;
}

if (DiscardEvent) {
if (PI_SUCCESS != EnqueueKernel())
throw runtime_error("Enqueue process failed.",
PI_ERROR_INVALID_OPERATION);
} else {
NewEvent = std::make_shared<detail::event_impl>(MQueue);
NewEvent->setContextImpl(MQueue->getContextImplPtr());
NewEvent->setStateIncomplete();
OutEvent = &NewEvent->getHandleRef();

if (PI_SUCCESS != EnqueueKernel())
throw runtime_error("Enqueue process failed.",
PI_ERROR_INVALID_OPERATION);
else if (NewEvent->is_host() || NewEvent->getHandleRef() == nullptr)
NewEvent->setComplete();

MLastEvent = detail::createSyclObjFromImpl<event>(NewEvent);
if (DiscardEvent) {
if (PI_SUCCESS != EnqueueKernel())
throw runtime_error("Enqueue process failed.",
PI_ERROR_INVALID_OPERATION);
} else {
NewEvent = std::make_shared<detail::event_impl>(MQueue);
NewEvent->setContextImpl(MQueue->getContextImplPtr());
NewEvent->setStateIncomplete();
OutEvent = &NewEvent->getHandleRef();

if (PI_SUCCESS != EnqueueKernel())
throw runtime_error("Enqueue process failed.",
PI_ERROR_INVALID_OPERATION);
else if (NewEvent->is_host() || NewEvent->getHandleRef() == nullptr)
NewEvent->setComplete();

MLastEvent = detail::createSyclObjFromImpl<event>(NewEvent);
}
return MLastEvent;
}
return MLastEvent;
}

std::unique_ptr<detail::CG> CommandGroup;
Expand Down
15 changes: 15 additions & 0 deletions sycl/unittests/SYCL2020/SpecializationConstant.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -267,3 +267,18 @@ TEST(SpecializationConstant, UseKernelBundleAfterSetSpecConst) {
// Expected path
}
}

TEST(SpecializationConstant, NoKernel) {
sycl::unittest::PiMock Mock;
sycl::platform Plt = Mock.getPlatform();

const sycl::device Dev = Plt.get_devices()[0];
sycl::context Ctx{Dev};
sycl::queue Queue{Ctx, Dev};

Queue.submit([&](sycl::handler &CGH) {
int ExpectedValue = 42;
CGH.set_specialization_constant<SpecConst1>(ExpectedValue);
EXPECT_EQ(CGH.get_specialization_constant<SpecConst1>(), ExpectedValue);
});
}

0 comments on commit b333cee

Please sign in to comment.