@@ -251,8 +251,8 @@ Expected<OffloadFile> getInputBitcodeLibrary(StringRef Input) {
251251 Image.StringData [" arch"
252252 Image.Image = std::move (*ImageOrError);
253253
254- std::unique_ptr<MemoryBuffer> Binary =
255- MemoryBuffer::getMemBufferCopy ( OffloadBinary::write (Image));
254+ std::unique_ptr<MemoryBuffer> Binary = MemoryBuffer::getMemBufferCopy (
255+ OffloadBinary::write (Image), Image. Image -> getBufferIdentifier ( ));
256256 auto NewBinaryOrErr = OffloadBinary::create (*Binary);
257257 if (!NewBinaryOrErr)
258258 return NewBinaryOrErr.takeError ();
@@ -1358,6 +1358,135 @@ static Expected<StringRef> linkDevice(ArrayRef<StringRef> InputFiles,
13581358 return *DeviceLinkedFile;
13591359}
13601360
1361+ static bool isStaticArchiveFile (const StringRef Filename) {
1362+ if (!llvm::sys::path::has_extension (Filename))
1363+ // Any file with no extension should not be considered an Archive.
1364+ return false ;
1365+ llvm::file_magic Magic;
1366+ llvm::identify_magic (Filename, Magic);
1367+ // Only archive files are to be considered.
1368+ // TODO: .lib check to be added
1369+ return (Magic == llvm::file_magic::archive);
1370+ }
1371+
1372+ static Expected<StringRef> listSection (StringRef Filename,
1373+ const ArgList &Args) {
1374+ Expected<std::string> OffloadBundlerPath = findProgram (
1375+ " clang-offload-bundler" getMainExecutable (" clang-offload-bundler"
1376+ if (!OffloadBundlerPath)
1377+ return OffloadBundlerPath.takeError ();
1378+ BumpPtrAllocator Alloc;
1379+ StringSaver Saver (Alloc);
1380+
1381+ SmallVector<StringRef, 8 > CmdArgs;
1382+ CmdArgs.push_back (*OffloadBundlerPath);
1383+ bool IsArchive = isStaticArchiveFile (Filename);
1384+ CmdArgs.push_back (IsArchive ? " -type=aoo" " -type=o"
1385+ CmdArgs.push_back (Saver.save (" -input="
1386+ CmdArgs.push_back (" -list"
1387+ auto Output = createOutputFile (" bundled-targets" " list"
1388+ if (!Output)
1389+ return Output.takeError ();
1390+ SmallVector<std::optional<StringRef>> Redirects{std::nullopt , *Output,
1391+ std::nullopt };
1392+ int ErrCode = llvm::sys::ExecuteAndWait (*OffloadBundlerPath, CmdArgs,
1393+ std::nullopt , Redirects);
1394+ if (ErrCode != 0 )
1395+ return createStringError (inconvertibleErrorCode (),
1396+ " Failed to list targets"
1397+ return *Output;
1398+ }
1399+
1400+ // This routine is used to run the clang-offload-bundler tool and unbundle
1401+ // device inputs that have been created with an older compiler where the
1402+ // device object is bundled into a host object.
1403+ static Expected<StringRef> unbundle (StringRef Filename, const ArgList &Args,
1404+ llvm::Triple Triple) {
1405+ Expected<std::string> OffloadBundlerPath = findProgram (
1406+ " clang-offload-bundler" getMainExecutable (" clang-offload-bundler"
1407+ if (!OffloadBundlerPath)
1408+ return OffloadBundlerPath.takeError ();
1409+
1410+ // Create a new file to write the unbundled file to.
1411+ auto TempFileOrErr =
1412+ createOutputFile (sys::path::filename (ExecutableName), " ir"
1413+ if (!TempFileOrErr)
1414+ return TempFileOrErr.takeError ();
1415+
1416+ BumpPtrAllocator Alloc;
1417+ StringSaver Saver (Alloc);
1418+
1419+ SmallVector<StringRef, 8 > CmdArgs;
1420+ CmdArgs.push_back (*OffloadBundlerPath);
1421+ bool IsArchive = isStaticArchiveFile (Filename);
1422+ CmdArgs.push_back (IsArchive ? " -type=aoo" " -type=o"
1423+ auto *Target = Args.MakeArgString (Twine (" -targets=sycl-" str ());
1424+ CmdArgs.push_back (Target);
1425+ CmdArgs.push_back (Saver.save (" -input="
1426+ CmdArgs.push_back (Saver.save (" -output="
1427+ CmdArgs.push_back (" -unbundle"
1428+ CmdArgs.push_back (" -allow-missing-bundles"
1429+ if (Error Err = executeCommands (*OffloadBundlerPath, CmdArgs))
1430+ return std::move (Err);
1431+ return *TempFileOrErr;
1432+ }
1433+
1434+ Error extractBundledObjects (StringRef Filename, const ArgList &Args,
1435+ SmallVector<OffloadFile> &Binaries) {
1436+ auto List = listSection (Filename, Args);
1437+ if (!List)
1438+ return List.takeError ();
1439+ SmallVector<StringRef> TriplesInFile;
1440+ llvm::ErrorOr<std::unique_ptr<MemoryBuffer>> TripleList =
1441+ llvm::MemoryBuffer::getFileOrSTDIN (*List, /* isText=*/ true );
1442+ if (std::error_code EC = TripleList.getError ())
1443+ return createFileError (*List, EC);
1444+ (*TripleList)
1445+ ->getBuffer ()
1446+ .split (TriplesInFile, ' \n ' /* MaxSplit=*/ 1 , /* KeepEmpty=*/ false );
1447+ for (StringRef TripleStr : TriplesInFile) {
1448+ StringRef SYCLPrefix = " sycl-"
1449+ if (!TripleStr.starts_with (SYCLPrefix))
1450+ continue ;
1451+ llvm::Triple Triple (TripleStr.substr (SYCLPrefix.size ()));
1452+ auto UnbundledFile = unbundle (Filename, Args, Triple);
1453+ if (!UnbundledFile)
1454+ return UnbundledFile.takeError ();
1455+ if (*UnbundledFile == Filename)
1456+ continue ;
1457+
1458+ SmallVector<StringRef> ObjectFilePaths;
1459+ if (sycl::isStaticArchiveFile (Filename)) {
1460+ llvm::ErrorOr<std::unique_ptr<MemoryBuffer>> ObjList =
1461+ llvm::MemoryBuffer::getFileOrSTDIN (*UnbundledFile, /* isText=*/ true );
1462+ if (std::error_code EC = ObjList.getError ())
1463+ return createFileError (*UnbundledFile, EC);
1464+ (*ObjList)->getBuffer ().split (ObjectFilePaths, ' \n ' /* MaxSplit=*/ 1 ,
1465+ /* KeepEmpty=*/ false );
1466+ } else {
1467+ ObjectFilePaths.push_back (*UnbundledFile);
1468+ }
1469+ for (StringRef ObjectFilePath : ObjectFilePaths) {
1470+ llvm::file_magic Magic;
1471+ llvm::identify_magic (ObjectFilePath, Magic);
1472+ if (Magic == file_magic::spirv_object)
1473+ return createStringError (
1474+ " SPIR-V fat objects must be generated with --offload-new-driver"
1475+ auto Arg = Args.MakeArgString (
1476+ " sycl-"
1477+ (Triple.isSPIROrSPIRV () ? Triple.str () + " -" str ()) + " ="
1478+ ObjectFilePath);
1479+ auto Binary = getInputBitcodeLibrary (Arg);
1480+
1481+ if (!Binary)
1482+ return Binary.takeError ();
1483+
1484+ Binaries.push_back (std::move (*Binary));
1485+ }
1486+ }
1487+ return Error::success ();
1488+ }
1489+
13611490} // namespace sycl
13621491
13631492namespace generic {
@@ -2634,8 +2763,14 @@ getDeviceInput(const ArgList &Args) {
26342763 if (identify_magic (Buffer.getBuffer ()) == file_magic::elf_shared_object)
26352764 continue ;
26362765 SmallVector<OffloadFile> Binaries;
2766+ size_t OldSize = Binaries.size ();
26372767 if (Error Err = extractOffloadBinaries (Buffer, Binaries))
26382768 return std::move (Err);
2769+ if (Binaries.size () == OldSize) {
2770+ if (Error Err = sycl::extractBundledObjects (*Filename, Args, Binaries))
2771+ return std::move (Err);
2772+ }
2773+
26392774 for (auto &OffloadFile : Binaries) {
26402775 if (identify_magic (Buffer.getBuffer ()) == file_magic::archive &&
26412776 !WholeArchive)
0 commit comments