@@ -408,42 +408,40 @@ class kernel_bundle_impl {
408408
409409 std::string UserArgs = syclex::detail::userArgsAsString (BuildOptions);
410410
411- std::vector<ur_device_handle_t > DeviceHandles;
412- std::transform (Devices.begin (), Devices.end (),
413- std::back_inserter (DeviceHandles), [](const device &Dev) {
414- return getSyclObjImpl (Dev)->getHandleRef ();
415- });
416-
417- std::vector<const uint8_t *> Binaries;
418- std::vector<size_t > Lengths;
419- std::vector<std::vector<std::vector<char >>> PersistentBinaries;
420- for (size_t i = 0 ; i < Devices.size (); i++) {
421- std::vector<std::vector<char >> BinProg =
422- PersistentDeviceCodeCache::getCompiledKernelFromDisc (
423- Devices[i], UserArgs, SourceStr);
424-
425- // exit if any device binary is missing
426- if (BinProg.empty ()) {
427- return false ;
428- }
429- PersistentBinaries.push_back (BinProg);
430-
431- Binaries.push_back ((uint8_t *)(BinProg[0 ].data ()));
432- Lengths.push_back (BinProg[0 ].size ());
411+ std::vector<ur_device_handle_t > DeviceHandles;
412+ std::transform (
413+ Devices.begin (), Devices.end (), std::back_inserter (DeviceHandles),
414+ [](const device &Dev) { return getSyclObjImpl (Dev)->getHandleRef (); });
415+
416+ std::vector<const uint8_t *> Binaries;
417+ std::vector<size_t > Lengths;
418+ std::vector<std::vector<std::vector<char >>> PersistentBinaries;
419+ for (size_t i = 0 ; i < Devices.size (); i++) {
420+ std::vector<std::vector<char >> BinProg =
421+ PersistentDeviceCodeCache::getCompiledKernelFromDisc (
422+ Devices[i], UserArgs, SourceStr);
423+
424+ // exit if any device binary is missing
425+ if (BinProg.empty ()) {
426+ return false ;
433427 }
428+ PersistentBinaries.push_back (BinProg);
434429
435- ur_program_properties_t Properties = {};
436- Properties.stype = UR_STRUCTURE_TYPE_PROGRAM_PROPERTIES;
437- Properties.pNext = nullptr ;
438- Properties.count = 0 ;
439- Properties.pMetadatas = nullptr ;
430+ Binaries.push_back ((uint8_t *)(BinProg[0 ].data ()));
431+ Lengths.push_back (BinProg[0 ].size ());
432+ }
440433
441- Adapter->call <UrApiKind::urProgramCreateWithBinary>(
442- ContextImpl->getHandleRef (), DeviceHandles.size (),
443- DeviceHandles.data (), Lengths.data (), Binaries.data (), &Properties,
444- &UrProgram);
434+ ur_program_properties_t Properties = {};
435+ Properties.stype = UR_STRUCTURE_TYPE_PROGRAM_PROPERTIES;
436+ Properties.pNext = nullptr ;
437+ Properties.count = 0 ;
438+ Properties.pMetadatas = nullptr ;
445439
446- return true ;
440+ Adapter->call <UrApiKind::urProgramCreateWithBinary>(
441+ ContextImpl->getHandleRef (), DeviceHandles.size (), DeviceHandles.data (),
442+ Lengths.data (), Binaries.data (), &Properties, &UrProgram);
443+
444+ return true ;
447445 }
448446
449447 std::shared_ptr<kernel_bundle_impl>
@@ -480,8 +478,8 @@ class kernel_bundle_impl {
480478 // if successful, the log is empty. if failed, throws an error with
481479 // the compilation log.
482480 std::vector<uint32_t > IPVersionVec (Devices.size ());
483- std::transform (DeviceVec.begin (), DeviceVec.end (), IPVersionVec. begin (),
484- [&](ur_device_handle_t d) {
481+ std::transform (DeviceVec.begin (), DeviceVec.end (),
482+ IPVersionVec. begin (), [&](ur_device_handle_t d) {
485483 uint32_t ipVersion = 0 ;
486484 Adapter->call <UrApiKind::urDeviceGetInfo>(
487485 d, UR_DEVICE_INFO_IP_VERSION, sizeof (uint32_t ),
@@ -511,23 +509,23 @@ class kernel_bundle_impl {
511509 " languages at this time" );
512510 }();
513511
514- Adapter->call <UrApiKind::urProgramCreateWithIL>(ContextImpl-> getHandleRef (),
515- spirv.data (), spirv.size (),
516- nullptr , &UrProgram);
517- // program created by urProgramCreateWithIL is implicitly retained.
518- if (UrProgram == nullptr )
519- throw sycl::exception (
520- sycl::make_error_code (errc::invalid),
521- " urProgramCreateWithIL resulted in a null program handle." );
522-
523- std::string XsFlags = extractXsFlags (BuildOptions);
524- auto Res = Adapter->call_nocheck <UrApiKind::urProgramBuildExp>(
525- UrProgram, DeviceVec.size (), DeviceVec.data (), XsFlags.c_str ());
526- if (Res == UR_RESULT_ERROR_UNSUPPORTED_FEATURE) {
527- Res = Adapter->call_nocheck <UrApiKind::urProgramBuild>(
528- ContextImpl->getHandleRef (), UrProgram, XsFlags.c_str ());
529- }
530- Adapter->checkUrResult <errc::build>(Res);
512+ Adapter->call <UrApiKind::urProgramCreateWithIL>(
513+ ContextImpl-> getHandleRef (), spirv.data (), spirv.size (), nullptr ,
514+ &UrProgram);
515+ // program created by urProgramCreateWithIL is implicitly retained.
516+ if (UrProgram == nullptr )
517+ throw sycl::exception (
518+ sycl::make_error_code (errc::invalid),
519+ " urProgramCreateWithIL resulted in a null program handle." );
520+
521+ std::string XsFlags = extractXsFlags (BuildOptions);
522+ auto Res = Adapter->call_nocheck <UrApiKind::urProgramBuildExp>(
523+ UrProgram, DeviceVec.size (), DeviceVec.data (), XsFlags.c_str ());
524+ if (Res == UR_RESULT_ERROR_UNSUPPORTED_FEATURE) {
525+ Res = Adapter->call_nocheck <UrApiKind::urProgramBuild>(
526+ ContextImpl->getHandleRef (), UrProgram, XsFlags.c_str ());
527+ }
528+ Adapter->checkUrResult <errc::build>(Res);
531529
532530 } // if(!FetchedFromCache)
533531
0 commit comments