@@ -468,59 +468,60 @@ class kernel_bundle_impl {
468468 }
469469
470470 if (!FetchedFromCache) {
471- if (Language == syclex::source_language::sycl_jit) {
472- const auto &SourceStr = std::get<std::string>(this ->Source );
473- const auto &Img = syclex::detail::SYCL_JIT_to_SPIRV (
474- SourceStr, IncludePairs, BuildOptions, LogPtr,
475- RegisteredKernelNames);
476- UrProgram = ProgramManager::getInstance ().createURProgram (Img, MContext,
477- MDevices);
478- } else {
479- const auto spirv = [&]() -> std::vector<uint8_t > {
480- if (Language == syclex::source_language::opencl) {
481- // if successful, the log is empty. if failed, throws an error with
482- // the compilation log.
483- std::vector<uint32_t > IPVersionVec (Devices.size ());
484- std::transform (DeviceVec.begin (), DeviceVec.end (),
485- IPVersionVec.begin (), [&](ur_device_handle_t d) {
486- uint32_t ipVersion = 0 ;
487- Adapter->call <UrApiKind::urDeviceGetInfo>(
488- d, UR_DEVICE_INFO_IP_VERSION, sizeof (uint32_t ),
489- &ipVersion, nullptr );
490- return ipVersion;
491- });
492- return syclex::detail::OpenCLC_to_SPIRV (*SourceStrPtr, IPVersionVec,
493- BuildOptions, LogPtr);
494- }
495- if (Language == syclex::source_language::spirv) {
496- const auto &SourceBytes =
497- std::get<std::vector<std::byte>>(this ->Source );
498- std::vector<uint8_t > Result (SourceBytes.size ());
499- std::transform (SourceBytes.cbegin (), SourceBytes.cend (),
500- Result.begin (),
501- [](std::byte B) { return static_cast <uint8_t >(B); });
502- return Result;
503- }
504- if (Language == syclex::source_language::sycl) {
505- return syclex::detail::SYCL_to_SPIRV (*SourceStrPtr, IncludePairs,
506- BuildOptions, LogPtr,
507- RegisteredKernelNames);
508- }
509- throw sycl::exception (
510- make_error_code (errc::invalid),
511- " SYCL C++, OpenCL C and SPIR-V are the only supported "
512- " languages at this time"
513- }();
514-
515- Adapter->call <UrApiKind::urProgramCreateWithIL>(
516- ContextImpl->getHandleRef (), spirv.data (), spirv.size (), nullptr ,
517- &UrProgram);
518- // program created by urProgramCreateWithIL is implicitly retained.
519- if (UrProgram == nullptr )
520- throw sycl::exception (
521- sycl::make_error_code (errc::invalid),
522- " urProgramCreateWithIL resulted in a null program handle."
523- }
471+ const auto spirv = [&]() -> std::vector<uint8_t > {
472+ if (Language == syclex::source_language::opencl) {
473+ // if successful, the log is empty. if failed, throws an error with
474+ // the compilation log.
475+ std::vector<uint32_t > IPVersionVec (Devices.size ());
476+ std::transform (DeviceVec.begin (), DeviceVec.end (),
477+ IPVersionVec.begin (), [&](ur_device_handle_t d) {
478+ uint32_t ipVersion = 0 ;
479+ Adapter->call <UrApiKind::urDeviceGetInfo>(
480+ d, UR_DEVICE_INFO_IP_VERSION, sizeof (uint32_t ),
481+ &ipVersion, nullptr );
482+ return ipVersion;
483+ });
484+ return syclex::detail::OpenCLC_to_SPIRV (*SourceStrPtr, IPVersionVec,
485+ BuildOptions, LogPtr);
486+ }
487+ if (Language == syclex::source_language::spirv) {
488+ const auto &SourceBytes =
489+ std::get<std::vector<std::byte>>(this ->Source );
490+ std::vector<uint8_t > Result (SourceBytes.size ());
491+ std::transform (SourceBytes.cbegin (), SourceBytes.cend (),
492+ Result.begin (),
493+ [](std::byte B) { return static_cast <uint8_t >(B); });
494+ return Result;
495+ }
496+ if (Language == syclex::source_language::sycl) {
497+ return syclex::detail::SYCL_to_SPIRV (*SourceStrPtr, IncludePairs,
498+ BuildOptions, LogPtr,
499+ RegisteredKernelNames);
500+ }
501+ if (Language == syclex::source_language::sycl_jit) {
502+ auto *Binaries = syclex::detail::SYCL_JIT_to_SPIRV (
503+ *SourceStrPtr, IncludePairs, BuildOptions, LogPtr,
504+ RegisteredKernelNames);
505+ assert (Binaries->NumDeviceBinaries == 1 &&
506+ " Device code splitting is not yet supported"
507+ return std::vector<uint8_t >(Binaries->DeviceBinaries ->BinaryStart ,
508+ Binaries->DeviceBinaries ->BinaryEnd );
509+ }
510+ throw sycl::exception (
511+ make_error_code (errc::invalid),
512+ " SYCL C++, OpenCL C and SPIR-V are the only supported "
513+ " languages at this time"
514+ }();
515+
516+ Adapter->call <UrApiKind::urProgramCreateWithIL>(
517+ ContextImpl->getHandleRef (), spirv.data (), spirv.size (), nullptr ,
518+ &UrProgram);
519+ // program created by urProgramCreateWithIL is implicitly retained.
520+ if (UrProgram == nullptr )
521+ throw sycl::exception (
522+ sycl::make_error_code (errc::invalid),
523+ " urProgramCreateWithIL resulted in a null program handle."
524+
524525 } // if(!FetchedFromCache)
525526
526527 std::string XsFlags = extractXsFlags (BuildOptions);
0 commit comments