From 16952fca5c06c1ccc8290b641ad15503dcafbfd9 Mon Sep 17 00:00:00 2001 From: Arvind Sudarsanam Date: Wed, 24 Apr 2024 17:08:43 -0700 Subject: [PATCH 1/9] [NewOffloadModel][SYCL DeviceLib] Generate SYCL device library objects using new offload model Following are changes in this PR: 1. Generate new SYCL device library files using new offload model 2. Pass these new files to clang-linker-wrapper 3. Avoid use of these files in old offload model 4. Remove support for bundled objects in clang-linker-wrapper Signed-off-by: Arvind Sudarsanam --- clang/lib/Driver/Driver.cpp | 5 +- clang/lib/Driver/ToolChains/Clang.cpp | 6 +- clang/lib/Driver/ToolChains/SYCL.cpp | 28 ++- .../ClangLinkerWrapper.cpp | 202 +++++------------- libdevice/cmake/modules/SYCLLibdevice.cmake | 136 +++++++++++- 5 files changed, 218 insertions(+), 159 deletions(-) diff --git a/clang/lib/Driver/Driver.cpp b/clang/lib/Driver/Driver.cpp index 1a1d4db9a1e46..dd81d59c16eac 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; diff --git a/clang/lib/Driver/ToolChains/Clang.cpp b/clang/lib/Driver/ToolChains/Clang.cpp index 6fc3101a9d415..dccc82161512d 100644 --- a/clang/lib/Driver/ToolChains/Clang.cpp +++ b/clang/lib/Driver/ToolChains/Clang.cpp @@ -11049,7 +11049,11 @@ 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 = Args.hasFlag(options::OPT_offload_new_driver, + options::OPT_no_offload_new_driver, false); + 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..f2f06ac4204b0 100644 --- a/clang/lib/Driver/ToolChains/SYCL.cpp +++ b/clang/lib/Driver/ToolChains/SYCL.cpp @@ -291,12 +291,17 @@ SYCL::getDeviceLibraries(const Compilation &C, const llvm::Triple &TargetTriple, const SYCLDeviceLibsList SYCLDeviceSanitizerLibs = { {"libsycl-sanitizer", "internal"}}; #endif + bool IsWindowsMSVCEnv = + C.getDefaultToolChain().getTriple().isWindowsMSVCEnvironment(); + bool IsNewOffload = Args.hasFlag(options::OPT_offload_new_driver, + options::OPT_no_offload_new_driver, false); 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]) @@ -434,13 +439,11 @@ const char *SYCL::Linker::constructLLVMLinkCommand( auto isSYCLDeviceLib = [&](const InputInfo &II) { const ToolChain *HostTC = C.getSingleOffloadToolChain(); const bool IsNVPTX = this->getToolChain().getTriple().isNVPTX(); + bool IsWindowsMSVCEnv = HostTC->getTriple().isWindowsMSVCEnvironment(); StringRef LibPostfix = ".bc"; - if (IsNVPTX) { - LibPostfix = ".o"; - if (HostTC->getTriple().isWindowsMSVCEnvironment() && - C.getDriver().IsCLMode()) - LibPostfix = ".obj"; - } + if (IsNVPTX) + LibPostfix = (IsWindowsMSVCEnv && C.getDriver().IsCLMode()) ? ".obj" : ".o"; + StringRef NewLibPostfix = (IsWindowsMSVCEnv && C.getDriver().IsCLMode()) ? ".new.obj" : ".new.o"; std::string FileName = this->getToolChain().getInputFilename(II); StringRef InputFilename = llvm::sys::path::filename(FileName); if (IsNVPTX || IsSYCLNativeCPU) { @@ -448,12 +451,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/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp b/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp index c01effcc4bf24..79d173e3b002f 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,78 +492,6 @@ 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. static Expected convertSPIRVToIR(StringRef Filename, const ArgList &Args) { @@ -558,10 +511,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 +833,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 +851,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 +908,31 @@ 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); + 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); + } + } + } + for (auto &File : ExtractedDeviceLibFiles) InputFilesVec.emplace_back(File); // second llvm-link step auto DeviceLinkedFile = sycl::linkDeviceLibFiles(InputFilesVec, Args); @@ -1280,7 +1232,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 +1439,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 +2016,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::getFileOrSTDIN(*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..42dba45fe0d01 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}) @@ -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) From ee4227c6ee4524dc009b05bc14ff84447daa5d41 Mon Sep 17 00:00:00 2001 From: Arvind Sudarsanam Date: Fri, 26 Apr 2024 18:46:14 -0700 Subject: [PATCH 2/9] Remove --offload-new-driver when building SYCL device libs for nvidia targets Signed-off-by: Arvind Sudarsanam --- libdevice/cmake/modules/SYCLLibdevice.cmake | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) diff --git a/libdevice/cmake/modules/SYCLLibdevice.cmake b/libdevice/cmake/modules/SYCLLibdevice.cmake index 42dba45fe0d01..692deb852aa80 100644 --- a/libdevice/cmake/modules/SYCLLibdevice.cmake +++ b/libdevice/cmake/modules/SYCLLibdevice.cmake @@ -59,7 +59,8 @@ if ("NVPTX" IN_LIST LLVM_TARGETS_TO_BUILD) list(APPEND compile_opts "-fno-sycl-libspirv" "-fno-bundle-offload-arch" - "-nocudalib") + "-nocudalib" + "--no-offload-new-driver") endif() if (WIN32) From f1caa3780a64ec70ecf548ce9b84bfeba3139b0b Mon Sep 17 00:00:00 2001 From: Arvind Sudarsanam Date: Wed, 1 May 2024 11:40:37 -0700 Subject: [PATCH 3/9] Fix test failures and reintroduce --offload-new-driver for nvptx Signed-off-by: Arvind Sudarsanam --- clang/lib/Driver/Driver.cpp | 5 ++++ clang/lib/Driver/ToolChains/SYCL.cpp | 12 +++++---- clang/test/Driver/linker-wrapper-sycl.cpp | 9 +++---- clang/test/Driver/sycl-offload-new-driver.c | 27 +++++++++------------ libdevice/cmake/modules/SYCLLibdevice.cmake | 4 +-- 5 files changed, 28 insertions(+), 29 deletions(-) diff --git a/clang/lib/Driver/Driver.cpp b/clang/lib/Driver/Driver.cpp index dd81d59c16eac..70960136e50c7 100644 --- a/clang/lib/Driver/Driver.cpp +++ b/clang/lib/Driver/Driver.cpp @@ -7880,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/SYCL.cpp b/clang/lib/Driver/ToolChains/SYCL.cpp index f2f06ac4204b0..5e040c20fdf67 100644 --- a/clang/lib/Driver/ToolChains/SYCL.cpp +++ b/clang/lib/Driver/ToolChains/SYCL.cpp @@ -438,14 +438,16 @@ const char *SYCL::Linker::constructLLVMLinkCommand( }; auto isSYCLDeviceLib = [&](const InputInfo &II) { const ToolChain *HostTC = C.getSingleOffloadToolChain(); - const bool IsNVPTX = this->getToolChain().getTriple().isNVPTX(); - bool IsWindowsMSVCEnv = HostTC->getTriple().isWindowsMSVCEnvironment(); StringRef LibPostfix = ".bc"; - if (IsNVPTX) - LibPostfix = (IsWindowsMSVCEnv && C.getDriver().IsCLMode()) ? ".obj" : ".o"; - StringRef NewLibPostfix = (IsWindowsMSVCEnv && C.getDriver().IsCLMode()) ? ".new.obj" : ".new.o"; + StringRef NewLibPostfix = ".new.o"; + if (HostTC->getTriple().isWindowsMSVCEnvironment() && + C.getDriver().IsCLMode()) { + LibPostfix = ".obj"; + NewLibPostfix = ".new.obj"; + } std::string FileName = this->getToolChain().getInputFilename(II); StringRef InputFilename = llvm::sys::path::filename(FileName); + const bool IsNVPTX = this->getToolChain().getTriple().isNVPTX(); if (IsNVPTX || IsSYCLNativeCPU) { // Linking SYCL Device libs requires libclc as well as libdevice if ((InputFilename.find("libspirv") != InputFilename.npos || 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/libdevice/cmake/modules/SYCLLibdevice.cmake b/libdevice/cmake/modules/SYCLLibdevice.cmake index 692deb852aa80..8398cc267220f 100644 --- a/libdevice/cmake/modules/SYCLLibdevice.cmake +++ b/libdevice/cmake/modules/SYCLLibdevice.cmake @@ -58,9 +58,7 @@ 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" - "--no-offload-new-driver") + "-nocudalib") endif() if (WIN32) From c43cef5d2f9cc35246d5f9ec5ad0c307c326b9e5 Mon Sep 17 00:00:00 2001 From: Arvind Sudarsanam Date: Wed, 1 May 2024 17:29:49 -0700 Subject: [PATCH 4/9] Add dummy device arch for nvidia device lib gen Signed-off-by: Arvind Sudarsanam --- libdevice/cmake/modules/SYCLLibdevice.cmake | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) diff --git a/libdevice/cmake/modules/SYCLLibdevice.cmake b/libdevice/cmake/modules/SYCLLibdevice.cmake index 8398cc267220f..6de9124db0dba 100644 --- a/libdevice/cmake/modules/SYCLLibdevice.cmake +++ b/libdevice/cmake/modules/SYCLLibdevice.cmake @@ -58,7 +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" - "-nocudalib") + "-nocudalib" + "--cuda-gpu-arch=sm_50") endif() if (WIN32) From 41bc9e91597932728e34a3f4d77653eed8c7ff53 Mon Sep 17 00:00:00 2001 From: Arvind Sudarsanam Date: Wed, 1 May 2024 18:01:34 -0700 Subject: [PATCH 5/9] Adding updated copies of device library files Signed-off-by: Arvind Sudarsanam --- clang/test/Driver/Inputs/libsycl-complex.new.o | Bin 0 -> 70232 bytes clang/test/Driver/Inputs/libsycl-crt.new.o | Bin 0 -> 30224 bytes 2 files changed, 0 insertions(+), 0 deletions(-) create mode 100644 clang/test/Driver/Inputs/libsycl-complex.new.o create mode 100644 clang/test/Driver/Inputs/libsycl-crt.new.o 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 0000000000000000000000000000000000000000..c7501e92dc7a9c5bb419a30ec608578fc5bfd264 GIT binary patch literal 70232 zcmeHw3tSUt*7t;jOpGB3iW*9Y6K@-D7(kJ$27-!}ZV@l5-6{zo2rFPn5CQvj6M}?V zTJhG!t?#?RilwdXqU*NquI&O*mm;;)T3@QJKd4}7U)x&Sy7pV&^*hfbPZEel(Z1hm zo%#KeGjnq0nKS3PX6E@n|6k8bo$tqD5qM(}a|q87))It7?9r-dwTc)9mg1Y@*U>Mv z5l&mmDF}s2`Vky@FU5WA+^KaP6Rz`YuHzh+taiy7MR~!+uja4C-=+6t5ah}aI2hHeSr zvw}U@F5XL|@zRqN+{y9^B{t6@ zQdwrU%J2VDjhie=vXo~lR@kDL&v#4D z+Cy6`vJRG=EN!qEiKaCRidDbg_ulek@m=?XX~VA3E2g>UyH0Q zQDz3;iV5C^5?d79HIFPY+lmY}Q;xNmN0zBbv#h61*2bbPS!69*>6yBKu0+`dI8iI@ z>kjRy3piRQYj2i9I%hSqQ;B2SSWRq=b(I)haxUT5$Td$U3`9yrOshKjkUHiNEB=T& z_AonPhc%w)ZW3N5#9v+!=Bq?k#lkCu=(8^2ryblc#KLdQLL#A6ov?$gYG)@@IHLEf zLz=AdN8m(9bcZ9ZRUJ~LPN=oUKstw-|mQoB3$&nQg|6sxFYOH6J1Ud=GO>& zh$F1HBaY|_c3h1*w%rli+8k2hh(FDasj$YktD}#wAQ8Aw9GEp)6Ar0Ez@SkbTg68D zy^iP|j*x0c+#z;+qdKO`G2^^7zJeX!q>gE|hE&52c&MIJ&5k?N9M{B70Gm#AOe32j zo47CY%&UkI)R^i5o?oUllSfRJIK-Tz+)GGrZ0@V1s3=7%QBxi+NL1xfBUsd?fQkrW zZ=+=e&m0sbPG!}!a9{c!g_5;`tvcHj4SCuW>`Ek%NGDDGqo2?>oGvB}bh3p6X@o?&gis&Q6(Mojfk))2 z!kb;kp5EffGnT#c!rLw1#Kuhfoh_s1;M?s_ANcrM@uF*)FCXhUn2-?m*W+JB?Ot{K z*I!*#N>{(|m*a-(HGE~oGQtiqRl$ewYa%L_R!-QTWZu77cXa-xp0WGibt>|=^rKfBo(nKX@x)W2s#Nq6d};0 z{E@%4k>4z*PeOi3nrb$W5$Hw{RiN@*F*@b=6wU%jn*ue#(h%w?Ws2%$1jY7t{(EsS zI0l`A8ClVeE#0F8$DmvAIQDjO7ThB$O745~d-fA>EEb&vAKI~Hqdz#nZ??c8iU#di zTbM8lLIGa59b1->+B$_}gM=usayr(j^l|J>xsPKV4j;$sFu1iD6qtX3Kvm|9akY|?_`9CiA{G+u8S$AAC3X>XFf&!_F< z*T0>%FF^lQT8p?>vH1C*-i*a3$Mo)6{N*FPCl+s(>HlZ(9{I!>>1`A0bJMT>tbTWT z^)&x~!W9od`}!HAAIURIk9JG5gIEbw>gWnbb&{DZEg?;{g@Q|j$k~fU3X)_k&7mjj zq}fV;+RPdPJsNav)$D}!lIkJ@xsgYfrq~Pyvdq9PF0mOXTX1`z@N$i?TPZ3~iY{Ie zUP%*1ijRg~kV`w|p)Fcjw(^;{Bi6VM=ug@4M^Jwh*M23aSVL~iA!p~kjSU7ZMRr^8 z$t9vM3DMOwVNZ?lqDs_F386cVq!gj&n`JGDvQu&?3_s3yNqf3OKWE9ZC5xk>&)sc} zYfqctx)B&S!KkKoaq71p;xrrrJJnL?T)xcbpizB zMoWB;I<5*19o5J9aHnYWRRkj$1!=6cJb6(pyxb)0O%r~j9N&B@QP$Hvrn5Wr0yvNo zO!xj)7#l=at7F<7@rRpZS{-nY@$D&x3V9@WRR{rGVGBNYTzFX{`l5?_g%b5DMZJ{h zTAHvxEWD@`!r?bvTo}gmqzSJp$A5BJD{Zlac6N`sEDyaT4?QXmg&gdYkIkn3BkqVg z;V>Fw#I&zC0=W-2TLyln*ftn!!6!Zyex($3cL_gJi9WN#RU9m~sey!h$*6^*f+CED zowb&VL>QAm=n^2O+EqvGFzgIPBeZ5|_5__$h=y275hT*9RLiaFC8R0EmSML&WrlMP z>|n)LJL1uRN*xQusU0FFNRf|9(#pWqlg**&mZ0p^HA>-El<2E2?pJA|Z&6ah-Y)JX zO4Qu}H+xZKEv-T0PjD=F$s^8z(2b;+E;Lw*tyaf&IA*k~V-K-o+N(}?$u1-V zh5=Y?t2N;;YZ{E=D%5d1U}R{GuStr5p&AS|U3bJLN-2U^0Cm+UOQFFmF4#K zW?p#`35w8mr8qI%q2c{a;8)mA1+@yA!Tn)vRnNWm#(VW&Jh%V(*V~U(Prsa35X4$& zEVE8Eh3HJXtc$jr4ouzt;-ZbFkcnyA(D5HkyUJeM9x^p#;)-;w?hk8C3&&42y%aK0 zLlCb(do~?nG@Nf#Ks!N}tuWL=*+j)Ero&NwNqXG*C#e|~;XbNG!#!5R<}B!y2L^4gtoJl9S1HGA3#5j22*4^ z{0WHL4UH2!!ETL*hN{Ma`WP)mK+gbUVNT*SFs1vQ9zKis8kDH>7|A&%jN%B1$Uy4| z@+Fm8@~>e>Rf+JO9G<5NI!Xva#1N8TX#dv=X$-(B9m=eO*IW5F-Ho*#a`u0oCOaoks-n7*1?xXk4 z@1ym8z7K!V9p^{W{d^x2@BBUv|9kIa5tOhWt{?Nf?n5~;Tv4dxPY&8QPH>a+lARLV ztoTF{MF}#A%77}B#2mqm3l;EPx3V&TXOoz<>~UhQ#$h3*!1c1d z>Q7XCYWqd@HX>!os#UqOW2Vp8^4aT+H_q?*%{jv_{v0~T@!mp~z_Eq*%oF>zZ9Y|C zd}6ii!(VE@|I@q~Z*`BJmD;X*;{BuX?#fvCohz~JV*{Hqe<{{To=#3V9f9iU`>dek z$a%F1{=4*vBfq(BoVTJmfLkbQdvO z{*#Z{nnCMky>HB%9a)W98{2Z_ag~7IVxIt2bz&W=s#(EN)D9i0s&)x+v!VhW(lu8D zHP?C(YHsvu`O2V7I5S7ehwoaumR&xMtLZfJDp;Kcih!r%1hE{~*X{YUYH`0+sh#y3 zyg$G2mi73H$%j)|EUUA^~()=`}dC} z`~TcHK-0dn!1lhulukIVj8EG0!YliHy3~M?-$*(m!?$3kDoYb^rC-2llq@x`lyD@K zV*^M0F%@;YCc^)N)WGbF4zIRk%SzWhJ)KMp|KYYEa?m_F=9Nb^iQR!#_Q8cp)qJ^v zHv%Fo;h{LERBPEua_=~!=keL@rfiGStX=AE<)E2zH<{a4C7QL@1LccP&`lYcCVnqS z%UzaAhO@5COt$Hk^}C;#$u8rpbjQ(nxWrwyZ@pgkbfF*+DXMgfBn>pxnhfQ8F zqn|rWfaeY=r>N-aa!L*lgDmvZ+hg$bR)U}229SPe$_aPYjbCe)wYeUI>^gnrrq%P{ z;Bm)I_I&izb%pzR1Uy>xJgPuXT~~TPb=@*A;N}deDskqLijPOl^gNF_&i@a<-l|ORfbpgkLsGYean?{EwM>Rje?IEEA3f=8TwZQCt|&a4^;~?fS^RPaZ|@ zy3nSF2|RP*U<7-|74!@VX8jNu@S+P%j7fzLvReTkv>^$vnusfCngwiI!7dc!8zAR` zMDNwmhJQ!<$R2O_cO=myHSq-JIV0d0>5xQv3FzS!Bdx;{+sj;bSYms4Jx9-7#S+`k z0(~HSu*CNBUG`XFd-O~O*<*?Ap9CH95ld{J2|DCAme{@sbZ9Oume{@&bSQl+vHfP3 zJ(k#hi_0EMY;SkjV~OqSUG`XF`!_*P!Us!i|2F8*{8}ureFx~!{bGsj&wviyFP7N8 z&t;D#w*S&)k0rK81q-E*CAR0HvWU`hNo+51*<*?AWiC4`vAq)XDE?St`&po;%Y#c| z`}r<=EV2C}mmQYa{z;cTme^kFvd0qJ8$plm7fWnk0(x}6SYrFlphx+ICAQxRdX!&S zV*9`HM=!xq`dDK7GcJ29vHjm&_E=*3E1*aBizT+d26}YA zSYmq~v>`muV~Op>pbr8)me_uR%N|Q?Kh0&2CAL?Aek`PeCDQYtok80wXaoNW657Bo zP?^!?=4;nvXJlmQi*t4QyxhzTr$%SgYxDZ+X^l>69AJ(VI&G$@zcEti490c8>aV6X zI(>E#-jc|&>Ed7w;*SL9%*#Cb*2G& zTB9q{4z#B=C>@_<=v4Y^A(3p9Mt?0*=(Om!ei5fNI<2pZP8I#jfHROAbV>u9M@hk5 zX?@*ys_0(^q~&xoe|?@w7wPR+MuuCLTkP{&mQ(jccyvYv0$>zptkdOX7>qC{C_8Og zWa5&DMY=N2Y@j=x=`$#nV@BROoi;Bk*R(dmGnEL6u0GE-p=bt_U}sUD0n?2D@lleo z!KgJD^u{oq!GJ0NY#5hP72!OloZ_~)lp2@a^>F7-h3h!SC7*H0fw+`4dN&y5mY|fR zJHVru_zWI}!J{yE6b6rShj^4j9z04kOhC3p!3^0IFp<`_vHyO+J2VmqU?^pPN1?$Y z#WXquKqLSbaRE?hq=*xhQUsGAot8BuvZTlo`m9pQgFZo!6gMy>4cU~@-;kygg92a| z1WD=OehH}o@P&rFpk(cVEj?F6t1tf}x_#2ipRWCQYW(lFJFb>qpBHxS#N{zB=C6Az zYew-CsbS|nym)o(jRMEV`WycU6FIu|e_AWBiC!VpfG2>a(la|=pV9)*3eZrfb!H*V zN4pp%+QS%${=-M3OiH&6gpcrrd;qMVmR$@VGXktE3r+S(YKa12UtpZ40xQA~VP9So z``}tini1vz4f((lhJG-${~4Zcwq)72o3_=T>eA_j-fHC+p+JCRRHEd0V)(B8oN^--~)!xhxo!r5`0h^J2(2G zG{nAOh)Lsozz~L`J}8avt9&9>zTGEc%YNn)v9_hY5qs4)Vyk^|E#1DjmY!oij&(ol z<5>DGpNNf$^ns6T&hw4fS>K4g_oBBF<6wL$Bex$W3VVgj~RnS8G6hZEXvSh z#$Zu~9y11u;)xl9MfrDQYzirFEJ`z}Py!aEo|j_)EXp&(8Z>}Kd5Pag02bxr;fij+ zqMQ@tSOAOiT|k2zuqazZea(PHc}1+q1T2apFvkzDD1RT(um`Xx{}|bK>QOY$25t9> zML}_T4QA*tSQO_Rdp!3YV^FX}dW5~f7!)kAJ;tD5iS0241xsvihq$A8^H^g0de9-P z1(w(zV^FZf_85bLCARN?G#D%jV3MGspm<`5%2{_V>DYCZKy9p+dv#jwqcDH15mK^Wk(^1JHew| z6W$>nMeO89sc`e7Bn`rkl7#qC0>^ptqa^j`M?rIyUHmA+5I@Q=%#V@`{3uD2crwI~ zVh$WPpO`efDk3dW-azmZM=ch-T2&R}r>vbOe$^f)2A&i&M;vjZ90tj7qcGekKVfbZ zzkB992?Wj)EpVPVRQ)+m-t=Hhlw?V?&EQRP;=TV9B|L$)!!_^SO?ZMd{RvNyhQ^P$ z2~UuQpy$SW2v3lT28MVDPmqdE%0qa9R09Z4kZJ(o2~s()=q5ZtD(9qqH{l7b@6n}FgjUgxxd~5b z744XZ@C4~-$2^25NJTs5Av{4U+A$B|2~yFHc?eICigwIHc!E?;$J~S`w2F4jO?X19 zXvaK+CrC#-<{>;mDyL&^!V{#T9rF;LAeA#K-GnDd<;+Sq;R#YX9di?&&??$7H{l7b zq8;-Po**6Vn1}ELsc6SMgeOQvJLVxgK`Lif-jeX-y-=F)1g%LwfbZng2t|caQSNsp zOL5L>T%cI8>C|e)hm{96D700^9f}2;ul`=~>9bec6&s$r+N+qi<>0>*AJl9cGg-gw z)Xd4>Z%_aES`rU`-M z2N>|atr33J1=KV;ornBn1+W7l_Cmy_QA%wxN^Bc|*azreO2`ehB+YNK9O~um_)T)C zHHhEj*f(JN@8mei&ReG~T9;ptYgkvLFDRlA-O7?&w2A?xo;sCUnv|GG+>!^Snu>ly z=`e6w1?q8%NS-aiYMvI!aR@BojZi|*qK5L|&Hb=UMH!TfSNRr$MaWt+KN@-&y7 zctfzDoa^9LqzRfDS%Ct!GBa&v)G%&R7$vP>vxtdsji`^3k1l0BOE1P=jQJK|;jn#( zJ7vX=z`!N`({`?}|9j7&Bk}v*YPo89c6EVc)n~=m1#`{iIs15#Q-1ahYay}s-7P## z>E+3-`(J(6xsd#LzZLxde)KQz=9ZSGuG`f6>FBr5SN^4^R~fQ8eev1HetzPFu59YG zxo`FLMxEck+eyJPR!dw&of-?)xbOoBJS&YL=uJxz#NaVOJS%q}la}c}#h4(R%-zQ% z=_eVJiT4zfmT3=WIe8jt(GRyL2)~2M3F2D1MSIfjX-~HGZ%;JQ(4JIIlUDsszBq`V zc|CY^Wq_ipdyQH8xA(;f>y&va-Z8}0)T|_)<_;@IKiMb5pjN7Rk7BN4*+^Kf04>+722S2a*vnsn6R3uMFR54d;S9;SySVNX?{~j~MNIqtJm?A0{Pj>N0Wc|`j zt^kw_kLTW!^y2k4PeUae{;Z3yf^NoswQ+FzirWA3rmxsJ4--8AeMS6R&$$l~h0U9V z-sBdwo!DO>u6WjRD~5KO%a$pY#9aLEHa&-~7iIu>-&BR`5~FzKu<{ z>}lI7(BLv#%SOr<_I(|xJQ)Hld{PAA;3gXc=HFRLWpmZ+W2j}9(k=TjqFN%|@>qFu z;_&w_UJte__CWie_}f0ZSK0sS4Z*JUq}M35YfZ-+&!TIDeMn#9Nd7mGPdzynT_YS_ zLn&yggwZQO)lsK&fx|YsDmo&-AysdnKxIu`-jy5mk_NPYGAcG@4jM7t`}L7 z<(Bg2 z`o({}yYok7^96UGk48M5EAWU1<+5oky+-rOrQxARIg$Z`q0{2f=kxoK=0$Bvwq0a9 z#2^2{JCP+J{e~V1Zoa8ZkyjyE4(ms=o&UekJRtP{>rFwjDnEF$`D}!Q^+$cQA9{C* zHZ4UEgO4eCecP1DL7xApw;5A%r!alW{obbxp8p8Tcn|+6CNuf}VCVnDd&++e0qrv5+J^R_4kzjTqge?@cMu3p7Os^ z#?=1@yZ-myQ~q@@_$^Qqu{?EIo{I8TU_5Z=H z|N48%e=d{%_kI0;3f=?#@SLJeuKe%!zOLT^nr&2rUfyY^q4&f=ZD^)YGJ21E)GQ!^ z!&|Eo=L;e-t0fb*K31MMj`R=bS$@~NVl;a3T*DR8lZpQGZ||oin)=T(nz|YN)HEp$ zz3@wXZOU9ncCYsfKQzxs1-?i1dv{#>1b$%#?~adKL(Y1?N(Qdfv9&KmxL8uJ^?O|> z8GZGpRU}Yn{i5~eaP-dav4ER%SGZr)p?7#o&TaQ|%~S#s*SVUNj-&CDuBCeB`Dk<9 zZxiA5;rsb(+=nn1b4$L)`$cbdzZbpBTW^x+qFdcB6wwRaNwaSJ)j;@vkQ$h+b-&<7 zNiCD#6kTspCd2L1^DNQ(Jm-79`S70ad3eu9!F#?Gzh9lOb56gY9o8_NJu>Z=m6Pn~ z1lY6`K@6Vn^*?gH`!V@$p7}kLihpaq`>ai(sWNj;@N$S?yEJ_}+hcxQ zEU~@BWsfDc$NachVtXanqje^*#P$qNu9JrvuStR>j*kQ4_Av6~z5`_i)+YIHW^EFN zCl`2;5rHl+=ejv@-Hgdzy8rJyxwl%=<40MX1j;^MpTvBZJh_)$Jh^5_3vGCP61-l? zfb~i6k|bFAEKPDJ>yi}0V0yTpoW?E%!!=3E_^pYpoP}I*M9el;gqy^Gs)+S-rCiw}k zO>(v=nx*151xp-a&Qb2AKq8$u*(9l?s3>IzQR8Re9#?W$F#I6yc@2|7Si|INSi|I^ zrT-cx``v4pFbug6XXgSA3_~tNg2OQ6E`X{>K`5N1f?>$zK`5N-jr^}OJF_VMfAPTLot|0=CT+^bmpd{A%3;*(=~cP;+%k=_%F zH_P<@vv`ku;*9jR3H7<@SASN&JH2|E|3BgR3|hD3Gf1CVw}e@@q?cK@Jc(IG$-68B%sxv*ym82_dg(a|J30an6+hDMj)*kp{b9|FJw$++&m^H1! z8egG~+rduQVU4dzim7(QwW^~#)UijbaR3M^r3hjH)YaRqWfCo>&{`({)^$v7xkgF5 zQk)p>(D42y@GESmf?5U5)OgFVwyNjed*i+OFP_`~{Oj$j-=*9c>jjh#ftJ;Z)7=I zZzSlWfpEh5bKQ|Q2>za3>w4HBbtz7oikJua>V2%7im>q7d{=;oSioV_JqEcjz%=_;EckWB76JFMiy=k_s{KUeprCmZFEbbvA`{|(IP)+D0`uZxiRbuZUR*4({bo29&40%d+hbl_EU~@a zWsfDcuLnKibHx(dV_sY=u|4L+#S+_hKpF@mf+e;;13Gk{SYmsOJ;4&&W9$i**dAj~ zu*CKldx9ml$Ji4ru|39~V2SN9_5@38&tOm(3<~1OMdg7v0d`3DrS{ zAD7|By$>-c7_Wilj|D$*<4+8Ci9hLb;ZH26GQo!NCzZ~_{`|PbQ8V!kU^x(gawqtc z0zQL3xlk9-mndst$-3mB=ewn6?V&9eSqIBbmNwXoL{+KXW{}v{=dep^Y#Z!0ql7Gj zS#de+!Uoc$v>9t`QMIv?lJisa-K*E9UQI)8n z9&4UOOlXnz$falNq-P1)1-M^}tSwPy2H%Pa-i8ue6x=nBEHT@P3^r4awHPMjo57+rEM z;n&DDPbLgRN<|Dz5Im%gImC)TqK-XGPaGt=n}nAM@t0SG`6>~xrdJJ_mrc0z?CdcQiP$r^tIPIN?fpy`GoRqBLVYYa?4Jk%V&*E(kh+e{MS zA_jl*zlcA1!-+rH*vO9cX3X{8?@Sxs+rj-zLl!}CRFU}_QK5z`6O+Xq+-nq4nkb}% z?=mH61f_~DxU^*{ZD};qXVs8KJ84XFowy!;qDc6LbJvT}Pv{#?7n24$*+POeLLvr! zp*^sr=Za|c<$pxCPkQ;&wI5H7|NVBy)za(p!p@zzJm$sxb#G* zJvx!E8%4B&%I(-ki;|B|;kXeL1iHsY#q;18+uJdHU??~S37O$!_y)%=f1(1%(7zeC zW8WRC0LK)daywQY;27wbK4jZK$M^|?fsXCmnBwD@IJ|lpLG^QN(zsx7491;ixT@1J z!%-A57?pY)`@YI2V&&U?BDU;jJ`rnM>Kn0FeIwQ%b3&l3>>lCcSkEyZ$GV^OaV&k8 zPsBz=ZXX!2&3V2NJL?;zkGRFefk_*eo9< zSp4o34hi{2p**6?SVYn`C}WUsDq4jimX|ycvJ&#mOsh~CQ+=jeE~k&s7m)APbW&-j z7qJLxIII}RU{28J0=oX<>;A!dw( z(sug$yjI#KmmYG;~Ukn)#}&|$BcG$>>+kcd)4VK*@Z;eneH(sVW3zi>*^kJ zNh`g;f)OKwIf+!xnyH)2U`{%<%`DNkl&CvR1S9?7RqG|BDaGdhel!4xI-wDAz_>oH zS{+l(PH1Jvp;rF`7^|RXsbo70_TzR#5L$}CoT!41q80HhN@6^B zo@u*=VlXE^3RXi2bi$}UIBwSL4WaJ zJxBuxdec(3!|80gi*m+bPCWApTy>MdoIG9d!HN9C3xV>I!JO!p^~2D_xcM5>VdYp< zG4Bpj*LxRr_0fmCPw>AB{KWi+JZ22$Wa$0qyt^o84CchAADtX@8;mf6IT?D)7|hAg zW5!@kh8{Bpb29XpF_@E~$Be<83_WHH=49wGV=yN}j~RnG@x+Y5oZMfSlOv>}>uEE` z&dZ?yb5b#^K@6Ca|K#^I0_NoOaK%=@oSYHlgahW}NZQ&wYmtME<7S4Gh3S#rHzh?h?%XA ziHeGj)H`P{p~nA*H?)Qthxq!Yk40FKzNCmim(JH0z)qIlke8cjB0yot)0bpyFlr43 zy>WV`E_`;>tT3IyK!g>nE7FH8`1!KwrXsCwEfI!X4O^F;owrV#m0PgJsmRb~Wx)-@ zk*Ua7q|F3bnvdjp+;{%M9awIj-u@q|mRQ>W literal 0 HcmV?d00001 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 0000000000000000000000000000000000000000..2fd53a58c4aa84ee257aeb21d43d28a8b1a8eece GIT binary patch literal 30224 zcmeHQ4Omm>y??_AIRWFzQBgw)@kFepogV=#4G_9~po^Av0qursHxj;tm5&f20^Qvb z5=pGsMe90x+wDSi_wK`Wv1LB)UY;!w>QK=!=UQsLE)1~R*-Dl6ad-9h{@;@v2+>-t z?Yi~meV&v5d;aHr|L@0n&-=dlonQX>&g%R%F+3jP0uSAfrq4JnLPHJHaWxxNqc})O ze?b2nr>|TPyv_s#N)5>v#Ajprsqu6?YwZc8mEk;zh8iL90?b9H$zgI4D|b8iAdv2a=t#LDldFh4DBFeLKah|8E_4$}YKY4V1 zTw#0ropf@xbpWAl3_`f87FY1;7qzWx=M=)n&iX!sK0a>g>O-;G!llCJ54RehLH8qA^s2}19H zOd4B!=h+DL;z5n9uPgaT*X+PZ@`>K$b6wIgo{(%bk~XtzOC{0bahYvo zlb+n7bY)h$%$=kaozuwP?2@7q**UrNkSX~Tl8uevjEC&9GaA_`we-zi=^)tYntfKD zJP9dP$lW~PsyAW_TNqc1(PgQ0SQ#=fdw3-IkX$-!m!2@mUeQoz80pzA*??R+W|H-d z%pT~HUg(kzkIWvnCtv7FKBcD4*pg3Zq*yOcKGY>WW0a2BW#1l=jmf0}p6rN5cG`#| zP`67a>2TNVGd$Tjd-9Q9*%5m(%0ij67J2#vB{V7=R?~z!KIt1>!44_gK6>blC@q+}#cp8u3vV5dHQ9RihSPqNxcaJ|Cdo7!tfi zQ6x?$agt9I773=vNN;a|U!tZe4DNBMaF#eQa=LqSd_=BS zmo==*@)8W`6a1OMozKKV=ezrn-bBCvxx1QKz-PA3=0An@dhL}7wgl*<%WN;H(Z3@f zkC}6p5swT}<3oap2@1Mm(m-88P;D=vKIqTs=0S%jeFQW%=Q&+cw@x`kWTQz&{H}q5 zhELMesF$ii;!6`aB%?i6!k_n_%{G_SL*{fkihZQY z>nQdT*_DnKvuiU<1qaD3;ql?XEVvivjij}cY*dpqMsdB7Y@x}Z>hY%ULJ=_D={o>yYlt2E=)smOqdh6*6 zFF&;XaQWjckA8c{_~gZ}%i8=$Ew4(bKWTQXkqAiy9V~5R$@o-1Yln*tG_)B{Ul2iGI?b zm;DvF5~soIqYzy}Z5ym5%+HV-&TXKA{LM&rPwr07!$JylYOIT~sciE3KU;O=9yA1} z4A$|~GN0$3uf9+|ztsV!xM&+<#72FrllM|Q+5qQ5&<3j)?@$=O!n4KWd{9EL+~;Y* z;62Cyg)qn$GWboY&$clZ-2(~vpbj=<2VbIq8p5h+!)qs7!=LU)M8$mcV@OCt?oiPy z|LvuCuon!tXlB2rQ!$r^)C+Kzis49?bx0ED39I6-cG z&hY%K)qg4Y;cJUt{>9aTQ>PdG=r30bb_W)X%l;Re-4D0sJje{8P1-*qtw`YIx^+sY zzg=r1o14gnj(X8qBo5Az86vH$uQ&O4m$X6=%O-Oq!F&l5XFHKQ*wk({laC9@X1&X7 zCR@xzW0T9wxDp5FP!~MZhyor^@#q9KVW2YT{^T>TM#+==OtK2awwxYE&JautgsKP6 zc{zg<39wo|UP<0p3CkLJpUq{4HPc8orjm_UV!`Ydo2yvps+YJd9j@jM`!bl+-wsgk z!UV4%N?<7)>k6)x($gmC5xKM<)+>|X6%S0Mx~vZN^B_#LI;ES)+3jBmg}&dKa}s7g zzv#DBE_0E~Vs=<7$tI)ADsfpmU5Urni{FN+&OnX9+@ltcFceJo8BA*O>0Vi%Ms`9j zJ#Ckr4oG3qev>DwkQBg#m;Ib0XV9=D)DWX99YFr%#yS z4Af;sO795`Y#A29;t%U{g+7s8TO6?RDBE>egMQV)-mC#Xv{5yv@2+PzBlXaMMXtox z4^bD4;;Hb>ML`QJCtgZz7wy z&S~!0|I1!jS`%Ia6<1}1~=Xy-3YP07ssugvcsxbMa61L+Gh4tQVsRRMp= z;q$X!7_4bOR!*&W@c7FwsZZh_1X%2e!Tq6pFdQ)uX?v}?baQ~UO0B=j(vkW zYIoas=qE6jErmv}If85{GIzkZ#6v1;$lrQadZ=qQ?6hGDgE0|&_7tCh1U@bRhVUn0k3)*es2}II~f++`+*b`#U8JP1v>#_~B8>fhQB@%>g0@s*HXQ ztBn98Imqs~U}7S2I~^*Rq&$8+Pn($F_nCQojSBBo@cK88?>|K=5lRub>Yp+iY`qL$ z>LhrGglm~vq3^{G{g70xS+(A|ae!IfYNrYLgK!D)(t`OP;jc2|a~wdx$ziYCMm;CI!fc0>2$jPFz7ynJVs{ffZtQ*iE=&UH5YLAL&j&?t z;Q68xjyZ5B@nC__3m20?o&kv`mNBe^z+K02oSqBvPy|aZt__QqLY#pE2L~JsjzHK0 z2ZJLJ_`G0UIElYkd6q!CK|dF+a}w(%pqGZybCOHHIFz1~T>4d^^qjq<=3Yik?pN=;VS3#=_GZfg991kE3npamzmk1HY{Okbyj8(e1# zxWRR>AF94~b5*r_7TAMT7#_~htWSTSv}Jntb+bFG8{)Go zt8ssqRW&@8KD~8=W|minHgGGUwZUeuEQN}1v6#%}a?6TRvl$N-2prB_bc;B%Czu@# zgm7kO=)5u{w}oUR&b+C73S7P}$gI5ye0d!V>?po$jN;3Hn132jdlX-e;>-V-_;U9& zzN`e)%9RPQN)ezhuE%G_XaAagIRUV83*gHv_}j>$y#NXW_!|P0S>!i}GF#!>b1oFiOto_OryydVx9t;4LD5;`w2FKm@$YvpIY_@aKM$6}v4d{2y zQKvW$pO^IvU##zV@S)$m`r=oczHygp{PM*1M-P1ewUwE#pQ%0kjSJl!j~@7s-;Tdp z^w0-q#!bqeWm|TPOeCMIblT0x#zQt(TWAk_>&;2MgoqJ#!wg#}hJh~|Z5x)9BwT|9 z&)@mPHCXU>CBFse7ECd41hL?C`Fl&}x9Z@OhQr2U!K?2Qych{ppO_c7F%qiYzitme z)p#4qZVy?g+A{t%fU1SqJRCvDSRDogV;~uz>Yplv(?GB~RPLSeCLb8Y%WDk0tq5Y@ zhj^lnhfyb#pvGzz)c&L1HjP#r-~tBI@0b;aR*OPKPh86S8E#lG;G*Xy@#7c4*Ml-U z*4-?&JT$y${HIq7+D|Wf`OMXVLw{Nndq)&oj$+GEY&nW8N3rE7wv0?n6kARaM6qSO zr@&+3TFg3%E#pn%JpS!q%dZLaG+@ha(ml~^i$6vl3;?z~Ne*`cw*0K{3jN7o%XdIM zUx(dm6k85r%ouCMm@S7Xa}w(@){TQaD5m7+30uav9$?Fve2^aj9Y?X{e-do@I+*Nj z1Irw)%wPgA1cv=YH_cym!uw1dwgv^_aAo{%1BXalxx&23!~x2jyq3L;(9Ph>e}%#Q zS>Vg~yZ(*%%bg#GzYI1pcljd7D0ex^UH&KGE)Rsb%U%KEb;Kajme&!aN&!*Z52PH- z9~>Z9LL3^QE)EGW`38w6Jk;A}s%D59QH#%^A&2UuE(;jRLpo)4BRZf2r@YgUo+n?kmfbNsiZI=Sud#^*u>VZm@)zhmSAadXD3vsN0?RBtM31sUj zwsa4n8UhXw%M?czRL#f`khukucw2#K?)+QeEc0%SvwUD0Q&x~oN|!mDq#PdqoD-H| zf{Y|AV^t(!8S5em%UBmlSjM_Y!ZOxH5SFnjhyqU&mRVf{VVTuM5SFpd#)8z-gk@al zHH2kcLL^}sAzujf4kIk%W<(N}p^$3{%Q$za=<5=ey}2R6GRB)D_{#AqsBB#)eWXhm zU-VShUU7U@SKh)WUfH91>5ESk|7}nAufDkM)xYgY?fv4B-1iB{&T^K=Kt?&sM3l1} zrnbslXDzR_GMF^I zzN!X(;mGK|@C9arR->t|tK9^&ofKgkkeTySgo1lpjT_+X-c~v;WutPF7g*-Wr!qa2Net`I;)=iFy8YU*XKQ>bu)R-%(`&i zvpVK;&wI{*F*TDX3%q1A%>Q$o@|}J8e3O2!VxC4a_wKO=7$6 zRJxRRpj;jE*ijGvm8Ugv3b>AgiGjLlAK7rwoQ`hHYprjn*z#@+olCt zXxwYtet!kpw{)(sqD06x2#{M4CkN(r^8O}^t}+$VJRber(L zA^bkADPix~4*oG-LSDw|j;+GNn!DDF!o5kHg&Ek3Tg#q3EPoA}vh#b=!nZ)ryIy;4 zc#kmGc4m8YA9jyu7QIsl5)Y8m2!JwpFLc1PpX~Nz%f`=iLrI?f6v9^$z zHG`3o zN<;l-m+#aHoK+%cqE5J!>azpC9r7#HGntynWt%64*EmgE9JUB>HXy^X^0MM(zXnIQnh`eqHc+;N4&Ey~1t39D(t{ zy&o}rOx}G@h<9%TTR3p+Sp_BRHfpf8tK1^S6vo1l$!1GF4}lsxDIzWpU~*am$2 zXM|TK-~|DM+x6vwv5r@xm~zyY3ycGdDZ4=e-#$HxDF;2)d=4;W4uj?7Cjlo<;mPWc zi6?tQcrv`AWdnyNw}sAU;K_|2?J4|b@Z{A|JUQwo{8Mz99`)lA_2cr7>Br@Tkbl@4 zxd%6VrtZPb@%5jA*B09qwY?BKzB zrq>q4`b3%OwFR*b#%FqM0hfBz zgEfrs+7c}K`d(XR-6Dqk;vW~qHeW4x8(g>SyISz$(~IsrcNO!fqW-|6{=lRDz_Hiv zs6X(iKk%qO@R=Ocs6TMbLy!6cpTbik{ERzp$v3u%3Hq3uImY&|HvWJ313xLy`v60B zknT>vkO#9<+OQb z4b7$1%cmTqFR!d?u!4j8zsEnvz#zRrz`uRsqjx^ICdmqGd6O04Kdq^OUvZ)pW#wk@ zb>4u&+?+43D9coq-OGAAUr`Ic_FVC$haO(uU^SIKhF0J|xLHwGQBhrIDyyp96jT(O z%F5txa^S4NYO$J1KsMK4dE22FEa3N&k-^vy)YL^>z~9Zo13t$G{tt-opbkF?i{ocK zJgDJhR{jUL7#kc@K^=;jEmJBsf`11X9`HVh5BU}0N{D+y`LK_-@W9odllXUEI3Ah5 zKa_t{c;--Atm5+Hd|Vvsx$uQh|8eJ>j7(bq7l`Y$pX+}vAJ^{4{J4JjcWIG95h|Z+ P7tV*%adtv(jL`o9zXso7 literal 0 HcmV?d00001 From ff0fd350eb132ae4d634e81dafa00d9d274e7a9a Mon Sep 17 00:00:00 2001 From: Arvind Sudarsanam Date: Thu, 2 May 2024 08:20:24 -0700 Subject: [PATCH 6/9] Fix clang format issue Signed-off-by: Arvind Sudarsanam --- clang/lib/Driver/ToolChains/SYCL.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/clang/lib/Driver/ToolChains/SYCL.cpp b/clang/lib/Driver/ToolChains/SYCL.cpp index 5e040c20fdf67..4ed419fd92748 100644 --- a/clang/lib/Driver/ToolChains/SYCL.cpp +++ b/clang/lib/Driver/ToolChains/SYCL.cpp @@ -441,7 +441,7 @@ const char *SYCL::Linker::constructLLVMLinkCommand( StringRef LibPostfix = ".bc"; StringRef NewLibPostfix = ".new.o"; if (HostTC->getTriple().isWindowsMSVCEnvironment() && - C.getDriver().IsCLMode()) { + C.getDriver().IsCLMode()) { LibPostfix = ".obj"; NewLibPostfix = ".new.obj"; } From 07144b56587ca39841be87012ccef9acae05755d Mon Sep 17 00:00:00 2001 From: Arvind Sudarsanam Date: Thu, 2 May 2024 08:27:58 -0700 Subject: [PATCH 7/9] Fix some merge issues Signed-off-by: Arvind Sudarsanam --- clang/lib/Driver/ToolChains/SYCL.cpp | 12 ++++++++---- 1 file changed, 8 insertions(+), 4 deletions(-) diff --git a/clang/lib/Driver/ToolChains/SYCL.cpp b/clang/lib/Driver/ToolChains/SYCL.cpp index 4ed419fd92748..79a089a651490 100644 --- a/clang/lib/Driver/ToolChains/SYCL.cpp +++ b/clang/lib/Driver/ToolChains/SYCL.cpp @@ -438,16 +438,20 @@ const char *SYCL::Linker::constructLLVMLinkCommand( }; auto isSYCLDeviceLib = [&](const InputInfo &II) { const ToolChain *HostTC = C.getSingleOffloadToolChain(); + const bool IsNVPTX = this->getToolChain().getTriple().isNVPTX(); StringRef LibPostfix = ".bc"; + if (IsNVPTX) { + LibPostfix = ".o"; + if (HostTC->getTriple().isWindowsMSVCEnvironment() && + C.getDriver().IsCLMode()) + LibPostfix = ".obj"; + } StringRef NewLibPostfix = ".new.o"; if (HostTC->getTriple().isWindowsMSVCEnvironment() && - C.getDriver().IsCLMode()) { - LibPostfix = ".obj"; + C.getDriver().IsCLMode()) NewLibPostfix = ".new.obj"; - } std::string FileName = this->getToolChain().getInputFilename(II); StringRef InputFilename = llvm::sys::path::filename(FileName); - const bool IsNVPTX = this->getToolChain().getTriple().isNVPTX(); if (IsNVPTX || IsSYCLNativeCPU) { // Linking SYCL Device libs requires libclc as well as libdevice if ((InputFilename.find("libspirv") != InputFilename.npos || From 37009916eee45f53dc5e497f7c728d66f653b685 Mon Sep 17 00:00:00 2001 From: Sudarsanam Date: Mon, 6 May 2024 11:23:22 -0700 Subject: [PATCH 8/9] Add windows test and address review comments Signed-off-by: Sudarsanam --- clang/lib/Driver/ToolChains/Clang.cpp | 3 +-- clang/lib/Driver/ToolChains/SYCL.cpp | 3 +-- .../test/Driver/Inputs/libsycl-complex.new.obj | Bin 0 -> 60739 bytes clang/test/Driver/Inputs/libsycl-crt.new.obj | Bin 0 -> 37521 bytes clang/test/Driver/linker-wrapper-sycl-win.cpp | 12 ++++++++++++ .../clang-linker-wrapper/ClangLinkerWrapper.cpp | 11 ++++++++++- 6 files changed, 24 insertions(+), 5 deletions(-) create mode 100644 clang/test/Driver/Inputs/libsycl-complex.new.obj create mode 100644 clang/test/Driver/Inputs/libsycl-crt.new.obj create mode 100644 clang/test/Driver/linker-wrapper-sycl-win.cpp diff --git a/clang/lib/Driver/ToolChains/Clang.cpp b/clang/lib/Driver/ToolChains/Clang.cpp index dccc82161512d..3e6905e3da4fb 100644 --- a/clang/lib/Driver/ToolChains/Clang.cpp +++ b/clang/lib/Driver/ToolChains/Clang.cpp @@ -11049,8 +11049,7 @@ void LinkerWrapper::ConstructJob(Compilation &C, const JobAction &JA, SYCLInstallationDetector SYCLInstallation(D); SYCLInstallation.getSYCLDeviceLibPath(LibLocCandidates); SmallString<128> LibName("libsycl-crt"); - bool IsNewOffload = Args.hasFlag(options::OPT_offload_new_driver, - options::OPT_no_offload_new_driver, false); + bool IsNewOffload = D.getUseNewOffloadingDriver(); StringRef LibSuffix = TheTriple.isWindowsMSVCEnvironment() ? (IsNewOffload ? ".new.obj" : ".obj") : (IsNewOffload ? ".new.o" : ".o"); diff --git a/clang/lib/Driver/ToolChains/SYCL.cpp b/clang/lib/Driver/ToolChains/SYCL.cpp index 79a089a651490..7024e79c4aca5 100644 --- a/clang/lib/Driver/ToolChains/SYCL.cpp +++ b/clang/lib/Driver/ToolChains/SYCL.cpp @@ -293,8 +293,7 @@ SYCL::getDeviceLibraries(const Compilation &C, const llvm::Triple &TargetTriple, #endif bool IsWindowsMSVCEnv = C.getDefaultToolChain().getTriple().isWindowsMSVCEnvironment(); - bool IsNewOffload = Args.hasFlag(options::OPT_offload_new_driver, - options::OPT_no_offload_new_driver, false); + bool IsNewOffload = C.getDriver().getUseNewOffloadingDriver(); StringRef LibSuffix = ".bc"; if (TargetTriple.isNVPTX()) // For NVidia, we are unbundling objects. 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 0000000000000000000000000000000000000000..197cbbacec8e2e61e9e72e2e3f6e3700fbb66408 GIT binary patch literal 60739 zcmeHw4OmlGw)PGQNq`W70*xi$34#+9HJ~C0Xdr;tQj7S9tv?9?gbr#5h={#5AwaOw zia+gO$4-NlmO8eB)@hxNw*v$$0=1OuwA4CYP_WdomRiSad;RKN`y@LAVl!Tx=ew=V zd7fnDti0K0?RECqd*|%)u9aEAMhE0`ve8w91jxM2D9kD@w0tIl#UP)|l0s6&j7(*r za-tShaYjZ#!6=yBt`=}JXOC3(TtRl_A+F42My4i9Ralfo(82KG{NKr*MY;>n z%gv3iRE_&%YKeUR#{}O zMWOrF5 zuihP$P+ePk&}D*&)8XruR#g=u2$)ychR_ZQq4996plp`lSbXUs?=*PZo!^|A?Cv%> z@sw+P+GH;E5p8G*gxfb)H#bH;#XU0@f($k|$N(4fk#MuS>t*qp;Gt~THw3OOakoUJ*` z{CZA-6v{C~)EEl806B!H%i@WrbOBJp1)b=ME}*^Evs))ZXa)FbxKcf+OQ=J)aezLkgZ;An;$SH2Wl-DGTa^ zn$}{&diw&pn*s)FMg0se2M(=O!%>(e)dsbeQw%POG(_YW)J9Gr0_R+*6`_p4L9ytR zGTV$G?Juiy`ZU8a{MVxj02A>8dg?^1{To6Z_P@|72 zlSl89`!|}R59y=3Z59)9Jzm|eBK}+i;>{lYOhO;xP-4|{C_m^ z2fEl_%J^u?V3YqLd2AW4)$nu!5o@ByFg=G>jJ~%cI6I zsB+J;aCE3am(JDtL={;JQ&)^s?DFw`|z6->ZMD%$$}5z+p5d(Edz~zZ>!c&Sj~wx>en@LVjA@8 z8kqBP^d)M;W@tMnI9qH#_w;&T-zzwUMoy88vq9mNui)UO-o?I2VPye@mGE97)7EdR`}?=Yp8mo8qrdkzJRP{b>F*guCATx){_3k2ueL7Uc{*!b$?EOT z{pW_rKUdwkFrbV)8M@`Uz9F%%g8O?U1J`Q7t%#*%@ZR7pN-yulBOdPCDa2DGmrrP( zic;x4cw;D#eCp%lR)`}p0Dg+Nh$&2KiW47kL7uRSdt3oxOAj5^Cyy>F%e}n89pT}O z3ojWGHda678>I9tmD4}Mc1X2i-vDt zo^5OVb#XWt=8k^^FWRs-5|m&V9s;a}mF#E(!z7@x*3s3k^{6{4q<>^F>~QEom2W8% z3c}u?4SQjYs0~~J1yP`)4LkgPP4#jvABG|sR?&tvDD4c(liL|qYp^p+<8|83u%HSs zEE9^hdaOMm7z}$AR93@$pT{1n23sjm(}qo*xZBQS>NGpUt_0f|w*L`3!=8A-&SS09 z%10YEb#>5a!(KRPZ`jlJhD|N7Gb|(9&amCj*%_ARZ|AYrxrWh(Jykf_&akgUqdj)` z!hSoC4XNx5ySmWMuv)<$Xk*jC^VT+2DxSKkyjKA#Ya2^UNKGU~U zZ`;OvCI8kr75~P(F%)Q=KkM1|LlMfB6Jk`S3C`z#np;h`b612jMz$fWSOZdUd!29%qghJ=Uql#UkUg( zwES7T_`v3-fHSqCY^f^^)BATaW837hFo<;)R~4!`+qj&q$p*EWQ=(=T6&uu)VM3=j z|7tbAPwKT%>UC*|Ka|R!C1?q_AQttA1DchA+0q@+C+Q&+1|=8?%b3xfLq7EMfipM9 zpvmFDv;yWAY6YiA!YR7t3dL^G8rE?Q`9g!H#;~NvJw?g~2c{b)e13|5RpIqjFMEjc z8jyMoP+m7v`5OiNOHw}Ue%s6bvX|eV%D*L@bn&WE)T|5W>3is^H~=O`E#d$;gM;FU z+0@I?C*`px@DwAeGyNo-eW+{+*qLnDqBczE_=5kn)T^(T|E0|9OCub`#9$WGbD%Ek zYM`lL5929kjjk*XrX=7xn9f9X%36#tvkT~r3%t-I$_`dZ`F;5OdO;xrQn^^aS;#3! zHmoxmo`(5j&Cz{~m?~4uK2v0ud>%BXPVg8nS?n9H3IS|lmin|!-Rzocj{wc5_HlO*;#W|5qMoGf?BsTCidzA zx|O1CUEp?QV0Udmn=>(pBx={c(a(Ajymg3y*)kTKsCHAeiyUUOw8Q z%28_8737v=WowihvZkvy}s+Y^|DpHE*L2W09srKdr!DRj^22B#!lT?1_u zH?ycNQ|V>VDf1tuCfuaW*Q4L;OA7uMsc@Nl_2Y}(>x)$WjRvibvsFU(Ld6ESg77t% zu0CLjJkEfB0I~Z_F29BjO$2|r&Qqd(C4L{&1vV>1UA3Y%^M#{qVtC`9#Cj3VmrCGa5`rAH!G=Qys+Vgf06HsXtWh5W9apsph7Y8q z5W;`A$Ft%>@w5Qz`rR4m8>x9v8qaFwN3i&4mbZQ!=T(_p_=8)EEDqii!t!K3Er{nY z&`HrGc4EO^1=SVb6c;H#4?!6bf6bgvxoKyA$7KpvByt;%`TX^TqP_od{d&y>_l}RI zm;A|dOWo&vvu-2zgZqOz50rKue5sQY%&4BcAV^zhJb3)K2VVTm2RBxvZ{PTF$A%M& zqV{_wyHJnKZ$_t>hyk}D@PPSb*=#ikB!+z$hL8tWj1+SE-tK1`A?b)Vw%S`hCJ>s2GgyHj>i7(wwe zP=#luOsMr{nvO1#O4rs(2f3yNj71W7uU%E)$u$VIN@kFNtuX1(R5)I?&Gi^jm(qEO zxd$aLU9%>4Zq)RdyS{w0;r981zxiCf^1XohrVke}c&1(49Z$Tr=h^O!nkUu<{>PKb z@BWlH^W8qbIVqj0Cq8P4v0mI2|F9IcccOP=#*+etaC<`X*>HT_f5h-fn3Y%)>$*QH zZv3~mw85yme-cj-lSFLsW`wR7815&$z|Ol}+qnP2*LJ3o;f>-$LxnEAzJ5GLFaE!V zHzkFTW1nV?!KY6+eG0z)@pf812A{6@#rrhhp-=1VeQF+ftL4gh%|nmNc<#-{VCV`$ zYH?SPIU$1Dr@~!%Qd)LHimAAabOFHN- zm6Iy?&_~H$p27)Z+?<_YP%R%(@2DvjaAumRB?b<0*Sgbls7fTr9Ioa)^sNIn^}N^d z>XBc3Kj3)0a`hRWyWOi#lbk+vcs+)U!Kd*~pE~qu;uw7TgV8xxbu?~djlrkiJALY~ zotBTmr#GEGb?DQ&G5GYl)29x7+BycGUUT}?p-%_K;L|HkpE~rZtMC7QetOyIQ-?kc z8G}zRI(_QUr-@_m=>?}x9r`qD3_d;gi}z`iL!Xx0_>{g^6b$!@lDnzMs!~b}H~Mt+ zo&Q$2^DiWK{yjM^4?~&kIk%rtEN{2m-U`k~K>rMFVWz9$#^HY{^dR;Y~F9<3_tsF_?Gg{W$+#KWcT5hddz2X%uufu&1y{8rI zxt+10{@BO`2bkl3_C{ITjy(MWAb;0f9a;`_|Bsa#_?@7^x{?2!B z&l+-^Ec^^K?2KdX8^ZS@@i|@xi>=`xQ{aVh>F~lEmaz1ShAfXA?FL)`?7syP-&e-# zHh3f7Yu=zLRU?1zS#CjpDBzU6*Pb9(DS@1*>{5zoMz>7$N|98-hhZm6q-vv5cUm^+q6VT!M5=rob7Ceyz zf6am?lHhT}!sQc5@NC=`aak5g@InioNP-Wv;D{u6Dd@5Pi6rHS@1*>{1E7IeTgLao1n+_C6eH|(1+l) z7$ON?0D2$L6G`yF7CeyzA8Nr9N$@hzPlR%a#Ck6DGk9AAec<~bp%1)@YRgK~ZL zx^!PtO3i3JR;dc|HH9O;pf#$3+>JS-^jM=(Z7LY0r!}fVz)XUT(=IBm5SDFQsu2vYv8%L?9}D6;+BR# zswx?NV(wm_gmcJyN|-k@cXI(ejW_(ToqBUdUardW%-u|Az~+`b6Q0Nea#ZoUEgGd- zou!$fQmb+MC;Y7wc#wpRqTdmcSmclKx2i4Q2U~tuX4%iN$Q>3rn!mL^%Szh1D|sva z9(Y^xoxCk4Z_CNsa`LwBg|~Hln70)P59%8t;Q8uwc;MWyZRBUq_fz>Bs`)_IDgoXW zO{yxQSyVu_0#cQQs6{iX%vGvf0L z8&9vjdc0=!JFkB*@ae6;RlMDC>+3av6{Zf&r$YY3G)4tbsPNPYPpfDOReN^5( zCc_=G3+6`Eh=GuT=R7byGPAL+2iwuCL_R?5!hZxPcp(N5yAIE*v@7bhRLl{B|BqUX zVquQh3h}tn{I73HF*gU#z^z9d$Fbw)BzoDCh9=sRh92>;SRw+$U^1L*Nd*-PUNpze)!&|G6=6@aDzRrX z{g*v4jmoei4W%U5k%syz?buA$z3n{a`YSt+9bRnju{Z2JwtR=3$4=VwznVkr3_HEk z&amc(>{au|L}Lzn-?|eqOZ$N&ej>Z}5J6w>=4_~By@N#mtoSZEuXUoaia&oquoGmA33qLpK1v-hc!vt^-eJ2_iU z&eku5W9&MH=Zu}4Eyv?k$QXR;f~%W_G#7_eCp(EIreG!7<}sFY&rI6 z-57l8f~%W_UXVFeCp(EIrgcm&2#tn8dJc@*>ddDkTLkw$=P!3)5J0O z)XCX$?9;3<_|(bSa&orJ4?m4()!DS4VSbvEv-Q8>Y_)JC4Zztt#?7e)&X&=wJ{34y z?e2p^f6@x}@FY#Z*{b!-Q37YHgI~`8&elJ?25Wgb>J@KEJaD#7kIT6ToUL!i*Ovlk z>t)};(tjCe3-btn4^K2XIa_8f3gL(lt`?D4k2z|Dt3@Qi6RsAK1W&kHL=wCa{EnZs zB$D9kK!=}*B$D6>SBpr3CtNKe2_8S9i8(h!5n>cU z`zLJSJ{UGDZ)=fxpE8So%G;{3;3&)Q$}Rgj7Kz6*yp85b0C-3xy! z79PI8r@UJNbKb2oYu>GR2YI*Rao#QOAltlK@gwtYVeY9V@0J_RyX8jmZY4n8t@z2@ zK%93=>m8JYCVNzdr^bovk$ar)65ip;$|x6UO{n0oFhjQtQwU}3#;g&hLdSw)u?1zST!n{7FL;$XicU?t7xK&HJKKzVquoo za562d89SY`H@HJKJxnLTDrrbVl0!>q})XccYP za562dqYWEQriE3sVZ+I^u!=TpIGGk!(S{8t)50pV$E?Y;u*z(hHJKKzq7Abq)1pWw)?`{(Wj4&3Obe@M!-kV-(JI<7Ycefb zMH@DpObhGGr_!2C3#(|uhLdSw6>ZpXGA*p44I56Tg;nNLdDmoG@5s!_v~YftQM|0~ zagtJvq|D`$Ov#{LlOoA356_fbtYp_TzyXSm85ak z9`=;9-Q5$XT-&oJc1rucV=Jfhy?D^TFz3_y0*>L#r^R&U({kq1`WNQY3XZ73xm>q4 zaOP6w8b~7xIpH94STScyjbXx*!-=&T<0B!f)^SMh1WA>cQy?W6omCE!pGKcB(ed-Q zQV6YtYxJMx&dQ;FjdN$Uehb+D>-<^Sd7G4ln>K9BRc|WH+E_^8bY4qyaR?AfK5ZJc zEIuv{-8IKn6&3l6)TCy$@#I0&EG{QpZw#HqGVyd_rqtw49y2U#7SE@(Em=_-o`F+m zL;dVAN?lx~>Dci&@OXW2-FuZw8X-4Yo2sd{83`$hSWY7fmh`%PZEMx4X z6K52W+&d5m@HL#5D}A50_fpr;7dF@Zt^fGRnAhHIzFx3%?MBm@FN<#R7HCUzUgOT1 z`m1jli_oFdySR$2SEsZcIeglj^kb4sy8GWB`Sa=Aty@zzl@EO3_wVN`{@gzx^TAydasE}bP0wy#0WF5EkO>i2H2%@c)5V|s6fT|{vFGu}mlrqI5hCcHVSlRP-$@OPG$+j$N-=Wt;j(@4sU2 zcO=`G@H6IcR-lR{Q!F`yGM}_2$N`GR9i0%fT+U7V_QO&2e zsfu1mB*}UfvKnIEea^az)1keSZ=1-dria)INs4yr?v~}pH}ezM$IKbr_CkiBZ)%oX zstbiL`($*Q#nT@*uf?=pNbJ|2l3HS@+c@*v>XVCvbYO+C=$% zwTX&p&VB+X_4>KBi8mdN9d-6?f^Z~{zS1Xg$D_g?COs;o?kU3)+iWEhD_?x=n_1E` z{xCYkha(d^LCw>?qT3p{K+bH%qk@PY6%CV3Zt$u_NxKL3%_w za`-lH|7MQO47qoG*IPUBF=qS+eTWTOmVfYwQUSkE!oDtdpuOhbXI3xU= zbk2}suAWL$^N%bz-B-QJ% zO{#a;E~#G2-IMAKk5u)elj=PN1kIYyY)7g!QRb|FF!*$+GOc05VT^q3CK?`bGnO#2 zH+tLT`h0w2&WJ$bZ(4nWhqx`}%8`jihiV>O%YVDioVIQS+jcY)t|%}M-r&@;pWbye zI~UX_yoiS zHUAIu#`#z+Z8z?l21Xp+1bL^Bz*}D_oW$m2?p0`w%@k`p`(h-X%1Fn}p{Fvd{1%u@ z$8D!F&{AKk#Vxg@mu{(3pS{g{VYY2cZRnHD1KUoG$Oovjb^ERj!^v3T93|D^1c39U zyxlg9am0DqY@JZxHv6J=9)V|3J?Uf?h0lJ$L^=-qs%2pTEiB=v3D#Lb!1EhMBz27_ zPcU3!ngpM%v@K+*|A^^KtThi|hL=s_SphNXVFH5ZCHQ$F{Gae0UV?1HglDzq!i9`K zW<|Q-^;I!^PtN00_&gqm^Qc_~-`vYw#T7UO_4iSqku6KrRjcmVr}y8CJq6%(EP#DmvQ zwBA>nU_0By1K3Xt+*g~Z2z2%n4_-gv3LMM(xIf`+6AxfN5prK`q9VZAPds@2MB;t5 z3AVFMJb?W~)_t{!3X!v)c<}m(^80ENY-gMJS@si{yZPg5nRS*nF=8!q#138^tHukP z%22$v?$eH+%u2v(?=4y$nh@4j8JENh&!`dx?|!T_E{NkA#?`&jl#Ny(`r)nL zeeh9gT*1iax~4sgziONujTb}Xem`}ADSN=o@xFt*V%@IqPLci;G^jZdcur#@LR*$6niCJL#n=x4Vycun*KUK4#XUCPo<4;8u?{op&iDJg^=fyMD3 zZyT-u;>v`+ki5gZ>MoKzb8NOLfYigeqT?Dj&s+lx7F zzbW34K`DYJXJ$;!IXF4*;$(HkqJ0OtcDyXB(*LTY`B&RrLYw15PnCuq7JUe8Q zW+$j-C%-y7<%mb zjeKsAer}0r?sjI(1*7MADyBsh(tp?iaEoK?Uu)WqL00z;13#O&tHn|NBoOu z{zdux%li3OO!KcX7u=-0Zcqzus1^)0@~>$Zd~aIt12Z8$At9+OA->N`R-2I0pKx32 z^*u^VOGsQ%miR~wAN_$Jn8P;RjV!urBV`noTQ#ARm)jf8u<*^H_#TOG6UoKmmr}KNN zcJbBz#Y4f5-^hIY)}hC@*FN-g)8jv|mc=J7ODcb;Q2CIi&rhve_DJip#jeZK5|^(i zU*4`1wbd?PJ+QomA!=r<$Vgm~UB2R}))gCFSAHT6I4fRRn6eEm1}zH)?BL%yxO$pg6rBFjEUF9Yx}d- zey^N(OGoFK@5g!O)wt_;7e5w*yAIj((0Y$K>C6apuF>-X_*TZhtmT)^YtTnG>SMa} zvAz0QFV=!v%v7%H5=oMdDo>M^FO`)~^5rd0cU_U=x-xk!Lsy<&UB1d#UgziYo4mL8cG@r z+ZqkiCWii74`=%z=b5%27|b2jf*rf59d8ED-fM2dy(>yXo@fqP>oWQMni{u1>~(vo z-tD(1+zxQ{2N3h+9=G2Oy1gRQzan5Za6OI+J&sBA$0#N=3#}@>Zp;nHhw^Gacl$(QI(sI9atM7eF%lp=2Q2UWq z@H$_z`#Y`A?*utv59@^WBUKzw1KJ zi3Z1Az|Se+Eemo@FMu7$H9a%|_Te!Yv<`?6;UH7sg>mWd!W)**-J>D=5DaX)0p|}A z@)SB_H-R5;EC&)RhQw?llEjuFF`I}a_-Em_n0G=X!IPLxL=wEwf+v#T>p<@fFCqz^ z#B3sx;7QCTA_=|=%5cVRLg-Pr0I)xaB>q_qdVIc!B>q%d@I(^42K3mUL=t>4=yAIt zlHi{OJ)Tz)N$|TtKOXc%68s@c`9u=@VN3Z$68uR^`9u=D$x=R%1W#f%5lQeQW)qPF zPhvI^N$@0Q6OjZ@Vm1*;@FZpvkpxd-HW5kiBxVzl1TTX=>K_}VWa4G^iSh9Ram zb159J2_SV2({(k@i(qGq3xvKTP$SR_Prg81=e zvKorl%+d)a^Wu~#!U|K3R8StKquiPMWjuo}g{ll%GFf2SEg8={n7prizk;{$U}z_+ zHXq_TeFJfwF6ll9JX*^9tpn*Q|c$^$!Lw>xfqy(X~2)S>xQ$e)z93SVKQk(@H&!w$L9LKTCf}ZGQmp6XuMEkt) zkJ#sp|4pS`-gw0l5f}#3Xe}HS(T24L%>%;(uxmB!vsN7#76K}(VOMs#tt#(L#Xo`< z9k|EU-#$HignfGU68pUI|5a^g7?mLe!{Gmb7K%r-VJQiAdE@&k?efN7_qOwx>#yuQ zc6hP9$KJ5_*zz5AhMi2aGpsqp&al%v?F?&v$j-33aKlk|lnE8Iw!^8}lU2T@uYwAn z0duIM&+I++$K7^@JzZdDSkihs!#+D>XV{gO?F@5WV1FvN+n-8z@o8xM&?YFbm2Uh~ zg$MV*bp+>4hE*>3I?5~yhEoZ%5iM5Hr?RK=5olx3nGLrwUH1n^;xT~UAUG6tG|njL za6vLZF2eW7N5ykW3rS@Sz}>AQr&GH)12w6g?pbZAJ)Q&Kr)s@JrY||~leJ;V5n)LE zk{ zauP;p%a<0rE`=F|kh79&D3TZ^qz#8~s*aDUGDSlar!M)tlltgRQ`Lu_ml=WGad0Qx zQy~`hi$&*ZMT51Xew`28Z@&;1*sKfe?h81t48(oJRtlkHkpC>cQ=}jm$9MWakL`5V zxJ{i>L0p(g!ENQaEa{{~ZM-IGlFqHY^0^P+`mpY+=Z-x8W@l^F^s9LreHe>0CHiRv z{;Gog`bYN`9G$lJw~uZs@DEAdgZJMq*kAJdz5dhuL(&AIOLzFD`C`@#J|od-%g558n2IM*3RbGR2p?WZF<;T3Wwy?%_RQLuij z20^b6!!U$AxMHNhi#%^(t$97zaouoirF*Pf+t00TuE6QO$GSEC-0Bu`?{z!z)2o}~ zJVuE&b+g7GdM$`|hxMwF;@v5`D2$+Z8K}auQYO@TGfhVqNu_IRrGs440>&ZM4jackfpzryv12#@aYpypF+F&@o}Mi3_e})i}z{1 zL!Z{!`;^3xQO;`~dR)eHZ#D))R}fN*yMoLK5!5~v?h1@Tbf>Hg@6tVk0(u7hCFmI< z*NRv9WWX55Ew_x(qo&Mj>MKZpuy@0p8ZypSC~d;H97go@@LyHF7O*!Fzj=OeO5#6nqS`wX77>)yL&#B+D% zAT#XcI5f)5*$D>K@)7lpvc@DygC_vt5O=LRJ%=Lk(QM{Wj>oG3n|j{sc=gCHz8`Sx zQ&-^5k-6Kw`ZUSuQ-{}M$QXPY@ARodpC*pMr#~2-b5%#!KW9U zK6U8R#4-5vg43rCeVR1}pPu{0`!vd-Ps?q5I-GmHn~JO|rNnTfPe)VF>Iw`+1GQTc`U$=thtaaf# zj+pvEZ~2rhSF4#LseVSwYT)LmpVxAK=Trsr%b zD4YQwW)x-><8cgc&`lyHezR`oof(fpz z$z=3gG4H+iIl4WNp z3uj32lmX@=*rplTxp{Q4p!oqb>s?fva4yN>bz3w_wK_{PT{$-_J3MQKO0CAfGH;VW f%*e|t+Aw2NHqJ+xnY(elS+P!;nW-trU61|;L{nr# literal 0 HcmV?d00001 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 0000000000000000000000000000000000000000..511fa246ef1d5f0838f90dfeef2fb15de68e862a GIT binary patch literal 37521 zcmeHw4O|mfw)o8l$p9f45Hv`@380U*is3^9MV%19ShbCc7Hj)R_!72!M2LuX+m|m0 zwpimwUHZA8w$FC$E^V!Uck8Z%P?svUXl)l;w+k&=+qEtAYwgx;-#L>Eg6+25-M;_6 zw~zbF%(-*s+;i_e_s+d%a_&7-+!#b2*UczB(@e-@;tPqEl6s3jj^L+o>=qEB6Z()+ z01Po06N^n2lMgj~-sq>{5iHSkBo@}z^4AC{X5^JGZY69-V*O}_SO(TbdW2uUpRJF#SrPGx=$uO2Ch0r>7?KBplq+M94ii@FNr} zOf0S`DYDd+@Wk;D83%G36@qAzXdYcY4G_!(y48MKEHBn$`izWGNhNUgha34(G|@ZG zC+IhYAniFkx%G~&1aE1gVv-Xt$#NiI^_wcHw(qu=x+mdmnd)O{$=h=O2um;i^E zV8xVNObh2oE~bZb5Er{0&J(%Vf5170i#-hI@mx$G77%e3qzV${1c?G_1O`C+gFr4M zLCF$BY{0fiY_ySqZ>MxW67PcqzQRk!AD@R85LUo}_xQZR;W`Qq5m#yz7n=&_Nn9)w z&I}jBqC|7CJUC+zk-K~mh!qdV3hqo_YN{%}ZS>>)7RnUDlO z6oLdXpu{i;eSoz1VV*yo%YR+whkTDI_g|A*b9qf#YC(N^T0vUsw6#_DR8_C5;vq5B z6qON6O}V+k-;uL)9q-NnQDDa?lxgA@ul!Bv@4wi0dsKIPwjeoxsG(Uapl=2X)m)jd zU8T0TE!uAJ0tua*u~p!dO%9+UW5f>@TElaMSzCj%cGJn>>6CI|Q1a3yy}3&rmmTJ6 zrnOsz#og5N;fyY~iFs9!ywsU5cW;T!cDLDg3M0MJqY;5io0})o(X%Yw@aQu>8O4^0 z=@+u>3qqH|Y0H{UgFYy5O3vYctff<^{jb`M50UW3&_+fO+PR=psNedB(1=*N|2=Vs z-LBmhkgaY`GCf3OK_*#H1AEvR)*5O}ZZ|y8VNkau?U~x6EY70T&M;@Plb#anyty@l z*37r|go*f_awF9e=+Xzgkvn~2aTS)&v^a)S& zxi;0HfRff3rB<_RU72Wu+hw*&>-Exg>8`Xgm$_AHA?LI)AGN7SVaz$T>aZ#LZ4xs$ zfC&%VV&2upyv?dU@~C=%Pg}@Yb@U~m%p^_gyIpIH$YCAh+F*3mlsPPn)Ea)NKjcJP z%vm5d#k|X?&bGyLt1;d=5Ypck(`Ac3VN$(ojqdZryxkTu*dNlTj_$Ilj(B1Ke`Fx! zur}u0z=Thr3WAt^VB!fm3=siyRz2ykE&A>L=zd#FUz_TDo2pOX78ylqF}bONkqDTE z5OEUO-98ZU0E%N0NF%;eHKL}?3R%(qIY>PPp+1X+`O{%v*hT52SA)!fiZXg<}-_;h=<52;hw~a9;w3Chq zIz&eADkWw)j|h!HL`A|(@(Z;5V}rbsC_Z(|KVcPT7zLZ&E}Bi8vk&T`dL5ZgQD%p3 z)&XzEKJ9F$Xilq$Ra`JA&Jg+{s~FJ8|JW-TVicc3LXG^=ko-cgwAwAdtien#4oLSS4FJw&rExRNn{g~Nfk_(AV2M3!v!P} zwM(ow5i`*{GrD!@`*a!m1epgx2%K@?opANI>X=#_j-%P(2w zWKOqk&UO*oBbwv%rtj57bvQB)LPBqPuXlF0E~-g4r^S&0eD--VcRS{87g?pl1^##l zYNt34&Eb|061P3(pmx?tZzd4;c%z)&%#)%FC&YE>2L+iY9qH}5S)8WzaKiWio4c8vUBU-6?rbOP2PrY#{l+ipQ{y*4-ngWU5q|5#b zN`O2r?3Z>-B4Gw8J4w8sR17_f=t=x&vq!p4@2cP9A8q<4SOkhp3UB4W-VUEC@V|%ND3f-!0Yz~hUq-M1| z{*hqGpKkr#2PL0hi+J*PubMv3yk`I1E9Wm=eWmoBccz_s@QLx;j$XNR;M&BaAO7Lx zOquH^XMUcf`i*1aCJikopl8w`bKwO-jMAK#N6f@3AXdQ8ZqU(Y@A;CY;zSfGYGk_= zw3|Nd!<3SeqVfY+B3Ms|X7|nG-p>#4YitB`lx~r7>2n(`d$Y_)gE|RK1_iQb{@kTs zA+X{#G!dk~KpOLQ#ZLYitpc1CP>{`fqI>qvr0vT|D_pX05TP~Gs($n5`Xe`sQH@~` zCtzbb+|SOkv2Lge_LB$!~;*XwMGZs?dYUPTh^SLiXTtDHaTIf(OxrtO2Vo_`wOX4622%B z0-xuh&Cdf&p{%N|0aaRo2W71xiG`=DcXFHBCN%!4>j&+Zut2hHi;GU?1R7<8mFtu8a@PoP62Nb9Z! z0NZs|R{`Z(qjc4@xYoDWrotTdNx%Fu%x4-=VK-=wF(rVBPPjPFr#$q zGFrIv0dM+tU3wEt8{X-V={5Uon9ke1vfq}w%=xYwv%^v*tv9+XN|&Y875UmS#U}*R z1^FOM*q|=Jj1IM65~5FgVmh@kCvY^3Io+=U1@0q3OsR4SPHB%jX7?E4{S`5~*40bv zTb3-kd z34)rs)H7iT_}J)a(268&OQgswdpW} z_IP2+%;@$)9W#6M`_^!Z=^7AVzAN&b!}9Y+#mD`q{hZ-6m@5W(6)k^XBZte2{Swd# zKvTJ@iF@z7N!4kK?i-kJUL6f8X_q=0TJW5DQYrHU=vZ?Oh(H_2=*d3_tq+ym0Oi!X z)|p+AM?a8%s!?L;>Vx>FBGOr}>y zWSL^1sIPgVm)gQh7p>OFKLsVPU-GFzaS2(;&-6>qGKzs-sO&D*QC9})sVj2+#2!Z` z=y?ad={+qcjWHKYDya1l!K8jplRjaKx!)9XqAj}Hl)F#zCr}pK<%i7j3q$hDL~&@i zlPlA#4a@2SmSIh$cTldn1XtuzcwRtT%ItGQxwA5wQIWb{;A0_Y1)$)@ocDwn)qXV^ zUnFedpbEn)L8mR|#6a|E6TUE$^$a0*0IZ48GrA^OWMECQo{6(HZ`pD@yOmBT&P+~=N@|KKOjWCyx{{jO^6Dxk zDI-yns99ZZVG_zL7IW>a>C@q&Y;9p;QFY~XQ|;PfQ%$XDdPPND<+S8Pq|?kyN=+@C zkv1cJW|1i+BRw^#uq1t^CMmh7sH8N*RAegA%$PZ2M)6G*)sPRKsw=C?&DEBYDhq=} zT2@{O@1TtCmRp#-EUlKrQTjBywJeV624IG(nRdF#BlCtbDYA{$9!AwIv%kE0Bc`IJOSo_F^pKXh}C2Dg0Qd7~3_tY+kyQOwl)MP6m&%oF?4dfOa zRFeQ>j&wu5xdr+;cF~Mpul=m*a9aqx#lesWee4|fERRop-zfjU07Hag5C;my2L}12 zcB@UgK7kv->Rs^k;q#kse%zbB4>es-sfG8zV5mAhpgPUKYa5==aqmLCZK`hHLuc{` z6qlf*Lcg`jgPVMvKj8Ir=<0nipn2iZn+;DojKkfa1VcCu^Bw+;B<+FzHv4fH=-{2l zkqLvI+lw=ciShm4GES_`gx7kQ-#|wbcfCXB6A~x+ZOt~LdAElVtDGVM0iA)I&7wM9 z$Y0p2VzsHuoRLSF9Q$*$NWBZjF@)g%Or!Y^j^OWpoac|_9L@h! zP!fHjc&hDh;VqV>iO2iP-nxAx@3+EIr3U&~xM_c^{-s~i!uF|$WA$l&op+(H+*46y zz#}}!-eG765PK#X5@?|=VmvLb)a022Y`X+5h`5q%BkChxwTXjv6(4!p^5LYT1a-or^Daq1kabJM)NqQ*Y)Dv);2(fm-7O~junKGebkHBsX zPCvt_Kt&)79FkPxHD-!*(^f1e#|#&*VI3fY>9x_H@Yw z2yG-yMWNpHbPC8+2ZDPz)GwXV6wS%PEEMILaM?>IUw)1b)5^~3G$-woxM$^*0t8pZ zIyy2B>p-%rdP%CtCmR5M*(_nPuqYRl7s#dG0dhI%vBwgL8ZbUNTm?8zynvGf$}c>~ z-Fi)`-%Jb%!V_mEvHPzjY-dB<(E=s~#mv=x9KhK5-R>np*)Whw#LT z>S}Xh<=R?{xu$w`O$qn^{5Y{wwt{-(|EW9DWN&p;`djGZW-q-(8#eywDB*&k6^7j=QDLaCKRmEB&~6ghw+M}09Y&UxJvz7d$(8oP z)}2F5A1U`x=d61kvITh+O@qoR`UH`cM$&PN_Q8S1R*_?+WRN--(NuN7bgr8T42A~$ zxB6QW-fb5=PYKKqNM97c(k}4S3p9mn>dZh%=uOQ^`q*N-HajRzl3lyaNN?$E5*$Ar zugu>4O(YK$Z8MVRI;E0IjU-nq%hD7YH4G)$bW6Kd+3XX+`$6!ZSd<(ZuomQ0w;aaH zb2B^gr7zI;_sd!~x7=Q|VA-I^OMkv^bIXbA@ss-(M~l=~n<}Fs)Z}PkdbVwjy^@5X&@jVC9H=WI#MIpi#MDrCCBH|Exe%ZFZn>WQzu3SY8`Eb6Kbl+u6c6l5Pt z&AF-kDdE|16O!*HTz2+1WOEC%UGPMF5_d;0*7rB?ENs5?R+m3({S&KB#4kQRT^%{# ztNzKCo)2FE1t^Jtq4;1%#&BgXS551(r#hD-{;D86J0+)O!5shHeP1-){n!#mY9~v5 znUeGJ=03MtbNeT8J9JR~C#)}!WBv18gK_)!Q9bX^d~Go9@_mvXVX9+N^w!NM4?Gt> ze$AXMbJW6%gL9W&JnYr;4bgXdtY-5c(zk8S-=a=G)*m;%fx0c_YSZG!mP1=?eP>A$ zZQtyBR?%_+<5>58*~0e9$fvFEB%_V72EKzn zzRvIFi?YIegJo(4^#BYe=9^?74#D}|`8UYm>t*HUaoK1Y#B+SVj|_&(O0m>iSWpNi z#%JmJ2ASNDPqr_gzeOe^@>v7g4y$?0)RgQcpmcffFgH8pKk>Sef|=Vl7c&v(+djF zAX-;2zoNR(RI#9{uDZx%fp^{o#Rc_s-$Ahcy9hQQ!Aw4jMZo$Sa`=|}Tn68AANc!= zpS8BU!m^<18*1@Q1k2!cpSMNdz~O6!tSqT4GOsTvC@!fhFDj`hFD&q3;9Of`xd9It zyjy9$0S}lOuzmZq%kt6^OL=8UVy(qgV@Wh?Mq^8=iUC)Psf!Rq9?cXT}DH{2Q&Fz;^wJ3nWHuj8!lV4Mddyx*|TS>S!vehd7|@1t}} zk&=K``v8<1uLcLtt5+deEJKvRFOX(p`?no15O8I|ZH+&Iytqp|;R1w@ z7Y*zca)t%&C59vGkgNx8pzn|1R#vNp1qficuwqJ|>7E&3kcU4VaLK@OSW7Z&kZ1h0 z2y4(Z0Yj{+0N#mx#eqLY9Ik_U zKW5_}v+?_E#o<9?HvX@#CE#T8-)Q6iv$Tu>8~-!Z9s)N0-GS$fVB_x(R)&I&-z7^A z1snehvV2buC@&EHQ5!$*$DIiZ;+T!!XX{5>GTOFz+ddCD8#CCpdD}h@c|331=ONFh zRWN4br^MH7`-9Nm(ALk%4e-2epNEKtKhK73``=;X{|rdSZ2Y+2e9XrG&ASi$HvVt0 zjsL?2e%`)MZh}q-f#J^Ie?Nb5f-gaXKOJx1f53lzqkX^ByxPPY_IddA-Tbib=X==s zAAx`G{~30EHPr3f_VXY3d-wAj7lkTs7NpS3 z#Lf3kZu5OTxB0#oH{Zu*;pTfixA}enY`)jsWk4CtBli&c|Dub6n+ZtiN#DMC-p#{h z*da+PF8jSET(c+k!y0+VY`5cZt#e?ZGC4T?|N<**y~ zPyM6fQ9lNUAWZams~r4ToNEm^ua=)@74LP(-|mPR1f)?u2uPD0w#=Pn5qXaDc)?LB zkdl=Taj7n_@=L7zBVWQ+WHJb6xYZ^HnoF&ChUD)NIkIKs{VZm88u!j2F)+shupbK4 zS;^hta3;>0OKn>=hxzE=Km7B-sXsl`d2v?8)StQvhJMxZM9iVSihEz4eDCvb6xN)2 zGUg`*MHzp&_b0A_p%+fiX!Nc#ydNs*MEfuXmW_pN)UEpXaopy6LVqXMLuY=rId%eOqVoRaf6p*0(<)xS`_ z3LD&!83q9x+u{DvUAv1T>^dKYSP2tq)Gwdj_(Bl30vXX@nx{0o-=45wWkWx2aAt^J ztX%q%?0GPzI3=o+&&$%I1z_0%QMPp(X-e4zX-Lve!nN=8^2<#h(E;)afz! z$XKZjEa`F|3M}big2uc(I<%yVICgM@tHlHNC6-6nSGtcwJaH*O|wE^A2b6JTnG z5sdgtIl^?}lgy^xJD$tpFnGec=qu=isT&D#2M9UIeur&7p_$M(*EQPf5-gjRN+&0r z+?I9trdJ|=1>5u=82U+TuPyL_l`XAh8v`FW(b9Tub0D`5e?HWIY#+X8Y#;vEKKvhW zAO6&|7T1rJ;5njUZGv>Y*;Qk771X=x^wJGi8^EtbZHp^%Rn%VCMD2h*(zCB+rNeIO zYxeY$4$f7CcY3=Qc2@Jd^!rX5qrnr^X>f%#FacaYyf!caoL7PO2XF^)1Tf$r0CsWX zw(Z<2l0O>c7y9MkwCNBnzhIUhLN^NFIYM#KEe9tHhhUCl6e&~F!QX^ix(yO#!ip-Dx>jjwI?Va8M z+s?hRZuwi_t-up~#1?%&_%9ex-S3I+1hJSJEz*lb@gbwQ(l6O<09RP(tm~3n{%JqB zt|;<4e2e++HH-2gcXX71JLK0h4K+&FI&k~}?jh=>>sq9o&w~OOkH5?3K@)Qy`aC## zF<$gf_jzEdTF?0XEPT7ygJI_gu$KYq7)LdM^Mi^wikfM+THp$NAXq&zU06bj?7`L^ zCjN*__y+F-L7R~x%RbYi$={LUHRru16Ey-t4~Viy|C&<-&-!n6ieSCLDMHgIrwHzU zd7lUwAbZ%$H^L_Z_$)BvzGunKEuD*;wzA!taVBQsO+DG-rre+f%9u2Ua+1c7rUAxk z5@l&9E2V4FIGfm9nf&k+PyZf=mP)6jsB?g>M@rww?#dH9$nCv-S_f~nEzO~!%L1m3 zALsW`khV?cOvR15xeoo7e%LS@*#o{7evQ5sRurYq)yP7*ef6M9;(pZr{jwZmWn^>5 z+sP`y{t{Mr&r!GdtsUAx4cy25iuk`M>?hxY`^kGfu+RQ1ARw3+LcU(tc-Xh88eLZ# z^LE15>#Dq7SJuJTm3idvt?QL6sPxKB$FD-YCGg5pL~#{`atfdhH`JAVXGC4!{+@MB z2VSGwHQQfTYUifEqg~mJ*XxQ}^~iReQpC0KE8=6mJ57_xF@g0X_y|d)Y2`<3Olrmq4p2`^1>6 z4|hiDsKeG=ZA+4~^;c>-Ily$7w*Rq_)kol+v^()y_(>!C8f7&i5s_6<#y;x-p-`Em znMjKrl8Ek(*p+mtOx-Hl$x){qP}9HSVyOqzZj-1pm;%25W#AWJV!qMw&_T8_N;rdd ztZZbd1ex%c#t_R6V`~ysyu%2-0-9JF9?B#NlC+vD!eo!%Stg2gu-BB%B~$`cd_bKL zplK;#9`N5*-R&0Linqmi6r^JmlyOI^zhs9eTw+6VWkJ84;_3LXVHT4{b#2z>H(2d# zV+b{qqO7MGckl%gC>Q>c?T0Zy4*t9{s6pY}^qmeik*@EvwmYn(8*cY2+(SSshC$;J z!;ljW7$rEiYoyCs=OWww1`rBG37L7oV3|93ksyUNdu}?{WIJwUl&P0B!MnisKq~ki zSUjgk2wn`7A>e*tnc_tQI5Nl;*^Rv?y;EC+(gBu`TxkC*=mR7de1MG4g~3Mox{+}_ z94fqcZoJZfK0!)XpqGH$JJ2)8m2r1XnF-E8*e&fetDEyBtKT+#+H1x2oUafgI=DE6 zUP7)!DD?y5jv2L!vlm6QU-h{Tc<}jjLHo6^t)W-$(9h5Nv@QqSUpzL>V|egBQ4=J* zJ=L)wv_o_H>tOm!FE(sgd1QpW4*3TUQdRT)(l=ru{(n8}oN5El5 zOgyln-4{%Z&%VI#JBLS@IHAF>KgZ4+$8zx30i82}vWj(?k@lVaoTC_g`(Kfg*HM@< zqB)atv1yfylcJmQ-U)`Y-`7j>s}^P}xQBpA1?7iw^vO?W#QfEO!+8$o8#ms=U}6Yq z1~^Wf&5cPoAB`T9MwWT^w~aksiTHV0hyOvSZk6qpb~dZUQ6JkhMUis=U{hFg(r(?X5 zuJOZ^do+@Xl{VQ-v55f}CN|nVp*x_k^+mDBre$6fd)XXTY}Zt50KOMR7T7A zLmzhH6Tv+=rO?$I@xc*AL04}J5mpXAIGjHWu75w8v%u|5Xwu>opcZ$yrNKdJBFW4* zLPw|9q3<8gcNJvA^_%qo`TmXl3jEzzdq6hV`*_!leOxonr>~SXjSor*ow&0hE2wdq zCbNQ7$po?}`fwI>fb6#6USxLD#me_CYs{ycJf1#q{!+<^kNSI9O?&uuxX0GN>%Ujc zh&TOC_t|~_u8=CNf@x*oL1yz}yGJaIlXQW57>&o%B=`)NEHZ-!nas36$vD+Fkk~}0 z;+6V!Zf1T-co!q~58CH#0h@e7`HRvO`vzF<}k0z7&ccT2Otw zL2sFFtc-Mn-V$%yUy={H=n&|lwxCp?!?6LMeHa_s{l1OVJ}36RwnqOH}t7_RNn=R)k??9*saHieonaWCD`0iEXE#(^(>6C%1J zBdRe;DO&9{EPoB02m#+wdcIHZsc4!K1jEJv-O(;cN)^LnPRRsrx;tPI^r&o>yeJlg z;?Q8kp`nwe3#our#n!?D>ZVR3sv-w2F|>1~BsMTt)+`ZwT2F}m1oHMEAnknnG? zFK9n_n0k~YdrrlQ!o4AaQeg8+D7`O7lF%4smxYu@HpjJ=KPXG?34^hwo6VV`tTWT4 zO|NIUrCx~>2B~&|8eGV%WG2)@;VJ|IzgDm3`5Xt$h_JtrcVHS2oPIhFvlx%VQj0ux@fPTNmVjfx;gD)8Ov`HmH?C!ZQH#j5 zl^7aeg!7F|rSpSMP~Z*e_3&t0q2HI$YB&1iS(Y7oHbtLj)HqUeV1n+crTluunfM*4 zIjBl64}q@E^~Y$azRS&_K6%P||KFq}Z&^HGH4;zmf z1K!#p2;MtQH@?}kLmLd3_JE@HlB@QlzT{@DMuG~#-f51BKyL~WCif+`SxRysHuybA24&Rf6R>gM$mN`8b20_t@c*Rg4dJ(6GzPA54tTx5wyr67)rb?_Z!~ zXGt1wlB4PzFYmO0zaK~^8S(CTr4ZjSj&d&;U;^UmjbkK^&j7e!Vtn6O;KxvWfWmJb zIewx~n(O&3;PRt}^BX@RKMD6v^WFKOtAbJaJuouA;vXzO98d6h4G#AyOc3RQiTNhE zI!EREYRG@LTz8Mk??W!Xe@w3E5AUzbwd)7VZ&-(;U}okd3Txm%rf9?Q`5x+_%FHcR zsxxp;$4{-YKLwwok60DI2%*|OvW9)2NEcAH@W<)=+AC!NFhzf)@@3BM@=~K7T|!#8+l_yzg5OxKDE-+((Fb^v`p*pPq+2 zeUYDzhdljSKRpi-k82O__0#i^r{4s4G33WXo_-79ac<-x;?WNe28KIdhsggwfgTNH zJml$L20U8Wc*yhb06dOwJmm4+fCqMjhlnQ(7(@Y&{)>2sc$5#^;@X9WK*sqI()#f{ zgy!b(e0%W_y9Jh?Z#y1}{rFt~88~=|bwPTsAJ0SVhlod8B#v7=#Bx%Ae;;saIC#k6 zArtI>T>tb#9v=e*-;yfGEq#{3$MeYde-(dNzjGvXm=rj=FKS`y*Q zN7!F+jK#pfQ49x;2WB{&aH!xwJ3sd43OM*6@~;WNm0(EkTN4Hyp9><`bI-?cH1gfW zb&eQB99LhA0pPnEj|d2&;dKWh0t1@=?)az}^@PYo3wX%lsC<8^iv(CMTyu8tS?}-R zk%W*~Y_gcRnD6ae$s!8mdd33wUq@BiXI5Cb&k>m>~2Xf$xfQocHId=g@O zfnT_f>-#t?=1ZvQYT<~_f=w$-n^`n7b>!E~nK!SX zAP11}`1*&Zg4Nvj?+R*4R+raWYD_ikJgufaD@T)+o;_m*TU=s+53aEVE0S1D#Ih^3 z`T2K}c=&{yrDS>4y7H>x#U^w7ym^{$<%%EH`wqrSa`$;yW#o4u3FE4EbE#3!@!H1mEq*|Dj5Et$uE z-;ZTeR?LeZxq6U|kI!1EK{*#JWLY5n);9aDo{+NQTa$mi$&->(Fw=PjI<`Qwn7!>* mefIpt3$xjRJJb0x&x?n`L9f--*XP!lDogbE?ZA?n6!Kp$*H!%h literal 0 HcmV?d00001 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..77331c55e4bc7 --- /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-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.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/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp b/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp index 79d173e3b002f..37ecc006a797b 100644 --- a/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp +++ b/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp @@ -493,6 +493,9 @@ static Error getSYCLDeviceLibs(SmallVector &DeviceLibFiles, } // 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( @@ -922,6 +925,7 @@ static Expected linkDevice(ArrayRef InputFiles, 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()) { @@ -929,9 +933,14 @@ static Expected linkDevice(ArrayRef InputFiles, 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 @@ -2017,7 +2026,7 @@ getDeviceInput(const ArgList &Args) { continue; ErrorOr> BufferOrErr = - MemoryBuffer::getFileOrSTDIN(*Filename); + MemoryBuffer::getFile(*Filename); if (std::error_code EC = BufferOrErr.getError()) return createFileError(*Filename, EC); From 66de5b9a47d56def8639ab79a5040959d731587e Mon Sep 17 00:00:00 2001 From: Sudarsanam Date: Mon, 6 May 2024 12:04:19 -0700 Subject: [PATCH 9/9] Fix incorrect host triple Signed-off-by: Sudarsanam --- clang/test/Driver/linker-wrapper-sycl-win.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/clang/test/Driver/linker-wrapper-sycl-win.cpp b/clang/test/Driver/linker-wrapper-sycl-win.cpp index 77331c55e4bc7..1854dee476641 100644 --- a/clang/test/Driver/linker-wrapper-sycl-win.cpp +++ b/clang/test/Driver/linker-wrapper-sycl-win.cpp @@ -1,7 +1,7 @@ // 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-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 +// 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