diff --git a/llvm/include/llvm/SYCLPostLink/Utils.h b/llvm/include/llvm/SYCLPostLink/Utils.h new file mode 100644 index 0000000000000..42eddf9b3af2d --- /dev/null +++ b/llvm/include/llvm/SYCLPostLink/Utils.h @@ -0,0 +1,78 @@ +//===------------ Utils.h - Utility functions for SYCL Offloading ---------===// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// +// Low-level utility functions for SYCL post-link processing. +//===----------------------------------------------------------------------===// + +#ifndef LLVM_SYCL_POST_LINK_UTILS_H +#define LLVM_SYCL_POST_LINK_UTILS_H + +#include "llvm/ADT/StringRef.h" +#include "llvm/IR/Module.h" +#include "llvm/SYCLPostLink/ComputeModuleRuntimeInfo.h" +#include "llvm/SYCLPostLink/ModuleSplitter.h" +#include "llvm/Support/Error.h" + +#include + +namespace llvm { +namespace sycl_post_link { + +/// \brief Saves an LLVM module to a file in either bitcode or LLVM assembly +/// format. +/// +/// \param M The LLVM module to be saved. +/// \param Filename The path where the module should be saved. If the file +/// exists, it will be overwritten. +/// \param OutputAssembly If true, saves the module as human-readable LLVM IR +/// assembly (.ll format). If false, saves as bitcode +/// (.bc format). +/// +/// \return Error::success() on successful write, or a StringError containing +/// details about the failure (typically file I/O errors). +llvm::Error saveModuleIR(Module &M, StringRef Filename, bool OutputAssembly); + +/// Checks if the given target and module are compatible. +/// A target and module are compatible if all the optional kernel features +/// the module uses are supported by that target (i.e. that module can be +/// compiled for that target and then be executed on that target). This +/// information comes from the device config file (DeviceConfigFile.td). +/// For example, the intel_gpu_tgllp target does not support fp64 - therefore, +/// a module using fp64 would *not* be compatible with intel_gpu_tgllp. +bool isTargetCompatibleWithModule(const std::string &Target, + module_split::ModuleDesc &IrMD); + +/// \brief Saves module properties to a file with optional target specification. +/// +/// \param MD Module descriptor containing the module and entry points +/// \param GlobProps Global binary image properties to include +/// \param Filename Output file path for the properties +/// \param Target Optional target name to add as compile_target requirement +/// \param AllowDeviceImageDependencies If true, preserves inter-module +/// dependencies +/// \param SplitMode The module splitting mode used +/// +/// \return Error::success() on success, or error details on failure +llvm::Error saveModuleProperties(const module_split::ModuleDesc &MD, + const sycl::GlobalBinImageProps &GlobProps, + StringRef Filename, StringRef Target, + bool AllowDeviceImageDependencies, + module_split::IRSplitMode SplitMode); + +/// \brief Saves the symbol table (entry point names) for a module to a file. +/// +/// \param MD Module descriptor containing the module and entry points +/// \param Filename Output file path for the symbol table +/// +/// \return Error::success() on success, or error details on failure +llvm::Error saveModuleSymbolTable(const module_split::ModuleDesc &MD, + StringRef Filename); + +} // namespace sycl_post_link +} // namespace llvm + +#endif // LLVM_SYCL_POST_LINK_UTILS_H diff --git a/llvm/lib/SYCLPostLink/CMakeLists.txt b/llvm/lib/SYCLPostLink/CMakeLists.txt index 0ab3edfd535bc..6c4040798c96e 100644 --- a/llvm/lib/SYCLPostLink/CMakeLists.txt +++ b/llvm/lib/SYCLPostLink/CMakeLists.txt @@ -4,6 +4,7 @@ add_llvm_component_library(LLVMSYCLPostLink ModuleSplitter.cpp SpecializationConstants.cpp SYCLPostLink.cpp + Utils.cpp ADDITIONAL_HEADER_DIRS ${LLVM_MAIN_INCLUDE_DIR}/llvm/SYCLPostLink diff --git a/llvm/lib/SYCLPostLink/Utils.cpp b/llvm/lib/SYCLPostLink/Utils.cpp new file mode 100644 index 0000000000000..12ee791202611 --- /dev/null +++ b/llvm/lib/SYCLPostLink/Utils.cpp @@ -0,0 +1,141 @@ +//===------------ Utils.cpp - Utility functions for SYCL Offloading -------===// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// +// See comments in the header. +//===----------------------------------------------------------------------===// + +#include "llvm/SYCLPostLink/Utils.h" +#include "llvm/Bitcode/BitcodeWriterPass.h" +#include "llvm/IRPrinter/IRPrintingPasses.h" +#include "llvm/Passes/PassBuilder.h" +#include "llvm/SYCLLowerIR/DeviceConfigFile.hpp" +#include "llvm/SYCLPostLink/ComputeModuleRuntimeInfo.h" +#include "llvm/Support/FileSystem.h" +#include "llvm/Support/PropertySetIO.h" +#include "llvm/Support/raw_ostream.h" + +using namespace llvm; +using PropSetRegTy = llvm::util::PropertySetRegistry; + +namespace { + +PropSetRegTy +computeModulePropertiesHelper(const module_split::ModuleDesc &MD, + const sycl::GlobalBinImageProps &GlobProps, + bool AllowDeviceImageDependencies, + module_split::IRSplitMode SplitMode) { + PropSetRegTy PropSet; + // For bf16 devicelib module, no kernel included and no specialization + // constant used, skip regular Prop emit. However, we have fallback and + // native version of bf16 devicelib and we need new property values to + // indicate all exported function. + if (!MD.isSYCLDeviceLib()) + PropSet = sycl::computeModuleProperties( + MD.getModule(), MD.entries(), GlobProps, AllowDeviceImageDependencies); + else + PropSet = sycl::computeDeviceLibProperties(MD.getModule(), MD.Name); + + // When the split mode is none, the required work group size will be added + // to the whole module, which will make the runtime unable to + // launch the other kernels in the module that have different + // required work group sizes or no required work group sizes. So we need to + // remove the required work group size metadata in this case. + if (SplitMode == module_split::SPLIT_NONE) + PropSet.remove(PropSetRegTy::SYCL_DEVICE_REQUIREMENTS, + PropSetRegTy::PROPERTY_REQD_WORK_GROUP_SIZE); + return PropSet; +} + +} // end anonymous namespace + +Error llvm::sycl_post_link::saveModuleIR(Module &M, const StringRef Filename, + bool OutputAssembly) { + std::error_code EC; + raw_fd_ostream Out{Filename, EC, sys::fs::OF_None}; + if (EC) + return createStringError(EC, "error opening the file '" + Filename + "'"); + + ModulePassManager MPM; + ModuleAnalysisManager MAM; + PassBuilder PB; + PB.registerModuleAnalyses(MAM); + if (OutputAssembly) + MPM.addPass(PrintModulePass(Out)); + else + MPM.addPass(BitcodeWriterPass(Out)); + MPM.run(M, MAM); + return Error::success(); +} + +bool llvm::sycl_post_link::isTargetCompatibleWithModule( + const std::string &Target, module_split::ModuleDesc &IrMD) { + // When the user does not specify a target, + // (e.g. -o out.table compared to -o intel_gpu_pvc,out-pvc.table) + // Target will be empty and we will not want to perform any filtering, so + // we return true here. + if (Target.empty()) + return true; + + // TODO: If a target not found in the device config file is passed, + // to SYCLPostLink, then we should probably throw an error. However, + // since not all the information for all the targets is filled out + // right now, we return true, having the affect that unrecognized + // targets have no filtering applied to them. + if (!is_contained(DeviceConfigFile::TargetTable, Target)) + return true; + + const DeviceConfigFile::TargetInfo &TargetInfo = + DeviceConfigFile::TargetTable[Target]; + const SYCLDeviceRequirements &ModuleReqs = + IrMD.getOrComputeDeviceRequirements(); + + // Check to see if all the requirements of the input module + // are compatbile with the target. + for (const auto &Aspect : ModuleReqs.Aspects) { + if (!is_contained(TargetInfo.aspects, Aspect.Name)) + return false; + } + + // Check if module sub group size is compatible with the target. + // For ESIMD, the reqd_sub_group_size will be 1; this is not + // a supported by any backend (e.g. no backend can support a kernel + // with sycl::reqd_sub_group_size(1)), but for ESIMD, this is + // a special case. + if (!IrMD.isESIMD() && ModuleReqs.SubGroupSize.has_value() && + !is_contained(TargetInfo.subGroupSizes, *ModuleReqs.SubGroupSize)) + return false; + + return true; +} + +Error llvm::sycl_post_link::saveModuleProperties( + const module_split::ModuleDesc &MD, + const sycl::GlobalBinImageProps &GlobProps, StringRef Filename, + StringRef Target, bool AllowDeviceImageDependencies, + module_split::IRSplitMode SplitMode) { + PropSetRegTy PropSet = computeModulePropertiesHelper( + MD, GlobProps, AllowDeviceImageDependencies, SplitMode); + + if (!Target.empty()) + PropSet.add(PropSetRegTy::SYCL_DEVICE_REQUIREMENTS, "compile_target", + Target); + + return writeToOutput(Filename, [&](raw_ostream &OS) -> Error { + PropSet.write(OS); + return Error::success(); + }); +} + +Error llvm::sycl_post_link::saveModuleSymbolTable( + const module_split::ModuleDesc &MD, StringRef Filename) { + std::string SymT = + sycl::computeModuleSymbolTable(MD.getModule(), MD.entries()); + return writeToOutput(Filename, [&](raw_ostream &OS) -> Error { + OS << SymT; + return Error::success(); + }); +} diff --git a/llvm/tools/sycl-post-link/sycl-post-link.cpp b/llvm/tools/sycl-post-link/sycl-post-link.cpp index 48b83bd1f86d4..b4b34dbbdbb1d 100644 --- a/llvm/tools/sycl-post-link/sycl-post-link.cpp +++ b/llvm/tools/sycl-post-link/sycl-post-link.cpp @@ -14,13 +14,9 @@ //===----------------------------------------------------------------------===// #include "llvm/ADT/StringRef.h" -#include "llvm/Bitcode/BitcodeWriterPass.h" #include "llvm/IR/LLVMContext.h" #include "llvm/IR/Module.h" -#include "llvm/IRPrinter/IRPrintingPasses.h" #include "llvm/IRReader/IRReader.h" -#include "llvm/Passes/PassBuilder.h" -#include "llvm/SYCLLowerIR/DeviceConfigFile.hpp" #include "llvm/SYCLLowerIR/ESIMD/ESIMDUtils.h" #include "llvm/SYCLLowerIR/SYCLDeviceLibBF16.h" #include "llvm/SYCLLowerIR/SYCLUtils.h" @@ -29,8 +25,8 @@ #include "llvm/SYCLPostLink/ESIMDPostSplitProcessing.h" #include "llvm/SYCLPostLink/ModuleSplitter.h" #include "llvm/SYCLPostLink/SpecializationConstants.h" +#include "llvm/SYCLPostLink/Utils.h" #include "llvm/Support/CommandLine.h" -#include "llvm/Support/Error.h" #include "llvm/Support/FileSystem.h" #include "llvm/Support/InitLLVM.h" #include "llvm/Support/Path.h" @@ -47,7 +43,6 @@ using namespace llvm; using namespace llvm::sycl; -using namespace llvm::sycl_post_link; namespace { @@ -267,80 +262,6 @@ unsigned getOptLevel() { return 2; // default value } -Error saveModuleIR(Module &M, const StringRef Filename) { - std::error_code EC; - raw_fd_ostream Out{Filename, EC, sys::fs::OF_None}; - if (EC) - return createStringError(EC, "error opening the file '" + Filename + "'"); - - ModulePassManager MPM; - ModuleAnalysisManager MAM; - PassBuilder PB; - PB.registerModuleAnalyses(MAM); - if (OutputAssembly) - MPM.addPass(PrintModulePass(Out)); - else if (Force || !CheckBitcodeOutputToConsole(Out)) - MPM.addPass(BitcodeWriterPass(Out)); - MPM.run(M, MAM); - return Error::success(); -} - -PropSetRegTy computeModuleProperties(const module_split::ModuleDesc &MD, - const GlobalBinImageProps &GlobProps) { - PropSetRegTy PropSet; - // For bf16 devicelib module, no kernel included and no specialization - // constant used, skip regular Prop emit. However, we have fallback and - // native version of bf16 devicelib and we need new property values to - // indicate all exported function. - if (!MD.isSYCLDeviceLib()) - PropSet = computeModuleProperties(MD.getModule(), MD.entries(), GlobProps, - AllowDeviceImageDependencies); - else - PropSet = computeDeviceLibProperties(MD.getModule(), MD.Name); - - // When the split mode is none, the required work group size will be added - // to the whole module, which will make the runtime unable to - // launch the other kernels in the module that have different - // required work group sizes or no required work group sizes. So we need to - // remove the required work group size metadata in this case. - if (SplitMode == module_split::SPLIT_NONE) - PropSet.remove(PropSetRegTy::SYCL_DEVICE_REQUIREMENTS, - PropSetRegTy::PROPERTY_REQD_WORK_GROUP_SIZE); - return PropSet; -} - -Error writePropertiesToFile(const StringRef Filename, - const util::PropertySetRegistry &PropSet) { - return writeToOutput(Filename, [&](raw_ostream &OS) -> Error { - PropSet.write(OS); - return Error::success(); - }); -} - -Error saveModuleProperties(const module_split::ModuleDesc &MD, - const GlobalBinImageProps &GlobProps, - const StringRef Filename, StringRef Target = "") { - PropSetRegTy PropSet = computeModuleProperties(MD, GlobProps); - - if (!Target.empty()) - PropSet.add(PropSetRegTy::SYCL_DEVICE_REQUIREMENTS, "compile_target", - Target); - - return writePropertiesToFile(Filename, PropSet); -} - -// Saves specified collection of symbols to a file. -Error saveModuleSymbolTable(const module_split::ModuleDesc &MD, - const StringRef Filename) { - return writeToOutput(Filename, [&](raw_ostream &OS) -> Error { - OS << computeModuleSymbolTable(MD.getModule(), MD.entries()); - return Error::success(); - }); -} - -bool isTargetCompatibleWithModule(const std::string &Target, - module_split::ModuleDesc &IrMD); - void addTableRow(util::SimpleTable &Table, const IrPropSymFilenameTriple &RowData); @@ -372,17 +293,18 @@ void saveModule( StringRef IRExtension = OutputAssembly ? ".ll" : ".bc"; BaseTriple.Ir = (OutputPrefix + Suffix + "_" + Twine(I) + IRExtension).str(); - ExitOnErr(saveModuleIR(MD.getModule(), BaseTriple.Ir)); + ExitOnErr(sycl_post_link::saveModuleIR(MD.getModule(), BaseTriple.Ir, + OutputAssembly)); } if (DoSymGen) { // Save the names of the entry points - the symbol table. BaseTriple.Sym = (OutputPrefix + Suffix + "_" + Twine(I) + ".sym").str(); - ExitOnErr(saveModuleSymbolTable(MD, BaseTriple.Sym)); + ExitOnErr(sycl_post_link::saveModuleSymbolTable(MD, BaseTriple.Sym)); } for (const auto &[Table, OutputFile] : zip_equal(OutTables, OutputFiles)) { - if (!isTargetCompatibleWithModule(OutputFile.Target, MD)) + if (!sycl_post_link::isTargetCompatibleWithModule(OutputFile.Target, MD)) continue; auto CopyTriple = BaseTriple; if (DoPropGen) { @@ -396,7 +318,9 @@ void saveModule( CopyTriple.Prop = (OutputPrefix + NewSuff + "_" + Twine(I) + ".prop").str(); - ExitOnErr(saveModuleProperties(MD, Props, CopyTriple.Prop, Target)); + ExitOnErr(sycl_post_link::saveModuleProperties( + MD, Props, CopyTriple.Prop, Target, AllowDeviceImageDependencies, + SplitMode)); } addTableRow(*Table, CopyTriple); } @@ -439,54 +363,6 @@ void addTableRow(util::SimpleTable &Table, Table.addRow(Row); } -// Checks if the given target and module are compatible. -// A target and module are compatible if all the optional kernel features -// the module uses are supported by that target (i.e. that module can be -// compiled for that target and then be executed on that target). This -// information comes from the device config file (DeviceConfigFile.td). -// For example, the intel_gpu_tgllp target does not support fp64 - therefore, -// a module using fp64 would *not* be compatible with intel_gpu_tgllp. -bool isTargetCompatibleWithModule(const std::string &Target, - module_split::ModuleDesc &IrMD) { - // When the user does not specify a target, - // (e.g. -o out.table compared to -o intel_gpu_pvc,out-pvc.table) - // Target will be empty and we will not want to perform any filtering, so - // we return true here. - if (Target.empty()) - return true; - - // TODO: If a target not found in the device config file is passed, - // to sycl-post-link, then we should probably throw an error. However, - // since not all the information for all the targets is filled out - // right now, we return true, having the affect that unrecognized - // targets have no filtering applied to them. - if (!is_contained(DeviceConfigFile::TargetTable, Target)) - return true; - - const DeviceConfigFile::TargetInfo &TargetInfo = - DeviceConfigFile::TargetTable[Target]; - const SYCLDeviceRequirements &ModuleReqs = - IrMD.getOrComputeDeviceRequirements(); - - // Check to see if all the requirements of the input module - // are compatbile with the target. - for (const auto &Aspect : ModuleReqs.Aspects) { - if (!is_contained(TargetInfo.aspects, Aspect.Name)) - return false; - } - - // Check if module sub group size is compatible with the target. - // For ESIMD, the reqd_sub_group_size will be 1; this is not - // a supported by any backend (e.g. no backend can support a kernel - // with sycl::reqd_sub_group_size(1)), but for ESIMD, this is - // a special case. - if (!IrMD.isESIMD() && ModuleReqs.SubGroupSize.has_value() && - !is_contained(TargetInfo.subGroupSizes, *ModuleReqs.SubGroupSize)) - return false; - - return true; -} - std::vector> processInputModule(std::unique_ptr M, const StringRef OutputPrefix) { // Construct the resulting table which will accumulate all the outputs. @@ -559,15 +435,16 @@ processInputModule(std::unique_ptr M, const StringRef OutputPrefix) { MDesc->fixupLinkageOfDirectInvokeSimdTargets(); - ESIMDProcessingOptions Options = {SplitMode, - EmitOnlyKernelsAsEntryPoints, - AllowDeviceImageDependencies, - LowerEsimd, - SplitEsimd, - getOptLevel(), - ForceDisableESIMDOpt}; - auto ModulesOrErr = - handleESIMD(std::move(MDesc), Options, Modified, SplitOccurred); + llvm::sycl_post_link::ESIMDProcessingOptions Options = { + SplitMode, + EmitOnlyKernelsAsEntryPoints, + AllowDeviceImageDependencies, + LowerEsimd, + SplitEsimd, + getOptLevel(), + ForceDisableESIMDOpt}; + auto ModulesOrErr = llvm::sycl_post_link::handleESIMD( + std::move(MDesc), Options, Modified, SplitOccurred); CHECK_AND_EXIT(ModulesOrErr.takeError()); SmallVector, 2> &MMs = *ModulesOrErr; @@ -584,8 +461,8 @@ processInputModule(std::unique_ptr M, const StringRef OutputPrefix) { "' can't be used"); } MMs.front()->cleanup(AllowDeviceImageDependencies); - ExitOnErr( - saveModuleIR(MMs.front()->getModule(), OutputFiles[0].Filename)); + ExitOnErr(sycl_post_link::saveModuleIR( + MMs.front()->getModule(), OutputFiles[0].Filename, OutputAssembly)); return Tables; } // Empty IR file name directs saveModule to generate one and save IR to