diff --git a/clang/lib/Driver/Driver.cpp b/clang/lib/Driver/Driver.cpp index 1a1d4db9a1e46..70960136e50c7 100644 --- a/clang/lib/Driver/Driver.cpp +++ b/clang/lib/Driver/Driver.cpp @@ -3683,8 +3683,11 @@ getLinkerArgs(Compilation &C, DerivedArgList &Args, bool IncludeObj = false) { static bool IsSYCLDeviceLibObj(std::string ObjFilePath, bool isMSVCEnv) { StringRef ObjFileName = llvm::sys::path::filename(ObjFilePath); StringRef ObjSuffix = isMSVCEnv ? ".obj" : ".o"; + StringRef NewObjSuffix = isMSVCEnv ? ".new.obj" : ".new.o"; bool Ret = - (ObjFileName.starts_with("libsycl-") && ObjFileName.ends_with(ObjSuffix)) + (ObjFileName.starts_with("libsycl-") && + ObjFileName.ends_with(ObjSuffix) && + !ObjFileName.ends_with(NewObjSuffix)) // Avoid new-offload-driver objs ? true : false; return Ret; @@ -7877,6 +7880,11 @@ Action *Driver::BuildOffloadingActions(Compilation &C, break; } + // Backend/Assemble actions are not used for the SYCL device side + if (Kind == Action::OFK_SYCL && + (Phase == phases::Backend || Phase == phases::Assemble)) + continue; + auto TCAndArch = TCAndArchs.begin(); for (Action *&A : DeviceActions) { if (A->getType() == types::TY_Nothing) diff --git a/clang/lib/Driver/ToolChains/Clang.cpp b/clang/lib/Driver/ToolChains/Clang.cpp index 6fc3101a9d415..3e6905e3da4fb 100644 --- a/clang/lib/Driver/ToolChains/Clang.cpp +++ b/clang/lib/Driver/ToolChains/Clang.cpp @@ -11049,7 +11049,10 @@ void LinkerWrapper::ConstructJob(Compilation &C, const JobAction &JA, SYCLInstallationDetector SYCLInstallation(D); SYCLInstallation.getSYCLDeviceLibPath(LibLocCandidates); SmallString<128> LibName("libsycl-crt"); - StringRef LibSuffix = TheTriple.isWindowsMSVCEnvironment() ? ".obj" : ".o"; + bool IsNewOffload = D.getUseNewOffloadingDriver(); + StringRef LibSuffix = TheTriple.isWindowsMSVCEnvironment() + ? (IsNewOffload ? ".new.obj" : ".obj") + : (IsNewOffload ? ".new.o" : ".o"); llvm::sys::path::replace_extension(LibName, LibSuffix); for (const auto &LibLoc : LibLocCandidates) { SmallString<128> FullLibName(LibLoc); diff --git a/clang/lib/Driver/ToolChains/SYCL.cpp b/clang/lib/Driver/ToolChains/SYCL.cpp index fa5bc219c06ca..7024e79c4aca5 100644 --- a/clang/lib/Driver/ToolChains/SYCL.cpp +++ b/clang/lib/Driver/ToolChains/SYCL.cpp @@ -291,12 +291,16 @@ SYCL::getDeviceLibraries(const Compilation &C, const llvm::Triple &TargetTriple, const SYCLDeviceLibsList SYCLDeviceSanitizerLibs = { {"libsycl-sanitizer", "internal"}}; #endif + bool IsWindowsMSVCEnv = + C.getDefaultToolChain().getTriple().isWindowsMSVCEnvironment(); + bool IsNewOffload = C.getDriver().getUseNewOffloadingDriver(); StringRef LibSuffix = ".bc"; if (TargetTriple.isNVPTX()) // For NVidia, we are unbundling objects. - LibSuffix = C.getDefaultToolChain().getTriple().isWindowsMSVCEnvironment() - ? ".obj" - : ".o"; + LibSuffix = IsWindowsMSVCEnv ? ".obj" : ".o"; + if (IsNewOffload) + // For new offload model, we use packaged .bc files. + LibSuffix = IsWindowsMSVCEnv ? ".new.obj" : ".new.o"; auto addLibraries = [&](const SYCLDeviceLibsList &LibsList) { for (const DeviceLibOptInfo &Lib : LibsList) { if (!DeviceLibLinkInfo[Lib.DeviceLibOption]) @@ -441,6 +445,10 @@ const char *SYCL::Linker::constructLLVMLinkCommand( C.getDriver().IsCLMode()) LibPostfix = ".obj"; } + StringRef NewLibPostfix = ".new.o"; + if (HostTC->getTriple().isWindowsMSVCEnvironment() && + C.getDriver().IsCLMode()) + NewLibPostfix = ".new.obj"; std::string FileName = this->getToolChain().getInputFilename(II); StringRef InputFilename = llvm::sys::path::filename(FileName); if (IsNVPTX || IsSYCLNativeCPU) { @@ -448,12 +456,15 @@ const char *SYCL::Linker::constructLLVMLinkCommand( if ((InputFilename.find("libspirv") != InputFilename.npos || InputFilename.find("libdevice") != InputFilename.npos)) return true; - if (IsNVPTX) + if (IsNVPTX) { LibPostfix = ".cubin"; + NewLibPostfix = ".new.cubin"; + } } StringRef LibSyclPrefix("libsycl-"); if (!InputFilename.starts_with(LibSyclPrefix) || - !InputFilename.ends_with(LibPostfix)) + !InputFilename.ends_with(LibPostfix) || + InputFilename.ends_with(NewLibPostfix)) return false; // Skip the prefix "libsycl-" std::string PureLibName = diff --git a/clang/test/Driver/Inputs/libsycl-complex.new.o b/clang/test/Driver/Inputs/libsycl-complex.new.o new file mode 100644 index 0000000000000..c7501e92dc7a9 Binary files /dev/null and b/clang/test/Driver/Inputs/libsycl-complex.new.o differ diff --git a/clang/test/Driver/Inputs/libsycl-complex.new.obj b/clang/test/Driver/Inputs/libsycl-complex.new.obj new file mode 100644 index 0000000000000..197cbbacec8e2 Binary files /dev/null and b/clang/test/Driver/Inputs/libsycl-complex.new.obj differ diff --git a/clang/test/Driver/Inputs/libsycl-crt.new.o b/clang/test/Driver/Inputs/libsycl-crt.new.o new file mode 100644 index 0000000000000..2fd53a58c4aa8 Binary files /dev/null and b/clang/test/Driver/Inputs/libsycl-crt.new.o differ diff --git a/clang/test/Driver/Inputs/libsycl-crt.new.obj b/clang/test/Driver/Inputs/libsycl-crt.new.obj new file mode 100644 index 0000000000000..511fa246ef1d5 Binary files /dev/null and b/clang/test/Driver/Inputs/libsycl-crt.new.obj differ diff --git a/clang/test/Driver/linker-wrapper-sycl-win.cpp b/clang/test/Driver/linker-wrapper-sycl-win.cpp new file mode 100644 index 0000000000000..1854dee476641 --- /dev/null +++ b/clang/test/Driver/linker-wrapper-sycl-win.cpp @@ -0,0 +1,12 @@ +// REQUIRES: system-windows + +/// Check for list of commands for standalone clang-linker-wrapper run for sycl +// RUN: clang-linker-wrapper -sycl-device-library-location=%S/Inputs -sycl-device-libraries=libsycl-crt.new.obj,libsycl-complex.new.obj -sycl-post-link-options="SYCL_POST_LINK_OPTIONS" -llvm-spirv-options="LLVM_SPIRV_OPTIONS" "--host-triple=x86_64-pc-windows-msvc" "--triple=spir64" "--linker-path=/usr/bin/ld" "--" HOST_LINKER_FLAGS "-dynamic-linker" HOST_DYN_LIB "-o" "a.out" HOST_LIB_PATH HOST_STAT_LIB %S/Inputs/test-sycl.o --dry-run 2>&1 | FileCheck -check-prefix=CHK-CMDS %s +// CHK-CMDS: "{{.*}}spirv-to-ir-wrapper.exe" {{.*}} -o [[FIRSTLLVMLINKIN:.*]].bc --llvm-spirv-opts=--spirv-preserve-auxdata --llvm-spirv-opts=--spirv-target-env=SPV-IR --llvm-spirv-opts=--spirv-builtin-format=global +// CHK-CMDS-NEXT: "{{.*}}llvm-link.exe" [[FIRSTLLVMLINKIN:.*]].bc -o [[FIRSTLLVMLINKOUT:.*]].bc --suppress-warnings +// CHK-CMDS-NEXT: "{{.*}}llvm-link.exe" -only-needed [[FIRSTLLVMLINKOUT]].bc {{.*}}.bc {{.*}}.bc -o [[SECONDLLVMLINKOUT:.*]].bc --suppress-warnings +// CHK-CMDS-NEXT: "{{.*}}sycl-post-link.exe" SYCL_POST_LINK_OPTIONS -o [[SYCLPOSTLINKOUT:.*]].table [[SECONDLLVMLINKOUT]].bc +// LLVM-SPIRV is not called in dry-run +// CHK-CMDS-NEXT: offload-wrapper: input: [[LLVMSPIRVOUT:.*]].table, output: [[WRAPPEROUT:.*]].bc +// CHK-CMDS-NEXT: "{{.*}}llc.exe" -filetype=obj -o [[LLCOUT:.*]].o [[WRAPPEROUT]].bc +// CHK-CMDS-NEXT: "{{.*}}/ld" -- HOST_LINKER_FLAGS -dynamic-linker HOST_DYN_LIB -o a.out [[LLCOUT]].o HOST_LIB_PATH HOST_STAT_LIB {{.*}}test-sycl.o diff --git a/clang/test/Driver/linker-wrapper-sycl.cpp b/clang/test/Driver/linker-wrapper-sycl.cpp index 9d65f80f18c71..bf0f7306b070f 100644 --- a/clang/test/Driver/linker-wrapper-sycl.cpp +++ b/clang/test/Driver/linker-wrapper-sycl.cpp @@ -1,11 +1,10 @@ // REQUIRES: system-linux /// Check for list of commands for standalone clang-linker-wrapper run for sycl -// RUN: clang-linker-wrapper -sycl-device-library-location=%S/Inputs -sycl-device-libraries=libsycl-crt.o,libsycl-complex.o -sycl-post-link-options="SYCL_POST_LINK_OPTIONS" -llvm-spirv-options="LLVM_SPIRV_OPTIONS" "--host-triple=x86_64-unknown-linux-gnu" "--triple=spir64" "--linker-path=/usr/bin/ld" "--" HOST_LINKER_FLAGS "-dynamic-linker" HOST_DYN_LIB "-o" "a.out" HOST_LIB_PATH HOST_STAT_LIB %S/Inputs/test-sycl.o --dry-run 2>&1 | FileCheck -check-prefix=CHK-CMDS %s -// CHK-CMDS: "{{.*}}llvm-link" [[INPUT:.*]].bc -o [[FIRSTLLVMLINKOUT:.*]].bc --suppress-warnings -// CHK-CMDS-NEXT: "{{.*}}clang-offload-bundler" -type=o -targets=sycl-spir64-unknown-unknown -input={{.*}}libsycl-crt.o -output=[[FIRSTUNBUNDLEDLIB:.*]].bc -unbundle -allow-missing-bundles -// CHK-CMDS-NEXT: "{{.*}}clang-offload-bundler" -type=o -targets=sycl-spir64-unknown-unknown -input={{.*}}libsycl-complex.o -output=[[SECONDUNBUNDLEDLIB:.*]].bc -unbundle -allow-missing-bundles -// CHK-CMDS-NEXT: "{{.*}}llvm-link" -only-needed [[FIRSTLLVMLINKOUT]].bc [[FIRSTUNBUNDLEDLIB]].bc [[SECONDUNBUNDLEDLIB]].bc -o [[SECONDLLVMLINKOUT:.*]].bc --suppress-warnings +// RUN: clang-linker-wrapper -sycl-device-library-location=%S/Inputs -sycl-device-libraries=libsycl-crt.new.o,libsycl-complex.new.o -sycl-post-link-options="SYCL_POST_LINK_OPTIONS" -llvm-spirv-options="LLVM_SPIRV_OPTIONS" "--host-triple=x86_64-unknown-linux-gnu" "--triple=spir64" "--linker-path=/usr/bin/ld" "--" HOST_LINKER_FLAGS "-dynamic-linker" HOST_DYN_LIB "-o" "a.out" HOST_LIB_PATH HOST_STAT_LIB %S/Inputs/test-sycl.o --dry-run 2>&1 | FileCheck -check-prefix=CHK-CMDS %s +// CHK-CMDS: "{{.*}}spirv-to-ir-wrapper" {{.*}} -o [[FIRSTLLVMLINKIN:.*]].bc --llvm-spirv-opts=--spirv-preserve-auxdata --llvm-spirv-opts=--spirv-target-env=SPV-IR --llvm-spirv-opts=--spirv-builtin-format=global +// CHK-CMDS-NEXT: "{{.*}}llvm-link" [[FIRSTLLVMLINKIN:.*]].bc -o [[FIRSTLLVMLINKOUT:.*]].bc --suppress-warnings +// CHK-CMDS-NEXT: "{{.*}}llvm-link" -only-needed [[FIRSTLLVMLINKOUT]].bc {{.*}}.bc {{.*}}.bc -o [[SECONDLLVMLINKOUT:.*]].bc --suppress-warnings // CHK-CMDS-NEXT: "{{.*}}sycl-post-link" SYCL_POST_LINK_OPTIONS -o [[SYCLPOSTLINKOUT:.*]].table [[SECONDLLVMLINKOUT]].bc // LLVM-SPIRV is not called in dry-run // CHK-CMDS-NEXT: offload-wrapper: input: [[LLVMSPIRVOUT:.*]].table, output: [[WRAPPEROUT:.*]].bc diff --git a/clang/test/Driver/sycl-offload-new-driver.c b/clang/test/Driver/sycl-offload-new-driver.c index 113d7ab3bfa95..ded8f86b42f2b 100644 --- a/clang/test/Driver/sycl-offload-new-driver.c +++ b/clang/test/Driver/sycl-offload-new-driver.c @@ -1,5 +1,4 @@ // REQUIRES: system-linux - /// Verify --offload-new-driver option phases // RUN: %clang --target=x86_64-unknown-linux-gnu -fsycl -fsycl-targets=nvptx64-nvidia-cuda,spir64 --offload-new-driver -ccc-print-phases %s 2>&1 \ // RUN: | FileCheck -check-prefix=OFFLOAD-NEW-DRIVER %s @@ -10,20 +9,16 @@ // OFFLOAD-NEW-DRIVER: 4: input, "[[INPUT]]", c++, (device-sycl) // OFFLOAD-NEW-DRIVER: 5: preprocessor, {4}, c++-cpp-output, (device-sycl) // OFFLOAD-NEW-DRIVER: 6: compiler, {5}, ir, (device-sycl) -// OFFLOAD-NEW-DRIVER: 7: backend, {6}, assembler, (device-sycl) -// OFFLOAD-NEW-DRIVER: 8: assembler, {7}, object, (device-sycl) -// OFFLOAD-NEW-DRIVER: 9: offload, "device-sycl (nvptx64-nvidia-cuda)" {8}, object -// OFFLOAD-NEW-DRIVER: 10: input, "[[INPUT]]", c++, (device-sycl) -// OFFLOAD-NEW-DRIVER: 11: preprocessor, {10}, c++-cpp-output, (device-sycl) -// OFFLOAD-NEW-DRIVER: 12: compiler, {11}, ir, (device-sycl) -// OFFLOAD-NEW-DRIVER: 13: backend, {12}, assembler, (device-sycl) -// OFFLOAD-NEW-DRIVER: 14: assembler, {13}, object, (device-sycl) -// OFFLOAD-NEW-DRIVER: 15: offload, "device-sycl (spir64-unknown-unknown)" {14}, object -// OFFLOAD-NEW-DRIVER: 16: clang-offload-packager, {9, 15}, image, (device-sycl) -// OFFLOAD-NEW-DRIVER: 17: offload, "host-sycl (x86_64-unknown-linux-gnu)" {3}, "device-sycl (x86_64-unknown-linux-gnu)" {16}, ir -// OFFLOAD-NEW-DRIVER: 18: backend, {17}, assembler, (host-sycl) -// OFFLOAD-NEW-DRIVER: 19: assembler, {18}, object, (host-sycl) -// OFFLOAD-NEW-DRIVER: 20: clang-linker-wrapper, {19}, image, (host-sycl) +// OFFLOAD-NEW-DRIVER: 7: offload, "device-sycl (nvptx64-nvidia-cuda)" {6}, ir +// OFFLOAD-NEW-DRIVER: 8: input, "[[INPUT]]", c++, (device-sycl) +// OFFLOAD-NEW-DRIVER: 9: preprocessor, {8}, c++-cpp-output, (device-sycl) +// OFFLOAD-NEW-DRIVER: 10: compiler, {9}, ir, (device-sycl) +// OFFLOAD-NEW-DRIVER: 11: offload, "device-sycl (spir64-unknown-unknown)" {10}, ir +// OFFLOAD-NEW-DRIVER: 12: clang-offload-packager, {7, 11}, image, (device-sycl) +// OFFLOAD-NEW-DRIVER: 13: offload, "host-sycl (x86_64-unknown-linux-gnu)" {3}, "device-sycl (x86_64-unknown-linux-gnu)" {12}, ir +// OFFLOAD-NEW-DRIVER: 14: backend, {13}, assembler, (host-sycl) +// OFFLOAD-NEW-DRIVER: 15: assembler, {14}, object, (host-sycl) +// OFFLOAD-NEW-DRIVER: 16: clang-linker-wrapper, {15}, image, (host-sycl) /// Check the toolflow for SYCL compilation using new offload model // RUN: %clangxx -### --target=x86_64-unknown-linux-gnu -fsycl -fsycl-targets=spir64 --offload-new-driver %s 2>&1 | FileCheck -check-prefix=CHK-FLOW %s @@ -38,7 +33,7 @@ // RUN: --sysroot=%S/Inputs/SYCL -### %s 2>&1 \ // RUN: | FileCheck -check-prefix WRAPPER_OPTIONS %s // WRAPPER_OPTIONS: clang-linker-wrapper{{.*}} "--triple=spir64" -// WRAPPER_OPTIONS-SAME: "-sycl-device-libraries=libsycl-crt.bc,libsycl-complex.bc,libsycl-complex-fp64.bc,libsycl-cmath.bc,libsycl-cmath-fp64.bc,libsycl-imf.bc,libsycl-imf-fp64.bc,libsycl-imf-bf16.bc,libsycl-itt-user-wrappers.bc,libsycl-itt-compiler-wrappers.bc,libsycl-itt-stubs.bc" +// WRAPPER_OPTIONS-SAME: "-sycl-device-libraries=libsycl-crt.new.o,libsycl-complex.new.o,libsycl-complex-fp64.new.o,libsycl-cmath.new.o,libsycl-cmath-fp64.new.o,libsycl-imf.new.o,libsycl-imf-fp64.new.o,libsycl-imf-bf16.new.o,libsycl-itt-user-wrappers.new.o,libsycl-itt-compiler-wrappers.new.o,libsycl-itt-stubs.new.o" // WRAPPER_OPTIONS-SAME: "-sycl-device-library-location={{.*}}/lib" // RUN: %clangxx --target=x86_64-unknown-linux-gnu -fsycl --offload-new-driver \ diff --git a/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp b/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp index c01effcc4bf24..37ecc006a797b 100644 --- a/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp +++ b/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp @@ -41,6 +41,7 @@ #include "llvm/Support/Errc.h" #include "llvm/Support/FileOutputBuffer.h" #include "llvm/Support/FileSystem.h" +#include "llvm/Support/FileUtilities.h" #include "llvm/Support/InitLLVM.h" #include "llvm/Support/LineIterator.h" #include "llvm/Support/MemoryBuffer.h" @@ -218,6 +219,30 @@ Expected createOutputFile(const Twine &Prefix, StringRef Extension) { return TempFiles.back(); } +Expected writeOffloadFile(const OffloadFile &File) { + const OffloadBinary &Binary = *File.getBinary(); + + StringRef Prefix = + sys::path::stem(Binary.getMemoryBufferRef().getBufferIdentifier()); + StringRef Suffix = getImageKindName(Binary.getImageKind()); + + auto TempFileOrErr = createOutputFile( + Prefix + "-" + Binary.getTriple() + "-" + Binary.getArch(), Suffix); + if (!TempFileOrErr) + return TempFileOrErr.takeError(); + + Expected> OutputOrErr = + FileOutputBuffer::create(*TempFileOrErr, Binary.getImage().size()); + if (!OutputOrErr) + return OutputOrErr.takeError(); + std::unique_ptr Output = std::move(*OutputOrErr); + llvm::copy(Binary.getImage(), Output->getBufferStart()); + if (Error E = Output->commit()) + return std::move(E); + + return *TempFileOrErr; +} + /// Execute the command \p ExecutablePath with the arguments \p Args. Error executeCommands(StringRef ExecutablePath, ArrayRef Args) { if (Verbose || DryRun) @@ -467,79 +492,10 @@ static Error getSYCLDeviceLibs(SmallVector &DeviceLibFiles, return Error::success(); } -static bool isStaticArchiveFile(const StringRef Filename) { - if (!llvm::sys::path::has_extension(Filename)) - // Any file with no extension should not be considered an Archive. - return false; - llvm::file_magic Magic; - llvm::identify_magic(Filename, Magic); - // Only archive files are to be considered. - // TODO: .lib check to be added - return (Magic == llvm::file_magic::archive); -} - -// Find if section related to triple is present in a bundled file -static Expected checkSection(StringRef Filename, llvm::Triple Triple, - const ArgList &Args) { - Expected OffloadBundlerPath = findProgram( - "clang-offload-bundler", {getMainExecutable("clang-offload-bundler")}); - if (!OffloadBundlerPath) - return OffloadBundlerPath.takeError(); - BumpPtrAllocator Alloc; - StringSaver Saver(Alloc); - - auto *Target = Args.MakeArgString(Twine("-targets=sycl-") + Triple.str()); - SmallVector CmdArgs; - CmdArgs.push_back(*OffloadBundlerPath); - CmdArgs.push_back(Target); - bool IsArchive = isStaticArchiveFile(Filename); - CmdArgs.push_back(IsArchive ? "-type=ao" : "-type=o"); - CmdArgs.push_back(Saver.save("-input=" + Filename)); - CmdArgs.push_back("-check-section"); - return !(llvm::sys::ExecuteAndWait(*OffloadBundlerPath, CmdArgs)); -} - -// This routine is used to run the clang-offload-bundler tool and unbundle -// device inputs that have been created with an older compiler where the -// device object is bundled into a host object. -static Expected unbundle(StringRef Filename, const ArgList &Args) { - Expected OffloadBundlerPath = findProgram( - "clang-offload-bundler", {getMainExecutable("clang-offload-bundler")}); - if (!OffloadBundlerPath) - return OffloadBundlerPath.takeError(); - - llvm::Triple Triple(Args.getLastArgValue(OPT_triple_EQ)); - // Check if section with Triple is available in input bundle - // If no section is available, then we assume it's not a valid bundle and - // return original file. - auto CheckSection = checkSection(Filename, Triple, Args); - if (!CheckSection) - return CheckSection.takeError(); - if (!(*CheckSection)) - return Filename; - // Create a new file to write the unbundled file to. - auto TempFileOrErr = - createOutputFile(sys::path::filename(ExecutableName), "bc"); - if (!TempFileOrErr) - return TempFileOrErr.takeError(); - - BumpPtrAllocator Alloc; - StringSaver Saver(Alloc); - - SmallVector CmdArgs; - CmdArgs.push_back(*OffloadBundlerPath); - CmdArgs.push_back("-type=o"); - CmdArgs.push_back(Saver.save("-targets=sycl-" + Triple.str())); - CmdArgs.push_back(Saver.save("-input=" + Filename)); - CmdArgs.push_back(Saver.save("-output=" + *TempFileOrErr)); - CmdArgs.push_back("-unbundle"); - CmdArgs.push_back("-allow-missing-bundles"); - if (Error Err = executeCommands(*OffloadBundlerPath, CmdArgs)) - return std::move(Err); - return *TempFileOrErr; -} - // This routine is used to convert SPIR-V input files into LLVM IR files. +// If input is not a SPIR-V file, then the original file is returned. +// TODO: Add a check to identify SPIR-V files and exit early if the input is +// not a SPIR-V file. static Expected convertSPIRVToIR(StringRef Filename, const ArgList &Args) { Expected SPIRVToIRWrapperPath = findProgram( @@ -558,10 +514,9 @@ static Expected convertSPIRVToIR(StringRef Filename, CmdArgs.push_back(Filename); CmdArgs.push_back("-o"); CmdArgs.push_back(*TempFileOrErr); - CmdArgs.push_back("-llvm-spirv-opts"); - CmdArgs.push_back("\"--spirv-preserve-auxdata"); - CmdArgs.push_back("--spirv-target-env=SPV-IR"); - CmdArgs.push_back("--spirv-builtin-format=global\""); + CmdArgs.push_back("--llvm-spirv-opts=--spirv-preserve-auxdata"); + CmdArgs.push_back("--llvm-spirv-opts=--spirv-target-env=SPV-IR"); + CmdArgs.push_back("--llvm-spirv-opts=--spirv-builtin-format=global"); if (Error Err = executeCommands(*SPIRVToIRWrapperPath, CmdArgs)) return std::move(Err); return *TempFileOrErr; @@ -881,25 +836,6 @@ static Expected runWrapperAndCompile(StringRef &InputFile, return *OutputFileOrErr; } -// This routine is used to unbundle all device library files that will be -// linked with input device codes. -static Error -unbundleSYCLDeviceLibs(const SmallVector &Files, - SmallVector &UnbundledFiles, - const ArgList &Args) { - for (auto &Filename : Files) { - assert(!sys::fs::is_directory(Filename) && "Filename cannot be directory"); - if (!sys::fs::exists(Filename)) - continue; - // Run unbundler - auto UnbundledFile = sycl::unbundle(Filename, Args); - if (!UnbundledFile) - return UnbundledFile.takeError(); - UnbundledFiles.push_back((*UnbundledFile).str()); - } - return Error::success(); -} - // Link all SYCL input files into one before adding device library files. Expected linkDeviceInputFiles(SmallVectorImpl &InputFiles, const ArgList &Args) { @@ -918,8 +854,12 @@ Expected linkDeviceInputFiles(SmallVectorImpl &InputFiles, SmallVector CmdArgs; CmdArgs.push_back(*LLVMLinkPath); - for (auto &File : InputFiles) - CmdArgs.push_back(File); + for (auto &File : InputFiles) { + auto IRFile = sycl::convertSPIRVToIR(File, Args); + if (!IRFile) + return IRFile.takeError(); + CmdArgs.push_back(*IRFile); + } CmdArgs.push_back("-o"); CmdArgs.push_back(*OutFileOrErr); CmdArgs.push_back("--suppress-warnings"); @@ -971,16 +911,37 @@ static Expected linkDevice(ArrayRef InputFiles, InputFilesVec.clear(); InputFilesVec.emplace_back(*LinkedFile); - // Get SYCL device library files // Gathering device library files SmallVector DeviceLibFiles; if (Error Err = sycl::getSYCLDeviceLibs(DeviceLibFiles, Args)) reportError(std::move(Err)); - SmallVector UnbundledDeviceLibFiles; - if (Error Err = sycl::unbundleSYCLDeviceLibs(DeviceLibFiles, - UnbundledDeviceLibFiles, Args)) - reportError(std::move(Err)); - for (auto &File : UnbundledDeviceLibFiles) + const llvm::Triple Triple(Args.getLastArgValue(OPT_triple_EQ)); + SmallVector ExtractedDeviceLibFiles; + for (auto &File : DeviceLibFiles) { + auto BufferOrErr = MemoryBuffer::getFile(File); + if (!BufferOrErr) + return createFileError(File, BufferOrErr.getError()); + auto Buffer = std::move(*BufferOrErr); + SmallVector Binaries; + if (Error Err = extractOffloadBinaries(Buffer->getMemBufferRef(), Binaries)) + return std::move(Err); + bool CompatibleBinaryFound = false; + for (auto &Binary : Binaries) { + auto BinTriple = Binary.getBinary()->getTriple(); + if (BinTriple == Triple.getTriple()) { + auto FileNameOrErr = writeOffloadFile(Binary); + if (!FileNameOrErr) + return FileNameOrErr.takeError(); + ExtractedDeviceLibFiles.emplace_back(*FileNameOrErr); + CompatibleBinaryFound = true; + } + } + if (!CompatibleBinaryFound) + WithColor::warning(errs(), LinkerExecutable) + << "Compatible SYCL device library binary not found\n"; + } + + for (auto &File : ExtractedDeviceLibFiles) InputFilesVec.emplace_back(File); // second llvm-link step auto DeviceLinkedFile = sycl::linkDeviceLibFiles(InputFilesVec, Args); @@ -1280,7 +1241,7 @@ Error linkBitcodeFiles(SmallVectorImpl &InputFiles, StringRef Arch = Args.getLastArgValue(OPT_arch_EQ); // Early exit for SPIR targets - if (Triple.isSPIR()) + if (Triple.isSPIROrSPIRV()) return Error::success(); SmallVector BitcodeInputFiles; @@ -1487,30 +1448,6 @@ Error linkBitcodeFiles(SmallVectorImpl &InputFiles, return Error::success(); } -Expected writeOffloadFile(const OffloadFile &File) { - const OffloadBinary &Binary = *File.getBinary(); - - StringRef Prefix = - sys::path::stem(Binary.getMemoryBufferRef().getBufferIdentifier()); - StringRef Suffix = getImageKindName(Binary.getImageKind()); - - auto TempFileOrErr = createOutputFile( - Prefix + "-" + Binary.getTriple() + "-" + Binary.getArch(), Suffix); - if (!TempFileOrErr) - return TempFileOrErr.takeError(); - - Expected> OutputOrErr = - FileOutputBuffer::create(*TempFileOrErr, Binary.getImage().size()); - if (!OutputOrErr) - return OutputOrErr.takeError(); - std::unique_ptr Output = std::move(*OutputOrErr); - llvm::copy(Binary.getImage(), Output->getBufferStart()); - if (Error E = Output->commit()) - return std::move(E); - - return *TempFileOrErr; -} - // Compile the module to an object file using the appropriate target machine for // the host triple. Expected compileModule(Module &M, OffloadKind Kind) { @@ -2088,31 +2025,17 @@ getDeviceInput(const ArgList &Args) { sys::fs::is_directory(*Filename)) continue; - // Some of the object files may be bundled using clang-offload-bundler - // Following code tries to unbundle these files. - auto UnbundledFile = sycl::unbundle(*Filename, Args); - if (!UnbundledFile) - return UnbundledFile.takeError(); - // In some cases, fat objects are created with SPIR-V files embedded. - // e.g. when fat object is created using `-fsycl-device-obj=spirv` option. - auto IRFile = (*UnbundledFile == *Filename) - ? *Filename - : sycl::convertSPIRVToIR(*UnbundledFile, Args); - if (!IRFile) - return IRFile.takeError(); ErrorOr> BufferOrErr = - MemoryBuffer::getFileOrSTDIN(*IRFile); + MemoryBuffer::getFile(*Filename); if (std::error_code EC = BufferOrErr.getError()) - return createFileError(*IRFile, EC); + return createFileError(*Filename, EC); MemoryBufferRef Buffer = **BufferOrErr; if (identify_magic(Buffer.getBuffer()) == file_magic::elf_shared_object) continue; - SmallVector Binaries; if (Error Err = extractOffloadBinaries(Buffer, Binaries)) return std::move(Err); - for (auto &OffloadFile : Binaries) { if (identify_magic(Buffer.getBuffer()) == file_magic::archive && !WholeArchive) diff --git a/libdevice/cmake/modules/SYCLLibdevice.cmake b/libdevice/cmake/modules/SYCLLibdevice.cmake index 1c4167e999b9b..6de9124db0dba 100644 --- a/libdevice/cmake/modules/SYCLLibdevice.cmake +++ b/libdevice/cmake/modules/SYCLLibdevice.cmake @@ -1,14 +1,19 @@ set(obj_binary_dir "${CMAKE_LIBRARY_OUTPUT_DIRECTORY}") +set(obj_new_offload_binary_dir "${CMAKE_LIBRARY_OUTPUT_DIRECTORY}") if (MSVC) set(lib-suffix obj) + set(new-offload-lib-suffix new.obj) set(spv_binary_dir "${CMAKE_RUNTIME_OUTPUT_DIRECTORY}") set(install_dest_spv bin) set(devicelib_host_static sycl-devicelib-host.lib) + set(devicelib_host_static_new_offload sycl-devicelib-host.new.lib) else() set(lib-suffix o) + set(new-offload-lib-suffix new.o) set(spv_binary_dir "${CMAKE_LIBRARY_OUTPUT_DIRECTORY}") set(install_dest_spv lib${LLVM_LIBDIR_SUFFIX}) set(devicelib_host_static libsycl-devicelib-host.a) + set(devicelib_host_static_new_offload libsycl-devicelib-host.new.a) endif() set(bc_binary_dir "${CMAKE_LIBRARY_OUTPUT_DIRECTORY}") set(install_dest_lib lib${LLVM_LIBDIR_SUFFIX}) @@ -53,8 +58,8 @@ if ("NVPTX" IN_LIST LLVM_TARGETS_TO_BUILD) string(APPEND sycl_targets_opt ",nvptx64-nvidia-cuda") list(APPEND compile_opts "-fno-sycl-libspirv" - "-fno-bundle-offload-arch" - "-nocudalib") + "-nocudalib" + "--cuda-gpu-arch=sm_50") endif() if (WIN32) @@ -63,13 +68,15 @@ if (WIN32) endif() add_custom_target(libsycldevice-obj) +add_custom_target(libsycldevice-obj-new-offload) add_custom_target(libsycldevice-spv) add_custom_target(libsycldevice-bc) add_custom_target(libsycldevice DEPENDS libsycldevice-obj - libsycldevice-spv - libsycldevice-bc) + libsycldevice-bc + libsycldevice-obj-new-offload + libsycldevice-spv) function(add_devicelib_obj obj_filename) cmake_parse_arguments(OBJ "" "" "SRC;DEP;EXTRA_ARGS" ${ARGN}) @@ -88,6 +95,22 @@ function(add_devicelib_obj obj_filename) install(FILES ${devicelib-obj-file} DESTINATION ${install_dest_lib} COMPONENT libsycldevice) + + set(devicelib-obj-file-new-offload ${obj_new_offload_binary_dir}/${obj_filename}.${new-offload-lib-suffix}) + add_custom_command(OUTPUT ${devicelib-obj-file-new-offload} + COMMAND ${clang} -fsycl -c --offload-new-driver + ${compile_opts} ${sycl_targets_opt} ${OBJ_EXTRA_ARGS} + ${CMAKE_CURRENT_SOURCE_DIR}/${OBJ_SRC} + -o ${devicelib-obj-file-new-offload} + MAIN_DEPENDENCY ${OBJ_SRC} + DEPENDS ${OBJ_DEP} + VERBATIM) + set(devicelib-obj-target-new-offload ${obj_filename}-new-offload-obj) + add_custom_target(${devicelib-obj-target-new-offload} DEPENDS ${devicelib-obj-file-new-offload}) + add_dependencies(libsycldevice-obj ${devicelib-obj-target-new-offload}) + install(FILES ${devicelib-obj-file-new-offload} + DESTINATION ${install_dest_lib} + COMPONENT libsycldevice) endfunction() function(add_devicelib_spv spv_filename) @@ -248,6 +271,14 @@ add_custom_command(OUTPUT ${obj_binary_dir}/libsycl-fallback-imf.${lib-suffix} DEPENDS ${imf_fallback_fp32_deps} get_imf_fallback_fp32 sycl-compiler VERBATIM) +add_custom_command(OUTPUT ${obj_binary_dir}/libsycl-fallback-imf.${new-offload-lib-suffix} + COMMAND ${clang} -fsycl -c --offload-new-driver + ${compile_opts} ${sycl_targets_opt} + ${imf_fp32_fallback_src} -I ${CMAKE_CURRENT_SOURCE_DIR}/imf + -o ${obj_binary_dir}/libsycl-fallback-imf.${new-offload-lib-suffix} + DEPENDS ${imf_fallback_fp32_deps} get_imf_fallback_fp32 sycl-compiler + VERBATIM) + add_custom_command(OUTPUT ${obj_binary_dir}/fallback-imf-fp32-host.${lib-suffix} COMMAND ${clang} ${imf_host_cxx_flags} -I ${CMAKE_CURRENT_SOURCE_DIR}/imf @@ -256,6 +287,14 @@ add_custom_command(OUTPUT ${obj_binary_dir}/fallback-imf-fp32-host.${lib-suffix} DEPENDS ${imf_fallback_fp32_deps} get_imf_fallback_fp32 sycl-compiler VERBATIM) +add_custom_command(OUTPUT ${obj_binary_dir}/fallback-imf-fp32-host.${new-offload-lib-suffix} + COMMAND ${clang} ${imf_host_cxx_flags} --offload-new-driver + -I ${CMAKE_CURRENT_SOURCE_DIR}/imf + ${imf_fp32_fallback_src} + -o ${obj_binary_dir}/fallback-imf-fp32-host.${new-offload-lib-suffix} + DEPENDS ${imf_fallback_fp32_deps} get_imf_fallback_fp32 sycl-compiler + VERBATIM) + add_custom_target(get_imf_fallback_fp64 DEPENDS ${imf_fp64_fallback_src}) add_custom_command(OUTPUT ${spv_binary_dir}/libsycl-fallback-imf-fp64.spv COMMAND ${clang} -fsycl-device-only -fsycl-device-obj=spirv @@ -282,6 +321,15 @@ add_custom_command(OUTPUT ${obj_binary_dir}/libsycl-fallback-imf-fp64.${lib-suff DEPENDS ${imf_fallback_fp64_deps} get_imf_fallback_fp64 sycl-compiler VERBATIM) +add_custom_command(OUTPUT ${obj_binary_dir}/libsycl-fallback-imf-fp64.${new-offload-lib-suffix} + COMMAND ${clang} -fsycl -c -I ${CMAKE_CURRENT_SOURCE_DIR}/imf + --offload-new-driver + ${compile_opts} ${sycl_targets_opt} + ${imf_fp64_fallback_src} + -o ${obj_binary_dir}/libsycl-fallback-imf-fp64.${new-offload-lib-suffix} + DEPENDS ${imf_fallback_fp64_deps} get_imf_fallback_fp64 sycl-compiler + VERBATIM) + add_custom_command(OUTPUT ${obj_binary_dir}/fallback-imf-fp64-host.${lib-suffix} COMMAND ${clang} ${imf_host_cxx_flags} -I ${CMAKE_CURRENT_SOURCE_DIR}/imf @@ -290,6 +338,14 @@ add_custom_command(OUTPUT ${obj_binary_dir}/fallback-imf-fp64-host.${lib-suffix} DEPENDS ${imf_fallback_fp64_deps} get_imf_fallback_fp64 sycl-compiler VERBATIM) +add_custom_command(OUTPUT ${obj_binary_dir}/fallback-imf-fp64-host.${new-offload-lib-suffix} + COMMAND ${clang} ${imf_host_cxx_flags} --offload-new-driver + -I ${CMAKE_CURRENT_SOURCE_DIR}/imf + ${imf_fp64_fallback_src} + -o ${obj_binary_dir}/fallback-imf-fp64-host.${new-offload-lib-suffix} + DEPENDS ${imf_fallback_fp64_deps} get_imf_fallback_fp64 sycl-compiler + VERBATIM) + add_custom_target(get_imf_fallback_bf16 DEPENDS ${imf_bf16_fallback_src}) add_custom_command(OUTPUT ${spv_binary_dir}/libsycl-fallback-imf-bf16.spv COMMAND ${clang} -fsycl-device-only -fsycl-device-obj=spirv @@ -316,6 +372,15 @@ add_custom_command(OUTPUT ${obj_binary_dir}/libsycl-fallback-imf-bf16.${lib-suff DEPENDS ${imf_fallback_bf16_deps} get_imf_fallback_bf16 sycl-compiler VERBATIM) +add_custom_command(OUTPUT ${obj_binary_dir}/libsycl-fallback-imf-bf16.${new-offload-lib-suffix} + COMMAND ${clang} -fsycl -c -I ${CMAKE_CURRENT_SOURCE_DIR}/imf + --offload-new-driver + ${compile_opts} ${sycl_targets_opt} + ${imf_bf16_fallback_src} + -o ${obj_binary_dir}/libsycl-fallback-imf-bf16.${new-offload-lib-suffix} + DEPENDS ${imf_fallback_bf16_deps} get_imf_fallback_bf16 sycl-compiler + VERBATIM) + add_custom_command(OUTPUT ${obj_binary_dir}/fallback-imf-bf16-host.${lib-suffix} COMMAND ${clang} ${imf_host_cxx_flags} -I ${CMAKE_CURRENT_SOURCE_DIR}/imf @@ -324,29 +389,46 @@ add_custom_command(OUTPUT ${obj_binary_dir}/fallback-imf-bf16-host.${lib-suffix} DEPENDS ${imf_fallback_bf16_deps} get_imf_fallback_bf16 sycl-compiler VERBATIM) +add_custom_command(OUTPUT ${obj_binary_dir}/fallback-imf-bf16-host.${new-offload-lib-suffix} + COMMAND ${clang} ${imf_host_cxx_flags} --offload-new-driver + -I ${CMAKE_CURRENT_SOURCE_DIR}/imf + ${imf_bf16_fallback_src} + -o ${obj_binary_dir}/fallback-imf-bf16-host.${new-offload-lib-suffix} + DEPENDS ${imf_fallback_bf16_deps} get_imf_fallback_bf16 sycl-compiler + VERBATIM) + add_custom_target(imf_fallback_fp32_spv DEPENDS ${spv_binary_dir}/libsycl-fallback-imf.spv) add_custom_target(imf_fallback_fp32_bc DEPENDS ${bc_binary_dir}/libsycl-fallback-imf.bc) add_custom_target(imf_fallback_fp32_obj DEPENDS ${obj_binary_dir}/libsycl-fallback-imf.${lib-suffix}) add_custom_target(imf_fallback_fp32_host_obj DEPENDS ${obj_binary_dir}/fallback-imf-fp32-host.${lib-suffix}) +add_custom_target(imf_fallback_fp32_new_offload_obj DEPENDS ${obj_binary_dir}/libsycl-fallback-imf.${new-offload-lib-suffix}) +add_custom_target(imf_fallback_fp32_host_new_offload_obj DEPENDS ${obj_binary_dir}/fallback-imf-fp32-host.${new-offload-lib-suffix}) add_dependencies(libsycldevice-spv imf_fallback_fp32_spv) add_dependencies(libsycldevice-bc imf_fallback_fp32_bc) add_dependencies(libsycldevice-obj imf_fallback_fp32_obj) +add_dependencies(libsycldevice-obj imf_fallback_fp32_new_offload_obj) add_custom_target(imf_fallback_fp64_spv DEPENDS ${spv_binary_dir}/libsycl-fallback-imf-fp64.spv) add_custom_target(imf_fallback_fp64_bc DEPENDS ${bc_binary_dir}/libsycl-fallback-imf-fp64.bc) add_custom_target(imf_fallback_fp64_obj DEPENDS ${obj_binary_dir}/libsycl-fallback-imf-fp64.${lib-suffix}) add_custom_target(imf_fallback_fp64_host_obj DEPENDS ${obj_binary_dir}/fallback-imf-fp64-host.${lib-suffix}) +add_custom_target(imf_fallback_fp64_new_offload_obj DEPENDS ${obj_binary_dir}/libsycl-fallback-imf-fp64.${new-offload-lib-suffix}) +add_custom_target(imf_fallback_fp64_host_new_offload_obj DEPENDS ${obj_binary_dir}/fallback-imf-fp64-host.${new-offload-lib-suffix}) add_dependencies(libsycldevice-spv imf_fallback_fp64_spv) add_dependencies(libsycldevice-bc imf_fallback_fp64_bc) add_dependencies(libsycldevice-obj imf_fallback_fp64_obj) +add_dependencies(libsycldevice-obj imf_fallback_fp64_new_offload_obj) add_custom_target(imf_fallback_bf16_spv DEPENDS ${spv_binary_dir}/libsycl-fallback-imf-bf16.spv) add_custom_target(imf_fallback_bf16_bc DEPENDS ${bc_binary_dir}/libsycl-fallback-imf-bf16.bc) add_custom_target(imf_fallback_bf16_obj DEPENDS ${obj_binary_dir}/libsycl-fallback-imf-bf16.${lib-suffix}) add_custom_target(imf_fallback_bf16_host_obj DEPENDS ${obj_binary_dir}/fallback-imf-bf16-host.${lib-suffix}) +add_custom_target(imf_fallback_bf16_new_offload_obj DEPENDS ${obj_binary_dir}/libsycl-fallback-imf-bf16.${new-offload-lib-suffix}) +add_custom_target(imf_fallback_bf16_host_new_offload_obj DEPENDS ${obj_binary_dir}/fallback-imf-bf16-host.${new-offload-lib-suffix}) add_dependencies(libsycldevice-spv imf_fallback_bf16_spv) add_dependencies(libsycldevice-bc imf_fallback_bf16_bc) add_dependencies(libsycldevice-obj imf_fallback_bf16_obj) +add_dependencies(libsycldevice-obj imf_fallback_bf16_new_offload_obj) add_custom_command(OUTPUT ${obj_binary_dir}/imf-fp32-host.${lib-suffix} COMMAND ${clang} ${imf_host_cxx_flags} @@ -356,6 +438,14 @@ add_custom_command(OUTPUT ${obj_binary_dir}/imf-fp32-host.${lib-suffix} DEPENDS ${imf_obj_deps} VERBATIM) +add_custom_command(OUTPUT ${obj_binary_dir}/imf-fp32-host.${new-offload-lib-suffix} + COMMAND ${clang} ${imf_host_cxx_flags} --offload-new-driver + ${CMAKE_CURRENT_SOURCE_DIR}/imf_wrapper.cpp + -o ${obj_binary_dir}/imf-fp32-host.${new-offload-lib-suffix} + MAIN_DEPENDENCY ${CMAKE_CURRENT_SOURCE_DIR}/imf_wrapper.cpp + DEPENDS ${imf_obj_deps} + VERBATIM) + add_custom_command(OUTPUT ${obj_binary_dir}/imf-fp64-host.${lib-suffix} COMMAND ${clang} ${imf_host_cxx_flags} ${CMAKE_CURRENT_SOURCE_DIR}/imf_wrapper_fp64.cpp @@ -364,6 +454,14 @@ add_custom_command(OUTPUT ${obj_binary_dir}/imf-fp64-host.${lib-suffix} DEPENDS ${imf_obj_deps} VERBATIM) +add_custom_command(OUTPUT ${obj_binary_dir}/imf-fp64-host.${new-offload-lib-suffix} + COMMAND ${clang} ${imf_host_cxx_flags} --offload-new-driver + ${CMAKE_CURRENT_SOURCE_DIR}/imf_wrapper_fp64.cpp + -o ${obj_binary_dir}/imf-fp64-host.${new-offload-lib-suffix} + MAIN_DEPENDENCY ${CMAKE_CURRENT_SOURCE_DIR}/imf_wrapper_fp64.cpp + DEPENDS ${imf_obj_deps} + VERBATIM) + add_custom_command(OUTPUT ${obj_binary_dir}/imf-bf16-host.${lib-suffix} COMMAND ${clang} ${imf_host_cxx_flags} ${CMAKE_CURRENT_SOURCE_DIR}/imf_wrapper_bf16.cpp @@ -372,10 +470,22 @@ add_custom_command(OUTPUT ${obj_binary_dir}/imf-bf16-host.${lib-suffix} DEPENDS ${imf_obj_deps} VERBATIM) +add_custom_command(OUTPUT ${obj_binary_dir}/imf-bf16-host.${new-offload-lib-suffix} + COMMAND ${clang} ${imf_host_cxx_flags} --offload-new-driver + ${CMAKE_CURRENT_SOURCE_DIR}/imf_wrapper_bf16.cpp + -o ${obj_binary_dir}/imf-bf16-host.${new-offload-lib-suffix} + MAIN_DEPENDENCY ${CMAKE_CURRENT_SOURCE_DIR}/imf_wrapper_bf16.cpp + DEPENDS ${imf_obj_deps} + VERBATIM) + add_custom_target(imf_fp32_host_obj DEPENDS ${obj_binary_dir}/imf-fp32-host.${lib-suffix}) add_custom_target(imf_fp64_host_obj DEPENDS ${obj_binary_dir}/imf-fp64-host.${lib-suffix}) add_custom_target(imf_bf16_host_obj DEPENDS ${obj_binary_dir}/imf-bf16-host.${lib-suffix}) +add_custom_target(imf_fp32_host_new_offload_obj DEPENDS ${obj_binary_dir}/imf-fp32-host.${new-offload-lib-suffix}) +add_custom_target(imf_fp64_host_new_offload_obj DEPENDS ${obj_binary_dir}/imf-fp64-host.${new-offload-lib-suffix}) +add_custom_target(imf_bf16_host_new_offload_obj DEPENDS ${obj_binary_dir}/imf-bf16-host.${new-offload-lib-suffix}) + add_custom_target(imf_host_obj DEPENDS ${obj_binary_dir}/${devicelib_host_static}) add_custom_command(OUTPUT ${obj_binary_dir}/${devicelib_host_static} COMMAND ${llvm-ar} rcs ${obj_binary_dir}/${devicelib_host_static} @@ -390,7 +500,22 @@ add_custom_command(OUTPUT ${obj_binary_dir}/${devicelib_host_static} DEPENDS imf_bf16_host_obj imf_fallback_bf16_host_obj DEPENDS sycl-compiler VERBATIM) +add_custom_target(imf_host_new_offload_obj DEPENDS ${obj_binary_dir}/${devicelib_host_static_new_offload}) +add_custom_command(OUTPUT ${obj_binary_dir}/${devicelib_host_static_new_offload} + COMMAND ${llvm-ar} rcs ${obj_binary_dir}/${devicelib_host_static_new_offload} + ${obj_binary_dir}/imf-fp32-host.${new-offload-lib-suffix} + ${obj_binary_dir}/fallback-imf-fp32-host.${new-offload-lib-suffix} + ${obj_binary_dir}/imf-fp64-host.${new-offload-lib-suffix} + ${obj_binary_dir}/fallback-imf-fp64-host.${new-offload-lib-suffix} + ${obj_binary_dir}/imf-bf16-host.${new-offload-lib-suffix} + ${obj_binary_dir}/fallback-imf-bf16-host.${new-offload-lib-suffix} + DEPENDS imf_fp32_host_new_offload_obj imf_fallback_fp32_host_new_offload_obj + DEPENDS imf_fp64_host_new_offload_obj imf_fallback_fp64_host_new_offload_obj + DEPENDS imf_bf16_host_new_offload_obj imf_fallback_bf16_host_new_offload_obj + DEPENDS sycl-compiler + VERBATIM) add_dependencies(libsycldevice-obj imf_host_obj) +add_dependencies(libsycldevice-obj imf_host_new_offload_obj) install(FILES ${spv_binary_dir}/libsycl-fallback-imf.spv ${spv_binary_dir}/libsycl-fallback-imf-fp64.spv ${spv_binary_dir}/libsycl-fallback-imf-bf16.spv @@ -409,3 +534,10 @@ install(FILES ${obj_binary_dir}/libsycl-fallback-imf.${lib-suffix} ${obj_binary_dir}/${devicelib_host_static} DESTINATION ${install_dest_lib} COMPONENT libsycldevice) + +install(FILES ${obj_binary_dir}/libsycl-fallback-imf.${new-offload-lib-suffix} + ${obj_binary_dir}/libsycl-fallback-imf-fp64.${new-offload-lib-suffix} + ${obj_binary_dir}/libsycl-fallback-imf-bf16.${new-offload-lib-suffix} + ${obj_binary_dir}/${devicelib_host_static_new_offload} + DESTINATION ${install_dest_lib} + COMPONENT libsycldevice)