diff --git a/clang/include/clang/Basic/DiagnosticDriverKinds.td b/clang/include/clang/Basic/DiagnosticDriverKinds.td index bc2f85d8094e..70e84fa10362 100644 --- a/clang/include/clang/Basic/DiagnosticDriverKinds.td +++ b/clang/include/clang/Basic/DiagnosticDriverKinds.td @@ -145,6 +145,9 @@ def warn_drv_unsupported_diag_option_for_flang : Warning< def warn_drv_unsupported_option_for_processor : Warning< "ignoring '%0' option as it is not currently supported for processor '%1'">, InGroup; +def warn_drv_unsupported_option_overrides_option : Warning< + "ignoring '%0' option as option '%1' overrides the behavior">, + InGroup; def warn_drv_unsupported_openmp_library : Warning< "the library '%0=%1' is not supported, OpenMP will not be enabled">, InGroup; diff --git a/clang/include/clang/CIR/Dialect/IR/CIRCUDAAttrs.td b/clang/include/clang/CIR/Dialect/IR/CIRCUDAAttrs.td index 504eb35c8357..f44023686e4e 100644 --- a/clang/include/clang/CIR/Dialect/IR/CIRCUDAAttrs.td +++ b/clang/include/clang/CIR/Dialect/IR/CIRCUDAAttrs.td @@ -30,7 +30,7 @@ def CIR_CUDAKernelNameAttr : CIR_Attr<"CUDAKernelName", "cu.kernel_name"> { names, we must record the corresponding device-side name for a stub. }]; - let parameters = (ins "std::string":$kernel_name); + let parameters = (ins StringRefParameter<"">:$kernel_name); let assemblyFormat = "`<` $kernel_name `>`"; } @@ -65,7 +65,7 @@ def CIR_CUDABinaryHandleAttr : CIR_Attr< and then generate various registration functions. }]; - let parameters = (ins "std::string":$name); + let parameters = (ins StringRefParameter<"">:$name); let assemblyFormat = "`<` $name `>`"; } diff --git a/clang/include/clang/CIR/Dialect/IR/CIROps.td b/clang/include/clang/CIR/Dialect/IR/CIROps.td index d796d8340ff1..bf1f000a1811 100644 --- a/clang/include/clang/CIR/Dialect/IR/CIROps.td +++ b/clang/include/clang/CIR/Dialect/IR/CIROps.td @@ -3992,6 +3992,58 @@ def CIR_DerivedMethodOp : CIR_Op<"derived_method", [Pure]> { let hasVerifier = 1; } + +//===----------------------------------------------------------------------===// +// Offload container +//===----------------------------------------------------------------------===// + +def CIR_OffloadContainerOp : CIR_Op<"offload.container", + [NoRegionArguments, NoTerminator]> { + let summary = "Container for host and device CIR modules"; + let description = [{ + `cir.offload.container` is a top-level container used to keep host and device + CIR modules together for joint analysis and transformation. + + The operation owns a single region. The region typically contains nested + `module` operations such as `module @host { ... }` and `module @device { ... }`, + each providing its own symbol table scope to avoid host/device symbol conflicts. + + Example: + + ```mlir + module { + cir.offload.container { + module @host { + // host CIR + } + module @device { + // device CIR + } + } + } + ``` + }]; + + let arguments = (ins); + + let regions = (region AnyRegion:$body); + + // Use generic region printer/parser: `cir.offload.container { ... }` + let assemblyFormat = [{ + $body attr-dict + }]; + + let hasVerifier = 1; + + let extraClassDeclaration = [{ + /// Returns the nested `module @host` if present, otherwise nullopt. + std::optional getHostModule(); + + /// Returns nested `module @device` module. + std::optional getDeviceModule(); + }]; +} + //===----------------------------------------------------------------------===// // FuncOp //===----------------------------------------------------------------------===// diff --git a/clang/include/clang/CIR/FrontendAction/CIRCombineAction.h b/clang/include/clang/CIR/FrontendAction/CIRCombineAction.h new file mode 100644 index 000000000000..fdcc1a520638 --- /dev/null +++ b/clang/include/clang/CIR/FrontendAction/CIRCombineAction.h @@ -0,0 +1,29 @@ + +#include "clang/Frontend/FrontendAction.h" + +#include + +namespace mlir { +class MLIRContext; +class ModuleOp; +} // namespace mlir + +namespace cir { +class CIRCombineAction : public clang::FrontendAction { +private: + mlir::MLIRContext *mlirContext; + +public: + CIRCombineAction(); + std::unique_ptr + CreateASTConsumer(clang::CompilerInstance &CI, + llvm::StringRef InFile) override { + return std::make_unique(); + } + + void ExecuteAction() override; + // We don't need a preprocessor-only mode. + bool usesPreprocessorOnly() const override { return false; } + virtual bool hasCIRSupport() const override { return true; } +}; +} // namespace cir diff --git a/clang/include/clang/CIR/FrontendAction/CIRGenAction.h b/clang/include/clang/CIR/FrontendAction/CIRGenAction.h index 218829aec063..2a5c05bbc74f 100644 --- a/clang/include/clang/CIR/FrontendAction/CIRGenAction.h +++ b/clang/include/clang/CIR/FrontendAction/CIRGenAction.h @@ -6,8 +6,8 @@ // //===----------------------------------------------------------------------===// -#ifndef LLVM_CLANG_CIR_CIRGENACTION_H -#define LLVM_CLANG_CIR_CIRGENACTION_H +#ifndef LLVM_CLANG_CIR_CIRCOMBINEACTION_H +#define LLVM_CLANG_CIR_CIRCOMBINEACTION_H #include "clang/CodeGen/CodeGenAction.h" #include "clang/Frontend/FrontendAction.h" @@ -49,8 +49,6 @@ class CIRGenAction : public clang::ASTFrontendAction { mlir::MLIRContext *mlirContext; - mlir::OwningOpRef loadModule(llvm::MemoryBufferRef mbRef); - protected: CIRGenAction(OutputType action, mlir::MLIRContext *_MLIRContext = nullptr); diff --git a/clang/include/clang/Driver/Action.h b/clang/include/clang/Driver/Action.h index dbf1187da4db..7866b86a331f 100644 --- a/clang/include/clang/Driver/Action.h +++ b/clang/include/clang/Driver/Action.h @@ -76,6 +76,8 @@ class Action { StaticLibJobClass, BinaryAnalyzeJobClass, BinaryTranslatorJobClass, + CIRCombineJobClass, + CIRSplitJobClass, ObjcopyJobClass, JobClassFirst = PreprocessJobClass, @@ -180,8 +182,7 @@ class Action { /// files for each offloading kind. By default, no prefix is used for /// non-device kinds, except if \a CreatePrefixForHost is set. static std::string - GetOffloadingFileNamePrefix(OffloadKind Kind, - StringRef NormalizedTriple, + GetOffloadingFileNamePrefix(OffloadKind Kind, StringRef NormalizedTriple, bool CreatePrefixForHost = false); /// Return a string containing a offload kind name. @@ -242,9 +243,7 @@ class InputAction : public Action { void setId(StringRef _Id) { Id = _Id.str(); } StringRef getId() const { return Id; } - static bool classof(const Action *A) { - return A->getKind() == InputClass; - } + static bool classof(const Action *A) { return A->getKind() == InputClass; } }; class BindArchAction : public Action { @@ -259,9 +258,7 @@ class BindArchAction : public Action { StringRef getArchName() const { return ArchName; } - static bool classof(const Action *A) { - return A->getKind() == BindArchClass; - } + static bool classof(const Action *A) { return A->getKind() == BindArchClass; } }; /// An offload action combines host or/and device actions according to the @@ -407,8 +404,7 @@ class JobAction : public Action { public: static bool classof(const Action *A) { - return (A->getKind() >= JobClassFirst && - A->getKind() <= JobClassLast); + return (A->getKind() >= JobClassFirst && A->getKind() <= JobClassLast); } }; @@ -511,9 +507,7 @@ class LinkJobAction : public JobAction { public: LinkJobAction(ActionList &Inputs, types::ID Type); - static bool classof(const Action *A) { - return A->getKind() == LinkJobClass; - } + static bool classof(const Action *A) { return A->getKind() == LinkJobClass; } }; class LipoJobAction : public JobAction { @@ -522,9 +516,7 @@ class LipoJobAction : public JobAction { public: LipoJobAction(ActionList &Inputs, types::ID Type); - static bool classof(const Action *A) { - return A->getKind() == LipoJobClass; - } + static bool classof(const Action *A) { return A->getKind() == LipoJobClass; } }; class DsymutilJobAction : public JobAction { @@ -644,6 +636,50 @@ class OffloadPackagerJobAction : public JobAction { } }; +class CombineCIRJobAction : public JobAction { + void anchor() override; + const ToolChain *HostToolChain; + const ToolChain *DeviceToolChain; + Action *HostAction; + Action *DeviceAction; + char *HostBoundArch; + const char *DeviceBoundArch; + unsigned HostOffloadKind; + +public: + CombineCIRJobAction(const ToolChain *HostToolChain, + const ToolChain *DeviceToolChain, Action *HostAction, + Action *DeviceAction, char *HostBoundArch, + const char *DeviceBoundArch, unsigned HostOffloadKind, + types::ID Type, OffloadKind OffloadDeviceKind); + + static bool classof(const Action *A) { + return A->getKind() == CIRCombineJobClass; + } + + Action *getHostAction() { return HostAction; } + Action *getDeviceAction() { return DeviceAction; } + + const ToolChain *getHostToolChain() const { return HostToolChain; } + const ToolChain *getDeviceToolChain() const { return DeviceToolChain; } + + const char *getHostBoundArch() const { return HostBoundArch; } + const char *getDeviceBoundArch() const { return DeviceBoundArch; } +}; + +class SplitCIRJobAction : public JobAction { + void anchor() override; + +public: + bool isHost; + SplitCIRJobAction(Action *Input, bool isHost, types::ID Type, + OffloadKind Kind = OFK_None); + + static bool classof(const Action *A) { + return A->getKind() == CIRSplitJobClass; + } +}; + class LinkerWrapperJobAction : public JobAction { void anchor() override; diff --git a/clang/include/clang/Driver/Options.td b/clang/include/clang/Driver/Options.td index eaa0ad5406e6..9f8750a8cdd6 100644 --- a/clang/include/clang/Driver/Options.td +++ b/clang/include/clang/Driver/Options.td @@ -3156,6 +3156,13 @@ defm clangir : BoolFOption<"clangir", PosFlag, NegFlag LLVM pipeline to compile">, BothFlags<[], [ClangOption, CC1Option], "">>; +defm clangir_offload : BoolFOption<"clangir-offload", + FrontendOpts<"UseClangIROffloadPipeline">, DefaultFalse, + PosFlag, + NegFlag, + BothFlags<[], [ClangOption, CC1Option], "">>; def fcir_output_EQ : Joined<["-"], "fcir-output=">, Group, Visibility<[ClangOption, CC1Option]>, Flags<[NoArgumentUnused]>, @@ -3208,6 +3215,7 @@ def fclangir_mem2reg : Flag<["-"], "fclangir-mem2reg">, HelpText<"Enable mem2reg on the flat ClangIR">, MarshallingInfoFlag>; +def CIR_Group : OptionGroup<"">; def clangir_disable_passes : Flag<["-"], "clangir-disable-passes">, Visibility<[ClangOption, CC1Option]>, HelpText<"Disable CIR transformations pipeline">, @@ -3256,6 +3264,25 @@ def emit_cir_only : Flag<["-"], "emit-cir-only">, def emit_cir_flat : Flag<["-"], "emit-cir-flat">, Visibility<[ClangOption, CC1Option]>, Group, Alias, AliasArgs<["cir-flat"]>, HelpText<"Similar to -emit-cir but also lowers structured CFG into basic blocks.">; +def cir_combine : Flag<["-"], "cir-combine">, + Visibility<[CC1Option]>, + Group, + HelpText<"Combine host and device CIR modules into a single offload container CIR module.">; +def cir_host_input : Separate<["-"], "cir-host-input">, + Visibility<[CC1Option]>, Group, + HelpText<"Host CIR input for -cir-combine.">; +def cir_device_input : Separate<["-"], "cir-device-input">, + Visibility<[CC1Option]>, Group, + HelpText<"Device CIR input for -cir-combine (may be repeated).">; +def cir_emit_split : Flag<["-"], "cir-emit-split">, + Visibility<[CC1Option]>, Group, + HelpText<"Emit split host/device CIR instead of a combined CIR container">; +def cir_host_output : Separate<["-"], "cir-host-output">, + Visibility<[CC1Option]>, Group, + HelpText<"Output path for host CIR when -cir-emit-split is used">; +def cir_device_output : Separate<["-"], "cir-device-output">, + Visibility<[CC1Option]>, Group, + HelpText<"Output path for device CIR when -cir-emit-split is used">; /// ClangIR-specific options - END def flto : Flag<["-"], "flto">, diff --git a/clang/include/clang/Frontend/FrontendOptions.h b/clang/include/clang/Frontend/FrontendOptions.h index ccc1ab710b63..e04c8f3363f4 100644 --- a/clang/include/clang/Frontend/FrontendOptions.h +++ b/clang/include/clang/Frontend/FrontendOptions.h @@ -68,6 +68,10 @@ enum ActionKind { /// Generate CIR, bud don't emit anything. EmitCIROnly, + /// Combine multiple CIR modules (e.g. host and device) into a single + /// container + CIRCombine, + /// Emit a .mlir file EmitMLIR, @@ -413,6 +417,11 @@ class FrontendOptions { LLVM_PREFERRED_TYPE(bool) unsigned UseClangIRPipeline : 1; + /// Use CIR-based offload pipeline (combine/split/fatbin/embed) when compiling + /// offload code. + LLVM_PREFERRED_TYPE(bool) + unsigned UseClangIROffloadPipeline : 1; + /// Lower directly from ClangIR to LLVM unsigned ClangIRDirectLowering : 1; @@ -454,6 +463,11 @@ class FrontendOptions { std::string ClangIRIdiomRecognizerOpts; std::string ClangIRLibOptOpts; std::string ClangIRFile; + std::string CIRHostInput; + std::string CIRDeviceInput; + bool EmitSplit; + std::string CIRHostOutput; + std::string CIRDeviceOutput; frontend::MLIRDialectKind MLIRTargetDialect = frontend::MLIR_CORE; @@ -532,7 +546,6 @@ class FrontendOptions { /// should only be used for debugging and experimental features. std::vector MLIRArgs; - /// File name of the file that will provide record layouts /// (in the format produced by -fdump-record-layouts). std::string OverrideRecordLayoutsFile; @@ -587,7 +600,7 @@ class FrontendOptions { ClangIRVerifyDiags(false), ClangIRLifetimeCheck(false), ClangIRIdiomRecognizer(false), ClangIRLibOpt(false), ClangIRCallConvLowering(true), ClangIREnableMem2Reg(false), - ClangIRAnalysisOnly(false), EmitClangIRFile(false), + ClangIRAnalysisOnly(false), EmitClangIRFile(false), EmitSplit(false), TimeTraceGranularity(500), TimeTraceVerbose(false) {} /// getInputKindForExtension - Return the appropriate input kind for a file diff --git a/clang/lib/CIR/CodeGen/TargetInfo.cpp b/clang/lib/CIR/CodeGen/TargetInfo.cpp index 5c9be81f8212..269fee58da91 100644 --- a/clang/lib/CIR/CodeGen/TargetInfo.cpp +++ b/clang/lib/CIR/CodeGen/TargetInfo.cpp @@ -386,6 +386,29 @@ class AMDGPUTargetCIRGenInfo : public TargetCIRGenInfo { ft = getABIInfo().getContext().adjustFunctionType( ft, ft->getExtInfo().withCallingConv(CC_DeviceKernel)); } + + void setTargetAttributes(const clang::Decl *decl, mlir::Operation *global, + CIRGenModule &cgm) const override { + + if (const auto *fd = clang::dyn_cast_or_null(decl)) { + cir::FuncOp func = mlir::cast(global); + if (func.isDeclaration()) + return; + + if (cgm.getLangOpts().HIP) { + if (fd->hasAttr()) { + func.setCallingConv(cir::CallingConv::AMDGPUKernel); + func.setLinkageAttr(cir::GlobalLinkageKindAttr::get( + func.getContext(), cir::GlobalLinkageKind::ExternalLinkage)); + func.setVisibility(mlir::SymbolTable::Visibility::Public); + func.setGlobalVisibility(cir::VisibilityKind::Protected); + } + } + + if (fd->getAttr()) + llvm_unreachable("NYI"); + } + } }; } // namespace diff --git a/clang/lib/CIR/Dialect/IR/CIRDialect.cpp b/clang/lib/CIR/Dialect/IR/CIRDialect.cpp index 8a472617c7e2..84c0e071c2cb 100644 --- a/clang/lib/CIR/Dialect/IR/CIRDialect.cpp +++ b/clang/lib/CIR/Dialect/IR/CIRDialect.cpp @@ -2642,6 +2642,49 @@ LogicalResult cir::VTTAddrPointOp::verify() { return success(); } +//===----------------------------------------------------------------------===// +// Offload container +//===----------------------------------------------------------------------===// + +std::optional cir::OffloadContainerOp::getHostModule() { + if (getBody().empty()) + return std::nullopt; + + for (mlir::Operation &op : getBody().front()) { + auto mod = llvm::dyn_cast(op); + if (!mod) + continue; + if (auto name = mod.getSymNameAttr(); name && name.getValue() == "host") + return mod; + } + return std::nullopt; +} + +std::optional cir::OffloadContainerOp::getDeviceModule() { + if (getBody().empty()) + return std::nullopt; + + for (mlir::Operation &op : getBody().front()) { + auto mod = llvm::dyn_cast(op); + if (!mod) + continue; + if (auto name = mod.getSymNameAttr(); name && name.getValue() == "device") + return mod; + } + return std::nullopt; +} + +mlir::LogicalResult cir::OffloadContainerOp::verify() { + auto host = getHostModule(); + auto dev = getDeviceModule(); + + if (!host) + return emitOpError() << "expects nested module @host"; + if (!dev) + return emitOpError() << "expects nested module @device"; + return mlir::success(); +} + //===----------------------------------------------------------------------===// // FuncOp //===----------------------------------------------------------------------===// @@ -3798,31 +3841,47 @@ ::mlir::Attribute cir::ConstArrayAttr::parse(::mlir::AsmParser &parser, return {}; } - // ArrayAttrrs have per-element type, not the type of the array... - if (llvm::dyn_cast(*resultVal)) { - // Array has implicit type: infer from const array type. - if (parser.parseOptionalColon().failed()) { - resultTy = type; - } else { // Array has explicit type: parse it. - resultTy = ::mlir::FieldParser<::mlir::Type>::parse(parser); - if (failed(resultTy)) { - parser.emitError( - parser.getCurrentLocation(), - "failed to parse ConstArrayAttr parameter 'type' which is " - "to be a `::mlir::Type`"); - return {}; - } + // Case 1: array attribute => infer from outer `type` (since printer never + // prints `: ty`). + if (llvm::isa(*resultVal)) { + resultTy = type; + if (!type) { + parser.emitError(parser.getCurrentLocation(), + "missing outer type for ConstArrayAttr"); + return {}; } - } else { - assert(mlir::isa(*resultVal) && "IDK"); - auto ta = mlir::cast(*resultVal); - resultTy = ta.getType(); - if (mlir::isa(resultTy.value())) { + } + // Case 2: string literal => also infer from outer `type`. + else if (llvm::isa(*resultVal)) { + resultTy = type; + if (!type) { parser.emitError(parser.getCurrentLocation(), - "expected type declaration for string literal"); + "expected outer type declaration for string literal"); return {}; } } + // Case 3: typed attribute => use its type, but allow fallback to outer type + // if NoneType. + else { + auto ta = llvm::dyn_cast(*resultVal); + if (!ta) { + parser.emitError( + parser.getCurrentLocation(), + "expected array, string, or typed attribute for ConstArrayAttr"); + return {}; + } + resultTy = ta.getType(); + if (mlir::isa(resultTy.value())) { + // For safety: prefer outer type if provided. + if (type) + resultTy = type; + else { + parser.emitError(parser.getCurrentLocation(), + "expected type declaration for attribute"); + return {}; + } + } + } auto zeros = 0; if (parser.parseOptionalComma().succeeded()) { diff --git a/clang/lib/CIR/Dialect/Transforms/LoweringPrepare.cpp b/clang/lib/CIR/Dialect/Transforms/LoweringPrepare.cpp index d7359bf960e4..4202837a777c 100644 --- a/clang/lib/CIR/Dialect/Transforms/LoweringPrepare.cpp +++ b/clang/lib/CIR/Dialect/Transforms/LoweringPrepare.cpp @@ -1014,7 +1014,7 @@ void LoweringPreparePass::buildCUDAModuleCtor() { assert(!cir::MissingFeatures::hipModuleCtor()); return; } - std::string cudaGPUBinaryName = + llvm::StringRef cudaGPUBinaryName = cast(cudaBinaryHandleAttr).getName(); constexpr unsigned cudaFatMagic = 0x466243b1; @@ -1830,8 +1830,7 @@ void LoweringPreparePass::runOnOp(Operation *op) { if (auto attr = fnOp.getExtraAttrs().getElements().get( CUDAKernelNameAttr::getMnemonic())) { auto cudaBinaryAttr = dyn_cast(attr); - std::string kernelName = cudaBinaryAttr.getKernelName(); - cudaKernelMap[kernelName] = fnOp; + cudaKernelMap[cudaBinaryAttr.getKernelName()] = fnOp; } if (std::optional annotations = fnOp.getAnnotations()) addGlobalAnnotations(fnOp, annotations.value()); diff --git a/clang/lib/CIR/Dialect/Transforms/TargetLowering/Targets/AMDGPU.cpp b/clang/lib/CIR/Dialect/Transforms/TargetLowering/Targets/AMDGPU.cpp index 76d457f1607b..0f2b2e7209be 100644 --- a/clang/lib/CIR/Dialect/Transforms/TargetLowering/Targets/AMDGPU.cpp +++ b/clang/lib/CIR/Dialect/Transforms/TargetLowering/Targets/AMDGPU.cpp @@ -34,9 +34,15 @@ class AMDGPUABIInfo : public ABIInfo { AMDGPUABIInfo(LowerTypes <) : ABIInfo(lt) {} private: - void computeInfo(LowerFunctionInfo &fi) const override { - llvm_unreachable("NYI"); - } + static const unsigned MaxNumRegsForArgsRet = 16; + + ABIArgInfo classifyReturnType(mlir::Type ty) const; + ABIArgInfo classifyArgumentType(mlir::Type Ty, bool Variadic, + unsigned &NumRegsLeft) const; + + ABIArgInfo classifyKernelArgumentType(mlir::Type ty) const; + + void computeInfo(LowerFunctionInfo &fi) const override; }; class AMDGPUTargetLoweringInfo : public TargetLoweringInfo { @@ -64,6 +70,77 @@ class AMDGPUTargetLoweringInfo : public TargetLoweringInfo { }; } // namespace + +ABIArgInfo AMDGPUABIInfo::classifyReturnType(mlir::Type ty) const { + if (llvm::isa(ty)) + return ABIArgInfo::getIgnore(); + + if (getContext().getLangOpts().OpenMP) + llvm_unreachable("NYI"); + + if (!isScalarType(ty)) + return ABIArgInfo::getDirect(); + + // OG treats enums as their underlying type. + // This has already been done for CIR. + + // Integers with size < 32 must be extended to 32 bits. + // (See Section 3.3 of PTX ABI.) + return (isPromotableIntegerTypeForABI(ty) ? ABIArgInfo::getExtend(ty) + : ABIArgInfo::getDirect()); +} + +/// For kernels all parameters are really passed in a special buffer. It doesn't +/// make sense to pass anything byval, so everything must be direct. +ABIArgInfo AMDGPUABIInfo::classifyKernelArgumentType(mlir::Type ty) const { + return ABIArgInfo::getDirect(); +} + +ABIArgInfo AMDGPUABIInfo::classifyArgumentType(mlir::Type ty, bool variadic, + unsigned &numRegsLeft) const { + assert(numRegsLeft <= MaxNumRegsForArgsRet && "register estimate underflow"); + + ty = useFirstFieldIfTransparentUnion(ty); + + // Variadic arguments: always direct. + if (variadic) { + return ABIArgInfo::getDirect(); + } + + // Aggregate (struct/array) handling + if (isAggregateTypeForABI(ty)) { + llvm_unreachable("NYI"); + } + + // === Non-aggregate fallback === + ABIArgInfo Info = isPromotableIntegerTypeForABI(ty) + ? ABIArgInfo::getExtend(ty) + : ABIArgInfo::getDirect(); + + return Info; +} + +void AMDGPUABIInfo::computeInfo(LowerFunctionInfo &fi) const { + llvm::CallingConv::ID cc = fi.getCallingConvention(); + + if (!getCXXABI().classifyReturnType(fi)) + fi.getReturnInfo() = classifyReturnType(fi.getReturnType()); + + unsigned argumentIndex = 0; + const unsigned numFixedArguments = fi.getNumRequiredArgs(); + + unsigned numRegsLeft = MaxNumRegsForArgsRet; + for (auto &argument : fi.arguments()) { + if (cc == llvm::CallingConv::AMDGPU_KERNEL) { + argument.info = classifyKernelArgumentType(argument.type); + } else { + bool fixedArgument = argumentIndex++ < numFixedArguments; + argument.info = + classifyArgumentType(argument.type, !fixedArgument, numRegsLeft); + } + } +} + std::unique_ptr createAMDGPUTargetLoweringInfo(LowerModule &lowerModule) { return std::make_unique(lowerModule.getTypes()); diff --git a/clang/lib/CIR/FrontendAction/CIRCombineAction.cpp b/clang/lib/CIR/FrontendAction/CIRCombineAction.cpp new file mode 100644 index 000000000000..e303a19ae8d4 --- /dev/null +++ b/clang/lib/CIR/FrontendAction/CIRCombineAction.cpp @@ -0,0 +1,153 @@ +#include "clang/CIR/FrontendAction/CIRCombineAction.h" +#include "mlir/Dialect/DLTI/DLTI.h" +#include "mlir/Dialect/Func/IR/FuncOps.h" +#include "mlir/Dialect/LLVMIR/LLVMDialect.h" +#include "mlir/Dialect/MemRef/IR/MemRef.h" +#include "mlir/Dialect/OpenMP/OpenMPDialect.h" +#include "mlir/IR/BuiltinDialect.h" +#include "mlir/IR/MLIRContext.h" +#include "mlir/IR/OwningOpRef.h" +#include "mlir/Parser/Parser.h" +#include "clang/Basic/DiagnosticDriver.h" +#include "clang/Basic/DiagnosticFrontend.h" +#include "clang/CIR/Dialect/IR/CIRDialect.h" +#include "clang/CodeGen/BackendUtil.h" + +using namespace cir; +using namespace clang; + +static mlir::FailureOr> +loadModule(llvm::MemoryBufferRef MbRef, mlir::MLIRContext &mlirContext) { + auto Module = + mlir::parseSourceString(MbRef.getBuffer(), &mlirContext); + if (!Module) + return mlir::failure(); + return Module; +} + +CIRCombineAction::CIRCombineAction() : mlirContext(new mlir::MLIRContext) {} + +void CIRCombineAction::ExecuteAction() { + auto &Ci = getCompilerInstance(); + auto &Diags = Ci.getDiagnostics(); + const clang::FrontendOptions &Fo = Ci.getFrontendOpts(); + + // Expect ParseFrontendArgs already validated these, but keep it defensive. + if (Fo.CIRHostInput.empty() || Fo.CIRDeviceInput.empty()) { + Diags.Report(clang::diag::err_fe_error_reading) + << "missing -cir-host-input/-cir-device-input"; + return; + } + + if (!Fo.EmitSplit && Fo.OutputFile.empty()) { + Diags.Report(clang::diag::err_drv_missing_arg_mtp) << "-o"; + return; + } else if (Fo.EmitSplit) { + if (Fo.CIRHostOutput.empty()) + Diags.Report(clang::diag::err_drv_missing_arg_mtp) << "-cir-host-output"; + if (Fo.CIRDeviceOutput.empty()) + Diags.Report(clang::diag::err_drv_missing_arg_mtp) + << "-cir-device-output"; + } + + mlirContext->getOrLoadDialect(); + mlirContext->getOrLoadDialect(); + mlirContext->getOrLoadDialect(); + mlirContext->getOrLoadDialect(); + mlirContext->getOrLoadDialect(); + mlirContext->getOrLoadDialect(); + mlirContext->getOrLoadDialect(); + mlirContext->getOrLoadDialect(); + + llvm::ErrorOr> HostInputOrErr = + Ci.getFileManager().getBufferForFile(Fo.CIRHostInput); + if (!HostInputOrErr) { + std::error_code Ec = HostInputOrErr.getError(); + Diags.Report(clang::diag::err_fe_error_reading) << Fo.CIRHostInput; + Diags.Report(clang::diag::note_drv_command_failed_diag_msg) << Ec.message(); + return; + } + + llvm::ErrorOr> DeviceInputOrErr = + Ci.getFileManager().getBufferForFile(Fo.CIRDeviceInput); + if (!DeviceInputOrErr) { + std::error_code Ec = DeviceInputOrErr.getError(); + Diags.Report(clang::diag::err_fe_error_reading) << Fo.CIRDeviceInput; + Diags.Report(clang::diag::note_drv_command_failed_diag_msg) << Ec.message(); + return; + } + + std::unique_ptr HostInput = std::move(*HostInputOrErr); + std::unique_ptr DeviceInput = + std::move(*DeviceInputOrErr); + + auto HostCirModuleOr = loadModule(*HostInput, *mlirContext); + if (mlir::failed(HostCirModuleOr)) { + Diags.Report(clang::diag::err_fe_error_reading) + << "failed to parse CIR module" << Fo.CIRHostInput; + return; + } + mlir::OwningOpRef HostCirModule = std::move(*HostCirModuleOr); + + auto DeviceCirModuleOr = loadModule(*DeviceInput, *mlirContext); + if (mlir::failed(DeviceCirModuleOr)) { + Diags.Report(clang::diag::err_fe_error_reading) + << "failed to parse CIR module" << Fo.CIRDeviceInput; + return; + } + mlir::OwningOpRef DeviceCirModule = + std::move(*DeviceCirModuleOr); + + // Rename nested modules so the container verifier can find them. + HostCirModule->setSymNameAttr(mlir::StringAttr::get(mlirContext, "host")); + DeviceCirModule->setSymNameAttr(mlir::StringAttr::get(mlirContext, "device")); + + // Create a new top-level module which will hold the container. + auto Loc = mlir::UnknownLoc::get(mlirContext); + mlir::ModuleOp Combined = mlir::ModuleOp::create(Loc); + + mlir::OpBuilder Builder = mlir::OpBuilder::atBlockBegin(Combined.getBody()); + // Insert the container op into the top-level module body. + Builder.setInsertionPointToStart(Combined.getBody()); + auto Container = cir::OffloadContainerOp::create(Builder, Loc); + + // Ensure the container region has a block. + mlir::Region &Body = Container.getBody(); + if (Body.empty()) + Body.push_back(new mlir::Block()); + + mlir::Block &Blk = Body.front(); + mlir::OpBuilder IB(&Blk, Blk.end()); + + // Clone the parsed host/device modules into the container body. + // (Cloning is simplest/robust for PR2.) + IB.insert(HostCirModule->getOperation()->clone()); + IB.insert(DeviceCirModule->getOperation()->clone()); + auto EmitCIR = [&](mlir::ModuleOp &mOp, StringRef Output) { + mlir::OpPrintingFlags Flags; + Flags.enableDebugInfo(/*enable=*/true, /*prettyForm=*/true); + + std::error_code EC; + llvm::raw_fd_ostream OS(Output, EC, llvm::sys::fs::OF_Text); + + if (EC) { + Diags.Report(clang::diag::err_fe_error_opening) << Output << EC.message(); + return; + } + + mOp.print(OS, Flags); + }; + + if (!Fo.EmitSplit) { + EmitCIR(Combined, Fo.OutputFile); + return; + } + auto devModOr = Container.getDeviceModule(); + if (!devModOr) + Container.emitError("missing device module in offload container"); + EmitCIR(*devModOr, Fo.CIRDeviceOutput); + auto hostModOr = Container.getHostModule(); + if (!hostModOr) + Container.emitError("missing host module in offload container"); + EmitCIR(*hostModOr, Fo.CIRHostOutput); +} diff --git a/clang/lib/CIR/FrontendAction/CIRGenAction.cpp b/clang/lib/CIR/FrontendAction/CIRGenAction.cpp index 6064cebb8e73..982c6c07f440 100644 --- a/clang/lib/CIR/FrontendAction/CIRGenAction.cpp +++ b/clang/lib/CIR/FrontendAction/CIRGenAction.cpp @@ -7,8 +7,14 @@ //===----------------------------------------------------------------------===// #include "clang/CIR/FrontendAction/CIRGenAction.h" +#include "mlir/Bytecode/BytecodeReader.h" +#include "mlir/Bytecode/BytecodeWriter.h" +#include "mlir/Dialect/DLTI/DLTI.h" #include "mlir/Dialect/Func/IR/FuncOps.h" +#include "mlir/Dialect/LLVMIR/LLVMDialect.h" #include "mlir/Dialect/MemRef/IR/MemRef.h" +#include "mlir/Dialect/OpenMP/OpenMPDialect.h" +#include "mlir/IR/BuiltinDialect.h" #include "mlir/IR/BuiltinOps.h" #include "mlir/IR/MLIRContext.h" #include "mlir/IR/OperationSupport.h" @@ -195,99 +201,62 @@ class CIRGenConsumer : public clang::ASTConsumer { llvm_unreachable("NYI"); } - void HandleTranslationUnit(ASTContext &C) override { - llvm::TimeTraceScope Scope("CIR Gen"); - - // Note that this method is called after `HandleTopLevelDecl` has already - // ran all over the top level decls. Here clang mostly wraps defered and - // global codegen, followed by running CIR passes. - Gen->HandleTranslationUnit(C); - - if (!FeOptions.ClangIRDisableCIRVerifier) - if (!Gen->verifyModule()) { - llvm::report_fatal_error( - "CIR codegen: module verification error before running CIR passes"); - return; - } - - auto MlirMod = Gen->getModule(); - auto MlirCtx = Gen->takeContext(); - - auto SetupCirPipelineAndExecute = [&] { - // Sanitize passes options. MLIR uses spaces between pass options - // and since that's hard to fly in clang, we currently use ';'. - std::string LifetimeOpts, IdiomRecognizerOpts, LibOptOpts; - if (FeOptions.ClangIRLifetimeCheck) - LifetimeOpts = sanitizePassOptions(FeOptions.ClangIRLifetimeCheckOpts); - if (FeOptions.ClangIRIdiomRecognizer) - IdiomRecognizerOpts = - sanitizePassOptions(FeOptions.ClangIRIdiomRecognizerOpts); - if (FeOptions.ClangIRLibOpt) - LibOptOpts = sanitizePassOptions(FeOptions.ClangIRLibOptOpts); - - bool EnableCcLowering = - FeOptions.ClangIRCallConvLowering && - !(Action == CIRGenAction::OutputType::EmitMLIR && - FeOptions.MLIRTargetDialect == frontend::MLIR_CIR); - bool FlattenCir = - Action == CIRGenAction::OutputType::EmitMLIR && - FeOptions.MLIRTargetDialect == clang::frontend::MLIR_CIR_FLAT; - - // Setup and run CIR pipeline. - std::string PassOptParsingFailure; - if (runCIRToCIRPasses( - MlirMod, MlirCtx.get(), C, !FeOptions.ClangIRDisableCIRVerifier, - FeOptions.ClangIRLifetimeCheck, LifetimeOpts, - FeOptions.ClangIRIdiomRecognizer, IdiomRecognizerOpts, - FeOptions.ClangIRLibOpt, LibOptOpts, PassOptParsingFailure, - CodeGenOpts.OptimizationLevel > 0, FlattenCir, - !FeOptions.ClangIRDirectLowering, EnableCcLowering, - FeOptions.ClangIREnableMem2Reg) - .failed()) { - if (!PassOptParsingFailure.empty()) { - auto D = Diags.Report(diag::err_drv_cir_pass_opt_parsing); - D << PassOptParsingFailure; - } else - llvm::report_fatal_error("CIR codegen: MLIR pass manager fails " - "when running CIR passes!"); - return; - } - }; - - if (!FeOptions.ClangIRDisablePasses) { - // Handle source manager properly given that lifetime analysis - // might emit warnings and remarks. - auto &ClangSourceMgr = C.getSourceManager(); - FileID MainFileID = ClangSourceMgr.getMainFileID(); - - std::unique_ptr FileBuf = - llvm::MemoryBuffer::getMemBuffer( - ClangSourceMgr.getBufferOrFake(MainFileID)); - - llvm::SourceMgr MlirSourceMgr; - MlirSourceMgr.AddNewSourceBuffer(std::move(FileBuf), llvm::SMLoc()); - - if (FeOptions.ClangIRVerifyDiags) { - mlir::SourceMgrDiagnosticVerifierHandler SourceMgrHandler( - MlirSourceMgr, MlirCtx.get()); - MlirCtx->printOpOnDiagnostic(false); - SetupCirPipelineAndExecute(); + void SetupCirPipelineAndExecute(mlir::ModuleOp MlirMod, + mlir::MLIRContext &MlirCtx, ASTContext &C) { + // Sanitize passes options. MLIR uses spaces between pass options + // and since that's hard to fly in clang, we currently use ';'. + std::string LifetimeOpts, IdiomRecognizerOpts, LibOptOpts; + if (FeOptions.ClangIRLifetimeCheck) + LifetimeOpts = sanitizePassOptions(FeOptions.ClangIRLifetimeCheckOpts); + if (FeOptions.ClangIRIdiomRecognizer) + IdiomRecognizerOpts = + sanitizePassOptions(FeOptions.ClangIRIdiomRecognizerOpts); + if (FeOptions.ClangIRLibOpt) + LibOptOpts = sanitizePassOptions(FeOptions.ClangIRLibOptOpts); + + bool EnableCcLowering = + FeOptions.ClangIRCallConvLowering && + !(Action == CIRGenAction::OutputType::EmitMLIR && + FeOptions.MLIRTargetDialect == frontend::MLIR_CIR); + bool FlattenCir = + Action == CIRGenAction::OutputType::EmitMLIR && + FeOptions.MLIRTargetDialect == clang::frontend::MLIR_CIR_FLAT; + + // Setup and run CIR pipeline. + std::string PassOptParsingFailure; + if (runCIRToCIRPasses( + MlirMod, &MlirCtx, C, !FeOptions.ClangIRDisableCIRVerifier, + FeOptions.ClangIRLifetimeCheck, LifetimeOpts, + FeOptions.ClangIRIdiomRecognizer, IdiomRecognizerOpts, + FeOptions.ClangIRLibOpt, LibOptOpts, PassOptParsingFailure, + CodeGenOpts.OptimizationLevel > 0, FlattenCir, + !FeOptions.ClangIRDirectLowering, EnableCcLowering, + FeOptions.ClangIREnableMem2Reg) + .failed()) { + if (!PassOptParsingFailure.empty()) { + auto D = Diags.Report(diag::err_drv_cir_pass_opt_parsing); + D << PassOptParsingFailure; + } else + llvm::report_fatal_error("CIR codegen: MLIR pass manager fails " + "when running CIR passes!"); + return; + } + } - // Verify the diagnostic handler to make sure that each of the - // diagnostics matched. - if (SourceMgrHandler.verify().failed()) { - // FIXME: we fail ungracefully, there's probably a better way - // to communicate non-zero return so tests can actually fail. - llvm::sys::RunInterruptHandlers(); - exit(1); - } - } else { - mlir::SourceMgrDiagnosticHandler SourceMgrHandler(MlirSourceMgr, - MlirCtx.get()); - SetupCirPipelineAndExecute(); - } + mlir::LogicalResult writeModuleBytecode(mlir::Operation *op, + raw_ostream &os) { + // If you have a ModuleOp, pass module.getOperation(). + // writeBytecodeToFile takes an Operation* and a raw_ostream. + if (mlir::failed(mlir::writeBytecodeToFile(op, os))) { + return mlir::failure(); } + os.flush(); + return mlir::success(); + } + + void GenerateOutput(mlir::ModuleOp MlirMod, + std::unique_ptr MlirCtx) { bool EmitCIR = LangOpts.EmitCIRToFile || FeOptions.EmitClangIRFile || !LangOpts.CIRFile.empty() || !FeOptions.ClangIRFile.empty(); if (EmitCIR) { @@ -332,12 +301,8 @@ class CIRGenConsumer : public clang::ASTConsumer { assert(MlirMod && "MLIR module does not exist, but lowering did not fail?"); assert(OutputStream && "Why are we here without an output stream?"); - // FIXME: we cannot roundtrip prettyForm=true right now. - mlir::OpPrintingFlags Flags; - Flags.enableDebugInfo(/*enable=*/true, /*prettyForm=*/false); - if (!Verify) - Flags.assumeVerified(); - MlirMod->print(*OutputStream, Flags); + // FIXME: Think of a better error handling mechanism + (void)writeModuleBytecode(MlirMod, *OutputStream); }; switch (Action) { @@ -382,15 +347,15 @@ class CIRGenConsumer : public clang::ASTConsumer { !FeOptions.ClangIRCallConvLowering, DisableDebugInfo); LlvmModule->setTargetTriple(llvm::Triple(CI.getTargetOpts().Triple)); - LlvmModule->setDataLayout(C.getTargetInfo().getDataLayoutString()); + LlvmModule->setDataLayout(CI.getTarget().getDataLayoutString()); LinkInModules(*LlvmModule); BackendAction BackendAction = getBackendActionFromOutputType(Action); - emitBackendOutput( - CI, CodeGenOpts, C.getTargetInfo().getDataLayoutString(), - LlvmModule.get(), BackendAction, FS, std::move(OutputStream)); + emitBackendOutput(CI, CodeGenOpts, CI.getTarget().getDataLayoutString(), + LlvmModule.get(), BackendAction, FS, + std::move(OutputStream)); break; } case CIRGenAction::OutputType::None: @@ -398,6 +363,61 @@ class CIRGenConsumer : public clang::ASTConsumer { } } + void HandleTranslationUnit(ASTContext &C) override { + llvm::TimeTraceScope Scope("CIR Gen"); + + // Note that this method is called after `HandleTopLevelDecl` has already + // ran all over the top level decls. Here clang mostly wraps defered and + // global codegen, followed by running CIR passes. + Gen->HandleTranslationUnit(C); + + if (!FeOptions.ClangIRDisableCIRVerifier) + if (!Gen->verifyModule()) { + llvm::report_fatal_error( + "CIR codegen: module verification error before running CIR passes"); + return; + } + + auto MlirMod = Gen->getModule(); + auto MlirCtx = Gen->takeContext(); + + if (!FeOptions.ClangIRDisablePasses) { + // Handle source manager properly given that lifetime analysis + // might emit warnings and remarks. + auto &ClangSourceMgr = C.getSourceManager(); + FileID MainFileID = ClangSourceMgr.getMainFileID(); + + std::unique_ptr FileBuf = + llvm::MemoryBuffer::getMemBuffer( + ClangSourceMgr.getBufferOrFake(MainFileID)); + + llvm::SourceMgr MlirSourceMgr; + MlirSourceMgr.AddNewSourceBuffer(std::move(FileBuf), llvm::SMLoc()); + + if (FeOptions.ClangIRVerifyDiags) { + mlir::SourceMgrDiagnosticVerifierHandler SourceMgrHandler( + MlirSourceMgr, MlirCtx.get()); + MlirCtx->printOpOnDiagnostic(false); + SetupCirPipelineAndExecute(MlirMod, *MlirCtx, C); + + // Verify the diagnostic handler to make sure that each of the + // diagnostics matched. + if (SourceMgrHandler.verify().failed()) { + // FIXME: we fail ungracefully, there's probably a better way + // to communicate non-zero return so tests can actually fail. + llvm::sys::RunInterruptHandlers(); + exit(1); + } + } else { + mlir::SourceMgrDiagnosticHandler SourceMgrHandler(MlirSourceMgr, + MlirCtx.get()); + SetupCirPipelineAndExecute(MlirMod, *MlirCtx, C); + } + } + + GenerateOutput(MlirMod, std::move(MlirCtx)); + } + void LoadLinkModules(llvm::LLVMContext &LlvmCtx) { for (const CodeGenOptions::BitcodeFileToLink &F : CI.getCodeGenOpts().LinkBitcodeFiles) { @@ -526,12 +546,15 @@ CIRGenAction::CreateASTConsumer(CompilerInstance &Ci, StringRef InputFile) { return std::move(Result); } -mlir::OwningOpRef -CIRGenAction::loadModule(llvm::MemoryBufferRef MbRef) { - auto Module = - mlir::parseSourceString(MbRef.getBuffer(), mlirContext); - assert(Module && "Failed to parse ClangIR module"); - return Module; +static mlir::FailureOr> +loadModule(std::unique_ptr buf, + mlir::MLIRContext &mlirContext) { + llvm::SourceMgr sm; + sm.AddNewSourceBuffer(std::move(buf), llvm::SMLoc()); + auto module = mlir::parseSourceFile(sm, &mlirContext); + if (!module) + return mlir::failure(); + return module; } void CIRGenAction::ExecuteAction() { @@ -540,44 +563,66 @@ void CIRGenAction::ExecuteAction() { return; } - // If this is a CIR file we have to treat it specially. - // TODO: This could be done more logically. This is just modeled at the moment - // mimicing CodeGenAction but this is clearly suboptimal. auto &Ci = getCompilerInstance(); - std::unique_ptr Outstream = - getOutputStream(Ci, getCurrentFile(), action); - if (action != OutputType::None && !Outstream) - return; + auto &Diags = Ci.getDiagnostics(); + const clang::FrontendOptions &Fo = Ci.getFrontendOpts(); + + if (Fo.Inputs.size() > 1) + llvm_unreachable("NYI: Missing support of 'linking CIR files'"); + const FrontendInputFile &Input = Fo.Inputs.front(); + StringRef InputFile = Input.getFile(); + InputKind Kind = Input.getKind(); + assert(Kind.getFormat() == InputKind::Source && + "Loading CIR files only support source code formats"); + auto Out = Ci.takeOutputStream(); auto &SourceManager = Ci.getSourceManager(); auto FileId = SourceManager.getMainFileID(); - auto MainFile = SourceManager.getBufferOrNone(FileId); - if (!MainFile) + if (!Out) + Out = getOutputStream(Ci, InputFile, action); + + auto Result = std::make_unique( + action, Ci, Ci.getDiagnostics(), &Ci.getVirtualFileSystem(), + Ci.getHeaderSearchOpts(), Ci.getCodeGenOpts(), Ci.getTargetOpts(), + Ci.getLangOpts(), Ci.getFrontendOpts(), InputFile, std::move(Out)); + cgConsumer = Result.get(); + + std::unique_ptr MlirContext{new mlir::MLIRContext}; + MlirContext->getOrLoadDialect(); + MlirContext->getOrLoadDialect(); + MlirContext->getOrLoadDialect(); + MlirContext->getOrLoadDialect(); + MlirContext->getOrLoadDialect(); + MlirContext->getOrLoadDialect(); + + llvm::ErrorOr> InputOrErr = + Ci.getFileManager().getBufferForFile(InputFile); + if (!InputOrErr) { + std::error_code Ec = InputOrErr.getError(); + Diags.Report(clang::diag::err_fe_error_reading) << InputFile; + Diags.Report(clang::diag::note_drv_command_failed_diag_msg) << Ec.message(); return; + } + std::unique_ptr InputBuf = std::move(*InputOrErr); - mlirContext->getOrLoadDialect(); - mlirContext->getOrLoadDialect(); - mlirContext->getOrLoadDialect(); + auto MlirModuleOr = loadModule(std::move(InputBuf), *MlirContext); - // TODO: unwrap this -- this exists because including the `OwningModuleRef` in - // CIRGenAction's header would require linking the Frontend against MLIR. - // Let's avoid that for now. - auto MlirModule = loadModule(*MainFile); - if (!MlirModule) + if (mlir::failed(MlirModuleOr)) { + Diags.Report(clang::diag::err_fe_error_reading) + << "failed to parse CIR module" << InputFile; return; + } + + // FIXME: This introduces a leak. The ownership model of "GenerateOutput" is + // puzzling. We give ownership of MLIRContext, but not of the MLIRModule. + // If the lifetime of the ModuleOp exceeds the lifetime of the context there + // are crashes. + // I need to check this with CIR team. + mlir::ModuleOp MlirModule = std::move(*MlirModuleOr).release(); - // FIXME(cir): This compilation path does not account for some flags. - llvm::LLVMContext LlvmCtx; - bool DisableDebugInfo = - Ci.getCodeGenOpts().getDebugInfo() == llvm::codegenoptions::NoDebugInfo; - auto LlvmModule = lowerFromCIRToLLVMIR( - Ci.getFrontendOpts(), MlirModule.release(), - std::unique_ptr(mlirContext), LlvmCtx, - /*disableVerifier=*/false, /*disableCCLowering=*/true, DisableDebugInfo); - - if (Outstream) - LlvmModule->print(*Outstream, nullptr); + assert(MlirModule && "Could not load module"); + cgConsumer->GenerateOutput(std::move(MlirModule), std::move(MlirContext)); } namespace cir { diff --git a/clang/lib/CIR/FrontendAction/CMakeLists.txt b/clang/lib/CIR/FrontendAction/CMakeLists.txt index 686754cf087a..cc7ece6d04e0 100644 --- a/clang/lib/CIR/FrontendAction/CMakeLists.txt +++ b/clang/lib/CIR/FrontendAction/CMakeLists.txt @@ -8,6 +8,7 @@ get_property(dialect_libs GLOBAL PROPERTY MLIR_DIALECT_LIBS) add_clang_library(clangCIRFrontendAction CIRGenAction.cpp + CIRCombineAction.cpp DEPENDS MLIRCIROpsIncGen @@ -36,4 +37,6 @@ add_clang_library(clangCIRFrontendAction MLIRTransforms MLIRSupport MLIRMemRefDialect + MLIRBytecodeWriter + MLIRBytecodeReader ) diff --git a/clang/lib/Driver/Action.cpp b/clang/lib/Driver/Action.cpp index 72a42a6f957e..ff36827d6e8f 100644 --- a/clang/lib/Driver/Action.cpp +++ b/clang/lib/Driver/Action.cpp @@ -19,25 +19,38 @@ Action::~Action() = default; const char *Action::getClassName(ActionClass AC) { switch (AC) { - case InputClass: return "input"; - case BindArchClass: return "bind-arch"; + case InputClass: + return "input"; + case BindArchClass: + return "bind-arch"; case OffloadClass: return "offload"; - case PreprocessJobClass: return "preprocessor"; - case PrecompileJobClass: return "precompiler"; + case PreprocessJobClass: + return "preprocessor"; + case PrecompileJobClass: + return "precompiler"; case ExtractAPIJobClass: return "api-extractor"; case AnalyzeJobClass: return "analyzer"; - case CompileJobClass: return "compiler"; - case BackendJobClass: return "backend"; - case AssembleJobClass: return "assembler"; - case IfsMergeJobClass: return "interface-stub-merger"; - case LinkJobClass: return "linker"; - case LipoJobClass: return "lipo"; - case DsymutilJobClass: return "dsymutil"; - case VerifyDebugInfoJobClass: return "verify-debug-info"; - case VerifyPCHJobClass: return "verify-pch"; + case CompileJobClass: + return "compiler"; + case BackendJobClass: + return "backend"; + case AssembleJobClass: + return "assembler"; + case IfsMergeJobClass: + return "interface-stub-merger"; + case LinkJobClass: + return "linker"; + case LipoJobClass: + return "lipo"; + case DsymutilJobClass: + return "dsymutil"; + case VerifyDebugInfoJobClass: + return "verify-debug-info"; + case VerifyPCHJobClass: + return "verify-pch"; case OffloadBundlingJobClass: return "clang-offload-bundler"; case OffloadUnbundlingJobClass: @@ -54,6 +67,10 @@ const char *Action::getClassName(ActionClass AC) { return "binary-translator"; case ObjcopyJobClass: return "objcopy"; + case CIRSplitJobClass: + return "splitcir"; + case CIRCombineJobClass: + return "comebinecir"; } llvm_unreachable("invalid class"); @@ -64,17 +81,26 @@ void Action::propagateDeviceOffloadInfo(OffloadKind OKind, const char *OArch, // Offload action set its own kinds on their dependences. if (Kind == OffloadClass) return; + // Unbundling actions use the host kinds. if (Kind == OffloadUnbundlingJobClass) return; - assert((OffloadingDeviceKind == OKind || OffloadingDeviceKind == OFK_None) && + assert((Kind == CIRCombineJobClass || OffloadingDeviceKind == OKind || + OffloadingDeviceKind == OFK_None) && "Setting device kind to a different device??"); assert(!ActiveOffloadKindMask && "Setting a device kind in a host action??"); OffloadingDeviceKind = OKind; OffloadingArch = OArch; OffloadingToolChain = OToolChain; + if (Kind == CIRCombineJobClass) { + auto *CIRCombineAction = dyn_cast(this); + CIRCombineAction->getDeviceAction()->propagateDeviceOffloadInfo( + OffloadingDeviceKind, OArch, OToolChain); + return; + } + for (auto *A : Inputs) A->propagateDeviceOffloadInfo(OffloadingDeviceKind, OArch, OToolChain); } @@ -83,12 +109,18 @@ void Action::propagateHostOffloadInfo(unsigned OKinds, const char *OArch) { // Offload action set its own kinds on their dependences. if (Kind == OffloadClass) return; - - assert(OffloadingDeviceKind == OFK_None && + assert((Kind == CIRCombineJobClass || OffloadingDeviceKind == OFK_None) && "Setting a host kind in a device action."); ActiveOffloadKindMask |= OKinds; OffloadingArch = OArch; + if (Kind == CIRCombineJobClass) { + auto *CIRCombineAction = dyn_cast(this); + CIRCombineAction->getHostAction()->propagateHostOffloadInfo( + ActiveOffloadKindMask, OArch); + return; + } + for (auto *A : Inputs) A->propagateHostOffloadInfo(ActiveOffloadKindMask, OArch); } @@ -144,10 +176,9 @@ std::string Action::getOffloadingKindPrefix() const { /// Return a string that can be used as prefix in order to generate unique files /// for each offloading kind. -std::string -Action::GetOffloadingFileNamePrefix(OffloadKind Kind, - StringRef NormalizedTriple, - bool CreatePrefixForHost) { +std::string Action::GetOffloadingFileNamePrefix(OffloadKind Kind, + StringRef NormalizedTriple, + bool CreatePrefixForHost) { // Don't generate prefix for host actions unless required. if (!CreatePrefixForHost && (Kind == OFK_None || Kind == OFK_Host)) return {}; @@ -217,8 +248,9 @@ OffloadAction::OffloadAction(const DeviceDependences &DDeps, types::ID Ty) OffloadingArch = BArchs.front(); // Propagate info to the dependencies. - for (unsigned i = 0, e = getInputs().size(); i != e; ++i) + for (unsigned i = 0, e = getInputs().size(); i != e; ++i) { getInputs()[i]->propagateDeviceOffloadInfo(OKinds[i], BArchs[i], OTCs[i]); + } } OffloadAction::OffloadAction(const HostDependence &HDep, @@ -364,7 +396,7 @@ PrecompileJobAction::PrecompileJobAction(Action *Input, types::ID OutputType) PrecompileJobAction::PrecompileJobAction(ActionClass Kind, Action *Input, types::ID OutputType) : JobAction(Kind, Input, OutputType) { - assert(isa((Action*)this) && "invalid action kind"); + assert(isa((Action *)this) && "invalid action kind"); } void ExtractAPIJobAction::anchor() {} @@ -448,6 +480,36 @@ OffloadPackagerJobAction::OffloadPackagerJobAction(ActionList &Inputs, types::ID Type) : JobAction(OffloadPackagerJobClass, Inputs, Type) {} +void CombineCIRJobAction::anchor() {} + +CombineCIRJobAction::CombineCIRJobAction( + const ToolChain *HostToolChain, const ToolChain *DeviceToolChain, + Action *HostAction, Action *DeviceAction, char *HostBoundArch, + const char *DeviceBoundArch, unsigned HostOffloadKind, types::ID Type, + OffloadKind OffloadDeviceKind) + : JobAction(CIRCombineJobClass, {HostAction, DeviceAction}, Type), + HostToolChain(HostToolChain), DeviceToolChain(DeviceToolChain), + HostAction(HostAction), DeviceAction(DeviceAction), + HostBoundArch(HostBoundArch), DeviceBoundArch(DeviceBoundArch), + HostOffloadKind(HostOffloadKind) { + + OffloadingDeviceKind = OffloadDeviceKind; + ActiveOffloadKindMask = 0; + + // Propagate info to the dependencies. + // NOTE: THIS IS LIKELY THE LAST STEP OF MAKING THE -print-passes-work + // properly. I need to extent the constructors and get access to the + // toolchains +} + +void SplitCIRJobAction::anchor() {} + +SplitCIRJobAction::SplitCIRJobAction(Action *Input, bool isHost, types::ID Type, + OffloadKind Kind) + : JobAction(CIRSplitJobClass, Input, Type), isHost(isHost) { + OffloadingDeviceKind = Kind; +} + void LinkerWrapperJobAction::anchor() {} LinkerWrapperJobAction::LinkerWrapperJobAction(ActionList &Inputs, diff --git a/clang/lib/Driver/Driver.cpp b/clang/lib/Driver/Driver.cpp index cfc365ed0cba..da7bdd8ae5c1 100644 --- a/clang/lib/Driver/Driver.cpp +++ b/clang/lib/Driver/Driver.cpp @@ -364,8 +364,7 @@ phases::ID Driver::getFinalPhase(const DerivedArgList &DAL, if (CCCIsCPP() || (PhaseArg = DAL.getLastArg(options::OPT_E)) || (PhaseArg = DAL.getLastArg(options::OPT__SLASH_EP)) || (PhaseArg = DAL.getLastArg(options::OPT_M, options::OPT_MM)) || - (PhaseArg = DAL.getLastArg(options::OPT__SLASH_P)) || - CCGenDiagnostics) { + (PhaseArg = DAL.getLastArg(options::OPT__SLASH_P)) || CCGenDiagnostics) { FinalPhase = phases::Preprocess; // --precompile only runs up to precompilation. @@ -391,18 +390,18 @@ phases::ID Driver::getFinalPhase(const DerivedArgList &DAL, (PhaseArg = DAL.getLastArg(options::OPT_emit_ast))) { FinalPhase = phases::Compile; - // -S only runs up to the backend. + // -S only runs up to the backend. } else if ((PhaseArg = DAL.getLastArg(options::OPT_S))) { FinalPhase = phases::Backend; - // -c compilation only runs up to the assembler. + // -c compilation only runs up to the assembler. } else if ((PhaseArg = DAL.getLastArg(options::OPT_c))) { FinalPhase = phases::Assemble; } else if ((PhaseArg = DAL.getLastArg(options::OPT_emit_interface_stubs))) { FinalPhase = phases::IfsMerge; - // Otherwise do everything. + // Otherwise do everything. } else FinalPhase = phases::Link; @@ -639,8 +638,7 @@ static void setZosTargetVersion(const Driver &D, llvm::Triple &Target, /// /// This routine provides the logic to compute a target triple from various /// args passed to the driver and the default triple string. -static llvm::Triple computeTargetTriple(const Driver &D, - StringRef TargetTriple, +static llvm::Triple computeTargetTriple(const Driver &D, StringRef TargetTriple, const ArgList &Args, StringRef DarwinArchName = "") { // FIXME: Already done in Compilation *Driver::BuildCompilation @@ -774,8 +772,8 @@ static llvm::Triple computeTargetTriple(const Driver &D, // Handle -miamcu flag. if (Args.hasFlag(options::OPT_miamcu, options::OPT_mno_iamcu, false)) { if (Target.get32BitArchVariant().getArch() != llvm::Triple::x86) - D.Diag(diag::err_drv_unsupported_opt_for_target) << "-miamcu" - << Target.str(); + D.Diag(diag::err_drv_unsupported_opt_for_target) + << "-miamcu" << Target.str(); if (A && !A->getOption().matches(options::OPT_m32)) D.Diag(diag::err_drv_argument_not_allowed_with) @@ -1687,14 +1685,13 @@ Compilation *Driver::BuildCompilation(ArrayRef ArgList) { if (Arg *A = Args.getLastArg(options::OPT_fembed_bitcode_EQ)) { StringRef Name = A->getValue(); unsigned Model = llvm::StringSwitch(Name) - .Case("off", EmbedNone) - .Case("all", EmbedBitcode) - .Case("bitcode", EmbedBitcode) - .Case("marker", EmbedMarker) - .Default(~0U); + .Case("off", EmbedNone) + .Case("all", EmbedBitcode) + .Case("bitcode", EmbedBitcode) + .Case("marker", EmbedMarker) + .Default(~0U); if (Model == ~0U) { - Diags.Report(diag::err_drv_invalid_value) << A->getAsString(Args) - << Name; + Diags.Report(diag::err_drv_invalid_value) << A->getAsString(Args) << Name; } else BitcodeEmbed = static_cast(Model); } @@ -1923,7 +1920,7 @@ bool Driver::getCrashDiagnosticFile(StringRef ReproCrashFilename, size_t LineEnd = Data.find_first_of("\n", ParentProcPos); if (LineEnd == StringRef::npos) continue; - StringRef ParentProcess = Data.slice(ParentProcPos+15, LineEnd).trim(); + StringRef ParentProcess = Data.slice(ParentProcPos + 15, LineEnd).trim(); int OpenBracket = -1, CloseBracket = -1; for (size_t i = 0, e = ParentProcess.size(); i < e; ++i) { if (ParentProcess[i] == '[') @@ -1936,7 +1933,8 @@ bool Driver::getCrashDiagnosticFile(StringRef ReproCrashFilename, int CrashPID; if (OpenBracket < 0 || CloseBracket < 0 || ParentProcess.slice(OpenBracket + 1, CloseBracket) - .getAsInteger(10, CrashPID) || CrashPID != PID) { + .getAsInteger(10, CrashPID) || + CrashPID != PID) { continue; } @@ -2192,8 +2190,7 @@ void Driver::generateCompilationDiagnostics( CrashDiagDir += "__.crash"; Diag(clang::diag::note_drv_command_failed_diag_msg) << "Crash backtrace is located in"; - Diag(clang::diag::note_drv_command_failed_diag_msg) - << CrashDiagDir.str(); + Diag(clang::diag::note_drv_command_failed_diag_msg) << CrashDiagDir.str(); Diag(clang::diag::note_drv_command_failed_diag_msg) << "(choose the .crash file that corresponds to your crash)"; } @@ -2305,8 +2302,7 @@ void Driver::PrintHelp(bool ShowHidden) const { std::string Usage = llvm::formatv("{0} [options] file...", Name).str(); getOpts().printHelp(llvm::outs(), Usage.c_str(), DriverTitle.c_str(), - ShowHidden, /*ShowAllAliases=*/false, - VisibilityMask); + ShowHidden, /*ShowAllAliases=*/false, VisibilityMask); } void Driver::PrintVersion(const Compilation &C, raw_ostream &OS) const { @@ -2483,11 +2479,11 @@ bool Driver::HandleImmediateArgs(Compilation &C) { if (C.getArgs().hasArg(options::OPT_v)) { if (!SystemConfigDir.empty()) - llvm::errs() << "System configuration file directory: " - << SystemConfigDir << "\n"; + llvm::errs() << "System configuration file directory: " << SystemConfigDir + << "\n"; if (!UserConfigDir.empty()) - llvm::errs() << "User configuration file directory: " - << UserConfigDir << "\n"; + llvm::errs() << "User configuration file directory: " << UserConfigDir + << "\n"; } const ToolChain &TC = C.getDefaultToolChain(); @@ -2570,7 +2566,7 @@ bool Driver::HandleImmediateArgs(Compilation &C) { StringRef ProgName = A->getValue(); // Null program name cannot have a path. - if (! ProgName.empty()) + if (!ProgName.empty()) llvm::outs() << GetProgramPath(ProgName, TC); llvm::outs() << "\n"; @@ -2709,6 +2705,35 @@ static unsigned PrintActions1(const Compilation &C, Action *A, IsFirst = false; SibKind = OtherSibAction; }); + } else if (CombineCIRJobAction *COA = dyn_cast(A)) { + bool IsFirst = true; + auto printAction = [&](Action *A, const ToolChain *TC, + const char *BoundArch) { + assert(TC && "Unknown host toolchain"); + // E.g. for two CUDA device dependences whose bound arch is sm_20 and + // sm_35 this will generate: + // "cuda-device" (nvptx64-nvidia-cuda:sm_20) {#ID}, "cuda-device" + // (nvptx64-nvidia-cuda:sm_35) {#ID} + if (!IsFirst) + os << ", "; + os << '"'; + os << A->getOffloadingKindPrefix(); + os << " ("; + os << TC->getTriple().normalize(); + if (BoundArch) + os << ":" << BoundArch; + os << ")"; + os << '"'; + os << " {" << PrintActions1(C, A, Ids, SibIndent, SibKind) << "}"; + IsFirst = false; + SibKind = OtherSibAction; + }; + printAction(COA->getHostAction(), COA->getHostToolChain(), + COA->getHostBoundArch()); + + printAction(COA->getDeviceAction(), COA->getOffloadingToolChain(), + COA->getDeviceBoundArch()); + } else { const ActionList *AL = &A->getInputs(); @@ -2728,7 +2753,7 @@ static unsigned PrintActions1(const Compilation &C, Action *A, // itself (e.g. (cuda-device, sm_20) or (cuda-host)). std::string offload_str; llvm::raw_string_ostream offload_os(offload_str); - if (!isa(A)) { + if (!isa(A) && !isa(A)) { auto S = A->getOffloadingKindPrefix(); if (!S.empty()) { offload_os << ", (" << S; @@ -2804,7 +2829,7 @@ void Driver::BuildUniversalActions(Compilation &C, const ToolChain &TC, // Add in arch bindings for every top level action, as well as lipo and // dsymutil steps if needed. - for (Action* Act : SingleActions) { + for (Action *Act : SingleActions) { // Make sure we can lipo this kind of output. If not (and it is an actual // output) then we disallow, since we can't create an output file with the // right name without overwriting it. We could remove this oddity by just @@ -3027,7 +3052,8 @@ void Driver::BuildInputs(const ToolChain &TC, DerivedArgList &Args, Ty = TC.LookupTypeForExtension(Ext + 1); if (Ty == types::TY_INVALID) { - if (IsCLMode() && (Args.hasArgNoClaim(options::OPT_E) || CCGenDiagnostics)) + if (IsCLMode() && + (Args.hasArgNoClaim(options::OPT_E) || CCGenDiagnostics)) Ty = types::TY_CXX; else if (CCCIsCPP() || CCGenDiagnostics) Ty = types::TY_C; @@ -3233,7 +3259,7 @@ class OffloadingActionBuilder final { virtual void appendLinkDeviceActions(ActionList &AL) {} /// Append linker host action generated by the builder. - virtual Action* appendLinkHostActions(ActionList &AL) { return nullptr; } + virtual Action *appendLinkHostActions(ActionList &AL) { return nullptr; } /// Append linker actions generated by the builder. virtual void appendLinkDependences(OffloadAction::DeviceDependences &DA) {} @@ -3253,6 +3279,15 @@ class OffloadingActionBuilder final { Action::OffloadKind getAssociatedOffloadKind() { return AssociatedOffloadKind; } + + // Return true if this builder can support combining host and device modules + // into a single CIR module and co-optimize them. + virtual bool hasCIRCombineSupport() { return false; }; + + virtual void addCIRCombineSplitActions(const ToolChain *HostToolChainconst, + Action *&HostAction, + char *HostBoundArch, + unsigned HostOffloadKind) {}; }; /// Base class for CUDA/HIP action builder. It injects device code in @@ -3466,6 +3501,38 @@ class OffloadingActionBuilder final { return false; } + + // Return true if this builder can support combining host and device modules + // into a single CIR module and co-optimize them. + bool hasCIRCombineSupport() override { return true; }; + + void addCIRCombineSplitActions(const ToolChain *HostToolChain, + Action *&HostAction, char *HostBoundArch, + unsigned HostOffloadKind) override { + // This assumes that there are multiple actions. one per architecture + // based on my understanding of the driver her. I am pretty sure this is + // not needed for CIR combination. + // TODO: Revisit this decision here once I have a working prototype for a + // single arch. + + for (unsigned I = 0, E = GpuArchList.size(); I != E; ++I) { + const char *GPUArch = GpuArchList[I].ID; + for (Action *&A : CudaDeviceActions) { + if (A->getType() != types::TY_CIR) + continue; + Action *CirCombineAction = C.MakeAction( + HostToolChain, ToolChains.front(), HostAction, A, HostBoundArch, + GPUArch, HostOffloadKind, types::TY_CIR, + A->getOffloadingDeviceKind()); + A = C.MakeAction(CirCombineAction, false, + types::TY_CIR, + A->getOffloadingDeviceKind()); + HostAction = C.MakeAction( + CirCombineAction, true, types::TY_CIR, + Action::OffloadKind::OFK_None); + } + } + } }; /// \brief CUDA action builder. It injects device code in the host backend @@ -3765,9 +3832,10 @@ class OffloadingActionBuilder final { } // By default, we produce an action for each device arch. - for (Action *&A : CudaDeviceActions) + for (Action *&A : CudaDeviceActions) { A = C.getDriver().ConstructPhaseAction(C, Args, CurPhase, A, AssociatedOffloadKind); + } if (CompileDeviceOnly && CurPhase == FinalPhase && BundleOutput && *BundleOutput) { @@ -3804,15 +3872,15 @@ class OffloadingActionBuilder final { for (auto &LI : DeviceLinkerInputs) { types::ID Output = Args.hasArg(options::OPT_emit_llvm) - ? types::TY_LLVM_BC - : types::TY_Image; + ? types::TY_LLVM_BC + : types::TY_Image; auto *DeviceLinkAction = C.MakeAction(LI, Output); // Linking all inputs for the current GPU arch. // LI contains all the inputs for the linker. OffloadAction::DeviceDependences DeviceLinkDeps; - DeviceLinkDeps.add(*DeviceLinkAction, *ToolChains[0], - GpuArchList[I], AssociatedOffloadKind); + DeviceLinkDeps.add(*DeviceLinkAction, *ToolChains[0], GpuArchList[I], + AssociatedOffloadKind); Actions.push_back(C.MakeAction( DeviceLinkDeps, DeviceLinkAction->getType())); ++I; @@ -3821,8 +3889,8 @@ class OffloadingActionBuilder final { // If emitting LLVM, do not generate final host/device compilation action if (Args.hasArg(options::OPT_emit_llvm)) { - AL.append(Actions); - return; + AL.append(Actions); + return; } // Create a host object from all the device images by embedding them @@ -3843,7 +3911,7 @@ class OffloadingActionBuilder final { } } - Action* appendLinkHostActions(ActionList &AL) override { return AL.back(); } + Action *appendLinkHostActions(ActionList &AL) override { return AL.back(); } void appendLinkDependences(OffloadAction::DeviceDependences &DA) override {} }; @@ -3922,7 +3990,8 @@ class OffloadingActionBuilder final { Action * addDeviceDependencesToHostAction(Action *HostAction, const Arg *InputArg, phases::ID CurPhase, phases::ID FinalPhase, - DeviceActionBuilder::PhasesTy &Phases) { + DeviceActionBuilder::PhasesTy &Phases, + bool combineCIR = false) { if (!IsValid) return nullptr; @@ -3932,6 +4001,42 @@ class OffloadingActionBuilder final { assert(HostAction && "Invalid host action!"); recordHostAction(HostAction, InputArg); + if (combineCIR && CurPhase == phases::Backend) { + // TODO: I ignore the multiple device architectures at this point. + // In principle a single CIR file should be able to represent all of them. + // However backend action may need to invoke toolchains in some other way. + // I need to investigate this and act accordingly here. + unsigned ActiveOffloadKinds = 0u; + for (auto &I : InputArgToOffloadKindMap) + ActiveOffloadKinds |= I.second; + + for (auto *SB : SpecializedBuilders) { + if (!SB->isValid()) + continue; + if (!SB->hasCIRCombineSupport()) + continue; + + HostAction->setHostOffloadInfo(ActiveOffloadKinds, + /*BoundArch=*/nullptr); + for (auto *A : HostAction->inputs()) { + auto ArgLoc = HostActionToInputArgMap.find(A); + if (ArgLoc == HostActionToInputArgMap.end()) + continue; + auto OFKLoc = InputArgToOffloadKindMap.find(ArgLoc->second); + if (OFKLoc == InputArgToOffloadKindMap.end()) + continue; + A->propagateHostOffloadInfo(OFKLoc->second, /*BoundArch=*/nullptr); + } + + if (HostAction->getType() != types::TY_CIR) + continue; + + SB->addCIRCombineSplitActions( + C.getSingleOffloadToolChain(), HostAction, + nullptr, ActiveOffloadKinds); + } + } + OffloadAction::DeviceDependences DDeps; // Check if all the programming models agree we should not emit the host // action. Also, keep track of the offloading kinds employed. @@ -3947,19 +4052,19 @@ class OffloadingActionBuilder final { SB->getDeviceDependences(DDeps, CurPhase, FinalPhase, Phases); // If the builder explicitly says the host action should be ignored, - // we need to increment the variable that tracks the builders that request - // the host object to be ignored. + // we need to increment the variable that tracks the builders that + // request the host object to be ignored. if (RetCode == DeviceActionBuilder::ABRT_Ignore_Host) ++IgnoringBuilders; - // Unless the builder was inactive for this action, we have to record the - // offload kind because the host will have to use it. + // Unless the builder was inactive for this action, we have to record + // the offload kind because the host will have to use it. if (RetCode != DeviceActionBuilder::ABRT_Inactive) OffloadKind |= SB->getAssociatedOffloadKind(); } - // If all builders agree that the host object should be ignored, just return - // nullptr. + // If all builders agree that the host object should be ignored, just + // return nullptr. if (IgnoringBuilders && SpecializedBuilders.size() == (InactiveBuilders + IgnoringBuilders)) return nullptr; @@ -3967,8 +4072,8 @@ class OffloadingActionBuilder final { if (DDeps.getActions().empty()) return HostAction; - // We have dependences we need to bundle together. We use an offload action - // for that. + // We have dependences we need to bundle together. We use an offload + // action for that. OffloadAction::HostDependence HDep( *HostAction, *C.getSingleOffloadToolChain(), /*BoundArch=*/nullptr, DDeps); @@ -4019,8 +4124,8 @@ class OffloadingActionBuilder final { assert(RetCode != DeviceActionBuilder::ABRT_Ignore_Host && "Host dependence not expected to be ignored.!"); - // Unless the builder was inactive for this action, we have to record the - // offload kind because the host will have to use it. + // Unless the builder was inactive for this action, we have to record + // the offload kind because the host will have to use it. if (RetCode != DeviceActionBuilder::ABRT_Inactive) OffloadKind |= SB->getAssociatedOffloadKind(); } @@ -4055,7 +4160,8 @@ class OffloadingActionBuilder final { // therefore only do this when HostAction is not a null pointer. if (CanUseBundler && ShouldUseBundler && HostAction && HostAction->getType() != types::TY_Nothing && !OffloadAL.empty()) { - // Add the host action to the list in order to create the bundling action. + // Add the host action to the list in order to create the bundling + // action. OffloadAL.push_back(HostAction); // We expect that the host action was just appended to the action list @@ -4091,7 +4197,7 @@ class OffloadingActionBuilder final { return nullptr; // Let builders add host linking actions. - Action* HA = nullptr; + Action *HA = nullptr; for (DeviceActionBuilder *SB : SpecializedBuilders) { if (!SB->isValid()) continue; @@ -4105,10 +4211,10 @@ class OffloadingActionBuilder final { return HA; } - /// Processes the host linker action. This currently consists of replacing it - /// with an offload action if there are device link objects and propagate to - /// the host action all the offload kinds used in the current compilation. The - /// resulting action is returned. + /// Processes the host linker action. This currently consists of replacing + /// it with an offload action if there are device link objects and propagate + /// to the host action all the offload kinds used in the current + /// compilation. The resulting action is returned. Action *processHostLinkAction(Action *HostAction) { // Add all the dependences from the device linking actions. OffloadAction::DeviceDependences DDeps; @@ -4124,8 +4230,8 @@ class OffloadingActionBuilder final { for (auto &I : InputArgToOffloadKindMap) ActiveOffloadKinds |= I.second; - // If we don't have device dependencies, we don't have to create an offload - // action. + // If we don't have device dependencies, we don't have to create an + // offload action. if (DDeps.getActions().empty()) { // Set all the active offloading kinds to the link action. Given that it // is a link action it is assumed to depend on all actions generated so @@ -4147,8 +4253,8 @@ class OffloadingActionBuilder final { } // Create the offload action with all dependences. When an offload action - // is created the kinds are propagated to the host action, so we don't have - // to do that explicitly here. + // is created the kinds are propagated to the host action, so we don't + // have to do that explicitly here. OffloadAction::HostDependence HDep( *HostAction, *C.getSingleOffloadToolChain(), /*BoundArch*/ nullptr, ActiveOffloadKinds); @@ -4224,7 +4330,8 @@ void Driver::handleArguments(Compilation &C, DerivedArgList &Args, getOpts().getOption(options::OPT_frtlib_add_rpath)); } // Emitting LLVM while linking disabled except in HIPAMD Toolchain - if (Args.hasArg(options::OPT_emit_llvm) && !Args.hasArg(options::OPT_hip_link)) + if (Args.hasArg(options::OPT_emit_llvm) && + !Args.hasArg(options::OPT_hip_link)) Diag(clang::diag::err_drv_emit_llvm_link); if (C.getDefaultToolChain().getTriple().isWindowsMSVCEnvironment() && LTOMode != LTOK_None && @@ -4455,11 +4562,13 @@ void Driver::BuildDefaultActions(Compilation &C, DerivedArgList &Args, break; for (phases::ID Phase : PL) { - // Add any offload action the host action depends on. - if (!UseNewOffloadingDriver) + if (!UseNewOffloadingDriver) { Current = OffloadBuilder->addDeviceDependencesToHostAction( - Current, InputArg, Phase, PL.back(), FullPL); + Current, InputArg, Phase, PL.back(), FullPL, + (Args.hasArg(options::OPT_fclangir_offload) && + Args.hasArg(options::OPT_fclangir))); + } if (!Current) break; @@ -4510,8 +4619,9 @@ void Driver::BuildDefaultActions(Compilation &C, DerivedArgList &Args, // Try to build the offloading actions and add the result as a dependency // to the host. - if (UseNewOffloadingDriver) + if (UseNewOffloadingDriver) { Current = BuildOffloadingActions(C, Args, I, CUID, Current); + } // Use the current host action in any of the offloading actions, if // required. else if (OffloadBuilder->addHostDependenceToDeviceActions(Current, @@ -5194,6 +5304,11 @@ Action *Driver::ConstructPhaseAction( return C.MakeAction(Input, types::TY_Nothing); if (Args.hasArg(options::OPT_extract_api)) return C.MakeAction(Input, types::TY_API_INFO); + if (Args.hasArg(options::OPT_fclangir_offload)) { + auto *Act = C.MakeAction(Input, types::TY_CIR); + Act->setCannotBeCollapsedWithNextDependentAction(); + return Act; + } return C.MakeAction(Input, types::TY_LLVM_BC); } case phases::Backend: { @@ -5495,6 +5610,10 @@ class ToolSelector final { return nullptr; Action *CurAction = *Inputs.begin(); + llvm::errs() << "Can be collapsed " << CanBeCollapsed + << " is action collapsible:" + << CurAction->isCollapsingWithNextDependentActionLegal() + << " Action type id:" << CurAction->getType() << "\n"; if (CanBeCollapsed && !CurAction->isCollapsingWithNextDependentActionLegal()) return nullptr; @@ -5702,8 +5821,8 @@ class ToolSelector final { continue; } - // This is legal to combine. Append any offload action we found and add the - // current input to preprocessor inputs. + // This is legal to combine. Append any offload action we found and add + // the current input to preprocessor inputs. CollapsedOffloadAction.append(PreprocessJobOffloadActions.begin(), PreprocessJobOffloadActions.end()); NewInputs.append(PJ->input_begin(), PJ->input_end()); @@ -5726,8 +5845,7 @@ class ToolSelector final { /// connected to collapsed actions are updated accordingly. The latter enables /// the caller of the selector to process them afterwards instead of just /// dropping them. If no suitable tool is found, null will be returned. - const Tool *getTool(ActionList &Inputs, - ActionList &CollapsedOffloadAction) { + const Tool *getTool(ActionList &Inputs, ActionList &CollapsedOffloadAction) { // // Get the largest chain of actions that we could combine. // @@ -5770,7 +5888,7 @@ class ToolSelector final { return T; } }; -} +} // namespace /// Return a string that uniquely identifies the result of a job. The bound arch /// is not necessarily represented in the toolchain's triple -- for example, @@ -5941,9 +6059,9 @@ InputInfoList Driver::BuildJobsForActionNoCache( StringRef ArchName = BAA->getArchName(); if (!ArchName.empty()) - TC = &getToolChain(C.getArgs(), - computeTargetTriple(*this, TargetTriple, - C.getArgs(), ArchName)); + TC = &getToolChain( + C.getArgs(), + computeTargetTriple(*this, TargetTriple, C.getArgs(), ArchName)); else TC = &C.getDefaultToolChain(); @@ -5952,7 +6070,6 @@ InputInfoList Driver::BuildJobsForActionNoCache( TargetDeviceOffloadKind); } - ActionList Inputs = A->getInputs(); const JobAction *JA = cast(A); @@ -6084,12 +6201,14 @@ InputInfoList Driver::BuildJobsForActionNoCache( std::string OffloadingPrefix = Action::GetOffloadingFileNamePrefix( A->getOffloadingDeviceKind(), EffectiveTriple.normalize(), /*CreatePrefixForHost=*/isa(A) || + /*CreatePrefixForHost=*/isa(A) || !(A->getOffloadingHostActiveKinds() == Action::OFK_None || AtTopLevel)); - Result = InputInfo(A, GetNamedOutputPath(C, *JA, BaseInput, BoundArch, - AtTopLevel, MultipleArchs, - OffloadingPrefix), - BaseInput); + Result = + InputInfo(A, + GetNamedOutputPath(C, *JA, BaseInput, BoundArch, AtTopLevel, + MultipleArchs, OffloadingPrefix), + BaseInput); if (T->canEmitIR() && OffloadingPrefix.empty()) handleTimeTrace(C, Args, JA, BaseInput, Result); } @@ -6434,12 +6553,10 @@ const char *Driver::GetNamedOutputPath(Compilation &C, const JobAction &JA, } } else if (JA.getType() == types::TY_PCH && IsCLMode()) { NamedOutput = C.getArgs().MakeArgString(GetClPchPath(C, BaseName)); - } else if ((JA.getType() == types::TY_Plist || JA.getType() == types::TY_AST) && + } else if ((JA.getType() == types::TY_Plist || + JA.getType() == types::TY_AST) && C.getArgs().hasArg(options::OPT__SLASH_o)) { - StringRef Val = - C.getArgs() - .getLastArg(options::OPT__SLASH_o) - ->getValue(); + StringRef Val = C.getArgs().getLastArg(options::OPT__SLASH_o)->getValue(); NamedOutput = MakeCLOutputFilename(C.getArgs(), Val, BaseName, types::TY_Object); } else { @@ -6852,15 +6969,15 @@ const ToolChain &Driver::getToolChain(const ArgList &Args, case llvm::Triple::Linux: case llvm::Triple::ELFIAMCU: if (Target.getArch() == llvm::Triple::hexagon) - TC = std::make_unique(*this, Target, - Args); + TC = + std::make_unique(*this, Target, Args); else if ((Target.getVendor() == llvm::Triple::MipsTechnologies) && !Target.hasEnvironment()) TC = std::make_unique(*this, Target, - Args); + Args); else if (Target.isPPC()) TC = std::make_unique(*this, Target, - Args); + Args); else if (Target.getArch() == llvm::Triple::ve) TC = std::make_unique(*this, Target, Args); else if (Target.isOHOSFamily()) @@ -6921,7 +7038,7 @@ const ToolChain &Driver::getToolChain(const ArgList &Args, break; case llvm::Triple::Itanium: TC = std::make_unique(*this, Target, - Args); + Args); break; case llvm::Triple::MSVC: case llvm::Triple::UnknownEnvironment: @@ -6930,8 +7047,7 @@ const ToolChain &Driver::getToolChain(const ArgList &Args, TC = std::make_unique( *this, Target, Args); else - TC = - std::make_unique(*this, Target, Args); + TC = std::make_unique(*this, Target, Args); break; } break; @@ -6965,8 +7081,8 @@ const ToolChain &Driver::getToolChain(const ArgList &Args, TC = std::make_unique(*this, Target, Args); break; case llvm::Triple::hexagon: - TC = std::make_unique(*this, Target, - Args); + TC = + std::make_unique(*this, Target, Args); break; case llvm::Triple::lanai: TC = std::make_unique(*this, Target, Args); @@ -7122,7 +7238,7 @@ Driver::getOptionVisibilityMask(bool UseDriverMode) const { return llvm::opt::Visibility(options::CLOption); if (IsDXCMode()) return llvm::opt::Visibility(options::DXCOption); - if (IsFlangMode()) { + if (IsFlangMode()) { return llvm::opt::Visibility(options::FlangOption); } return llvm::opt::Visibility(options::ClangOption); diff --git a/clang/lib/Driver/ToolChains/Clang.cpp b/clang/lib/Driver/ToolChains/Clang.cpp index 6cdd423c2274..0a6a4314248b 100644 --- a/clang/lib/Driver/ToolChains/Clang.cpp +++ b/clang/lib/Driver/ToolChains/Clang.cpp @@ -84,8 +84,8 @@ static void CheckCodeGenerationOptions(const Driver &D, const ArgList &Args) { if (Args.hasArg(options::OPT_static)) if (const Arg *A = Args.getLastArg(options::OPT_dynamic, options::OPT_mdynamic_no_pic)) - D.Diag(diag::err_drv_argument_not_allowed_with) << A->getAsString(Args) - << "-static"; + D.Diag(diag::err_drv_argument_not_allowed_with) + << A->getAsString(Args) << "-static"; } /// Apply \a Work on the current tool chain \a RegularToolChain and any other @@ -315,7 +315,7 @@ static void addMacroPrefixMapArg(const Driver &D, const ArgList &Args, /// Add a CC1 and CC1AS option to specify the coverage file path prefix map. static void addCoveragePrefixMapArg(const Driver &D, const ArgList &Args, - ArgStringList &CmdArgs) { + ArgStringList &CmdArgs) { for (const Arg *A : Args.filtered(options::OPT_ffile_prefix_map_EQ, options::OPT_fcoverage_prefix_map_EQ)) { StringRef Map = A->getValue(); @@ -374,13 +374,12 @@ static void addPGOAndCoverageFlags(const ToolChain &TC, Compilation &C, auto *CSPGOGenerateArg = getLastCSProfileGenerateArg(Args); - auto *ProfileGenerateArg = Args.getLastArg( - options::OPT_fprofile_instr_generate, - options::OPT_fprofile_instr_generate_EQ, - options::OPT_fno_profile_instr_generate); - if (ProfileGenerateArg && - ProfileGenerateArg->getOption().matches( - options::OPT_fno_profile_instr_generate)) + auto *ProfileGenerateArg = + Args.getLastArg(options::OPT_fprofile_instr_generate, + options::OPT_fprofile_instr_generate_EQ, + options::OPT_fno_profile_instr_generate); + if (ProfileGenerateArg && ProfileGenerateArg->getOption().matches( + options::OPT_fno_profile_instr_generate)) ProfileGenerateArg = nullptr; if (PGOGenerateArg && ProfileGenerateArg) @@ -1000,8 +999,8 @@ void Clang::AddPreprocessingOptions(Compilation &C, const JobAction &JA, } if (ThroughHeader.empty()) { - CmdArgs.push_back(Args.MakeArgString( - Twine("-pch-through-hdrstop-") + (YcArg ? "create" : "use"))); + CmdArgs.push_back(Args.MakeArgString(Twine("-pch-through-hdrstop-") + + (YcArg ? "create" : "use"))); } else { CmdArgs.push_back( Args.MakeArgString(Twine("-pch-through-header=") + ThroughHeader)); @@ -1040,8 +1039,8 @@ void Clang::AddPreprocessingOptions(Compilation &C, const JobAction &JA, continue; } else { // Ignore the PCH if not first on command line and emit warning. - D.Diag(diag::warn_drv_pch_not_first_include) << P - << A->getAsString(Args); + D.Diag(diag::warn_drv_pch_not_first_include) + << P << A->getAsString(Args); } } } else if (A->getOption().matches(options::OPT_isystem_after)) { @@ -1253,8 +1252,9 @@ static void renderRemarksOptions(const ArgList &Args, ArgStringList &CmdArgs, if (Arg *FinalOutput = Args.getLastArg(options::OPT_o)) F = FinalOutput->getValue(); } else { - if (Format != "yaml" && // For YAML, keep the original behavior. - Triple.isOSDarwin() && // Enable this only on darwin, since it's the only platform supporting .dSYM bundles. + if (Format != "yaml" && // For YAML, keep the original behavior. + Triple.isOSDarwin() && // Enable this only on darwin, since it's the + // only platform supporting .dSYM bundles. Output.isFilename()) F = Output.getFilename(); } @@ -1346,7 +1346,7 @@ void AddUnalignedAccessWarning(ArgStringList &CmdArgs) { StringRef(*StrictAlignIter) == "+strict-align") CmdArgs.push_back("-Wunaligned-access"); } -} +} // namespace static void CollectARMPACBTIOptions(const ToolChain &TC, const ArgList &Args, ArgStringList &CmdArgs, bool isAArch64) { @@ -1596,7 +1596,7 @@ void RenderAArch64ABI(const llvm::Triple &Triple, const ArgList &Args, CmdArgs.push_back("-target-abi"); CmdArgs.push_back(ABIName); } -} +} // namespace void Clang::AddAArch64TargetArgs(const ArgList &Args, ArgStringList &CmdArgs) const { @@ -1947,9 +1947,9 @@ void Clang::AddPPCTargetArgs(const ArgList &Args, ABIName = "elfv2"; A->claim(); } else if (V != "altivec") - // The ppc64 linux abis are all "altivec" abis by default. Accept and ignore - // the option if given as we don't have backend support for any targets - // that don't use the altivec abi. + // The ppc64 linux abis are all "altivec" abis by default. Accept and + // ignore the option if given as we don't have backend support for any + // targets that don't use the altivec abi. ABIName = A->getValue(); } if (IEEELongDouble) @@ -2100,7 +2100,7 @@ void Clang::AddSystemZTargetArgs(const ArgList &Args, if (HasBackchain && HasPackedStack && !HasSoftFloat) { const Driver &D = getToolChain().getDriver(); D.Diag(diag::err_drv_unsupported_opt) - << "-mpacked-stack -mbackchain -mhard-float"; + << "-mpacked-stack -mbackchain -mhard-float"; } if (HasBackchain) CmdArgs.push_back("-mbackchain"); @@ -2255,7 +2255,8 @@ void Clang::AddVETargetArgs(const ArgList &Args, ArgStringList &CmdArgs) const { void Clang::DumpCompilationDatabase(Compilation &C, StringRef Filename, StringRef Target, const InputInfo &Output, - const InputInfo &Input, const ArgList &Args) const { + const InputInfo &Input, + const ArgList &Args) const { // If this is a dry run, do not create the compilation database file. if (C.getArgs().hasArg(options::OPT__HASH_HASH_HASH)) return; @@ -2269,8 +2270,8 @@ void Clang::DumpCompilationDatabase(Compilation &C, StringRef Filename, Filename, EC, llvm::sys::fs::OF_TextWithCRLF | llvm::sys::fs::OF_Append); if (EC) { - D.Diag(clang::diag::err_drv_compilationdatabase) << Filename - << EC.message(); + D.Diag(clang::diag::err_drv_compilationdatabase) + << Filename << EC.message(); return; } CompilationDatabase = std::move(File); @@ -2296,7 +2297,7 @@ void Clang::DumpCompilationDatabase(Compilation &C, StringRef Filename, CDB << ", \"" << escape(Input.getFilename()) << "\""; if (Output.isFilename()) CDB << ", \"-o\", \"" << escape(Output.getFilename()) << "\""; - for (auto &A: Args) { + for (auto &A : Args) { auto &O = A->getOption(); // Skip language selection, which is positional. if (O.getID() == options::OPT_x) @@ -2315,7 +2316,7 @@ void Clang::DumpCompilationDatabase(Compilation &C, StringRef Filename, // All other arguments are quoted and appended. ArgStringList ASL; A->render(Args, ASL); - for (auto &it: ASL) + for (auto &it : ASL) CDB << ", \"" << escape(it) << "\""; } Buf = "--target="; @@ -2733,7 +2734,7 @@ static void RenderFloatingPointOptions(const ToolChain &TC, const Driver &D, bool AssociativeMath = false; bool ReciprocalMath = false; bool SignedZeros = true; - bool TrappingMath = false; // Implemented via -ffp-exception-behavior + bool TrappingMath = false; // Implemented via -ffp-exception-behavior bool TrappingMathPresent = false; // Is trapping-math in args, and not // overriden by ffp-exception-behavior? bool RoundingFPMath = false; @@ -2825,7 +2826,8 @@ static void RenderFloatingPointOptions(const ToolChain &TC, const Driver &D, switch (A->getOption().getID()) { // If this isn't an FP option skip the claim below - default: continue; + default: + continue; case options::OPT_fcx_limited_range: setComplexRange(D, A->getSpelling(), @@ -2925,20 +2927,48 @@ static void RenderFloatingPointOptions(const ToolChain &TC, const Driver &D, } // Options controlling individual features - case options::OPT_fhonor_infinities: HonorINFs = true; break; - case options::OPT_fno_honor_infinities: HonorINFs = false; break; - case options::OPT_fhonor_nans: HonorNaNs = true; break; - case options::OPT_fno_honor_nans: HonorNaNs = false; break; - case options::OPT_fapprox_func: ApproxFunc = true; break; - case options::OPT_fno_approx_func: ApproxFunc = false; break; - case options::OPT_fmath_errno: MathErrno = true; break; - case options::OPT_fno_math_errno: MathErrno = false; break; - case options::OPT_fassociative_math: AssociativeMath = true; break; - case options::OPT_fno_associative_math: AssociativeMath = false; break; - case options::OPT_freciprocal_math: ReciprocalMath = true; break; - case options::OPT_fno_reciprocal_math: ReciprocalMath = false; break; - case options::OPT_fsigned_zeros: SignedZeros = true; break; - case options::OPT_fno_signed_zeros: SignedZeros = false; break; + case options::OPT_fhonor_infinities: + HonorINFs = true; + break; + case options::OPT_fno_honor_infinities: + HonorINFs = false; + break; + case options::OPT_fhonor_nans: + HonorNaNs = true; + break; + case options::OPT_fno_honor_nans: + HonorNaNs = false; + break; + case options::OPT_fapprox_func: + ApproxFunc = true; + break; + case options::OPT_fno_approx_func: + ApproxFunc = false; + break; + case options::OPT_fmath_errno: + MathErrno = true; + break; + case options::OPT_fno_math_errno: + MathErrno = false; + break; + case options::OPT_fassociative_math: + AssociativeMath = true; + break; + case options::OPT_fno_associative_math: + AssociativeMath = false; + break; + case options::OPT_freciprocal_math: + ReciprocalMath = true; + break; + case options::OPT_fno_reciprocal_math: + ReciprocalMath = false; + break; + case options::OPT_fsigned_zeros: + SignedZeros = true; + break; + case options::OPT_fno_signed_zeros: + SignedZeros = false; + break; case options::OPT_ftrapping_math: if (!TrappingMathPresent && !FPExceptionBehavior.empty() && FPExceptionBehavior != "strict") @@ -3186,8 +3216,8 @@ static void RenderFloatingPointOptions(const ToolChain &TC, const Driver &D, << VecLibArg->getAsString(Args); } - if (AssociativeMath && ReciprocalMath && !SignedZeros && ApproxFunc && - !TrappingMath) + if (AssociativeMath && ReciprocalMath && !SignedZeros && ApproxFunc && + !TrappingMath) CmdArgs.push_back("-funsafe-math-optimizations"); if (!SignedZeros) @@ -3229,8 +3259,8 @@ static void RenderFloatingPointOptions(const ToolChain &TC, const Driver &D, CmdArgs.push_back(Args.MakeArgString("-fno-rounding-math")); if (!FPExceptionBehavior.empty()) - CmdArgs.push_back(Args.MakeArgString("-ffp-exception-behavior=" + - FPExceptionBehavior)); + CmdArgs.push_back( + Args.MakeArgString("-ffp-exception-behavior=" + FPExceptionBehavior)); if (!FPEvalMethod.empty()) CmdArgs.push_back(Args.MakeArgString("-ffp-eval-method=" + FPEvalMethod)); @@ -3334,8 +3364,7 @@ static void RenderAnalyzerOptions(const ArgList &Args, ArgStringList &CmdArgs, CmdArgs.push_back("-analyzer-checker=osx"); CmdArgs.push_back( "-analyzer-checker=security.insecureAPI.decodeValueOfObjCType"); - } - else if (Triple.isOSFuchsia()) + } else if (Triple.isOSFuchsia()) CmdArgs.push_back("-analyzer-checker=fuchsia"); CmdArgs.push_back("-analyzer-checker=deadcode"); @@ -3344,7 +3373,8 @@ static void RenderAnalyzerOptions(const ArgList &Args, ArgStringList &CmdArgs, CmdArgs.push_back("-analyzer-checker=cplusplus"); if (!Triple.isPS()) { - CmdArgs.push_back("-analyzer-checker=security.insecureAPI.UncheckedReturn"); + CmdArgs.push_back( + "-analyzer-checker=security.insecureAPI.UncheckedReturn"); CmdArgs.push_back("-analyzer-checker=security.insecureAPI.getpw"); CmdArgs.push_back("-analyzer-checker=security.insecureAPI.gets"); CmdArgs.push_back("-analyzer-checker=security.insecureAPI.mktemp"); @@ -3667,8 +3697,7 @@ static void RenderOpenCLOptions(const ArgList &Args, ArgStringList &CmdArgs, options::OPT_cl_mad_enable, options::OPT_cl_no_signed_zeros, options::OPT_cl_fp32_correctly_rounded_divide_sqrt, - options::OPT_cl_uniform_work_group_size - }; + options::OPT_cl_uniform_work_group_size}; if (Arg *A = Args.getLastArg(options::OPT_cl_std_EQ)) { std::string CLStdStr = std::string("-cl-std=") + A->getValue(); @@ -4043,10 +4072,9 @@ static bool RenderModulesOptions(Compilation &C, const Driver &D, static void RenderCharacterOptions(const ArgList &Args, const llvm::Triple &T, ArgStringList &CmdArgs) { // -fsigned-char is default. - if (const Arg *A = Args.getLastArg(options::OPT_fsigned_char, - options::OPT_fno_signed_char, - options::OPT_funsigned_char, - options::OPT_fno_unsigned_char)) { + if (const Arg *A = Args.getLastArg( + options::OPT_fsigned_char, options::OPT_fno_signed_char, + options::OPT_funsigned_char, options::OPT_fno_unsigned_char)) { if (A->getOption().matches(options::OPT_funsigned_char) || A->getOption().matches(options::OPT_fno_signed_char)) { CmdArgs.push_back("-fno-signed-char"); @@ -4140,9 +4168,8 @@ static void RenderObjCOptions(const ToolChain &TC, const Driver &D, auto *Arg = Args.getLastArg( options::OPT_fobjc_convert_messages_to_runtime_calls, options::OPT_fno_objc_convert_messages_to_runtime_calls); - if (Arg && - Arg->getOption().matches( - options::OPT_fno_objc_convert_messages_to_runtime_calls)) + if (Arg && Arg->getOption().matches( + options::OPT_fno_objc_convert_messages_to_runtime_calls)) CmdArgs.push_back("-fno-objc-convert-messages-to-runtime-calls"); } @@ -5045,8 +5072,8 @@ void Clang::ConstructJob(Compilation &C, const JobAction &JA, bool Failure = Triple.getArchName().substr(Offset).consumeInteger(10, Version); if (Failure || Version < 7) - D.Diag(diag::err_target_unsupported_arch) << Triple.getArchName() - << TripleStr; + D.Diag(diag::err_target_unsupported_arch) + << Triple.getArchName() << TripleStr; } // Push all default warning arguments that are specific to @@ -5077,6 +5104,7 @@ void Clang::ConstructJob(Compilation &C, const JobAction &JA, // -ffunction-sections not being used in -E mode either for example, even // though it's not really used either. if (!isa(JA)) { + llvm::errs() << "This is a Assemble Job Action\n"; // The args claimed here should match the args used in // CollectArgsForIntegratedAssembler(). if (TC.useIntegratedAs()) { @@ -5101,9 +5129,11 @@ void Clang::ConstructJob(Compilation &C, const JobAction &JA, } if (isa(JA)) { + llvm::errs() << "This is a AnalyzeJob Action\n"; assert(JA.getType() == types::TY_Plist && "Invalid output type."); CmdArgs.push_back("-analyze"); } else if (isa(JA)) { + llvm::errs() << "This is a PreprocessJobAction\n"; if (Output.getType() == types::TY_Dependencies) CmdArgs.push_back("-Eonly"); else { @@ -5115,6 +5145,7 @@ void Clang::ConstructJob(Compilation &C, const JobAction &JA, CmdArgs.push_back("-fdirectives-only"); } } else if (isa(JA)) { + llvm::errs() << "This is a AssembleJobAction\n"; CmdArgs.push_back("-emit-obj"); CollectArgsForIntegratedAssembler(C, Args, CmdArgs, D); @@ -5122,6 +5153,7 @@ void Clang::ConstructJob(Compilation &C, const JobAction &JA, // Also ignore explicit -force_cpusubtype_ALL option. (void)Args.hasArg(options::OPT_force__cpusubtype__ALL); } else if (isa(JA)) { + llvm::errs() << "This is a PrecompileJobAction\n"; if (JA.getType() == types::TY_Nothing) CmdArgs.push_back("-fsyntax-only"); else if (JA.getType() == types::TY_ModuleFile) @@ -5159,8 +5191,11 @@ void Clang::ConstructJob(Compilation &C, const JobAction &JA, } else { assert((isa(JA) || isa(JA)) && "Invalid action for clang tool."); + llvm::errs() << "Here accessing CompileJobAction and or BackEndJobAction\n"; if (JA.getType() == types::TY_Nothing) { CmdArgs.push_back("-fsyntax-only"); + } else if (JA.getType() == types::TY_CIR) { + CmdArgs.push_back("-emit-cir"); } else if (JA.getType() == types::TY_LLVM_IR || JA.getType() == types::TY_LTO_IR) { CmdArgs.push_back("-emit-llvm"); @@ -5242,8 +5277,7 @@ void Clang::ConstructJob(Compilation &C, const JobAction &JA, Twine("-flto=") + (LTOMode == LTOK_Thin ? "thin" : "full"))); // PS4 uses the legacy LTO API, which does not support some of the // features enabled by -flto-unit. - if (!RawTriple.isPS4() || - (D.getLTOMode() == LTOK_Full) || !UnifiedLTO) + if (!RawTriple.isPS4() || (D.getLTOMode() == LTOK_Full) || !UnifiedLTO) CmdArgs.push_back("-flto-unit"); } } @@ -5356,7 +5390,8 @@ void Clang::ConstructJob(Compilation &C, const JobAction &JA, // Render ABI arguments switch (TC.getArch()) { - default: break; + default: + break; case llvm::Triple::arm: case llvm::Triple::armeb: case llvm::Triple::thumbeb: @@ -5672,7 +5707,7 @@ void Clang::ConstructJob(Compilation &C, const JobAction &JA, if (Arg *A = Args.getLastArg(options::OPT_mabi_EQ_quadword_atomics)) { if (!Triple.isOSAIX() || Triple.isPPC32()) D.Diag(diag::err_drv_unsupported_opt_for_target) - << A->getSpelling() << RawTriple.str(); + << A->getSpelling() << RawTriple.str(); CmdArgs.push_back("-mabi=quadword-atomics"); } @@ -5747,7 +5782,7 @@ void Clang::ConstructJob(Compilation &C, const JobAction &JA, } CodeGenOptions::FramePointerKind FPKeepKind = - getFramePointerKind(Args, RawTriple); + getFramePointerKind(Args, RawTriple); const char *FPKeepKindStr = nullptr; switch (FPKeepKind) { case CodeGenOptions::FramePointerKind::None: @@ -5963,10 +5998,10 @@ void Clang::ConstructJob(Compilation &C, const JobAction &JA, // This is a coarse approximation of what llvm-gcc actually does, both // -fasynchronous-unwind-tables and -fnon-call-exceptions interact in more // complicated ways. - bool IsAsyncUnwindTablesDefault = - TC.getDefaultUnwindTableLevel(Args) == ToolChain::UnwindTableLevel::Asynchronous; - bool IsSyncUnwindTablesDefault = - TC.getDefaultUnwindTableLevel(Args) == ToolChain::UnwindTableLevel::Synchronous; + bool IsAsyncUnwindTablesDefault = TC.getDefaultUnwindTableLevel(Args) == + ToolChain::UnwindTableLevel::Asynchronous; + bool IsSyncUnwindTablesDefault = TC.getDefaultUnwindTableLevel(Args) == + ToolChain::UnwindTableLevel::Synchronous; bool AsyncUnwindTables = Args.hasFlag( options::OPT_fasynchronous_unwind_tables, @@ -5979,7 +6014,7 @@ void Clang::ConstructJob(Compilation &C, const JobAction &JA, if (AsyncUnwindTables) CmdArgs.push_back("-funwind-tables=2"); else if (UnwindTables) - CmdArgs.push_back("-funwind-tables=1"); + CmdArgs.push_back("-funwind-tables=1"); // Prepare `-aux-target-cpu` and `-aux-target-feature` unless // `--gpu-use-aux-triple-only` is specified. @@ -6355,8 +6390,7 @@ void Clang::ConstructJob(Compilation &C, const JobAction &JA, /*Joined=*/true); } else ImplyVCPPCVer = true; - } - else if (IsWindowsMSVC) + } else if (IsWindowsMSVC) ImplyVCPPCXXVer = true; Args.AddLastArg(CmdArgs, options::OPT_ftrigraphs, @@ -6432,7 +6466,7 @@ void Clang::ConstructJob(Compilation &C, const JobAction &JA, if (const Arg *A = Args.getLastArg(options::OPT_fcf_runtime_abi_EQ)) { static const char *kCFABIs[] = { - "standalone", "objc", "swift", "swift-5.0", "swift-4.2", "swift-4.1", + "standalone", "objc", "swift", "swift-5.0", "swift-4.2", "swift-4.1", }; if (!llvm::is_contained(kCFABIs, StringRef(A->getValue()))) @@ -6541,11 +6575,12 @@ void Clang::ConstructJob(Compilation &C, const JobAction &JA, } if (Args.hasFlag(options::OPT_fvisibility_inlines_hidden, - options::OPT_fno_visibility_inlines_hidden, false)) + options::OPT_fno_visibility_inlines_hidden, false)) CmdArgs.push_back("-fvisibility-inlines-hidden"); - Args.AddLastArg(CmdArgs, options::OPT_fvisibility_inlines_hidden_static_local_var, - options::OPT_fno_visibility_inlines_hidden_static_local_var); + Args.AddLastArg(CmdArgs, + options::OPT_fvisibility_inlines_hidden_static_local_var, + options::OPT_fno_visibility_inlines_hidden_static_local_var); // -fvisibility-global-new-delete-hidden is a deprecated spelling of // -fvisibility-global-new-delete=force-hidden. @@ -7042,8 +7077,8 @@ void Clang::ConstructJob(Compilation &C, const JobAction &JA, ToolChain::RTTIMode RTTIMode = TC.getRTTIMode(); - if (KernelOrKext || (types::isCXX(InputType) && - (RTTIMode == ToolChain::RM_Disabled))) + if (KernelOrKext || + (types::isCXX(InputType) && (RTTIMode == ToolChain::RM_Disabled))) CmdArgs.push_back("-fno-rtti"); // -fshort-enums=0 is default for all architectures except Hexagon and z/OS. @@ -7420,16 +7455,16 @@ void Clang::ConstructJob(Compilation &C, const JobAction &JA, if (Arg *inputCharset = Args.getLastArg(options::OPT_finput_charset_EQ)) { StringRef value = inputCharset->getValue(); if (!value.equals_insensitive("utf-8")) - D.Diag(diag::err_drv_invalid_value) << inputCharset->getAsString(Args) - << value; + D.Diag(diag::err_drv_invalid_value) + << inputCharset->getAsString(Args) << value; } // -fexec_charset=UTF-8 is default. Reject others if (Arg *execCharset = Args.getLastArg(options::OPT_fexec_charset_EQ)) { StringRef value = execCharset->getValue(); if (!value.equals_insensitive("utf-8")) - D.Diag(diag::err_drv_invalid_value) << execCharset->getAsString(Args) - << value; + D.Diag(diag::err_drv_invalid_value) + << execCharset->getAsString(Args) << value; } RenderDiagnosticsOptions(D, Args, CmdArgs); @@ -7607,8 +7642,7 @@ void Clang::ConstructJob(Compilation &C, const JobAction &JA, Arg->claim(); // -finclude-default-header flag is for preprocessor, // do not pass it to other cc1 commands when save-temps is enabled - if (C.getDriver().isSaveTempsEnabled() && - !isa(JA)) { + if (C.getDriver().isSaveTempsEnabled() && !isa(JA)) { if (StringRef(Arg->getValue()) == "-finclude-default-header") continue; } @@ -7832,7 +7866,7 @@ void Clang::ConstructJob(Compilation &C, const JobAction &JA, bool DefaultsSplitLTOUnit = ((WholeProgramVTables || SanitizeArgs.needsLTO()) && - (LTOMode == LTOK_Full || TC.canSplitThinLTOUnit())) || + (LTOMode == LTOK_Full || TC.canSplitThinLTOUnit())) || (!Triple.isPS4() && UnifiedLTO); bool SplitLTOUnit = Args.hasFlag(options::OPT_fsplit_lto_unit, @@ -7998,8 +8032,8 @@ void Clang::ConstructJob(Compilation &C, const JobAction &JA, if (Arg *A = Args.getLastArg(options::OPT_pg)) if (FPKeepKind == CodeGenOptions::FramePointerKind::None && !Args.hasArg(options::OPT_mfentry)) - D.Diag(diag::err_drv_argument_not_allowed_with) << "-fomit-frame-pointer" - << A->getAsString(Args); + D.Diag(diag::err_drv_argument_not_allowed_with) + << "-fomit-frame-pointer" << A->getAsString(Args); // Claim some arguments which clang supports automatically. @@ -8052,7 +8086,7 @@ ObjCRuntime Clang::AddObjCRuntimeArgs(const ArgList &args, !getToolChain().getTriple().isOSBinFormatCOFF()) { getToolChain().getDriver().Diag( diag::err_drv_gnustep_objc_runtime_incompatible_binary) - << runtime.getVersion().getMajor(); + << runtime.getVersion().getMajor(); } runtimeArg->render(args, cmdArgs); @@ -8282,37 +8316,36 @@ void Clang::AddClangCLArgs(const ArgList &Args, types::ID InputType, CmdArgs.push_back("-P"); } - if (Args.hasFlag(options::OPT__SLASH_Zc_dllexportInlines_, - options::OPT__SLASH_Zc_dllexportInlines, - false)) { - CmdArgs.push_back("-fno-dllexport-inlines"); - } - - if (Args.hasFlag(options::OPT__SLASH_Zc_wchar_t_, - options::OPT__SLASH_Zc_wchar_t, false)) { - CmdArgs.push_back("-fno-wchar"); - } - - if (Args.hasArg(options::OPT__SLASH_kernel)) { - llvm::Triple::ArchType Arch = getToolChain().getArch(); - std::vector Values = - Args.getAllArgValues(options::OPT__SLASH_arch); - if (!Values.empty()) { - llvm::SmallSet SupportedArches; - if (Arch == llvm::Triple::x86) - SupportedArches.insert("IA32"); - - for (auto &V : Values) - if (!SupportedArches.contains(V)) - D.Diag(diag::err_drv_argument_not_allowed_with) - << std::string("/arch:").append(V) << "/kernel"; - } - - CmdArgs.push_back("-fno-rtti"); - if (Args.hasFlag(options::OPT__SLASH_GR, options::OPT__SLASH_GR_, false)) - D.Diag(diag::err_drv_argument_not_allowed_with) << "/GR" - << "/kernel"; - } + if (Args.hasFlag(options::OPT__SLASH_Zc_dllexportInlines_, + options::OPT__SLASH_Zc_dllexportInlines, false)) { + CmdArgs.push_back("-fno-dllexport-inlines"); + } + + if (Args.hasFlag(options::OPT__SLASH_Zc_wchar_t_, + options::OPT__SLASH_Zc_wchar_t, false)) { + CmdArgs.push_back("-fno-wchar"); + } + + if (Args.hasArg(options::OPT__SLASH_kernel)) { + llvm::Triple::ArchType Arch = getToolChain().getArch(); + std::vector Values = + Args.getAllArgValues(options::OPT__SLASH_arch); + if (!Values.empty()) { + llvm::SmallSet SupportedArches; + if (Arch == llvm::Triple::x86) + SupportedArches.insert("IA32"); + + for (auto &V : Values) + if (!SupportedArches.contains(V)) + D.Diag(diag::err_drv_argument_not_allowed_with) + << std::string("/arch:").append(V) << "/kernel"; + } + + CmdArgs.push_back("-fno-rtti"); + if (Args.hasFlag(options::OPT__SLASH_GR, options::OPT__SLASH_GR_, false)) + D.Diag(diag::err_drv_argument_not_allowed_with) << "/GR" + << "/kernel"; + } Arg *MostGeneralArg = Args.getLastArg(options::OPT__SLASH_vmg); Arg *BestCaseArg = Args.getLastArg(options::OPT__SLASH_vmb); @@ -8494,7 +8527,7 @@ void ClangAs::AddLoongArchTargetArgs(const ArgList &Args, } void ClangAs::AddRISCVTargetArgs(const ArgList &Args, - ArgStringList &CmdArgs) const { + ArgStringList &CmdArgs) const { const llvm::Triple &Triple = getToolChain().getTriple(); StringRef ABIName = riscv::getRISCVABI(Args, Triple); @@ -8503,8 +8536,8 @@ void ClangAs::AddRISCVTargetArgs(const ArgList &Args, if (Args.hasFlag(options::OPT_mdefault_build_attributes, options::OPT_mno_default_build_attributes, true)) { - CmdArgs.push_back("-mllvm"); - CmdArgs.push_back("-riscv-add-build-attributes"); + CmdArgs.push_back("-mllvm"); + CmdArgs.push_back("-riscv-add-build-attributes"); } } @@ -8706,8 +8739,8 @@ void ClangAs::ConstructJob(Compilation &C, const JobAction &JA, // only, not C/C++. if (Args.hasFlag(options::OPT_mdefault_build_attributes, options::OPT_mno_default_build_attributes, true)) { - CmdArgs.push_back("-mllvm"); - CmdArgs.push_back("-arm-add-build-attributes"); + CmdArgs.push_back("-mllvm"); + CmdArgs.push_back("-arm-add-build-attributes"); } break; @@ -8770,12 +8803,12 @@ void ClangAs::ConstructJob(Compilation &C, const JobAction &JA, for (unsigned I = 0; I < JArgs.size(); ++I) { if (StringRef(JArgs[I]).starts_with("-object-file-name=") && Output.isFilename()) { - ArgStringList NewArgs(JArgs.begin(), JArgs.begin() + I); - addDebugObjectName(Args, NewArgs, DebugCompilationDir, - Output.getFilename()); - NewArgs.append(JArgs.begin() + I + 1, JArgs.end()); - J.replaceArguments(NewArgs); - break; + ArgStringList NewArgs(JArgs.begin(), JArgs.begin() + I); + addDebugObjectName(Args, NewArgs, DebugCompilationDir, + Output.getFilename()); + NewArgs.append(JArgs.begin() + I + 1, JArgs.end()); + J.replaceArguments(NewArgs); + break; } } } @@ -9029,7 +9062,7 @@ void OffloadPackager::ConstructJob(Compilation &C, const JobAction &JA, ? OffloadAction->getOffloadingArch() : TCArgs.getLastArgValue(options::OPT_march_EQ); StringRef Kind = - Action::GetOffloadKindName(OffloadAction->getOffloadingDeviceKind()); + Action::GetOffloadKindName(OffloadAction->getOffloadingDeviceKind()); ArgStringList Features; SmallVector FeatureArgs; diff --git a/clang/lib/Frontend/CompilerInvocation.cpp b/clang/lib/Frontend/CompilerInvocation.cpp index ca71f49aef13..32d2d75ec289 100644 --- a/clang/lib/Frontend/CompilerInvocation.cpp +++ b/clang/lib/Frontend/CompilerInvocation.cpp @@ -188,8 +188,7 @@ CompilerInvocation::operator=(const CowCompilerInvocation &X) { return *this; } -template -T &ensureOwned(std::shared_ptr &Storage) { +template T &ensureOwned(std::shared_ptr &Storage) { if (Storage.use_count() > 1) Storage = std::make_shared(*Storage); return *Storage; @@ -690,7 +689,8 @@ static bool FixupInvocation(CompilerInvocation &Invocation, emitError |= (DefaultCC == LangOptions::DCC_VectorCall || DefaultCC == LangOptions::DCC_RegCall) && !T.isX86(); - emitError |= DefaultCC == LangOptions::DCC_RtdCall && Arch != llvm::Triple::m68k; + emitError |= + DefaultCC == LangOptions::DCC_RtdCall && Arch != llvm::Triple::m68k; if (emitError) Diags.Report(diag::err_drv_argument_not_allowed_with) << A->getSpelling() << T.getTriple(); @@ -1074,13 +1074,12 @@ static bool ParseAnalyzerArgs(AnalyzerOptions &Opts, ArgList &Args, if (Arg *A = Args.getLastArg(OPT_analyzer_constraints)) { StringRef Name = A->getValue(); AnalysisConstraints Value = llvm::StringSwitch(Name) -#define ANALYSIS_CONSTRAINTS(NAME, CMDFLAG, DESC, CREATFN) \ - .Case(CMDFLAG, NAME##Model) +#define ANALYSIS_CONSTRAINTS(NAME, CMDFLAG, DESC, CREATFN) \ + .Case(CMDFLAG, NAME##Model) #include "clang/StaticAnalyzer/Core/Analyses.def" - .Default(NumConstraints); + .Default(NumConstraints); if (Value == NumConstraints) { - Diags.Report(diag::err_drv_invalid_value) - << A->getAsString(Args) << Name; + Diags.Report(diag::err_drv_invalid_value) << A->getAsString(Args) << Name; } else { #ifndef LLVM_WITH_Z3 if (Value == AnalysisConstraints::Z3ConstraintsModel) { @@ -1094,13 +1093,12 @@ static bool ParseAnalyzerArgs(AnalyzerOptions &Opts, ArgList &Args, if (Arg *A = Args.getLastArg(OPT_analyzer_output)) { StringRef Name = A->getValue(); AnalysisDiagClients Value = llvm::StringSwitch(Name) -#define ANALYSIS_DIAGNOSTICS(NAME, CMDFLAG, DESC, CREATFN) \ - .Case(CMDFLAG, PD_##NAME) +#define ANALYSIS_DIAGNOSTICS(NAME, CMDFLAG, DESC, CREATFN) \ + .Case(CMDFLAG, PD_##NAME) #include "clang/StaticAnalyzer/Core/Analyses.def" - .Default(NUM_ANALYSIS_DIAG_CLIENTS); + .Default(NUM_ANALYSIS_DIAG_CLIENTS); if (Value == NUM_ANALYSIS_DIAG_CLIENTS) { - Diags.Report(diag::err_drv_invalid_value) - << A->getAsString(Args) << Name; + Diags.Report(diag::err_drv_invalid_value) << A->getAsString(Args) << Name; } else { Opts.AnalysisDiagOpt = Value; } @@ -1109,13 +1107,11 @@ static bool ParseAnalyzerArgs(AnalyzerOptions &Opts, ArgList &Args, if (Arg *A = Args.getLastArg(OPT_analyzer_purge)) { StringRef Name = A->getValue(); AnalysisPurgeMode Value = llvm::StringSwitch(Name) -#define ANALYSIS_PURGE(NAME, CMDFLAG, DESC) \ - .Case(CMDFLAG, NAME) +#define ANALYSIS_PURGE(NAME, CMDFLAG, DESC) .Case(CMDFLAG, NAME) #include "clang/StaticAnalyzer/Core/Analyses.def" - .Default(NumPurgeModes); + .Default(NumPurgeModes); if (Value == NumPurgeModes) { - Diags.Report(diag::err_drv_invalid_value) - << A->getAsString(Args) << Name; + Diags.Report(diag::err_drv_invalid_value) << A->getAsString(Args) << Name; } else { Opts.AnalysisPurgeOpt = Value; } @@ -1124,13 +1120,11 @@ static bool ParseAnalyzerArgs(AnalyzerOptions &Opts, ArgList &Args, if (Arg *A = Args.getLastArg(OPT_analyzer_inlining_mode)) { StringRef Name = A->getValue(); AnalysisInliningMode Value = llvm::StringSwitch(Name) -#define ANALYSIS_INLINING_MODE(NAME, CMDFLAG, DESC) \ - .Case(CMDFLAG, NAME) +#define ANALYSIS_INLINING_MODE(NAME, CMDFLAG, DESC) .Case(CMDFLAG, NAME) #include "clang/StaticAnalyzer/Core/Analyses.def" - .Default(NumInliningModes); + .Default(NumInliningModes); if (Value == NumInliningModes) { - Diags.Report(diag::err_drv_invalid_value) - << A->getAsString(Args) << Name; + Diags.Report(diag::err_drv_invalid_value) << A->getAsString(Args) << Name; } else { Opts.InliningMode = Value; } @@ -1163,14 +1157,14 @@ static bool ParseAnalyzerArgs(AnalyzerOptions &Opts, ArgList &Args, StringRef key, val; std::tie(key, val) = configVal.split("="); if (val.empty()) { - Diags.Report(SourceLocation(), - diag::err_analyzer_config_no_value) << configVal; + Diags.Report(SourceLocation(), diag::err_analyzer_config_no_value) + << configVal; break; } if (val.contains('=')) { Diags.Report(SourceLocation(), diag::err_analyzer_config_multiple_values) - << configVal; + << configVal; break; } @@ -1208,9 +1202,8 @@ static StringRef getStringOption(AnalyzerOptions::ConfigTable &Config, } static void initOption(AnalyzerOptions::ConfigTable &Config, - DiagnosticsEngine *Diags, - StringRef &OptionField, StringRef Name, - StringRef DefaultVal) { + DiagnosticsEngine *Diags, StringRef &OptionField, + StringRef Name, StringRef DefaultVal) { // String options may be known to invalid (e.g. if the expected string is a // file name, but the file does not exist), those will have to be checked in // parseConfigs. @@ -1218,8 +1211,8 @@ static void initOption(AnalyzerOptions::ConfigTable &Config, } static void initOption(AnalyzerOptions::ConfigTable &Config, - DiagnosticsEngine *Diags, - bool &OptionField, StringRef Name, bool DefaultVal) { + DiagnosticsEngine *Diags, bool &OptionField, + StringRef Name, bool DefaultVal) { auto PossiblyInvalidVal = llvm::StringSwitch>( getStringOption(Config, Name, (DefaultVal ? "true" : "false"))) @@ -1230,7 +1223,7 @@ static void initOption(AnalyzerOptions::ConfigTable &Config, if (!PossiblyInvalidVal) { if (Diags) Diags->Report(diag::err_analyzer_config_invalid_input) - << Name << "a boolean"; + << Name << "a boolean"; else OptionField = DefaultVal; } else @@ -1238,16 +1231,15 @@ static void initOption(AnalyzerOptions::ConfigTable &Config, } static void initOption(AnalyzerOptions::ConfigTable &Config, - DiagnosticsEngine *Diags, - unsigned &OptionField, StringRef Name, - unsigned DefaultVal) { + DiagnosticsEngine *Diags, unsigned &OptionField, + StringRef Name, unsigned DefaultVal) { OptionField = DefaultVal; bool HasFailed = getStringOption(Config, Name, std::to_string(DefaultVal)) - .getAsInteger(0, OptionField); + .getAsInteger(0, OptionField); if (Diags && HasFailed) Diags->Report(diag::err_analyzer_config_invalid_input) - << Name << "an unsigned"; + << Name << "an unsigned"; } static void initOption(AnalyzerOptions::ConfigTable &Config, @@ -1401,12 +1393,12 @@ static bool parseDiagnosticLevelMask(StringRef FlagName, bool Success = true; for (const auto &Level : Levels) { DiagnosticLevelMask const PM = - llvm::StringSwitch(Level) - .Case("note", DiagnosticLevelMask::Note) - .Case("remark", DiagnosticLevelMask::Remark) - .Case("warning", DiagnosticLevelMask::Warning) - .Case("error", DiagnosticLevelMask::Error) - .Default(DiagnosticLevelMask::None); + llvm::StringSwitch(Level) + .Case("note", DiagnosticLevelMask::Note) + .Case("remark", DiagnosticLevelMask::Remark) + .Case("warning", DiagnosticLevelMask::Warning) + .Case("error", DiagnosticLevelMask::Error) + .Default(DiagnosticLevelMask::None); if (PM == DiagnosticLevelMask::None) { Success = false; Diags.Report(diag::err_drv_invalid_value) << FlagName << Level; @@ -1929,8 +1921,8 @@ bool CompilerInvocation::ParseCodeGenArgs(CodeGenOptions &Opts, ArgList &Args, .Case("unused-types", llvm::codegenoptions::UnusedTypeInfo) .Default(~0U); if (Val == ~0U) - Diags.Report(diag::err_drv_invalid_value) << A->getAsString(Args) - << A->getValue(); + Diags.Report(diag::err_drv_invalid_value) + << A->getAsString(Args) << A->getValue(); else Opts.setDebugInfo(static_cast(Val)); } @@ -1979,9 +1971,8 @@ bool CompilerInvocation::ParseCodeGenArgs(CodeGenOptions &Opts, ArgList &Args, Opts.SimplifyLibCalls = !LangOpts->NoBuiltin; if (Opts.SimplifyLibCalls) Opts.NoBuiltinFuncs = LangOpts->NoBuiltinFuncs; - Opts.UnrollLoops = - Args.hasFlag(OPT_funroll_loops, OPT_fno_unroll_loops, - (Opts.OptimizationLevel > 1)); + Opts.UnrollLoops = Args.hasFlag(OPT_funroll_loops, OPT_fno_unroll_loops, + (Opts.OptimizationLevel > 1)); Opts.InterchangeLoops = Args.hasFlag(OPT_floop_interchange, OPT_fno_loop_interchange, false); Opts.FuseLoops = Args.hasFlag(OPT_fexperimental_loop_fusion, @@ -1994,9 +1985,9 @@ bool CompilerInvocation::ParseCodeGenArgs(CodeGenOptions &Opts, ArgList &Args, Opts.DebugNameTable = static_cast( Args.hasArg(OPT_ggnu_pubnames) ? llvm::DICompileUnit::DebugNameTableKind::GNU - : Args.hasArg(OPT_gpubnames) - ? llvm::DICompileUnit::DebugNameTableKind::Default - : llvm::DICompileUnit::DebugNameTableKind::None); + : Args.hasArg(OPT_gpubnames) + ? llvm::DICompileUnit::DebugNameTableKind::Default + : llvm::DICompileUnit::DebugNameTableKind::None); if (const Arg *A = Args.getLastArg(OPT_gsimple_template_names_EQ)) { StringRef Value = A->getValue(); if (Value != "simple" && Value != "mangled") @@ -2227,7 +2218,7 @@ bool CompilerInvocation::ParseCodeGenArgs(CodeGenOptions &Opts, ArgList &Args, if (Arg *A = Args.getLastArg(OPT_mabi_EQ_quadword_atomics)) { if (!T.isOSAIX() || T.isPPC32()) Diags.Report(diag::err_drv_unsupported_opt_for_target) - << A->getSpelling() << T.str(); + << A->getSpelling() << T.str(); } bool NeedLocTracking = false; @@ -2321,12 +2312,12 @@ bool CompilerInvocation::ParseCodeGenArgs(CodeGenOptions &Opts, ArgList &Args, // Parse -fsanitize-recover= arguments. // FIXME: Report unrecoverable sanitizers incorrectly specified here. - parseSanitizerKinds("-fsanitize-recover=", - Args.getAllArgValues(OPT_fsanitize_recover_EQ), Diags, - Opts.SanitizeRecover); - parseSanitizerKinds("-fsanitize-trap=", - Args.getAllArgValues(OPT_fsanitize_trap_EQ), Diags, - Opts.SanitizeTrap); + parseSanitizerKinds( + "-fsanitize-recover=", Args.getAllArgValues(OPT_fsanitize_recover_EQ), + Diags, Opts.SanitizeRecover); + parseSanitizerKinds( + "-fsanitize-trap=", Args.getAllArgValues(OPT_fsanitize_trap_EQ), Diags, + Opts.SanitizeTrap); parseSanitizerKinds("-fsanitize-merge=", Args.getAllArgValues(OPT_fsanitize_merge_handlers_EQ), Diags, Opts.SanitizeMergeHandlers); @@ -2703,7 +2694,8 @@ bool clang::ParseDiagnosticArgs(DiagnosticOptions &Opts, ArgList &Args, Opts.DiagnosticSerializationFile = A->getValue(); Opts.ShowColors = parseShowColorsArgs(Args, DefaultDiagColor); - Opts.VerifyDiagnostics = Args.hasArg(OPT_verify) || Args.hasArg(OPT_verify_EQ); + Opts.VerifyDiagnostics = + Args.hasArg(OPT_verify) || Args.hasArg(OPT_verify_EQ); Opts.VerifyPrefixes = Args.getAllArgValues(OPT_verify_EQ); if (Args.hasArg(OPT_verify)) Opts.VerifyPrefixes.push_back("expected"); @@ -2739,21 +2731,22 @@ bool clang::ParseDiagnosticArgs(DiagnosticOptions &Opts, ArgList &Args, /// command-line argument. /// /// \returns true on error, false on success. -static bool parseTestModuleFileExtensionArg(StringRef Arg, - std::string &BlockName, - unsigned &MajorVersion, - unsigned &MinorVersion, - bool &Hashed, - std::string &UserInfo) { +static bool +parseTestModuleFileExtensionArg(StringRef Arg, std::string &BlockName, + unsigned &MajorVersion, unsigned &MinorVersion, + bool &Hashed, std::string &UserInfo) { SmallVector Args; Arg.split(Args, ':', 5); if (Args.size() < 5) return true; BlockName = std::string(Args[0]); - if (Args[1].getAsInteger(10, MajorVersion)) return true; - if (Args[2].getAsInteger(10, MinorVersion)) return true; - if (Args[3].getAsInteger(2, Hashed)) return true; + if (Args[1].getAsInteger(10, MajorVersion)) + return true; + if (Args[2].getAsInteger(10, MinorVersion)) + return true; + if (Args[3].getAsInteger(2, Hashed)) + return true; if (Args.size() > 4) UserInfo = std::string(Args[4]); return false; @@ -2790,6 +2783,7 @@ static const auto &getFrontendActionTable() { {frontend::EmitCodeGenOnly, OPT_emit_codegen_only}, {frontend::EmitObj, OPT_emit_obj}, {frontend::ExtractAPI, OPT_extract_api}, + {frontend::CIRCombine, OPT_cir_combine}, {frontend::FixIt, OPT_fixit_EQ}, {frontend::FixIt, OPT_fixit}, @@ -2901,6 +2895,20 @@ static void GenerateFrontendArgs(const FrontendOptions &Opts, }; } + if (Opts.ProgramAction == frontend::CIRCombine) { + if (!Opts.CIRHostInput.empty()) + GenerateArg(Consumer, OPT_cir_host_input, Opts.CIRHostInput); + if (!Opts.CIRDeviceInput.empty()) + GenerateArg(Consumer, OPT_cir_device_input, Opts.CIRDeviceInput); + if (Opts.EmitSplit) { + GenerateArg(Consumer, OPT_cir_emit_split); + if (!Opts.CIRHostOutput.empty()) + GenerateArg(Consumer, OPT_cir_host_output, Opts.CIRHostOutput); + if (!Opts.CIRDeviceOutput.empty()) + GenerateArg(Consumer, OPT_cir_device_output, Opts.CIRDeviceOutput); + } + } + if (Opts.ProgramAction == frontend::FixIt && !Opts.FixItSuffix.empty()) { GenerateProgramAction = [&]() { GenerateArg(Consumer, OPT_fixit_EQ, Opts.FixItSuffix); @@ -3001,6 +3009,8 @@ static void GenerateFrontendArgs(const FrontendOptions &Opts, Lang = "assembler-with-cpp"; break; case Language::Unknown: + if (Opts.ProgramAction == frontend::CIRCombine) + break; assert(Opts.DashX.getFormat() == InputKind::Precompiled && "Generating -x argument for unknown language (not precompiled)."); Lang = "ast"; @@ -3016,8 +3026,9 @@ static void GenerateFrontendArgs(const FrontendOptions &Opts, break; } - GenerateArg(Consumer, OPT_x, - Lang + HeaderUnit + Header + ModuleMap + Preprocessed); + if (Opts.ProgramAction != frontend::CIRCombine) + GenerateArg(Consumer, OPT_x, + Lang + HeaderUnit + Header + ModuleMap + Preprocessed); } // OPT_INPUT has a unique class, generate it directly. @@ -3115,7 +3126,7 @@ static bool ParseFrontendArgs(FrontendOptions &Opts, ArgList &Args, } } - if (const Arg* A = Args.getLastArg(OPT_plugin)) { + if (const Arg *A = Args.getLastArg(OPT_plugin)) { Opts.Plugins.emplace_back(A->getValue(0)); Opts.ProgramAction = frontend::PluginAction; Opts.ActionName = A->getValue(); @@ -3124,7 +3135,7 @@ static bool ParseFrontendArgs(FrontendOptions &Opts, ArgList &Args, Opts.PluginArgs[AA->getValue(0)].emplace_back(AA->getValue(1)); for (const std::string &Arg : - Args.getAllArgValues(OPT_ftest_module_file_extension_EQ)) { + Args.getAllArgValues(OPT_ftest_module_file_extension_EQ)) { std::string BlockName; unsigned MajorVersion; unsigned MinorVersion; @@ -3144,8 +3155,7 @@ static bool ParseFrontendArgs(FrontendOptions &Opts, ArgList &Args, } if (const Arg *A = Args.getLastArg(OPT_code_completion_at)) { - Opts.CodeCompletionAt = - ParsedSourceLocation::FromString(A->getValue()); + Opts.CodeCompletionAt = ParsedSourceLocation::FromString(A->getValue()); if (Opts.CodeCompletionAt.FileName.empty()) { Diags.Report(diag::err_drv_invalid_value) << A->getAsString(Args) << A->getValue(); @@ -3173,6 +3183,59 @@ static bool ParseFrontendArgs(FrontendOptions &Opts, ArgList &Args, if (Args.hasArg(OPT_fclangir_direct_lowering)) Opts.ClangIRDirectLowering = true; + if (Args.hasArg(OPT_cir_combine) && !Args.hasArg(OPT_fclangir)) + Diags.Report(diag::err_drv_argument_only_allowed_with) + << "-cir-combine" << "-fclangir"; + + if (Args.hasArg(OPT_fclangir_offload) && !Args.hasArg(OPT_fclangir)) + Diags.Report(diag::err_drv_argument_only_allowed_with) + << "-fclangir-offload" << "-fclangir"; + + if (Args.hasArg(OPT_cir_combine) && + Opts.ProgramAction == frontend::CIRCombine) { + if (!Opts.Inputs.empty()) { + Diags.Report(diag::err_drv_invalid_argument_to_option) + << "positional inputs" << "-cir-combine"; + } + + auto reqSingleInput = [&](llvm::opt::OptSpecifier Option, StringRef Name) { + auto args = Args.filtered(Option); + unsigned count = std::distance(args.begin(), args.end()); + if (count == 0) + Diags.Report(diag::err_drv_missing_argument) << Name << 1; + if (count > 1) + Diags.Report(diag::err_drv_invalid_argument_to_option) + << Twine("multiple ", Name).str() << "-cir-combine"; + return Args.getLastArgValue(Option).str(); + }; + + Opts.CIRHostInput = reqSingleInput(OPT_cir_host_input, "-cir-host-input"); + + if (!llvm::sys::fs::exists(Opts.CIRHostInput)) + Diags.Report(diag::err_drv_no_such_file) << Opts.CIRHostInput; + + Opts.CIRDeviceInput = + reqSingleInput(OPT_cir_device_input, "-cir-device-input"); + + if (!llvm::sys::fs::exists(Opts.CIRDeviceInput)) + Diags.Report(diag::err_drv_no_such_file) << Opts.CIRDeviceInput; + + if (!Args.hasArg(OPT_cir_emit_split)) { + if (!Args.hasArg(OPT_cir_combine) && !Args.getLastArg(OPT_o)) + Diags.Report(diag::err_drv_missing_argument) << "-o" << 1; + } else { + if (Args.hasArg(OPT_cir_combine) && Args.getLastArg(OPT_o)) + Diags.Report(diag::warn_drv_unsupported_option_overrides_option) + << "-o" << "-cir-emit-split"; + + Opts.EmitSplit = true; + Opts.CIRHostOutput = + reqSingleInput(OPT_cir_host_output, "-cir-host-output"); + Opts.CIRDeviceOutput = + reqSingleInput(OPT_cir_device_output, "-cir-device-output"); + } + } + if (Args.hasArg(OPT_clangir_disable_passes)) Opts.ClangIRDisablePasses = true; @@ -3284,7 +3347,7 @@ static bool ParseFrontendArgs(FrontendOptions &Opts, ArgList &Args, if (DashX.isUnknown()) Diags.Report(diag::err_drv_invalid_value) - << A->getAsString(Args) << A->getValue(); + << A->getAsString(Args) << A->getValue(); if (Preprocessed) DashX = DashX.getPreprocessed(); @@ -3301,6 +3364,14 @@ static bool ParseFrontendArgs(FrontendOptions &Opts, ArgList &Args, // '-' is the default input if none is given. std::vector Inputs = Args.getAllArgValues(OPT_INPUT); Opts.Inputs.clear(); + if (Opts.ProgramAction == frontend::CIRCombine) { + Opts.Inputs.emplace_back( + Opts.CIRHostInput, + clang::InputKind(clang::Language::CIR, clang::InputKind::Source), + false); + // CIRCombine does not require the -x flag for now. We skip this part + return Diags.getNumErrors() == NumErrorsBefore; + } if (Inputs.empty()) Inputs.push_back("-"); @@ -3312,7 +3383,7 @@ static bool ParseFrontendArgs(FrontendOptions &Opts, ArgList &Args, InputKind IK = DashX; if (IK.isUnknown()) { IK = FrontendOptions::getInputKindForExtension( - StringRef(Inputs[i]).rsplit('.').second); + StringRef(Inputs[i]).rsplit('.').second); // FIXME: Warn on this? if (IK.isUnknown()) IK = Language::C; @@ -3476,8 +3547,8 @@ static bool ParseHeaderSearchArgs(HeaderSearchOptions &Opts, ArgList &Args, StringRef Val = A->getValue(); if (Val.contains('=')) { auto Split = Val.split('='); - Opts.PrebuiltModuleFiles.insert_or_assign( - std::string(Split.first), std::string(Split.second)); + Opts.PrebuiltModuleFiles.insert_or_assign(std::string(Split.first), + std::string(Split.second)); } } for (const auto *A : Args.filtered(OPT_fprebuilt_module_path)) @@ -3551,7 +3622,7 @@ static bool ParseHeaderSearchArgs(HeaderSearchOptions &Opts, ArgList &Args, for (const auto *A : Args.filtered(OPT_cxx_isystem)) Opts.AddPath(A->getValue(), frontend::CXXSystem, false, true); for (const auto *A : Args.filtered(OPT_objc_isystem)) - Opts.AddPath(A->getValue(), frontend::ObjCSystem, false,true); + Opts.AddPath(A->getValue(), frontend::ObjCSystem, false, true); for (const auto *A : Args.filtered(OPT_objcxx_isystem)) Opts.AddPath(A->getValue(), frontend::ObjCXXSystem, false, true); @@ -3671,8 +3742,7 @@ static void ParsePointerAuthArgs(LangOptions &Opts, ArgList &Args, } /// Check if input file kind and language standard are compatible. -static bool IsInputCompatibleWithStandard(InputKind IK, - const LangStandard &S) { +static bool IsInputCompatibleWithStandard(InputKind IK, const LangStandard &S) { switch (IK.getLanguage()) { case Language::Unknown: case Language::LLVM_IR: @@ -4073,26 +4143,27 @@ bool CompilerInvocation::ParseLangArgs(LangOptions &Opts, ArgList &Args, LangStd = LangStandard::getLangKind(A->getValue()); if (LangStd == LangStandard::lang_unspecified) { Diags.Report(diag::err_drv_invalid_value) - << A->getAsString(Args) << A->getValue(); + << A->getAsString(Args) << A->getValue(); // Report supported standards with short description. - for (unsigned KindValue = 0; - KindValue != LangStandard::lang_unspecified; + for (unsigned KindValue = 0; KindValue != LangStandard::lang_unspecified; ++KindValue) { const LangStandard &Std = LangStandard::getLangStandardForKind( - static_cast(KindValue)); + static_cast(KindValue)); if (IsInputCompatibleWithStandard(IK, Std)) { auto Diag = Diags.Report(diag::note_drv_use_standard); Diag << Std.getName() << Std.getDescription(); unsigned NumAliases = 0; #define LANGSTANDARD(id, name, lang, desc, features, version) -#define LANGSTANDARD_ALIAS(id, alias) \ - if (KindValue == LangStandard::lang_##id) ++NumAliases; +#define LANGSTANDARD_ALIAS(id, alias) \ + if (KindValue == LangStandard::lang_##id) \ + ++NumAliases; #define LANGSTANDARD_ALIAS_DEPR(id, alias) #include "clang/Basic/LangStandards.def" Diag << NumAliases; #define LANGSTANDARD(id, name, lang, desc, features, version) -#define LANGSTANDARD_ALIAS(id, alias) \ - if (KindValue == LangStandard::lang_##id) Diag << alias; +#define LANGSTANDARD_ALIAS(id, alias) \ + if (KindValue == LangStandard::lang_##id) \ + Diag << alias; #define LANGSTANDARD_ALIAS_DEPR(id, alias) #include "clang/Basic/LangStandards.def" } @@ -4103,7 +4174,7 @@ bool CompilerInvocation::ParseLangArgs(LangOptions &Opts, ArgList &Args, const LangStandard &Std = LangStandard::getLangStandardForKind(LangStd); if (!IsInputCompatibleWithStandard(IK, Std)) { Diags.Report(diag::err_drv_argument_not_allowed_with) - << A->getAsString(Args) << GetInputKindName(IK); + << A->getAsString(Args) << GetInputKindName(IK); } } } @@ -4126,9 +4197,8 @@ bool CompilerInvocation::ParseLangArgs(LangOptions &Opts, ArgList &Args, if (OpenCLLangStd == LangStandard::lang_unspecified) { Diags.Report(diag::err_drv_invalid_value) - << A->getAsString(Args) << A->getValue(); - } - else + << A->getAsString(Args) << A->getValue(); + } else LangStd = OpenCLLangStd; } @@ -4225,7 +4295,7 @@ bool CompilerInvocation::ParseLangArgs(LangOptions &Opts, ArgList &Args, if (Args.hasArg(OPT_fobjc_subscripting_legacy_runtime)) Opts.ObjCSubscriptingLegacyRuntime = - (Opts.ObjCRuntime.getKind() == ObjCRuntime::FragileMacOSX); + (Opts.ObjCRuntime.getKind() == ObjCRuntime::FragileMacOSX); } if (Arg *A = Args.getLastArg(options::OPT_fgnuc_version_EQ)) { @@ -4251,8 +4321,7 @@ bool CompilerInvocation::ParseLangArgs(LangOptions &Opts, ArgList &Args, // Set the handler, if one is specified. Opts.OverflowHandler = std::string(Args.getLastArgValue(OPT_ftrapv_handler)); - } - else if (Args.hasArg(OPT_fwrapv)) + } else if (Args.hasArg(OPT_fwrapv)) Opts.setSignedOverflowBehavior(LangOptions::SOB_Defined); if (Args.hasArg(OPT_fwrapv_pointer)) Opts.PointerOverflowDefined = true; @@ -4261,8 +4330,8 @@ bool CompilerInvocation::ParseLangArgs(LangOptions &Opts, ArgList &Args, if (const Arg *A = Args.getLastArg(OPT_fms_compatibility_version)) { VersionTuple VT; if (VT.tryParse(A->getValue())) - Diags.Report(diag::err_drv_invalid_value) << A->getAsString(Args) - << A->getValue(); + Diags.Report(diag::err_drv_invalid_value) + << A->getAsString(Args) << A->getValue(); Opts.MSCompatibilityVersion = VT.getMajor() * 10000000 + VT.getMinor().value_or(0) * 100000 + VT.getSubminor().value_or(0); @@ -4281,8 +4350,8 @@ bool CompilerInvocation::ParseLangArgs(LangOptions &Opts, ArgList &Args, Opts.ZOSExt = Args.hasFlag(OPT_fzos_extensions, OPT_fno_zos_extensions, T.isOSzOS()); - Opts.Blocks = Args.hasArg(OPT_fblocks) || (Opts.OpenCL - && Opts.OpenCLVersion == 200); + Opts.Blocks = + Args.hasArg(OPT_fblocks) || (Opts.OpenCL && Opts.OpenCLVersion == 200); bool HasConvergentOperations = Opts.isTargetDevice() || Opts.OpenCL || Opts.HLSL || T.isAMDGPU() || T.isNVPTX(); @@ -4721,6 +4790,7 @@ static bool isStrictlyPreprocessorAction(frontend::ActionKind Action) { case frontend::ASTDump: case frontend::ASTPrint: case frontend::ASTView: + case frontend::CIRCombine: case frontend::EmitAssembly: case frontend::EmitBC: case frontend::EmitCIROnly: diff --git a/clang/lib/FrontendTool/ExecuteCompilerInvocation.cpp b/clang/lib/FrontendTool/ExecuteCompilerInvocation.cpp index 34c6ca11291c..3e262e0b23f4 100644 --- a/clang/lib/FrontendTool/ExecuteCompilerInvocation.cpp +++ b/clang/lib/FrontendTool/ExecuteCompilerInvocation.cpp @@ -35,6 +35,7 @@ #include "mlir/IR/MLIRContext.h" #include "mlir/Pass/PassManager.h" #include "clang/CIR/Dialect/Passes.h" +#include "clang/CIR/FrontendAction/CIRCombineAction.h" #include "clang/CIR/FrontendAction/CIRGenAction.h" #endif @@ -67,14 +68,26 @@ CreateFrontendBaseAction(CompilerInstance &CI) { "emitting of MLIR standard dialects"); switch (CI.getFrontendOpts().ProgramAction) { - case ASTDeclList: return std::make_unique(); - case ASTDump: return std::make_unique(); - case ASTPrint: return std::make_unique(); - case ASTView: return std::make_unique(); + case ASTDeclList: + return std::make_unique(); + case ASTDump: + return std::make_unique(); + case ASTPrint: + return std::make_unique(); + case ASTView: + return std::make_unique(); case DumpCompilerOptions: return std::make_unique(); - case DumpRawTokens: return std::make_unique(); - case DumpTokens: return std::make_unique(); + case DumpRawTokens: + return std::make_unique(); + case DumpTokens: + return std::make_unique(); + case CIRCombine: +#if CLANG_ENABLE_CIR + return std::make_unique(); +#else + llvm_unreachable("CIR suppport not built into clang"); +#endif case EmitAssembly: #if CLANG_ENABLE_CIR if (UseCIR) @@ -102,7 +115,8 @@ CreateFrontendBaseAction(CompilerInstance &CI) { case EmitMLIR: llvm_unreachable("CIR suppport not built into clang"); #endif - case EmitHTML: return std::make_unique(); + case EmitHTML: + return std::make_unique(); case EmitLLVM: { #if CLANG_ENABLE_CIR if (UseCIR) @@ -137,7 +151,8 @@ CreateFrontendBaseAction(CompilerInstance &CI) { } case ExtractAPI: return std::make_unique(); - case FixIt: return std::make_unique(); + case FixIt: + return std::make_unique(); case GenerateModule: return std::make_unique(); case GenerateModuleInterface: @@ -146,14 +161,20 @@ CreateFrontendBaseAction(CompilerInstance &CI) { return std::make_unique(); case GenerateHeaderUnit: return std::make_unique(); - case GeneratePCH: return std::make_unique(); + case GeneratePCH: + return std::make_unique(); case GenerateInterfaceStubs: return std::make_unique(); - case InitOnly: return std::make_unique(); - case ParseSyntaxOnly: return std::make_unique(); - case ModuleFileInfo: return std::make_unique(); - case VerifyPCH: return std::make_unique(); - case TemplightDump: return std::make_unique(); + case InitOnly: + return std::make_unique(); + case ParseSyntaxOnly: + return std::make_unique(); + case ModuleFileInfo: + return std::make_unique(); + case VerifyPCH: + return std::make_unique(); + case TemplightDump: + return std::make_unique(); case PluginAction: { for (const FrontendPluginRegistry::entry &Plugin : @@ -171,11 +192,12 @@ CreateFrontendBaseAction(CompilerInstance &CI) { } CI.getDiagnostics().Report(diag::err_fe_invalid_plugin_name) - << CI.getFrontendOpts().ActionName; + << CI.getFrontendOpts().ActionName; return nullptr; } - case PrintPreamble: return std::make_unique(); + case PrintPreamble: + return std::make_unique(); case PrintPreprocessedInput: { if (CI.getPreprocessorOutputOpts().RewriteIncludes || CI.getPreprocessorOutputOpts().RewriteImports) @@ -183,19 +205,28 @@ CreateFrontendBaseAction(CompilerInstance &CI) { return std::make_unique(); } - case RewriteMacros: return std::make_unique(); - case RewriteTest: return std::make_unique(); + case RewriteMacros: + return std::make_unique(); + case RewriteTest: + return std::make_unique(); #if CLANG_ENABLE_OBJC_REWRITER - case RewriteObjC: return std::make_unique(); + case RewriteObjC: + return std::make_unique(); #else - case RewriteObjC: Action = "RewriteObjC"; break; + case RewriteObjC: + Action = "RewriteObjC"; + break; #endif #if CLANG_ENABLE_STATIC_ANALYZER - case RunAnalysis: return std::make_unique(); + case RunAnalysis: + return std::make_unique(); #else - case RunAnalysis: Action = "RunAnalysis"; break; + case RunAnalysis: + Action = "RunAnalysis"; + break; #endif - case RunPreprocessorOnly: return std::make_unique(); + case RunPreprocessorOnly: + return std::make_unique(); case PrintDependencyDirectivesSourceMinimizerOutput: return std::make_unique(); } @@ -208,8 +239,7 @@ CreateFrontendBaseAction(CompilerInstance &CI) { #endif } -std::unique_ptr -CreateFrontendAction(CompilerInstance &CI) { +std::unique_ptr CreateFrontendAction(CompilerInstance &CI) { // Create the underlying action. std::unique_ptr Act = CreateFrontendBaseAction(CI); if (!Act) @@ -239,8 +269,8 @@ CreateFrontendAction(CompilerInstance &CI) { // If there are any AST files to merge, create a frontend action // adaptor to perform the merge. if (!FEOpts.ASTMergeFiles.empty()) - Act = std::make_unique(std::move(Act), - FEOpts.ASTMergeFiles); + Act = + std::make_unique(std::move(Act), FEOpts.ASTMergeFiles); return Act; } @@ -272,7 +302,7 @@ bool ExecuteCompilerInvocation(CompilerInstance *Clang) { // This should happen AFTER plugins have been loaded! if (!Clang->getFrontendOpts().LLVMArgs.empty()) { unsigned NumArgs = Clang->getFrontendOpts().LLVMArgs.size(); - auto Args = std::make_unique(NumArgs + 2); + auto Args = std::make_unique(NumArgs + 2); Args[0] = "clang (LLVM option parsing)"; for (unsigned i = 0; i != NumArgs; ++i) Args[i + 1] = Clang->getFrontendOpts().LLVMArgs[i].c_str(); diff --git a/clang/test/CIR/CodeGen/HIP/simple.cpp b/clang/test/CIR/CodeGen/HIP/simple.cpp index 2b17c3ef396b..4ed4013140d3 100644 --- a/clang/test/CIR/CodeGen/HIP/simple.cpp +++ b/clang/test/CIR/CodeGen/HIP/simple.cpp @@ -16,7 +16,8 @@ // RUN: FileCheck --check-prefix=LLVM-HOST --input-file=%t.ll %s // RUN: %clang_cc1 -triple=amdgcn-amd-amdhsa -x hip -fclangir \ -// RUN: -fcuda-is-device -fhip-new-launch-api \ +// RUN: -fcuda-is-device -fhip-new-launch-api \ +// RUN: -fvisibility=hidden \ // RUN: -I%S/../Inputs/ -emit-llvm %s -o %t.ll // RUN: FileCheck --check-prefix=LLVM-DEVICE --input-file=%t.ll %s @@ -27,6 +28,7 @@ // RUN: %clang_cc1 -triple=amdgcn-amd-amdhsa -x hip \ // RUN: -fcuda-is-device -fhip-new-launch-api \ +// RUN: -fvisibility=hidden \ // RUN: -I%S/../Inputs/ -emit-llvm %s -o %t.ll // RUN: FileCheck --check-prefix=OGCG-DEVICE --input-file=%t.ll %s @@ -65,8 +67,8 @@ __device__ void device_fn(int* a, double b, float c) {} __global__ void global_fn(int a) {} // CIR-DEVICE: @_Z9global_fni{{.*}} cc(amdgpu_kernel) -// LLVM-DEVICE: define dso_local amdgpu_kernel void @_Z9global_fni -// OGCG-DEVICE: define dso_local amdgpu_kernel void @_Z9global_fni +// LLVM-DEVICE: define protected amdgpu_kernel void @_Z9global_fni +// OGCG-DEVICE: define protected amdgpu_kernel void @_Z9global_fni // CIR-HOST: @_Z24__device_stub__global_fni{{.*}}extra([[Kernel]]) // CIR-HOST: %[[#CIRKernelArgs:]] = cir.alloca {{.*}}"kernel_args" diff --git a/clang/test/CIR/Driver/combine-cir.cpp b/clang/test/CIR/Driver/combine-cir.cpp new file mode 100644 index 000000000000..7503cf9db49f --- /dev/null +++ b/clang/test/CIR/Driver/combine-cir.cpp @@ -0,0 +1,138 @@ +// NOTE: This file exists only to host RUN lines; -cir-combine should reject +// positional inputs, so we never pass %s as an input. +// +// RUN: rm -rf %t && mkdir -p %t +// RUN: printf "module {}\n" > %t/host.cir +// RUN: printf "module {}\n" > %t/device.cir + +//------------------------------------------------------------------------------ +// Positive / baseline: all required args present. +//------------------------------------------------------------------------------ +// +// RUN: %clang_cc1 -fclangir -cir-combine \ +// RUN: -cir-host-input %t/host.cir \ +// RUN: -cir-device-input %t/device.cir \ +// RUN: -o %t/out.cir +// RUN: test -f %t/out.cir +// RUN: FileCheck %s --check-prefix=OK < %t/out.cir +// +// OK: cir.offload.container + +//------------------------------------------------------------------------------ +// Missing -o +//------------------------------------------------------------------------------ +// +// RUN: not %clang_cc1 -fclangir -cir-combine \ +// RUN: -cir-host-input %t/host.cir \ +// RUN: -cir-device-input %t/device.cir 2>&1 \ +// RUN: | FileCheck %s --check-prefix=NO_O +// +// NO_O: error: missing argument to '-o' + +//------------------------------------------------------------------------------ +// Missing host input +//------------------------------------------------------------------------------ +// +// RUN: not %clang_cc1 -fclangir -cir-combine \ +// RUN: -cir-device-input %t/device.cir \ +// RUN: -o %t/out.cir 2>&1 \ +// RUN: | FileCheck %s --check-prefix=NO_HOST +// +// NO_HOST: error: +// NO_HOST-SAME: -cir-host-input + +//------------------------------------------------------------------------------ +// Missing device input(s) +//------------------------------------------------------------------------------ +// +// RUN: not %clang_cc1 -fclangir -cir-combine \ +// RUN: -cir-host-input %t/host.cir \ +// RUN: -o %t/out.cir 2>&1 \ +// RUN: | FileCheck %s --check-prefix=NO_DEVICE +// +// NO_DEVICE: error: +// NO_DEVICE-SAME: -cir-device-input + +//------------------------------------------------------------------------------ +// Multiple host inputs (should be rejected) +//------------------------------------------------------------------------------ +// +// RUN: not %clang_cc1 -fclangir -cir-combine \ +// RUN: -cir-host-input %t/host.cir \ +// RUN: -cir-host-input %t/host.cir \ +// RUN: -cir-device-input %t/device.cir \ +// RUN: -o %t/out.cir 2>&1 \ +// RUN: | FileCheck %s --check-prefix=MULTI_HOST +// +// MULTI_HOST: error: +// MULTI_HOST-SAME: -cir-host-input + +//------------------------------------------------------------------------------ +// Non-existent host file +//------------------------------------------------------------------------------ +// +// RUN: not %clang_cc1 -fclangir -cir-combine \ +// RUN: -cir-host-input %t/does-not-exist.cir \ +// RUN: -cir-device-input %t/device.cir \ +// RUN: -o %t/out.cir 2>&1 \ +// RUN: | FileCheck %s --check-prefix=NO_SUCH_HOST +// +// NO_SUCH_HOST: error: no such file or directory: + +//------------------------------------------------------------------------------ +// Non-existent device file +//------------------------------------------------------------------------------ +// +// RUN: not %clang_cc1 -fclangir -cir-combine \ +// RUN: -cir-host-input %t/host.cir \ +// RUN: -cir-device-input %t/does-not-exist.cir \ +// RUN: -o %t/out.cir 2>&1 \ +// RUN: | FileCheck %s --check-prefix=NO_SUCH_DEVICE +// +// NO_SUCH_DEVICE: error: no such file or directory: + +//------------------------------------------------------------------------------ +// CIR split outputs to host/device +//------------------------------------------------------------------------------ +// RUN: %clang_cc1 -fclangir -cir-combine -cir-emit-split \ +// RUN: -cir-host-input %t/host.cir \ +// RUN: -cir-device-input %t/device.cir \ +// RUN: -cir-host-output %t/host.out.cir \ +// RUN: -cir-device-output %t/device.out.cir +// RUN: test -f %t/host.out.cir +// RUN: test -f %t/device.out.cir + +//------------------------------------------------------------------------------ +// CIR split outputs missing device output +//------------------------------------------------------------------------------ +// RUN: not %clang_cc1 -fclangir -cir-combine -cir-emit-split \ +// RUN: -cir-host-input %t/host.cir \ +// RUN: -cir-device-input %t/device.cir \ +// RUN: -cir-host-output %t/host.out.cir 2>&1 \ +// RUN: | FileCheck %s --check-prefix=SPLIT_MISSING_DEVOUT +// SPLIT_MISSING_DEVOUT: error: argument to '-cir-device-output' is missing (expected 1 value) + + +//------------------------------------------------------------------------------ +// CIR split outputs missing host output +//------------------------------------------------------------------------------ +// RUN: not %clang_cc1 -fclangir -cir-combine -cir-emit-split \ +// RUN: -cir-host-input %t/host.cir \ +// RUN: -cir-device-input %t/device.cir \ +// RUN: -cir-device-output %t/device.out.cir 2>&1 \ +// RUN: | FileCheck %s --check-prefix=SPLIT_MISSING_HOSTOUT +// SPLIT_MISSING_HOSTOUT: error: argument to '-cir-host-output' is missing (expected 1 value) + + +//------------------------------------------------------------------------------ +// CIR split outputs to host/device, pass unused '-o' +//------------------------------------------------------------------------------ +// RUN: %clang_cc1 -fclangir -cir-combine -cir-emit-split \ +// RUN: -cir-host-input %t/host.cir -o %t/out.cir \ +// RUN: -cir-device-input %t/device.cir \ +// RUN: -cir-host-output %t/host.out.cir \ +// RUN: -cir-device-output %t/device.out.cir 2>&1 \ +// RUN: | FileCheck %s --check-prefix=EXTRA_OUT +// RUN: test -f %t/host.out.cir +// RUN: test -f %t/device.out.cir +// EXTRA_OUT: warning: ignoring '-o' option as option '-cir-emit-split' overrides the behavior