@@ -2477,98 +2477,91 @@ static ur_result_t SetKernelParamsAndLaunch(
24772477 return Error;
24782478}
24792479
2480- ur_result_t enqueueImpCommandBufferKernel (
2481- context Ctx, DeviceImplPtr DeviceImpl,
2482- ur_exp_command_buffer_handle_t CommandBuffer,
2483- const CGExecKernel &CommandGroup,
2484- std::vector<ur_exp_command_buffer_sync_point_t > &SyncPoints,
2485- ur_exp_command_buffer_sync_point_t *OutSyncPoint,
2486- ur_exp_command_buffer_command_handle_t *OutCommand,
2487- const std::function<void *(Requirement *Req)> &getMemAllocationFunc) {
2488- auto ContextImpl = sycl::detail::getSyclObjImpl (Ctx);
2489- const sycl::detail::AdapterPtr &Adapter = ContextImpl->getAdapter ();
2490-
2491- const std::vector<std::weak_ptr<sycl::detail::CGExecKernel>>
2492- &AlternativeKernels = CommandGroup.MAlternativeKernels ;
2480+ namespace {
2481+ std::tuple<ur_kernel_handle_t , std::shared_ptr<device_image_impl>,
2482+ const KernelArgMask *>
2483+ getCGKernelInfo (const CGExecKernel &CommandGroup, ContextImplPtr ContextImpl,
2484+ DeviceImplPtr DeviceImpl,
2485+ std::vector<ur_kernel_handle_t > &UrKernelsToRelease,
2486+ std::vector<ur_program_handle_t > &UrProgramsToRelease) {
24932487
2494- // UR kernel and program for 'CommandGroup'
24952488 ur_kernel_handle_t UrKernel = nullptr ;
2496- ur_program_handle_t UrProgram = nullptr ;
2497-
2498- // Impl objects created when 'CommandGroup' is from a kernel bundle
2499- std::shared_ptr<kernel_impl> SyclKernelImpl = nullptr ;
25002489 std::shared_ptr<device_image_impl> DeviceImageImpl = nullptr ;
2501-
2502- // List of ur objects to be released after UR call
2503- std::vector<ur_kernel_handle_t > UrKernelsToRelease;
2504- std::vector<ur_program_handle_t > UrProgramsToRelease;
2505-
2506- auto Kernel = CommandGroup.MSyclKernel ;
2507- auto KernelBundleImplPtr = CommandGroup.MKernelBundle ;
25082490 const KernelArgMask *EliminatedArgMask = nullptr ;
25092491
25102492 // Use kernel_bundle if available unless it is interop.
25112493 // Interop bundles can't be used in the first branch, because the kernels
25122494 // in interop kernel bundles (if any) do not have kernel_id
25132495 // and can therefore not be looked up, but since they are self-contained
25142496 // they can simply be launched directly.
2515- if (KernelBundleImplPtr && !KernelBundleImplPtr->isInterop ()) {
2497+ if (auto KernelBundleImplPtr = CommandGroup.MKernelBundle ;
2498+ KernelBundleImplPtr && !KernelBundleImplPtr->isInterop ()) {
25162499 auto KernelName = CommandGroup.MKernelName ;
25172500 kernel_id KernelID =
25182501 detail::ProgramManager::getInstance ().getSYCLKernelID (KernelName);
2502+
25192503 kernel SyclKernel =
25202504 KernelBundleImplPtr->get_kernel (KernelID, KernelBundleImplPtr);
2521- SyclKernelImpl = detail::getSyclObjImpl (SyclKernel);
2505+
2506+ auto SyclKernelImpl = detail::getSyclObjImpl (SyclKernel);
25222507 UrKernel = SyclKernelImpl->getHandleRef ();
25232508 DeviceImageImpl = SyclKernelImpl->getDeviceImage ();
2524- UrProgram = DeviceImageImpl->get_ur_program_ref ();
25252509 EliminatedArgMask = SyclKernelImpl->getKernelArgMask ();
2526- } else if (Kernel != nullptr ) {
2510+ } else if (auto Kernel = CommandGroup. MSyclKernel ; Kernel != nullptr ) {
25272511 UrKernel = Kernel->getHandleRef ();
2528- UrProgram = Kernel->getProgramRef ();
25292512 EliminatedArgMask = Kernel->getKernelArgMask ();
25302513 } else {
2514+ ur_program_handle_t UrProgram = nullptr ;
25312515 std::tie (UrKernel, std::ignore, EliminatedArgMask, UrProgram) =
25322516 sycl::detail::ProgramManager::getInstance ().getOrCreateKernel (
25332517 ContextImpl, DeviceImpl, CommandGroup.MKernelName );
25342518 UrKernelsToRelease.push_back (UrKernel);
25352519 UrProgramsToRelease.push_back (UrProgram);
25362520 }
2521+ return std::make_tuple (UrKernel, DeviceImageImpl, EliminatedArgMask);
2522+ }
2523+ } // anonymous namespace
2524+
2525+ ur_result_t enqueueImpCommandBufferKernel (
2526+ context Ctx, DeviceImplPtr DeviceImpl,
2527+ ur_exp_command_buffer_handle_t CommandBuffer,
2528+ const CGExecKernel &CommandGroup,
2529+ std::vector<ur_exp_command_buffer_sync_point_t > &SyncPoints,
2530+ ur_exp_command_buffer_sync_point_t *OutSyncPoint,
2531+ ur_exp_command_buffer_command_handle_t *OutCommand,
2532+ const std::function<void *(Requirement *Req)> &getMemAllocationFunc) {
2533+ // List of ur objects to be released after UR call. We don't do anything
2534+ // with the ur_program_handle_t objects, but need to update their reference
2535+ // count.
2536+ std::vector<ur_kernel_handle_t > UrKernelsToRelease;
2537+ std::vector<ur_program_handle_t > UrProgramsToRelease;
2538+
2539+ ur_kernel_handle_t UrKernel = nullptr ;
2540+ std::shared_ptr<device_image_impl> DeviceImageImpl = nullptr ;
2541+ const KernelArgMask *EliminatedArgMask = nullptr ;
2542+
2543+ auto ContextImpl = sycl::detail::getSyclObjImpl (Ctx);
2544+ std::tie (UrKernel, DeviceImageImpl, EliminatedArgMask) =
2545+ getCGKernelInfo (CommandGroup, ContextImpl, DeviceImpl, UrKernelsToRelease,
2546+ UrProgramsToRelease);
25372547
25382548 // Build up the list of UR kernel handles that the UR command could be
25392549 // updated to use.
25402550 std::vector<ur_kernel_handle_t > AltUrKernels;
2551+ const std::vector<std::weak_ptr<sycl::detail::CGExecKernel>>
2552+ &AlternativeKernels = CommandGroup.MAlternativeKernels ;
25412553 for (const auto &AltCGKernelWP : AlternativeKernels) {
25422554 auto AltCGKernel = AltCGKernelWP.lock ();
25432555 assert (AltCGKernel != nullptr );
25442556
25452557 ur_kernel_handle_t AltUrKernel = nullptr ;
2546- if (auto KernelBundleImplPtr = AltCGKernel->MKernelBundle ;
2547- KernelBundleImplPtr && !KernelBundleImplPtr->isInterop ()) {
2548- auto KernelName = AltCGKernel->MKernelName ;
2549- kernel_id KernelID =
2550- detail::ProgramManager::getInstance ().getSYCLKernelID (KernelName);
2551- kernel SyclKernel =
2552- KernelBundleImplPtr->get_kernel (KernelID, KernelBundleImplPtr);
2553- AltUrKernel = detail::getSyclObjImpl (SyclKernel)->getHandleRef ();
2554- } else if (AltCGKernel->MSyclKernel != nullptr ) {
2555- AltUrKernel = Kernel->getHandleRef ();
2556- } else {
2557- ur_program_handle_t UrProgram = nullptr ;
2558- std::tie (AltUrKernel, std::ignore, std::ignore, UrProgram) =
2559- sycl::detail::ProgramManager::getInstance ().getOrCreateKernel (
2560- ContextImpl, DeviceImpl, AltCGKernel->MKernelName );
2561- UrKernelsToRelease.push_back (AltUrKernel);
2562- UrProgramsToRelease.push_back (UrProgram);
2563- }
2564-
2565- if (AltUrKernel != UrKernel) {
2566- // Don't include command-group 'CommandGroup' in the list to pass to UR,
2567- // as this will be used for the primary ur kernel parameter.
2568- AltUrKernels.push_back (AltUrKernel);
2569- }
2558+ std::tie (AltUrKernel, std::ignore, std::ignore) =
2559+ getCGKernelInfo (*AltCGKernel.get (), ContextImpl, DeviceImpl,
2560+ UrKernelsToRelease, UrProgramsToRelease);
2561+ AltUrKernels.push_back (AltUrKernel);
25702562 }
25712563
2564+ const sycl::detail::AdapterPtr &Adapter = ContextImpl->getAdapter ();
25722565 auto SetFunc = [&Adapter, &UrKernel, &DeviceImageImpl, &Ctx,
25732566 &getMemAllocationFunc](sycl::detail::ArgDesc &Arg,
25742567 size_t NextTrueIndex) {
0 commit comments