|
15 | 15 | //===---------------------------------------------------------------------===// |
16 | 16 |
|
17 | 17 | #include "OffloadWrapper.h" |
| 18 | +#include "SYCLOffloadWrapper.h" |
18 | 19 | #include "clang/Basic/Version.h" |
19 | 20 | #include "llvm/BinaryFormat/Magic.h" |
20 | 21 | #include "llvm/Bitcode/BitcodeWriter.h" |
|
46 | 47 | #include "llvm/Support/Path.h" |
47 | 48 | #include "llvm/Support/Program.h" |
48 | 49 | #include "llvm/Support/Signals.h" |
| 50 | +#include "llvm/Support/SimpleTable.h" |
49 | 51 | #include "llvm/Support/SourceMgr.h" |
50 | 52 | #include "llvm/Support/StringSaver.h" |
51 | 53 | #include "llvm/Support/TargetSelect.h" |
@@ -633,37 +635,142 @@ static Expected<StringRef> runLLVMToSPIRVTranslation(StringRef InputTable, |
633 | 635 | return *Output; |
634 | 636 | } |
635 | 637 |
|
636 | | -// Run clang-offload-wrapper |
637 | | -static Expected<StringRef> runWrapper(StringRef &InputFile, |
638 | | - const ArgList &Args) { |
639 | | - // Create a new file to write the wrapped file to. |
640 | | - auto TempFileOrErr = |
641 | | - createOutputFile(sys::path::filename(ExecutableName), "bc"); |
642 | | - if (!TempFileOrErr) |
643 | | - return TempFileOrErr.takeError(); |
644 | | - Expected<std::string> ClangOffloadWrapperPath = findProgram( |
645 | | - "clang-offload-wrapper", {getMainExecutable("clang-offload-wrapper")}); |
646 | | - if (!ClangOffloadWrapperPath) |
647 | | - return ClangOffloadWrapperPath.takeError(); |
| 638 | +Expected<std::vector<char>> readBinaryFile(StringRef File) { |
| 639 | + auto MBOrErr = MemoryBuffer::getFile(File, /*IsText*/ false, |
| 640 | + /*RequiresNullTerminator */ false); |
| 641 | + if (!MBOrErr) |
| 642 | + return createFileError(File, MBOrErr.getError()); |
648 | 643 |
|
649 | | - BumpPtrAllocator Alloc; |
650 | | - StringSaver Saver(Alloc); |
| 644 | + auto &MB = *MBOrErr; |
| 645 | + return std::vector<char>(MB->getBufferStart(), MB->getBufferEnd()); |
| 646 | +} |
651 | 647 |
|
652 | | - SmallVector<StringRef, 8> CmdArgs; |
653 | | - CmdArgs.push_back(*ClangOffloadWrapperPath); |
654 | | - CmdArgs.push_back(Saver.save("-o=" + *TempFileOrErr)); |
655 | | - llvm::Triple HostTriple( |
| 648 | +Expected<std::string> readTextFile(StringRef File) { |
| 649 | + auto MBOrErr = MemoryBuffer::getFile(File, /*IsText*/ true, |
| 650 | + /*RequiresNullTerminator */ true); |
| 651 | + if (!MBOrErr) |
| 652 | + return createFileError(File, MBOrErr.getError()); |
| 653 | + |
| 654 | + auto &MB = *MBOrErr; |
| 655 | + return std::string(MB->getBufferStart(), MB->getBufferEnd()); |
| 656 | +} |
| 657 | + |
| 658 | +Expected<std::unique_ptr<util::PropertySetRegistry>> |
| 659 | +readPropertyRegistryFromFile(StringRef File) { |
| 660 | + auto MBOrErr = MemoryBuffer::getFile(File, /*IsText*/ true); |
| 661 | + if (!MBOrErr) |
| 662 | + return createFileError(File, MBOrErr.getError()); |
| 663 | + |
| 664 | + auto &MB = *MBOrErr; |
| 665 | + return util::PropertySetRegistry::read(&*MB); |
| 666 | +} |
| 667 | + |
| 668 | +// The table format is the following: |
| 669 | +// [Code|Properties|Symbols] |
| 670 | +// a_0.bin|a_0.prop|a_0.sym |
| 671 | +// . |
| 672 | +// a_n.bin|a_n.prop|a_n.sym |
| 673 | +// |
| 674 | +// .bin extension might be a bc, spv or other native extension. |
| 675 | +Expected<SmallVector<SYCLImage>> readSYCLImagesFromTable(StringRef TableFile, |
| 676 | + const ArgList &Args) { |
| 677 | + auto TableOrErr = util::SimpleTable::read(TableFile); |
| 678 | + if (!TableOrErr) |
| 679 | + return TableOrErr.takeError(); |
| 680 | + |
| 681 | + std::unique_ptr<util::SimpleTable> Table = std::move(*TableOrErr); |
| 682 | + int CodeIndex = Table->getColumnId("Code"); |
| 683 | + int PropertiesIndex = Table->getColumnId("Properties"); |
| 684 | + int SymbolsIndex = Table->getColumnId("Symbols"); |
| 685 | + if (CodeIndex == -1 || PropertiesIndex == -1 || SymbolsIndex == -1) |
| 686 | + return createStringError( |
| 687 | + inconvertibleErrorCode(), |
| 688 | + "expected columns in the table: Code, Properties and Symbols"); |
| 689 | + |
| 690 | + SmallVector<SYCLImage> Images; |
| 691 | + for (const util::SimpleTable::Row &row : Table->rows()) { |
| 692 | + auto ImageOrErr = readBinaryFile(row.getCell("Code")); |
| 693 | + if (!ImageOrErr) |
| 694 | + return ImageOrErr.takeError(); |
| 695 | + |
| 696 | + auto PropertiesOrErr = |
| 697 | + readPropertyRegistryFromFile(row.getCell("Properties")); |
| 698 | + if (!PropertiesOrErr) |
| 699 | + return PropertiesOrErr.takeError(); |
| 700 | + |
| 701 | + auto SymbolsOrErr = readTextFile(row.getCell("Symbols")); |
| 702 | + if (!SymbolsOrErr) |
| 703 | + return SymbolsOrErr.takeError(); |
| 704 | + |
| 705 | + SYCLImage Image; |
| 706 | + Image.Image = std::move(*ImageOrErr); |
| 707 | + Image.PropertyRegistry = std::move(**PropertiesOrErr); |
| 708 | + Image.Entries = std::move(*SymbolsOrErr); |
| 709 | + Images.push_back(std::move(Image)); |
| 710 | + } |
| 711 | + |
| 712 | + return Images; |
| 713 | +} |
| 714 | + |
| 715 | +/// Reads device images from the given \p InputFile and wraps them |
| 716 | +/// in one LLVM IR Module as a constant data. |
| 717 | +/// |
| 718 | +/// \returns A path to the LLVM Module that contains wrapped images. |
| 719 | +Expected<StringRef> wrapSYCLBinariesFromFile(StringRef InputFile, |
| 720 | + const ArgList &Args) { |
| 721 | + auto OutputFileOrErr = createOutputFile( |
| 722 | + sys::path::filename(ExecutableName) + ".sycl.image.wrapper", "bc"); |
| 723 | + if (!OutputFileOrErr) |
| 724 | + return OutputFileOrErr.takeError(); |
| 725 | + |
| 726 | + StringRef OutputFilePath = *OutputFileOrErr; |
| 727 | + if (Verbose || DryRun) { |
| 728 | + errs() << formatv(" offload-wrapper: input: {0}, output: {1}\n", InputFile, |
| 729 | + OutputFilePath); |
| 730 | + if (DryRun) |
| 731 | + return OutputFilePath; |
| 732 | + } |
| 733 | + |
| 734 | + auto ImagesOrErr = readSYCLImagesFromTable(InputFile, Args); |
| 735 | + if (!ImagesOrErr) |
| 736 | + return ImagesOrErr.takeError(); |
| 737 | + |
| 738 | + auto &Images = *ImagesOrErr; |
| 739 | + StringRef Target = Args.getLastArgValue(OPT_triple_EQ); |
| 740 | + if (Target.empty()) |
| 741 | + return createStringError( |
| 742 | + inconvertibleErrorCode(), |
| 743 | + "can't wrap SYCL image. -triple argument is missed."); |
| 744 | + |
| 745 | + for (SYCLImage &Image : Images) |
| 746 | + Image.Target = Target; |
| 747 | + |
| 748 | + LLVMContext C; |
| 749 | + Module M("offload.wrapper.object", C); |
| 750 | + M.setTargetTriple( |
656 | 751 | Args.getLastArgValue(OPT_host_triple_EQ, sys::getDefaultTargetTriple())); |
657 | | - CmdArgs.push_back(Saver.save("-host=" + HostTriple.str())); |
658 | | - const llvm::Triple Triple(Args.getLastArgValue(OPT_triple_EQ)); |
659 | | - SmallString<128> TargetTripleOpt = Triple.getArchName(); |
660 | | - CmdArgs.push_back(Saver.save("-target=" + TargetTripleOpt)); |
661 | | - CmdArgs.push_back("-kind=sycl"); |
662 | | - CmdArgs.push_back("-batch"); |
663 | | - CmdArgs.push_back(InputFile); |
664 | | - if (Error Err = executeCommands(*ClangOffloadWrapperPath, CmdArgs)) |
665 | | - return std::move(Err); |
666 | | - return *TempFileOrErr; |
| 752 | + |
| 753 | + StringRef CompileOptions = |
| 754 | + Args.getLastArgValue(OPT_sycl_backend_compile_options_EQ); |
| 755 | + StringRef LinkOptions = Args.getLastArgValue(OPT_sycl_target_link_options_EQ); |
| 756 | + SYCLWrappingOptions WrappingOptions; |
| 757 | + WrappingOptions.CompileOptions = CompileOptions; |
| 758 | + WrappingOptions.LinkOptions = LinkOptions; |
| 759 | + if (Error E = wrapSYCLBinaries(M, Images, WrappingOptions)) |
| 760 | + return E; |
| 761 | + |
| 762 | + if (Args.hasArg(OPT_print_wrapped_module)) |
| 763 | + errs() << M; |
| 764 | + |
| 765 | + // TODO: Once "llc tool->runCompile" migration is finished we need to remove |
| 766 | + // this scope and use community flow. |
| 767 | + int FD = -1; |
| 768 | + if (std::error_code EC = sys::fs::openFileForWrite(OutputFilePath, FD)) |
| 769 | + return errorCodeToError(EC); |
| 770 | + |
| 771 | + raw_fd_ostream OS(FD, true); |
| 772 | + WriteBitcodeToFile(M, OS); |
| 773 | + return OutputFilePath; |
667 | 774 | } |
668 | 775 |
|
669 | 776 | // Run llc |
@@ -691,11 +798,10 @@ static Expected<StringRef> runCompile(StringRef &InputFile, |
691 | 798 | return *OutputFileOrErr; |
692 | 799 | } |
693 | 800 |
|
694 | | -// Run clang-offload-wrapper and llc |
| 801 | +// Run wrapping library and llc |
695 | 802 | static Expected<StringRef> runWrapperAndCompile(StringRef &InputFile, |
696 | 803 | const ArgList &Args) { |
697 | | - // call to clang-offload-wrapper |
698 | | - auto OutputFile = sycl::runWrapper(InputFile, Args); |
| 804 | + auto OutputFile = sycl::wrapSYCLBinariesFromFile(InputFile, Args); |
699 | 805 | if (!OutputFile) |
700 | 806 | return OutputFile.takeError(); |
701 | 807 | // call to llc |
|
0 commit comments