From d99d3a73fed79250860fca1db4b09ae8869f64e9 Mon Sep 17 00:00:00 2001 From: Zahira Ammarguellat Date: Wed, 31 Mar 2021 08:33:04 -0700 Subject: [PATCH 01/19] Implementing opt-report Signed-off-by: Zahira Ammarguellat --- clang/include/clang/Basic/Diagnostic.h | 3 + .../clang/Basic/SyclOptReportHandler.h | 59 + clang/lib/CodeGen/CGSYCLRuntime.cpp | 2 + clang/lib/CodeGen/CodeGenAction.cpp.mine | 1248 +++++++++++++++++ clang/lib/CodeGen/CodeGenFunction.cpp | 241 ++-- clang/lib/Sema/SemaSYCL.cpp | 45 +- 6 files changed, 1486 insertions(+), 112 deletions(-) create mode 100644 clang/include/clang/Basic/SyclOptReportHandler.h create mode 100644 clang/lib/CodeGen/CodeGenAction.cpp.mine diff --git a/clang/include/clang/Basic/Diagnostic.h b/clang/include/clang/Basic/Diagnostic.h index a3cc3af5a74a4..12bdf370773e4 100644 --- a/clang/include/clang/Basic/Diagnostic.h +++ b/clang/include/clang/Basic/Diagnostic.h @@ -18,6 +18,7 @@ #include "clang/Basic/DiagnosticOptions.h" #include "clang/Basic/SourceLocation.h" #include "clang/Basic/Specifiers.h" +#include "clang/Basic/SyclOptReportHandler.h" #include "llvm/ADT/ArrayRef.h" #include "llvm/ADT/DenseMap.h" #include "llvm/ADT/IntrusiveRefCntPtr.h" @@ -540,6 +541,8 @@ class DiagnosticsEngine : public RefCountedBase { LLVM_DUMP_METHOD void dump() const; LLVM_DUMP_METHOD void dump(StringRef DiagName) const; + SyclOptReportHandler OptReportHandler; + const IntrusiveRefCntPtr &getDiagnosticIDs() const { return Diags; } diff --git a/clang/include/clang/Basic/SyclOptReportHandler.h b/clang/include/clang/Basic/SyclOptReportHandler.h new file mode 100644 index 0000000000000..9d9bdbc8ba713 --- /dev/null +++ b/clang/include/clang/Basic/SyclOptReportHandler.h @@ -0,0 +1,59 @@ +//===---------------------- SyclOptReport.h ---------------------*- C++ -*-===// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// +/// +/// \file +/// Defines clang::SyclOptReport class. +/// +//===----------------------------------------------------------------------===// + +#ifndef LLVM_CLANG_BASIC_SYCLOPTREPORTHANDLER_H +#define LLVM_CLANG_BASIC_SYCLOPTREPORTHANDLER_H + +#include "clang/Basic/SourceLocation.h" +#include "llvm/ADT/DenseMap.h" +#include "llvm/ADT/StringRef.h" + +namespace clang { + +class FunctionDecl; + +class SyclOptReportHandler { +private: + struct OptReportInfo { + StringRef KernelArgName; + std::string KernelArgType; + SourceLocation KernelArgLoc; + + OptReportInfo(llvm::StringRef ArgName, std::string ArgType, + SourceLocation ArgLoc) : + KernelArgName(ArgName), + KernelArgType(ArgType), + KernelArgLoc(ArgLoc) {} + }; + llvm::DenseMap> Map; + +public: + void AddKernelArgs(const FunctionDecl *FD, StringRef ArgName, + std::string ArgType, SourceLocation ArgLoc) { + Map[FD].emplace_back(ArgName, ArgType, ArgLoc); + auto It = Map.find(FD); + } + SmallVector &getInfo(const FunctionDecl *FD) { + auto It = Map.find(FD); + assert(It != Map.end()); + return It->second; + } + + bool HasOptReportInfo(const FunctionDecl *FD) { + return Map.find(FD) != Map.end(); + } +}; + +} // namespace clang + +#endif // LLVM_CLANG_BASIC_SYCLOPTREPORTHANDLER_H diff --git a/clang/lib/CodeGen/CGSYCLRuntime.cpp b/clang/lib/CodeGen/CGSYCLRuntime.cpp index 1cdf825838547..03aebb500c52c 100644 --- a/clang/lib/CodeGen/CGSYCLRuntime.cpp +++ b/clang/lib/CodeGen/CGSYCLRuntime.cpp @@ -15,6 +15,8 @@ #include "clang/AST/Attr.h" #include "clang/AST/Decl.h" #include "llvm/IR/Instructions.h" +#include "llvm/Analysis/OptimizationRemarkEmitter.h" +#include "clang/Basic/SourceLocation.h" #include using namespace clang; diff --git a/clang/lib/CodeGen/CodeGenAction.cpp.mine b/clang/lib/CodeGen/CodeGenAction.cpp.mine new file mode 100644 index 0000000000000..7948c1244737d --- /dev/null +++ b/clang/lib/CodeGen/CodeGenAction.cpp.mine @@ -0,0 +1,1248 @@ +//===--- CodeGenAction.cpp - LLVM Code Generation Frontend Action ---------===// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// + +#include "clang/CodeGen/CodeGenAction.h" +#include "CodeGenModule.h" +#include "CoverageMappingGen.h" +#include "MacroPPCallbacks.h" +#include "clang/AST/ASTConsumer.h" +#include "clang/AST/ASTContext.h" +#include "clang/AST/DeclCXX.h" +#include "clang/AST/DeclGroup.h" +#include "clang/Basic/DiagnosticFrontend.h" +#include "clang/Basic/FileManager.h" +#include "clang/Basic/LangStandard.h" +#include "clang/Basic/SourceManager.h" +#include "clang/Basic/TargetInfo.h" +#include "clang/CodeGen/BackendUtil.h" +#include "clang/CodeGen/ModuleBuilder.h" +#include "clang/Driver/DriverDiagnostic.h" +#include "clang/Frontend/CompilerInstance.h" +#include "clang/Frontend/FrontendDiagnostic.h" +#include "clang/Lex/Preprocessor.h" +#include "llvm/Bitcode/BitcodeReader.h" +#include "llvm/CodeGen/MachineOptimizationRemarkEmitter.h" +#include "llvm/IR/DebugInfo.h" +#include "llvm/IR/DiagnosticInfo.h" +#include "llvm/IR/DiagnosticPrinter.h" +#include "llvm/IR/GlobalValue.h" +#include "llvm/IR/LLVMContext.h" +#include "llvm/IR/LLVMRemarkStreamer.h" +#include "llvm/IR/LegacyPassManager.h" +#include "llvm/IR/Module.h" +#include "llvm/IRReader/IRReader.h" +#include "llvm/LTO/LTOBackend.h" +#include "llvm/Linker/Linker.h" +#include "llvm/Pass.h" +#include "llvm/SYCLLowerIR/LowerWGScope.h" +#include "llvm/Support/MemoryBuffer.h" +#include "llvm/Support/SourceMgr.h" +#include "llvm/Support/TimeProfiler.h" +#include "llvm/Support/Timer.h" +#include "llvm/Support/ToolOutputFile.h" +#include "llvm/Support/YAMLTraits.h" +#include "llvm/Transforms/IPO/Internalize.h" + +#include +using namespace clang; +using namespace llvm; + +namespace clang { + class BackendConsumer; + class ClangDiagnosticHandler final : public DiagnosticHandler { + public: + ClangDiagnosticHandler(const CodeGenOptions &CGOpts, BackendConsumer *BCon) + : CodeGenOpts(CGOpts), BackendCon(BCon) {} + + bool handleDiagnostics(const DiagnosticInfo &DI) override; + + bool isAnalysisRemarkEnabled(StringRef PassName) const override { + return (CodeGenOpts.OptimizationRemarkAnalysisPattern && + CodeGenOpts.OptimizationRemarkAnalysisPattern->match(PassName)); + } + bool isMissedOptRemarkEnabled(StringRef PassName) const override { + return (CodeGenOpts.OptimizationRemarkMissedPattern && + CodeGenOpts.OptimizationRemarkMissedPattern->match(PassName)); + } + bool isPassedOptRemarkEnabled(StringRef PassName) const override { + return (CodeGenOpts.OptimizationRemarkPattern && + CodeGenOpts.OptimizationRemarkPattern->match(PassName)); + } + + bool isAnyRemarkEnabled() const override { + return (CodeGenOpts.OptimizationRemarkAnalysisPattern || + CodeGenOpts.OptimizationRemarkMissedPattern || + CodeGenOpts.OptimizationRemarkPattern); + } + + private: + const CodeGenOptions &CodeGenOpts; + BackendConsumer *BackendCon; + }; + + static void reportOptRecordError(Error E, DiagnosticsEngine &Diags, + const CodeGenOptions CodeGenOpts) { + handleAllErrors( + std::move(E), + [&](const LLVMRemarkSetupFileError &E) { + Diags.Report(diag::err_cannot_open_file) + << CodeGenOpts.OptRecordFile << E.message(); + }, + [&](const LLVMRemarkSetupPatternError &E) { + Diags.Report(diag::err_drv_optimization_remark_pattern) + << E.message() << CodeGenOpts.OptRecordPasses; + }, + [&](const LLVMRemarkSetupFormatError &E) { + Diags.Report(diag::err_drv_optimization_remark_format) + << CodeGenOpts.OptRecordFormat; + }); + } + + class BackendConsumer : public ASTConsumer { + using LinkModule = CodeGenAction::LinkModule; + + virtual void anchor(); + DiagnosticsEngine &Diags; + BackendAction Action; + const HeaderSearchOptions &HeaderSearchOpts; + const CodeGenOptions &CodeGenOpts; + const TargetOptions &TargetOpts; + const LangOptions &LangOpts; + std::unique_ptr AsmOutStream; + ASTContext *Context; + + Timer LLVMIRGeneration; + unsigned LLVMIRGenerationRefCount; + + /// True if we've finished generating IR. This prevents us from generating + /// additional LLVM IR after emitting output in HandleTranslationUnit. This + /// can happen when Clang plugins trigger additional AST deserialization. + bool IRGenFinished = false; + + bool TimerIsEnabled = false; + + std::unique_ptr Gen; + + SmallVector LinkModules; + + // This is here so that the diagnostic printer knows the module a diagnostic + // refers to. + llvm::Module *CurLinkModule = nullptr; + + public: + BackendConsumer(BackendAction Action, DiagnosticsEngine &Diags, + const HeaderSearchOptions &HeaderSearchOpts, + const PreprocessorOptions &PPOpts, + const CodeGenOptions &CodeGenOpts, + const TargetOptions &TargetOpts, + const LangOptions &LangOpts, const std::string &InFile, + SmallVector LinkModules, + std::unique_ptr OS, LLVMContext &C, + CoverageSourceInfo *CoverageInfo = nullptr) + : Diags(Diags), Action(Action), HeaderSearchOpts(HeaderSearchOpts), + CodeGenOpts(CodeGenOpts), TargetOpts(TargetOpts), LangOpts(LangOpts), + AsmOutStream(std::move(OS)), Context(nullptr), + LLVMIRGeneration("irgen", "LLVM IR Generation Time"), + LLVMIRGenerationRefCount(0), + Gen(CreateLLVMCodeGen(Diags, InFile, HeaderSearchOpts, PPOpts, + CodeGenOpts, C, CoverageInfo)), + LinkModules(std::move(LinkModules)) { + TimerIsEnabled = CodeGenOpts.TimePasses; + llvm::TimePassesIsEnabled = CodeGenOpts.TimePasses; + llvm::TimePassesPerRun = CodeGenOpts.TimePassesPerRun; + } + + // This constructor is used in installing an empty BackendConsumer + // to use the clang diagnostic handler for IR input files. It avoids + // initializing the OS field. + BackendConsumer(BackendAction Action, DiagnosticsEngine &Diags, + const HeaderSearchOptions &HeaderSearchOpts, + const PreprocessorOptions &PPOpts, + const CodeGenOptions &CodeGenOpts, + const TargetOptions &TargetOpts, + const LangOptions &LangOpts, + SmallVector LinkModules, LLVMContext &C, + CoverageSourceInfo *CoverageInfo = nullptr) + : Diags(Diags), Action(Action), HeaderSearchOpts(HeaderSearchOpts), + CodeGenOpts(CodeGenOpts), TargetOpts(TargetOpts), LangOpts(LangOpts), + Context(nullptr), + LLVMIRGeneration("irgen", "LLVM IR Generation Time"), + LLVMIRGenerationRefCount(0), + Gen(CreateLLVMCodeGen(Diags, "", HeaderSearchOpts, PPOpts, + CodeGenOpts, C, CoverageInfo)), + LinkModules(std::move(LinkModules)) { + TimerIsEnabled = CodeGenOpts.TimePasses; + llvm::TimePassesIsEnabled = CodeGenOpts.TimePasses; + llvm::TimePassesPerRun = CodeGenOpts.TimePassesPerRun; + } + llvm::Module *getModule() const { return Gen->GetModule(); } + std::unique_ptr takeModule() { + return std::unique_ptr(Gen->ReleaseModule()); + } + + CodeGenerator *getCodeGenerator() { return Gen.get(); } + + void HandleCXXStaticMemberVarInstantiation(VarDecl *VD) override { + Gen->HandleCXXStaticMemberVarInstantiation(VD); + } + + void Initialize(ASTContext &Ctx) override { + assert(!Context && "initialized multiple times"); + + Context = &Ctx; + + if (TimerIsEnabled) + LLVMIRGeneration.startTimer(); + + Gen->Initialize(Ctx); + + if (TimerIsEnabled) + LLVMIRGeneration.stopTimer(); + } + + bool HandleTopLevelDecl(DeclGroupRef D) override { + PrettyStackTraceDecl CrashInfo(*D.begin(), SourceLocation(), + Context->getSourceManager(), + "LLVM IR generation of declaration"); + + // Recurse. + if (TimerIsEnabled) { + LLVMIRGenerationRefCount += 1; + if (LLVMIRGenerationRefCount == 1) + LLVMIRGeneration.startTimer(); + } + + Gen->HandleTopLevelDecl(D); + + if (TimerIsEnabled) { + LLVMIRGenerationRefCount -= 1; + if (LLVMIRGenerationRefCount == 0) + LLVMIRGeneration.stopTimer(); + } + + return true; + } + + void HandleInlineFunctionDefinition(FunctionDecl *D) override { + PrettyStackTraceDecl CrashInfo(D, SourceLocation(), + Context->getSourceManager(), + "LLVM IR generation of inline function"); + if (TimerIsEnabled) + LLVMIRGeneration.startTimer(); + + Gen->HandleInlineFunctionDefinition(D); + + if (TimerIsEnabled) + LLVMIRGeneration.stopTimer(); + } + + void HandleInterestingDecl(DeclGroupRef D) override { + // Ignore interesting decls from the AST reader after IRGen is finished. + if (!IRGenFinished) + HandleTopLevelDecl(D); + } + + // Links each entry in LinkModules into our module. Returns true on error. + bool LinkInModules() { + for (auto &LM : LinkModules) { + if (LM.PropagateAttrs) + for (Function &F : *LM.Module) { + // Skip intrinsics. Keep consistent with how intrinsics are created + // in LLVM IR. + if (F.isIntrinsic()) + continue; + Gen->CGM().addDefaultFunctionDefinitionAttributes(F); + } + + CurLinkModule = LM.Module.get(); + + bool Err; + if (LM.Internalize) { + Err = Linker::linkModules( + *getModule(), std::move(LM.Module), LM.LinkFlags, + [](llvm::Module &M, const llvm::StringSet<> &GVS) { + internalizeModule(M, [&GVS](const llvm::GlobalValue &GV) { + return !GV.hasName() || (GVS.count(GV.getName()) == 0); + }); + }); + } else { + Err = Linker::linkModules(*getModule(), std::move(LM.Module), + LM.LinkFlags); + } + + if (Err) + return true; + } + return false; // success + } + + void HandleTranslationUnit(ASTContext &C) override { + { + llvm::TimeTraceScope TimeScope("Frontend"); + PrettyStackTraceString CrashInfo("Per-file LLVM IR generation"); + if (TimerIsEnabled) { + LLVMIRGenerationRefCount += 1; + if (LLVMIRGenerationRefCount == 1) + LLVMIRGeneration.startTimer(); + } + + Gen->HandleTranslationUnit(C); + + if (TimerIsEnabled) { + LLVMIRGenerationRefCount -= 1; + if (LLVMIRGenerationRefCount == 0) + LLVMIRGeneration.stopTimer(); + } + + IRGenFinished = true; + } + + // Silently ignore if we weren't initialized for some reason. + if (!getModule()) + return; + + // Install an inline asm handler so that diagnostics get printed through + // our diagnostics hooks. + LLVMContext &Ctx = getModule()->getContext(); + LLVMContext::InlineAsmDiagHandlerTy OldHandler = + Ctx.getInlineAsmDiagnosticHandler(); + void *OldContext = Ctx.getInlineAsmDiagnosticContext(); + Ctx.setInlineAsmDiagnosticHandler(InlineAsmDiagHandler, this); +#if !INTEL_CUSTOMIZATION + std::unique_ptr \ + OldDiagnosticHandler = + Ctx.getDiagnosticHandler(); + Ctx.setDiagnosticHandler(std::make_unique( + CodeGenOpts, this)); + + Expected> OptRecordFileOrErr = + setupLLVMOptimizationRemarks( + Ctx, CodeGenOpts.OptRecordFile, CodeGenOpts.OptRecordPasses, + CodeGenOpts.OptRecordFormat, CodeGenOpts.DiagnosticsWithHotness, + CodeGenOpts.DiagnosticsHotnessThreshold); + + if (Error E = OptRecordFileOrErr.takeError()) { + reportOptRecordError(std::move(E), Diags, CodeGenOpts); + return; + } + + std::unique_ptr OptRecordFile = + std::move(*OptRecordFileOrErr); + + if (OptRecordFile && + CodeGenOpts.getProfileUse() != CodeGenOptions::ProfileNone) + Ctx.setDiagnosticsHotnessRequested(true); +#endif // INTEL_CUSTOMIZATION + + // The parallel_for_work_group legalization pass can emit calls to + // builtins function. Definitions of those builtins can be provided in + // LinkModule. We force the pass to legalize the code before the link + // happens. + if (LangOpts.SYCLIsDevice) { + PrettyStackTraceString CrashInfo("Pre-linking SYCL passes"); + legacy::PassManager PreLinkingSyclPasses; + PreLinkingSyclPasses.add(llvm::createSYCLLowerWGScopePass()); + PreLinkingSyclPasses.run(*getModule()); + } + + // Link each LinkModule into our module. + if (LinkInModules()) + return; + + EmbedBitcode(getModule(), CodeGenOpts, llvm::MemoryBufferRef()); + + EmitBackendOutput(Diags, HeaderSearchOpts, CodeGenOpts, TargetOpts, + LangOpts, C.getTargetInfo().getDataLayout(), + getModule(), Action, std::move(AsmOutStream)); + + Ctx.setInlineAsmDiagnosticHandler(OldHandler, OldContext); +#if !INTEL_CUSTOMIZATION + Ctx.setDiagnosticHandler(std::move(OldDiagnosticHandler)); + + if (OptRecordFile) + OptRecordFile->keep(); +#endif + } + + void HandleTagDeclDefinition(TagDecl *D) override { + PrettyStackTraceDecl CrashInfo(D, SourceLocation(), + Context->getSourceManager(), + "LLVM IR generation of declaration"); + Gen->HandleTagDeclDefinition(D); + } + + void HandleTagDeclRequiredDefinition(const TagDecl *D) override { + Gen->HandleTagDeclRequiredDefinition(D); + } + + void CompleteTentativeDefinition(VarDecl *D) override { + Gen->CompleteTentativeDefinition(D); + } + + void CompleteExternalDeclaration(VarDecl *D) override { + Gen->CompleteExternalDeclaration(D); + } + + void AssignInheritanceModel(CXXRecordDecl *RD) override { + Gen->AssignInheritanceModel(RD); + } + + void HandleVTable(CXXRecordDecl *RD) override { + Gen->HandleVTable(RD); + } + + static void InlineAsmDiagHandler(const llvm::SMDiagnostic &SM,void *Context, + unsigned LocCookie) { + SourceLocation Loc = SourceLocation::getFromRawEncoding(LocCookie); + ((BackendConsumer*)Context)->InlineAsmDiagHandler2(SM, Loc); + } + + /// Get the best possible source location to represent a diagnostic that + /// may have associated debug info. + const FullSourceLoc + getBestLocationFromDebugLoc(const llvm::DiagnosticInfoWithLocationBase &D, + bool &BadDebugInfo, StringRef &Filename, + unsigned &Line, unsigned &Column) const; + + void InlineAsmDiagHandler2(const llvm::SMDiagnostic &, + SourceLocation LocCookie); + + void DiagnosticHandlerImpl(const llvm::DiagnosticInfo &DI); + /// Specialized handler for InlineAsm diagnostic. + /// \return True if the diagnostic has been successfully reported, false + /// otherwise. + bool InlineAsmDiagHandler(const llvm::DiagnosticInfoInlineAsm &D); + /// Specialized handler for StackSize diagnostic. + /// \return True if the diagnostic has been successfully reported, false + /// otherwise. + bool StackSizeDiagHandler(const llvm::DiagnosticInfoStackSize &D); + /// Specialized handler for unsupported backend feature diagnostic. + void UnsupportedDiagHandler(const llvm::DiagnosticInfoUnsupported &D); + /// Specialized handlers for optimization remarks. + /// Note that these handlers only accept remarks and they always handle + /// them. + void EmitOptimizationMessage(const llvm::DiagnosticInfoOptimizationBase &D, + unsigned DiagID); + void + OptimizationRemarkHandler(const llvm::DiagnosticInfoOptimizationBase &D); + void OptimizationRemarkHandler( + const llvm::OptimizationRemarkAnalysisFPCommute &D); + void OptimizationRemarkHandler( + const llvm::OptimizationRemarkAnalysisAliasing &D); + void OptimizationFailureHandler( + const llvm::DiagnosticInfoOptimizationFailure &D); + }; + + void BackendConsumer::anchor() {} +} + +bool ClangDiagnosticHandler::handleDiagnostics(const DiagnosticInfo &DI) { + BackendCon->DiagnosticHandlerImpl(DI); + return true; +} + +/// ConvertBackendLocation - Convert a location in a temporary llvm::SourceMgr +/// buffer to be a valid FullSourceLoc. +static FullSourceLoc ConvertBackendLocation(const llvm::SMDiagnostic &D, + SourceManager &CSM) { + // Get both the clang and llvm source managers. The location is relative to + // a memory buffer that the LLVM Source Manager is handling, we need to add + // a copy to the Clang source manager. + const llvm::SourceMgr &LSM = *D.getSourceMgr(); + + // We need to copy the underlying LLVM memory buffer because llvm::SourceMgr + // already owns its one and clang::SourceManager wants to own its one. + const MemoryBuffer *LBuf = + LSM.getMemoryBuffer(LSM.FindBufferContainingLoc(D.getLoc())); + + // Create the copy and transfer ownership to clang::SourceManager. + // TODO: Avoid copying files into memory. + std::unique_ptr CBuf = + llvm::MemoryBuffer::getMemBufferCopy(LBuf->getBuffer(), + LBuf->getBufferIdentifier()); + // FIXME: Keep a file ID map instead of creating new IDs for each location. + FileID FID = CSM.createFileID(std::move(CBuf)); + + // Translate the offset into the file. + unsigned Offset = D.getLoc().getPointer() - LBuf->getBufferStart(); + SourceLocation NewLoc = + CSM.getLocForStartOfFile(FID).getLocWithOffset(Offset); + return FullSourceLoc(NewLoc, CSM); +} + + +/// InlineAsmDiagHandler2 - This function is invoked when the backend hits an +/// error parsing inline asm. The SMDiagnostic indicates the error relative to +/// the temporary memory buffer that the inline asm parser has set up. +void BackendConsumer::InlineAsmDiagHandler2(const llvm::SMDiagnostic &D, + SourceLocation LocCookie) { + // There are a couple of different kinds of errors we could get here. First, + // we re-format the SMDiagnostic in terms of a clang diagnostic. + + // Strip "error: " off the start of the message string. + StringRef Message = D.getMessage(); + if (Message.startswith("error: ")) + Message = Message.substr(7); + + // If the SMDiagnostic has an inline asm source location, translate it. + FullSourceLoc Loc; + if (D.getLoc() != SMLoc()) + Loc = ConvertBackendLocation(D, Context->getSourceManager()); + + unsigned DiagID; + switch (D.getKind()) { + case llvm::SourceMgr::DK_Error: + DiagID = diag::err_fe_inline_asm; + break; + case llvm::SourceMgr::DK_Warning: + DiagID = diag::warn_fe_inline_asm; + break; + case llvm::SourceMgr::DK_Note: + DiagID = diag::note_fe_inline_asm; + break; + case llvm::SourceMgr::DK_Remark: + llvm_unreachable("remarks unexpected"); + } + // If this problem has clang-level source location information, report the + // issue in the source with a note showing the instantiated + // code. + if (LocCookie.isValid()) { + Diags.Report(LocCookie, DiagID).AddString(Message); + + if (D.getLoc().isValid()) { + DiagnosticBuilder B = Diags.Report(Loc, diag::note_fe_inline_asm_here); + // Convert the SMDiagnostic ranges into SourceRange and attach them + // to the diagnostic. + for (const std::pair &Range : D.getRanges()) { + unsigned Column = D.getColumnNo(); + B << SourceRange(Loc.getLocWithOffset(Range.first - Column), + Loc.getLocWithOffset(Range.second - Column)); + } + } + return; + } + + // Otherwise, report the backend issue as occurring in the generated .s file. + // If Loc is invalid, we still need to report the issue, it just gets no + // location info. + Diags.Report(Loc, DiagID).AddString(Message); +} + +#define ComputeDiagID(Severity, GroupName, DiagID) \ + do { \ + switch (Severity) { \ + case llvm::DS_Error: \ + DiagID = diag::err_fe_##GroupName; \ + break; \ + case llvm::DS_Warning: \ + DiagID = diag::warn_fe_##GroupName; \ + break; \ + case llvm::DS_Remark: \ + llvm_unreachable("'remark' severity not expected"); \ + break; \ + case llvm::DS_Note: \ + DiagID = diag::note_fe_##GroupName; \ + break; \ + } \ + } while (false) + +#define ComputeDiagRemarkID(Severity, GroupName, DiagID) \ + do { \ + switch (Severity) { \ + case llvm::DS_Error: \ + DiagID = diag::err_fe_##GroupName; \ + break; \ + case llvm::DS_Warning: \ + DiagID = diag::warn_fe_##GroupName; \ + break; \ + case llvm::DS_Remark: \ + DiagID = diag::remark_fe_##GroupName; \ + break; \ + case llvm::DS_Note: \ + DiagID = diag::note_fe_##GroupName; \ + break; \ + } \ + } while (false) + +bool +BackendConsumer::InlineAsmDiagHandler(const llvm::DiagnosticInfoInlineAsm &D) { + unsigned DiagID; + ComputeDiagID(D.getSeverity(), inline_asm, DiagID); + std::string Message = D.getMsgStr().str(); + + // If this problem has clang-level source location information, report the + // issue as being a problem in the source with a note showing the instantiated + // code. + SourceLocation LocCookie = + SourceLocation::getFromRawEncoding(D.getLocCookie()); + if (LocCookie.isValid()) + Diags.Report(LocCookie, DiagID).AddString(Message); + else { + // Otherwise, report the backend diagnostic as occurring in the generated + // .s file. + // If Loc is invalid, we still need to report the diagnostic, it just gets + // no location info. + FullSourceLoc Loc; + Diags.Report(Loc, DiagID).AddString(Message); + } + // We handled all the possible severities. + return true; +} + +bool +BackendConsumer::StackSizeDiagHandler(const llvm::DiagnosticInfoStackSize &D) { + if (D.getSeverity() != llvm::DS_Warning) + // For now, the only support we have for StackSize diagnostic is warning. + // We do not know how to format other severities. + return false; + + if (const Decl *ND = Gen->GetDeclForMangledName(D.getFunction().getName())) { + // FIXME: Shouldn't need to truncate to uint32_t + Diags.Report(ND->getASTContext().getFullLoc(ND->getLocation()), + diag::warn_fe_frame_larger_than) + << static_cast(D.getStackSize()) << Decl::castToDeclContext(ND); + return true; + } + + return false; +} + +const FullSourceLoc BackendConsumer::getBestLocationFromDebugLoc( + const llvm::DiagnosticInfoWithLocationBase &D, bool &BadDebugInfo, + StringRef &Filename, unsigned &Line, unsigned &Column) const { + SourceManager &SourceMgr = Context->getSourceManager(); + FileManager &FileMgr = SourceMgr.getFileManager(); + SourceLocation DILoc; + + if (D.isLocationAvailable()) { + D.getLocation(Filename, Line, Column); + if (Line > 0) { + auto FE = FileMgr.getFile(Filename); + if (!FE) + FE = FileMgr.getFile(D.getAbsolutePath()); + if (FE) { + // If -gcolumn-info was not used, Column will be 0. This upsets the + // source manager, so pass 1 if Column is not set. + DILoc = SourceMgr.translateFileLineCol(*FE, Line, Column ? Column : 1); + } + } + BadDebugInfo = DILoc.isInvalid(); + } + + // If a location isn't available, try to approximate it using the associated + // function definition. We use the definition's right brace to differentiate + // from diagnostics that genuinely relate to the function itself. + FullSourceLoc Loc(DILoc, SourceMgr); + if (Loc.isInvalid()) + if (const Decl *FD = Gen->GetDeclForMangledName(D.getFunction().getName())) + Loc = FD->getASTContext().getFullLoc(FD->getLocation()); + + if (DILoc.isInvalid() && D.isLocationAvailable()) + // If we were not able to translate the file:line:col information + // back to a SourceLocation, at least emit a note stating that + // we could not translate this location. This can happen in the + // case of #line directives. + Diags.Report(Loc, diag::note_fe_backend_invalid_loc) + << Filename << Line << Column; + + return Loc; +} + +void BackendConsumer::UnsupportedDiagHandler( + const llvm::DiagnosticInfoUnsupported &D) { + // We only support warnings or errors. + assert(D.getSeverity() == llvm::DS_Error || + D.getSeverity() == llvm::DS_Warning); + + StringRef Filename; + unsigned Line, Column; + bool BadDebugInfo = false; + FullSourceLoc Loc; + std::string Msg; + raw_string_ostream MsgStream(Msg); + + // Context will be nullptr for IR input files, we will construct the diag + // message from llvm::DiagnosticInfoUnsupported. + if (Context != nullptr) { + Loc = getBestLocationFromDebugLoc(D, BadDebugInfo, Filename, Line, Column); + MsgStream << D.getMessage(); + } else { + DiagnosticPrinterRawOStream DP(MsgStream); + D.print(DP); + } + + auto DiagType = D.getSeverity() == llvm::DS_Error + ? diag::err_fe_backend_unsupported + : diag::warn_fe_backend_unsupported; + Diags.Report(Loc, DiagType) << MsgStream.str(); + + if (BadDebugInfo) + // If we were not able to translate the file:line:col information + // back to a SourceLocation, at least emit a note stating that + // we could not translate this location. This can happen in the + // case of #line directives. + Diags.Report(Loc, diag::note_fe_backend_invalid_loc) + << Filename << Line << Column; +} + +void BackendConsumer::EmitOptimizationMessage( + const llvm::DiagnosticInfoOptimizationBase &D, unsigned DiagID) { + // We only support warnings and remarks. + assert(D.getSeverity() == llvm::DS_Remark || + D.getSeverity() == llvm::DS_Warning); + + StringRef Filename; + unsigned Line, Column; + bool BadDebugInfo = false; + FullSourceLoc Loc; + std::string Msg; + raw_string_ostream MsgStream(Msg); + + // Context will be nullptr for IR input files, we will construct the remark + // message from llvm::DiagnosticInfoOptimizationBase. + if (Context != nullptr) { + Loc = getBestLocationFromDebugLoc(D, BadDebugInfo, Filename, Line, Column); + MsgStream << D.getMsg(); + } else { + DiagnosticPrinterRawOStream DP(MsgStream); + D.print(DP); + } + + if (D.getHotness()) + MsgStream << " (hotness: " << *D.getHotness() << ")"; + + Diags.Report(Loc, DiagID) + << AddFlagValue(D.getPassName()) + << MsgStream.str(); + + if (BadDebugInfo) + // If we were not able to translate the file:line:col information + // back to a SourceLocation, at least emit a note stating that + // we could not translate this location. This can happen in the + // case of #line directives. + Diags.Report(Loc, diag::note_fe_backend_invalid_loc) + << Filename << Line << Column; +} + +void BackendConsumer::OptimizationRemarkHandler( + const llvm::DiagnosticInfoOptimizationBase &D) { + // Without hotness information, don't show noisy remarks. + if (D.isVerbose() && !D.getHotness()) + return; + + if (D.isPassed()) { + // Optimization remarks are active only if the -Rpass flag has a regular + // expression that matches the name of the pass name in \p D. + if (CodeGenOpts.OptimizationRemarkPattern && + CodeGenOpts.OptimizationRemarkPattern->match(D.getPassName())) + EmitOptimizationMessage(D, diag::remark_fe_backend_optimization_remark); + } else if (D.isMissed()) { + // Missed optimization remarks are active only if the -Rpass-missed + // flag has a regular expression that matches the name of the pass + // name in \p D. + if (CodeGenOpts.OptimizationRemarkMissedPattern && + CodeGenOpts.OptimizationRemarkMissedPattern->match(D.getPassName())) + EmitOptimizationMessage( + D, diag::remark_fe_backend_optimization_remark_missed); + } else { + assert(D.isAnalysis() && "Unknown remark type"); + + bool ShouldAlwaysPrint = false; + if (auto *ORA = dyn_cast(&D)) + ShouldAlwaysPrint = ORA->shouldAlwaysPrint(); + + if (ShouldAlwaysPrint || + (CodeGenOpts.OptimizationRemarkAnalysisPattern && + CodeGenOpts.OptimizationRemarkAnalysisPattern->match(D.getPassName()))) + EmitOptimizationMessage( + D, diag::remark_fe_backend_optimization_remark_analysis); + } +} + +void BackendConsumer::OptimizationRemarkHandler( + const llvm::OptimizationRemarkAnalysisFPCommute &D) { + // Optimization analysis remarks are active if the pass name is set to + // llvm::DiagnosticInfo::AlwasyPrint or if the -Rpass-analysis flag has a + // regular expression that matches the name of the pass name in \p D. + + if (D.shouldAlwaysPrint() || + (CodeGenOpts.OptimizationRemarkAnalysisPattern && + CodeGenOpts.OptimizationRemarkAnalysisPattern->match(D.getPassName()))) + EmitOptimizationMessage( + D, diag::remark_fe_backend_optimization_remark_analysis_fpcommute); +} + +void BackendConsumer::OptimizationRemarkHandler( + const llvm::OptimizationRemarkAnalysisAliasing &D) { + // Optimization analysis remarks are active if the pass name is set to + // llvm::DiagnosticInfo::AlwasyPrint or if the -Rpass-analysis flag has a + // regular expression that matches the name of the pass name in \p D. + + if (D.shouldAlwaysPrint() || + (CodeGenOpts.OptimizationRemarkAnalysisPattern && + CodeGenOpts.OptimizationRemarkAnalysisPattern->match(D.getPassName()))) + EmitOptimizationMessage( + D, diag::remark_fe_backend_optimization_remark_analysis_aliasing); +} + +void BackendConsumer::OptimizationFailureHandler( + const llvm::DiagnosticInfoOptimizationFailure &D) { + EmitOptimizationMessage(D, diag::warn_fe_backend_optimization_failure); +} + +/// This function is invoked when the backend needs +/// to report something to the user. +void BackendConsumer::DiagnosticHandlerImpl(const DiagnosticInfo &DI) { + unsigned DiagID = diag::err_fe_inline_asm; + llvm::DiagnosticSeverity Severity = DI.getSeverity(); + // Get the diagnostic ID based. + switch (DI.getKind()) { + case llvm::DK_InlineAsm: + if (InlineAsmDiagHandler(cast(DI))) + return; + ComputeDiagID(Severity, inline_asm, DiagID); + break; + case llvm::DK_StackSize: + if (StackSizeDiagHandler(cast(DI))) + return; + ComputeDiagID(Severity, backend_frame_larger_than, DiagID); + break; + case DK_Linker: + assert(CurLinkModule); + // FIXME: stop eating the warnings and notes. + if (Severity != DS_Error) + return; + DiagID = diag::err_fe_cannot_link_module; + break; + case llvm::DK_OptimizationRemark: + // Optimization remarks are always handled completely by this + // handler. There is no generic way of emitting them. + OptimizationRemarkHandler(cast(DI)); + return; + case llvm::DK_OptimizationRemarkMissed: + // Optimization remarks are always handled completely by this + // handler. There is no generic way of emitting them. + OptimizationRemarkHandler(cast(DI)); + return; + case llvm::DK_OptimizationRemarkAnalysis: + // Optimization remarks are always handled completely by this + // handler. There is no generic way of emitting them. + OptimizationRemarkHandler(cast(DI)); + return; + case llvm::DK_OptimizationRemarkAnalysisFPCommute: + // Optimization remarks are always handled completely by this + // handler. There is no generic way of emitting them. + OptimizationRemarkHandler(cast(DI)); + return; + case llvm::DK_OptimizationRemarkAnalysisAliasing: + // Optimization remarks are always handled completely by this + // handler. There is no generic way of emitting them. + OptimizationRemarkHandler(cast(DI)); + return; + case llvm::DK_MachineOptimizationRemark: + // Optimization remarks are always handled completely by this + // handler. There is no generic way of emitting them. + OptimizationRemarkHandler(cast(DI)); + return; + case llvm::DK_MachineOptimizationRemarkMissed: + // Optimization remarks are always handled completely by this + // handler. There is no generic way of emitting them. + OptimizationRemarkHandler(cast(DI)); + return; + case llvm::DK_MachineOptimizationRemarkAnalysis: + // Optimization remarks are always handled completely by this + // handler. There is no generic way of emitting them. + OptimizationRemarkHandler(cast(DI)); + return; + case llvm::DK_OptimizationFailure: + // Optimization failures are always handled completely by this + // handler. + OptimizationFailureHandler(cast(DI)); + return; + case llvm::DK_Unsupported: + UnsupportedDiagHandler(cast(DI)); + return; + default: + // Plugin IDs are not bound to any value as they are set dynamically. + ComputeDiagRemarkID(Severity, backend_plugin, DiagID); + break; + } + std::string MsgStorage; + { + raw_string_ostream Stream(MsgStorage); + DiagnosticPrinterRawOStream DP(Stream); + DI.print(DP); + } + + if (DiagID == diag::err_fe_cannot_link_module) { + Diags.Report(diag::err_fe_cannot_link_module) + << CurLinkModule->getModuleIdentifier() << MsgStorage; + return; + } + + // Report the backend message using the usual diagnostic mechanism. + FullSourceLoc Loc; + Diags.Report(Loc, DiagID).AddString(MsgStorage); +} +#undef ComputeDiagID + +CodeGenAction::CodeGenAction(unsigned _Act, LLVMContext *_VMContext) + : Act(_Act), VMContext(_VMContext ? _VMContext : new LLVMContext), + OwnsVMContext(!_VMContext) {} + +CodeGenAction::~CodeGenAction() { + TheModule.reset(); + if (OwnsVMContext) + delete VMContext; +} + +bool CodeGenAction::hasIRSupport() const { return true; } + +void CodeGenAction::EndSourceFileAction() { + // If the consumer creation failed, do nothing. + if (!getCompilerInstance().hasASTConsumer()) + return; + + // Steal the module from the consumer. + TheModule = BEConsumer->takeModule(); +} + +std::unique_ptr CodeGenAction::takeModule() { + return std::move(TheModule); +} + +llvm::LLVMContext *CodeGenAction::takeLLVMContext() { + OwnsVMContext = false; + return VMContext; +} + +static std::unique_ptr +GetOutputStream(CompilerInstance &CI, StringRef InFile, BackendAction Action) { + switch (Action) { + case Backend_EmitAssembly: + return CI.createDefaultOutputFile(false, InFile, "s"); + case Backend_EmitLL: + return CI.createDefaultOutputFile(false, InFile, "ll"); + case Backend_EmitBC: + return CI.createDefaultOutputFile(true, InFile, "bc"); + case Backend_EmitNothing: + return nullptr; + case Backend_EmitMCNull: + return CI.createNullOutputFile(); + case Backend_EmitObj: + return CI.createDefaultOutputFile(true, InFile, "o"); + } + + llvm_unreachable("Invalid action!"); +} + +std::unique_ptr +CodeGenAction::CreateASTConsumer(CompilerInstance &CI, StringRef InFile) { + BackendAction BA = static_cast(Act); + std::unique_ptr OS = CI.takeOutputStream(); + if (!OS) + OS = GetOutputStream(CI, InFile, BA); + + if (BA != Backend_EmitNothing && !OS) + return nullptr; + + // Load bitcode modules to link with, if we need to. + if (LinkModules.empty()) + for (const CodeGenOptions::BitcodeFileToLink &F : + CI.getCodeGenOpts().LinkBitcodeFiles) { + auto BCBuf = CI.getFileManager().getBufferForFile(F.Filename); + if (!BCBuf) { + CI.getDiagnostics().Report(diag::err_cannot_open_file) + << F.Filename << BCBuf.getError().message(); + LinkModules.clear(); + return nullptr; + } + + Expected> ModuleOrErr = + getOwningLazyBitcodeModule(std::move(*BCBuf), *VMContext); + if (!ModuleOrErr) { + handleAllErrors(ModuleOrErr.takeError(), [&](ErrorInfoBase &EIB) { + CI.getDiagnostics().Report(diag::err_cannot_open_file) + << F.Filename << EIB.message(); + }); + LinkModules.clear(); + return nullptr; + } + LinkModules.push_back({std::move(ModuleOrErr.get()), F.PropagateAttrs, + F.Internalize, F.LinkFlags}); + } + + CoverageSourceInfo *CoverageInfo = nullptr; + // Add the preprocessor callback only when the coverage mapping is generated. + if (CI.getCodeGenOpts().CoverageMapping) + CoverageInfo = CodeGen::CoverageMappingModuleGen::setUpCoverageCallbacks( + CI.getPreprocessor()); + + std::unique_ptr Result(new BackendConsumer( + BA, CI.getDiagnostics(), CI.getHeaderSearchOpts(), + CI.getPreprocessorOpts(), CI.getCodeGenOpts(), CI.getTargetOpts(), + CI.getLangOpts(), std::string(InFile), std::move(LinkModules), + std::move(OS), *VMContext, CoverageInfo)); + BEConsumer = Result.get(); + + // Enable generating macro debug info only when debug info is not disabled and + // also macro debug info is enabled. + if (CI.getCodeGenOpts().getDebugInfo() != codegenoptions::NoDebugInfo && + CI.getCodeGenOpts().MacroDebugInfo) { + std::unique_ptr Callbacks = + std::make_unique(BEConsumer->getCodeGenerator(), + CI.getPreprocessor()); + CI.getPreprocessor().addPPCallbacks(std::move(Callbacks)); + } + + return std::move(Result); +} + +static void BitcodeInlineAsmDiagHandler(const llvm::SMDiagnostic &SM, + void *Context, + unsigned LocCookie) { + SM.print(nullptr, llvm::errs()); + + auto Diags = static_cast(Context); + unsigned DiagID; + switch (SM.getKind()) { + case llvm::SourceMgr::DK_Error: + DiagID = diag::err_fe_inline_asm; + break; + case llvm::SourceMgr::DK_Warning: + DiagID = diag::warn_fe_inline_asm; + break; + case llvm::SourceMgr::DK_Note: + DiagID = diag::note_fe_inline_asm; + break; + case llvm::SourceMgr::DK_Remark: + llvm_unreachable("remarks unexpected"); + } + + Diags->Report(DiagID).AddString("cannot compile inline asm"); +} + +std::unique_ptr +CodeGenAction::loadModule(MemoryBufferRef MBRef) { + CompilerInstance &CI = getCompilerInstance(); + SourceManager &SM = CI.getSourceManager(); + + // For ThinLTO backend invocations, ensure that the context + // merges types based on ODR identifiers. We also need to read + // the correct module out of a multi-module bitcode file. + if (!CI.getCodeGenOpts().ThinLTOIndexFile.empty()) { + VMContext->enableDebugTypeODRUniquing(); + + auto DiagErrors = [&](Error E) -> std::unique_ptr { + unsigned DiagID = + CI.getDiagnostics().getCustomDiagID(DiagnosticsEngine::Error, "%0"); + handleAllErrors(std::move(E), [&](ErrorInfoBase &EIB) { + CI.getDiagnostics().Report(DiagID) << EIB.message(); + }); + return {}; + }; + + Expected> BMsOrErr = getBitcodeModuleList(MBRef); + if (!BMsOrErr) + return DiagErrors(BMsOrErr.takeError()); + BitcodeModule *Bm = llvm::lto::findThinLTOModule(*BMsOrErr); + // We have nothing to do if the file contains no ThinLTO module. This is + // possible if ThinLTO compilation was not able to split module. Content of + // the file was already processed by indexing and will be passed to the + // linker using merged object file. + if (!Bm) { + auto M = std::make_unique("empty", *VMContext); + M->setTargetTriple(CI.getTargetOpts().Triple); + return M; + } + Expected> MOrErr = + Bm->parseModule(*VMContext); + if (!MOrErr) + return DiagErrors(MOrErr.takeError()); + return std::move(*MOrErr); + } + + llvm::SMDiagnostic Err; + if (std::unique_ptr M = parseIR(MBRef, Err, *VMContext)) + return M; + + // Translate from the diagnostic info to the SourceManager location if + // available. + // TODO: Unify this with ConvertBackendLocation() + SourceLocation Loc; + if (Err.getLineNo() > 0) { + assert(Err.getColumnNo() >= 0); + Loc = SM.translateFileLineCol(SM.getFileEntryForID(SM.getMainFileID()), + Err.getLineNo(), Err.getColumnNo() + 1); + } + + // Strip off a leading diagnostic code if there is one. + StringRef Msg = Err.getMessage(); + if (Msg.startswith("error: ")) + Msg = Msg.substr(7); + + unsigned DiagID = + CI.getDiagnostics().getCustomDiagID(DiagnosticsEngine::Error, "%0"); + + CI.getDiagnostics().Report(Loc, DiagID) << Msg; + return {}; +} + +// zahira +namespace { +// Handles the initialization and cleanup of the OptRecordFile. This +// customization allows initialization before the clang codegen runs +// so it can also emit to the opt report. +struct OptRecordFileRAII { + std::unique_ptr OptRecordFile; + std::unique_ptr OldDiagnosticHandler; + llvm::LLVMContext &Ctx; + + OptRecordFileRAII(CodeGenAction &CGA, llvm::LLVMContext &Ctx, + BackendConsumer &BC) + : Ctx(Ctx) { + CompilerInstance &CI = CGA.getCompilerInstance(); + auto &CodeGenOpts = CI.getCodeGenOpts(); + + OldDiagnosticHandler = Ctx.getDiagnosticHandler(); + + Ctx.setDiagnosticHandler( + std::make_unique(CodeGenOpts, &BC)); + + Expected> OptRecordFileOrErr = + setupLLVMOptimizationRemarks( + Ctx, CodeGenOpts.OptRecordFile, CodeGenOpts.OptRecordPasses, + CodeGenOpts.OptRecordFormat, CodeGenOpts.DiagnosticsWithHotness, + CodeGenOpts.DiagnosticsHotnessThreshold); + + if (Error E = OptRecordFileOrErr.takeError()) + reportOptRecordError(std::move(E), CI.getDiagnostics(), CodeGenOpts); + else + OptRecordFile = std::move(*OptRecordFileOrErr); + + if (OptRecordFile && + CodeGenOpts.getProfileUse() != CodeGenOptions::ProfileNone) + Ctx.setDiagnosticsHotnessRequested(true); + } + ~OptRecordFileRAII() { + Ctx.setDiagnosticHandler(std::move(OldDiagnosticHandler)); + if (OptRecordFile) + OptRecordFile->keep(); + } +}; +} // namespace + +void CodeGenAction::ExecuteAction() { + if (getCurrentFileKind().getLanguage() != Language::LLVM_IR) { + OptRecordFileRAII ORF(*this, *VMContext, *BEConsumer); + this->ASTFrontendAction::ExecuteAction(); + return; + } + + // If this is an IR file, we have to treat it specially. + BackendAction BA = static_cast(Act); + CompilerInstance &CI = getCompilerInstance(); + auto &CodeGenOpts = CI.getCodeGenOpts(); + auto &Diagnostics = CI.getDiagnostics(); + std::unique_ptr OS = + GetOutputStream(CI, getCurrentFile(), BA); + if (BA != Backend_EmitNothing && !OS) + return; + + SourceManager &SM = CI.getSourceManager(); + FileID FID = SM.getMainFileID(); + Optional MainFile = SM.getBufferOrNone(FID); + if (!MainFile) + return; + + TheModule = loadModule(*MainFile); + if (!TheModule) + return; + + const TargetOptions &TargetOpts = CI.getTargetOpts(); + if (TheModule->getTargetTriple() != TargetOpts.Triple) { + Diagnostics.Report(SourceLocation(), diag::warn_fe_override_module) + << TargetOpts.Triple; + TheModule->setTargetTriple(TargetOpts.Triple); + } + + // EmbedBitcode(TheModule.get(), CodeGenOpts, *MainFile); + + LLVMContext &Ctx = TheModule->getContext(); + Ctx.setInlineAsmDiagnosticHandler(BitcodeInlineAsmDiagHandler, &Diagnostics); + + // Restore any diagnostic handler previously set before returning from this + // function. + struct RAII { + LLVMContext &Ctx; + std::unique_ptr PrevHandler = Ctx.getDiagnosticHandler(); + ~RAII() { Ctx.setDiagnosticHandler(std::move(PrevHandler)); } + } _{Ctx}; + + // Set clang diagnostic handler. To do this we need to create a fake + // BackendConsumer. + BackendConsumer Result(BA, CI.getDiagnostics(), CI.getHeaderSearchOpts(), + CI.getPreprocessorOpts(), CI.getCodeGenOpts(), + CI.getTargetOpts(), CI.getLangOpts(), + std::move(LinkModules), *VMContext, nullptr); + // PR44896: Force DiscardValueNames as false. DiscardValueNames cannot be + // true here because the valued names are needed for reading textual IR. + Ctx.setDiscardValueNames(false); + OptRecordFileRAII ORF(*this, Ctx, Result); +#if 0 + Ctx.setDiagnosticHandler( + std::make_unique(CodeGenOpts, &Result)); + + Expected> OptRecordFileOrErr = + setupLLVMOptimizationRemarks( + Ctx, CodeGenOpts.OptRecordFile, CodeGenOpts.OptRecordPasses, + CodeGenOpts.OptRecordFormat, CodeGenOpts.DiagnosticsWithHotness, + CodeGenOpts.DiagnosticsHotnessThreshold); + + if (Error E = OptRecordFileOrErr.takeError()) { + reportOptRecordError(std::move(E), Diagnostics, CodeGenOpts); + return; + } + std::unique_ptr OptRecordFile = + std::move(*OptRecordFileOrErr); +#endif + EmitBackendOutput(Diagnostics, CI.getHeaderSearchOpts(), CodeGenOpts, + TargetOpts, CI.getLangOpts(), + CI.getTarget().getDataLayout(), TheModule.get(), BA, + std::move(OS)); +#if 0 + if (OptRecordFile) + OptRecordFile->keep(); +#endif +} + +// + +void EmitAssemblyAction::anchor() { } +EmitAssemblyAction::EmitAssemblyAction(llvm::LLVMContext *_VMContext) + : CodeGenAction(Backend_EmitAssembly, _VMContext) {} + +void EmitBCAction::anchor() { } +EmitBCAction::EmitBCAction(llvm::LLVMContext *_VMContext) + : CodeGenAction(Backend_EmitBC, _VMContext) {} + +void EmitLLVMAction::anchor() { } +EmitLLVMAction::EmitLLVMAction(llvm::LLVMContext *_VMContext) + : CodeGenAction(Backend_EmitLL, _VMContext) {} + +void EmitLLVMOnlyAction::anchor() { } +EmitLLVMOnlyAction::EmitLLVMOnlyAction(llvm::LLVMContext *_VMContext) + : CodeGenAction(Backend_EmitNothing, _VMContext) {} + +void EmitCodeGenOnlyAction::anchor() { } +EmitCodeGenOnlyAction::EmitCodeGenOnlyAction(llvm::LLVMContext *_VMContext) + : CodeGenAction(Backend_EmitMCNull, _VMContext) {} + +void EmitObjAction::anchor() { } +EmitObjAction::EmitObjAction(llvm::LLVMContext *_VMContext) + : CodeGenAction(Backend_EmitObj, _VMContext) {} diff --git a/clang/lib/CodeGen/CodeGenFunction.cpp b/clang/lib/CodeGen/CodeGenFunction.cpp index 7343c96890e86..3b56c0cecfd3e 100644 --- a/clang/lib/CodeGen/CodeGenFunction.cpp +++ b/clang/lib/CodeGen/CodeGenFunction.cpp @@ -35,6 +35,7 @@ #include "clang/CodeGen/CGFunctionInfo.h" #include "clang/Frontend/FrontendDiagnostic.h" #include "llvm/ADT/ArrayRef.h" +#include "llvm/Analysis/OptimizationRemarkEmitter.h" #include "llvm/Frontend/OpenMP/OMPIRBuilder.h" #include "llvm/IR/DataLayout.h" #include "llvm/IR/Dominators.h" @@ -102,17 +103,20 @@ llvm::fp::ExceptionBehavior clang::ToConstrainedExceptMD(LangOptions::FPExceptionModeKind Kind) { switch (Kind) { - case LangOptions::FPE_Ignore: return llvm::fp::ebIgnore; - case LangOptions::FPE_MayTrap: return llvm::fp::ebMayTrap; - case LangOptions::FPE_Strict: return llvm::fp::ebStrict; + case LangOptions::FPE_Ignore: + return llvm::fp::ebIgnore; + case LangOptions::FPE_MayTrap: + return llvm::fp::ebMayTrap; + case LangOptions::FPE_Strict: + return llvm::fp::ebStrict; } llvm_unreachable("Unsupported FP Exception Behavior"); } void CodeGenFunction::SetFPModel() { llvm::RoundingMode RM = getLangOpts().getFPRoundingMode(); - auto fpExceptionBehavior = ToConstrainedExceptMD( - getLangOpts().getFPExceptionMode()); + auto fpExceptionBehavior = + ToConstrainedExceptMD(getLangOpts().getFPExceptionMode()); Builder.setDefaultConstrainedRounding(RM); Builder.setDefaultConstrainedExcept(fpExceptionBehavior); @@ -205,8 +209,8 @@ LValue CodeGenFunction::MakeNaturalAlignAddrLValue(llvm::Value *V, QualType T) { /// Given a value of type T* that may not be to a complete object, /// construct an l-value with the natural pointee alignment of T. -LValue -CodeGenFunction::MakeNaturalAlignPointeeAddrLValue(llvm::Value *V, QualType T) { +LValue CodeGenFunction::MakeNaturalAlignPointeeAddrLValue(llvm::Value *V, + QualType T) { LValueBaseInfo BaseInfo; TBAAAccessInfo TBAAInfo; CharUnits Align = CGM.getNaturalTypeAlignment(T, &BaseInfo, &TBAAInfo, @@ -214,7 +218,6 @@ CodeGenFunction::MakeNaturalAlignPointeeAddrLValue(llvm::Value *V, QualType T) { return MakeAddrLValue(Address(V, Align), T, BaseInfo, TBAAInfo); } - llvm::Type *CodeGenFunction::ConvertTypeForMem(QualType T) { return CGM.getTypes().ConvertTypeForMem(T); } @@ -303,7 +306,7 @@ llvm::DebugLoc CodeGenFunction::EmitReturnBlock() { // cleans up functions which started with a unified return block. if (ReturnBlock.getBlock()->hasOneUse()) { llvm::BranchInst *BI = - dyn_cast(*ReturnBlock.getBlock()->user_begin()); + dyn_cast(*ReturnBlock.getBlock()->user_begin()); if (BI && BI->isUnconditional() && BI->getSuccessor(0) == ReturnBlock.getBlock()) { // Record/return the DebugLoc of the simple 'return' expression to be used @@ -326,7 +329,8 @@ llvm::DebugLoc CodeGenFunction::EmitReturnBlock() { } static void EmitIfUsed(CodeGenFunction &CGF, llvm::BasicBlock *BB) { - if (!BB) return; + if (!BB) + return; if (!BB->use_empty()) return CGF.CurFn->getBasicBlockList().push_back(BB); delete BB; @@ -336,9 +340,9 @@ void CodeGenFunction::FinishFunction(SourceLocation EndLoc) { assert(BreakContinueStack.empty() && "mismatched push/pop in break/continue stack!"); - bool OnlySimpleReturnStmts = NumSimpleReturnExprs > 0 - && NumSimpleReturnExprs == NumReturnExprs - && ReturnBlock.getBlock()->use_empty(); + bool OnlySimpleReturnStmts = NumSimpleReturnExprs > 0 && + NumSimpleReturnExprs == NumReturnExprs && + ReturnBlock.getBlock()->use_empty(); // Usually the return expression is evaluated before the cleanup // code. If the function contains only a simple return statement, // such as a constant, the location before the cleanup code becomes @@ -403,8 +407,7 @@ void CodeGenFunction::FinishFunction(SourceLocation EndLoc) { EmitFunctionEpilog(*CurFnInfo, EmitRetDbgLoc, EndLoc); EmitEndEHSpec(CurCodeDecl); - assert(EHStack.empty() && - "did not remove all scopes from cleanup stack!"); + assert(EHStack.empty() && "did not remove all scopes from cleanup stack!"); // If someone did an indirect goto, emit the indirect goto block at the end of // the function. @@ -453,7 +456,7 @@ void CodeGenFunction::FinishFunction(SourceLocation EndLoc) { if (CGM.getCodeGenOpts().EmitDeclMetadata) EmitDeclMetadata(); - for (SmallVectorImpl >::iterator + for (SmallVectorImpl>::iterator I = DeferredReplacements.begin(), E = DeferredReplacements.end(); I != E; ++I) { @@ -582,8 +585,7 @@ CodeGenFunction::DecodeAddrUsedInPrologue(llvm::Value *F, } void CodeGenFunction::EmitOpenCLKernelMetadata(const FunctionDecl *FD, - llvm::Function *Fn) -{ + llvm::Function *Fn) { if (!FD->hasAttr()) return; @@ -616,7 +618,8 @@ void CodeGenFunction::EmitOpenCLKernelMetadata(const FunctionDecl *FD, llvm::ConstantAsMetadata::get(Builder.getInt32(A->getXDim())), llvm::ConstantAsMetadata::get(Builder.getInt32(A->getYDim())), llvm::ConstantAsMetadata::get(Builder.getInt32(A->getZDim()))}; - Fn->setMetadata("work_group_size_hint", llvm::MDNode::get(Context, AttrMDArgs)); + Fn->setMetadata("work_group_size_hint", + llvm::MDNode::get(Context, AttrMDArgs)); } if (const ReqdWorkGroupSizeAttr *A = FD->getAttr()) { @@ -660,7 +663,7 @@ void CodeGenFunction::EmitOpenCLKernelMetadata(const FunctionDecl *FD, } if (const SYCLIntelNumSimdWorkItemsAttr *A = - FD->getAttr()) { + FD->getAttr()) { const auto *CE = dyn_cast(A->getValue()); assert(CE && "Not an integer constant expression"); Optional ArgVal = CE->getResultAsAPSInt(); @@ -682,7 +685,7 @@ void CodeGenFunction::EmitOpenCLKernelMetadata(const FunctionDecl *FD, } if (const SYCLIntelMaxGlobalWorkDimAttr *A = - FD->getAttr()) { + FD->getAttr()) { const auto *CE = dyn_cast(A->getValue()); assert(CE && "Not an integer constant expression"); Optional ArgVal = CE->getResultAsAPSInt(); @@ -733,7 +736,7 @@ void CodeGenFunction::EmitOpenCLKernelMetadata(const FunctionDecl *FD, } /// Determine whether the function F ends with a return stmt. -static bool endsWithReturn(const Decl* F) { +static bool endsWithReturn(const Decl *F) { const Stmt *Body = nullptr; if (auto *FD = dyn_cast_or_null(F)) Body = FD->getBody(); @@ -846,7 +849,8 @@ void CodeGenFunction::StartFunction(GlobalDecl GD, QualType RetTy, // Apply sanitizer attributes to the function. if (SanOpts.hasOneOf(SanitizerKind::Address | SanitizerKind::KernelAddress)) Fn->addFnAttr(llvm::Attribute::SanitizeAddress); - if (SanOpts.hasOneOf(SanitizerKind::HWAddress | SanitizerKind::KernelHWAddress)) + if (SanOpts.hasOneOf(SanitizerKind::HWAddress | + SanitizerKind::KernelHWAddress)) Fn->addFnAttr(llvm::Attribute::SanitizeHWAddress); if (SanOpts.has(SanitizerKind::MemTag)) Fn->addFnAttr(llvm::Attribute::SanitizeMemTag); @@ -870,7 +874,8 @@ void CodeGenFunction::StartFunction(GlobalDecl GD, QualType RetTy, IdentifierInfo *II = OMD->getSelector().getIdentifierInfoForSlot(0); if (OMD->getMethodFamily() == OMF_dealloc || OMD->getMethodFamily() == OMF_initialize || - (OMD->getSelector().isUnarySelector() && II->isStr(".cxx_destruct"))) { + (OMD->getSelector().isUnarySelector() && + II->isStr(".cxx_destruct"))) { markAsIgnoreThreadCheckingAtRuntime(Fn); } } @@ -1012,9 +1017,8 @@ void CodeGenFunction::StartFunction(GlobalDecl GD, QualType RetTy, if (llvm::Constant *PrologueSig = getPrologueSignature(CGM, FD)) { // Remove any (C++17) exception specifications, to allow calling e.g. a // noexcept function through a non-noexcept pointer. - auto ProtoTy = - getContext().getFunctionTypeWithExceptionSpec(FD->getType(), - EST_None); + auto ProtoTy = getContext().getFunctionTypeWithExceptionSpec( + FD->getType(), EST_None); llvm::Constant *FTRTTIConst = CGM.GetAddrOfRTTIDescriptor(ProtoTy, /*ForEH=*/true); llvm::Constant *FTRTTIConstEncoded = @@ -1140,14 +1144,16 @@ void CodeGenFunction::StartFunction(GlobalDecl GD, QualType RetTy, if (CGM.getCodeGenOpts().MNopMCount) { if (!CGM.getCodeGenOpts().CallFEntry) CGM.getDiags().Report(diag::err_opt_not_valid_without_opt) - << "-mnop-mcount" << "-mfentry"; + << "-mnop-mcount" + << "-mfentry"; Fn->addFnAttr("mnop-mcount"); } if (CGM.getCodeGenOpts().RecordMCount) { if (!CGM.getCodeGenOpts().CallFEntry) CGM.getDiags().Report(diag::err_opt_not_valid_without_opt) - << "-mrecord-mcount" << "-mfentry"; + << "-mrecord-mcount" + << "-mfentry"; Fn->addFnAttr("mrecord-mcount"); } } @@ -1157,7 +1163,7 @@ void CodeGenFunction::StartFunction(GlobalDecl GD, QualType RetTy, if (getContext().getTargetInfo().getTriple().getArch() != llvm::Triple::systemz) CGM.getDiags().Report(diag::err_opt_not_valid_on_target) - << "-mpacked-stack"; + << "-mpacked-stack"; Fn->addFnAttr("packed-stack"); } @@ -1198,8 +1204,7 @@ void CodeGenFunction::StartFunction(GlobalDecl GD, QualType RetTy, // Tell the epilog emitter to autorelease the result. We do this // now so that various specialized functions can suppress it // during their IR-generation. - if (getLangOpts().ObjCAutoRefCount && - !CurFnInfo->isReturnsRetained() && + if (getLangOpts().ObjCAutoRefCount && !CurFnInfo->isReturnsRetained() && RetTy->isObjCRetainableType()) AutoreleaseResult = true; } @@ -1217,8 +1222,7 @@ void CodeGenFunction::StartFunction(GlobalDecl GD, QualType RetTy, if (D && isa(D) && cast(D)->isInstance()) { CGM.getCXXABI().EmitInstanceFunctionProlog(*this); const CXXMethodDecl *MD = cast(D); - if (MD->getParent()->isLambda() && - MD->getOverloadedOperator() == OO_Call) { + if (MD->getParent()->isLambda() && MD->getOverloadedOperator() == OO_Call) { // We're in a lambda; figure out the captures. MD->getParent()->getCaptureFields(LambdaCaptureFields, LambdaThisCaptureField); @@ -1229,21 +1233,24 @@ void CodeGenFunction::StartFunction(GlobalDecl GD, QualType RetTy, // Get the lvalue for the field (which is a copy of the enclosing object // or contains the address of the enclosing object). - LValue ThisFieldLValue = EmitLValueForLambdaField(LambdaThisCaptureField); + LValue ThisFieldLValue = + EmitLValueForLambdaField(LambdaThisCaptureField); if (!LambdaThisCaptureField->getType()->isPointerType()) { - // If the enclosing object was captured by value, just use its address. + // If the enclosing object was captured by value, just use its + // address. CXXThisValue = ThisFieldLValue.getAddress(*this).getPointer(); } else { // Load the lvalue pointed to by the field, since '*this' was captured // by reference. - CXXThisValue = - EmitLoadOfLValue(ThisFieldLValue, SourceLocation()).getScalarVal(); + CXXThisValue = EmitLoadOfLValue(ThisFieldLValue, SourceLocation()) + .getScalarVal(); } } for (auto *FD : MD->getParent()->fields()) { if (FD->hasCapturedVLAType()) { - auto *ExprArg = EmitLoadOfLValue(EmitLValueForLambdaField(FD), - SourceLocation()).getScalarVal(); + auto *ExprArg = + EmitLoadOfLValue(EmitLValueForLambdaField(FD), SourceLocation()) + .getScalarVal(); auto VAT = FD->getCapturedVLAType(); VLASizeMap[VAT->getSizeExpr()] = ExprArg; } @@ -1276,8 +1283,8 @@ void CodeGenFunction::StartFunction(GlobalDecl GD, QualType RetTy, // If any of the arguments have a variably modified type, make sure to // emit the type size. - for (FunctionArgList::const_iterator i = Args.begin(), e = Args.end(); - i != e; ++i) { + for (FunctionArgList::const_iterator i = Args.begin(), e = Args.end(); i != e; + ++i) { const VarDecl *VD = *i; // Dig out the type as written from ParmVarDecls; it's unclear whether @@ -1347,7 +1354,8 @@ void CodeGenFunction::EmitBlockWithFallThrough(llvm::BasicBlock *BB, static void TryMarkNoThrow(llvm::Function *F) { // LLVM treats 'nounwind' on a function as part of the type, so we // can't do this on functions that can be overwritten. - if (F->isInterposable()) return; + if (F->isInterposable()) + return; for (llvm::BasicBlock &BB : *F) for (llvm::Instruction &I : BB) @@ -1457,14 +1465,36 @@ void CodeGenFunction::GenerateCode(GlobalDecl GD, llvm::Function *Fn, // Emit the standard function prologue. StartFunction(GD, ResTy, Fn, FnInfo, Args, Loc, BodyRange.getBegin()); + SyclOptReportHandler &OptReportHandler = CGM.getDiags().OptReportHandler; + + // zahira + if (OptReportHandler.HasOptReportInfo(FD)) { + llvm::OptimizationRemarkEmitter ORE(Fn); + //llvm::DiagnosticLocation DL = SourceLocToDebugLoc(Body->getBeginLoc()); + for (auto &ORI : OptReportHandler.getInfo(FD)) { + llvm::DiagnosticLocation DL = SourceLocToDebugLoc(ORI.KernelArgLoc); + llvm::OptimizationRemarkMissed Remark("SYCL", "ZAHIRA_Region", DL, + &Fn->getEntryBlock()); +#if 1 + printf("Name of FD: %s \n", FD->getName().data()); + printf("ArgName: %s \n", ORI.KernelArgName.data()); + printf("ArgType: %s \n", ORI.KernelArgType.data()); + printf("ArgLoc: %d \n", ORI.KernelArgLoc.getRawEncoding()); +#endif + //llvm::OptimizationRemark Remark("sycl", "Kernelrgs", Fn); + Remark << llvm::ore::NV("KernelArg", ORI.KernelArgName); + Remark << llvm::ore::NV("ArgType", ORI.KernelArgType); + ORE.emit(Remark); + } + } + // Generate the body of the function. PGO.assignRegionCounters(GD, CurFn); if (isa(FD)) EmitDestructorBody(Args); else if (isa(FD)) EmitConstructorBody(Args); - else if (getLangOpts().CUDA && - !getLangOpts().CUDAIsDevice && + else if (getLangOpts().CUDA && !getLangOpts().CUDAIsDevice && FD->hasAttr()) CGM.getCUDARuntime().emitDeviceStub(*this, Args); else if (isa(FD) && @@ -1524,7 +1554,8 @@ void CodeGenFunction::GenerateCode(GlobalDecl GD, llvm::Function *Fn, /// that we can just remove the code. bool CodeGenFunction::ContainsLabel(const Stmt *S, bool IgnoreCaseStmts) { // Null statement, not a label! - if (!S) return false; + if (!S) + return false; // If this is a label, we have to emit the code, consider something like: // if (0) { ... foo: bar(); } goto foo; @@ -1556,7 +1587,8 @@ bool CodeGenFunction::ContainsLabel(const Stmt *S, bool IgnoreCaseStmts) { /// inside of it, this is fine. bool CodeGenFunction::containsBreak(const Stmt *S) { // Null statement, not a label! - if (!S) return false; + if (!S) + return false; // If this is a switch or loop that defines its own break scope, then we can // include it and anything inside of it. @@ -1576,7 +1608,8 @@ bool CodeGenFunction::containsBreak(const Stmt *S) { } bool CodeGenFunction::mightAddDeclToScope(const Stmt *S) { - if (!S) return false; + if (!S) + return false; // Some statement kinds add a scope and thus never add a decl to the current // scope. Note, this list is longer than the list of statements that might @@ -1622,11 +1655,11 @@ bool CodeGenFunction::ConstantFoldsToSimpleInteger(const Expr *Cond, // to bool. Expr::EvalResult Result; if (!Cond->EvaluateAsInt(Result, getContext())) - return false; // Not foldable, not integer or not fully evaluatable. + return false; // Not foldable, not integer or not fully evaluatable. llvm::APSInt Int = Result.Val.getInt(); if (!AllowLabels && CodeGenFunction::ContainsLabel(Cond)) - return false; // Contains a label. + return false; // Contains a label. ResultInt = Int; return true; @@ -1898,7 +1931,7 @@ void CodeGenFunction::EmitBranchOnBoolExpr(const Expr *Cond, // br(c ? throw x : y, t, f) -> br(c, br(throw x, t, f), br(y, t, f) // Fold this to: // br(c, throw x, br(y, t, f)) - EmitCXXThrowExpr(Throw, /*KeepInsertionPoint*/false); + EmitCXXThrowExpr(Throw, /*KeepInsertionPoint*/ false); return; } @@ -1949,13 +1982,12 @@ static void emitNonZeroVLAInit(CodeGenFunction &CGF, QualType baseType, CGBuilderTy &Builder = CGF.Builder; CharUnits baseSize = CGF.getContext().getTypeSizeInChars(baseType); - llvm::Value *baseSizeInChars - = llvm::ConstantInt::get(CGF.IntPtrTy, baseSize.getQuantity()); + llvm::Value *baseSizeInChars = + llvm::ConstantInt::get(CGF.IntPtrTy, baseSize.getQuantity()); - Address begin = - Builder.CreateElementBitCast(dest, CGF.Int8Ty, "vla.begin"); + Address begin = Builder.CreateElementBitCast(dest, CGF.Int8Ty, "vla.begin"); llvm::Value *end = - Builder.CreateInBoundsGEP(begin.getPointer(), sizeInChars, "vla.end"); + Builder.CreateInBoundsGEP(begin.getPointer(), sizeInChars, "vla.end"); llvm::BasicBlock *originBB = CGF.Builder.GetInsertBlock(); llvm::BasicBlock *loopBB = CGF.createBasicBlock("vla-init.loop"); @@ -1968,8 +2000,7 @@ static void emitNonZeroVLAInit(CodeGenFunction &CGF, QualType baseType, llvm::PHINode *cur = Builder.CreatePHI(begin.getType(), 2, "vla.cur"); cur->addIncoming(begin.getPointer(), originBB); - CharUnits curAlign = - dest.getAlignment().alignmentOfArrayElement(baseSize); + CharUnits curAlign = dest.getAlignment().alignmentOfArrayElement(baseSize); // memcpy the individual element bit-pattern. Builder.CreateMemCpy(Address(cur, curAlign), src, baseSizeInChars, @@ -1977,7 +2008,7 @@ static void emitNonZeroVLAInit(CodeGenFunction &CGF, QualType baseType, // Go to the next element. llvm::Value *next = - Builder.CreateInBoundsGEP(CGF.Int8Ty, cur, baseSizeInChars, "vla.next"); + Builder.CreateInBoundsGEP(CGF.Int8Ty, cur, baseSizeInChars, "vla.next"); // Leave if that's the end of the VLA. llvm::Value *done = Builder.CreateICmpEQ(next, end, "vla-init.isdone"); @@ -1987,8 +2018,7 @@ static void emitNonZeroVLAInit(CodeGenFunction &CGF, QualType baseType, CGF.EmitBlock(contBB); } -void -CodeGenFunction::EmitNullInitialization(Address DestPtr, QualType Ty) { +void CodeGenFunction::EmitNullInitialization(Address DestPtr, QualType Ty) { // Ignore empty classes in C++. if (getLangOpts().CPlusPlus) { if (const RecordType *RT = Ty->getAs()) { @@ -2010,9 +2040,8 @@ CodeGenFunction::EmitNullInitialization(Address DestPtr, QualType Ty) { // Don't bother emitting a zero-byte memset. if (size.isZero()) { // But note that getTypeInfo returns 0 for a VLA. - if (const VariableArrayType *vlaType = - dyn_cast_or_null( - getContext().getAsArrayType(Ty))) { + if (const VariableArrayType *vlaType = dyn_cast_or_null( + getContext().getAsArrayType(Ty))) { auto VlaSize = getVLASize(vlaType); SizeVal = VlaSize.NumElts; CharUnits eltSize = getContext().getTypeSizeInChars(VlaSize.Type); @@ -2033,21 +2062,22 @@ CodeGenFunction::EmitNullInitialization(Address DestPtr, QualType Ty) { // like -1, which happens to be the pattern used by member-pointers. if (!CGM.getTypes().isZeroInitializable(Ty)) { // For a VLA, emit a single element, then splat that over the VLA. - if (vla) Ty = getContext().getBaseElementType(vla); + if (vla) + Ty = getContext().getBaseElementType(vla); llvm::Constant *NullConstant = CGM.EmitNullConstant(Ty); - llvm::GlobalVariable *NullVariable = - new llvm::GlobalVariable(CGM.getModule(), NullConstant->getType(), - /*isConstant=*/true, - llvm::GlobalVariable::PrivateLinkage, - NullConstant, Twine()); + llvm::GlobalVariable *NullVariable = new llvm::GlobalVariable( + CGM.getModule(), NullConstant->getType(), + /*isConstant=*/true, llvm::GlobalVariable::PrivateLinkage, NullConstant, + Twine()); CharUnits NullAlign = DestPtr.getAlignment(); NullVariable->setAlignment(NullAlign.getAsAlign()); Address SrcPtr(Builder.CreateBitCast(NullVariable, Builder.getInt8PtrTy()), NullAlign); - if (vla) return emitNonZeroVLAInit(*this, Ty, DestPtr, SrcPtr, SizeVal); + if (vla) + return emitNonZeroVLAInit(*this, Ty, DestPtr, SrcPtr, SizeVal); // Get and call the appropriate llvm.memcpy overload. Builder.CreateMemCpy(DestPtr, SrcPtr, SizeVal, false); @@ -2074,13 +2104,14 @@ llvm::BlockAddress *CodeGenFunction::GetAddrOfLabel(const LabelDecl *L) { llvm::BasicBlock *CodeGenFunction::GetIndirectGotoBlock() { // If we already made the indirect branch for indirect goto, return its block. - if (IndirectBranch) return IndirectBranch->getParent(); + if (IndirectBranch) + return IndirectBranch->getParent(); CGBuilderTy TmpBuilder(*this, createBasicBlock("indirectgoto")); // Create the PHI node that indirect gotos will add entries to. - llvm::Value *DestVal = TmpBuilder.CreatePHI(Int8PtrTy, 0, - "indirect.goto.dest"); + llvm::Value *DestVal = + TmpBuilder.CreatePHI(Int8PtrTy, 0, "indirect.goto.dest"); // Create the indirect branch instruction. IndirectBranch = TmpBuilder.CreateIndirectBr(DestVal); @@ -2120,7 +2151,7 @@ llvm::Value *CodeGenFunction::emitArrayLength(const ArrayType *origArrayType, // We have some number of constant-length arrays, so addr should // have LLVM type [M x [N x [...]]]*. Build a GEP that walks // down to the first element of addr. - SmallVector gepIndices; + SmallVector gepIndices; // GEP down to the array type. llvm::ConstantInt *zero = Builder.getInt32(0); @@ -2130,18 +2161,17 @@ llvm::Value *CodeGenFunction::emitArrayLength(const ArrayType *origArrayType, QualType eltType; llvm::ArrayType *llvmArrayType = - dyn_cast(addr.getElementType()); + dyn_cast(addr.getElementType()); while (llvmArrayType) { assert(isa(arrayType)); - assert(cast(arrayType)->getSize().getZExtValue() - == llvmArrayType->getNumElements()); + assert(cast(arrayType)->getSize().getZExtValue() == + llvmArrayType->getNumElements()); gepIndices.push_back(zero); countFromCLAs *= llvmArrayType->getNumElements(); eltType = arrayType->getElementType(); - llvmArrayType = - dyn_cast(llvmArrayType->getElementType()); + llvmArrayType = dyn_cast(llvmArrayType->getElementType()); arrayType = getContext().getAsArrayType(arrayType->getElementType()); assert((!llvmArrayType || arrayType) && "LLVM and Clang types are out-of-synch"); @@ -2162,15 +2192,14 @@ llvm::Value *CodeGenFunction::emitArrayLength(const ArrayType *origArrayType, addr = Builder.CreateElementBitCast(addr, baseType, "array.begin"); } else { // Create the actual GEP. - addr = Address(Builder.CreateInBoundsGEP(addr.getPointer(), - gepIndices, "array.begin"), - addr.getAlignment()); + addr = Address( + Builder.CreateInBoundsGEP(addr.getPointer(), gepIndices, "array.begin"), + addr.getAlignment()); } baseType = eltType; - llvm::Value *numElements - = llvm::ConstantInt::get(SizeTy, countFromCLAs); + llvm::Value *numElements = llvm::ConstantInt::get(SizeTy, countFromCLAs); // If we had any VLA dimensions, factor them in. if (numVLAElements) @@ -2206,11 +2235,10 @@ CodeGenFunction::getVLASize(const VariableArrayType *type) { } } while ((type = getContext().getAsVariableArrayType(elementType))); - return { numElements, elementType }; + return {numElements, elementType}; } -CodeGenFunction::VlaSizePair -CodeGenFunction::getVLAElements1D(QualType type) { +CodeGenFunction::VlaSizePair CodeGenFunction::getVLAElements1D(QualType type) { const VariableArrayType *vla = getContext().getAsVariableArrayType(type); assert(vla && "type was not a variable array type!"); return getVLAElements1D(vla); @@ -2221,7 +2249,7 @@ CodeGenFunction::getVLAElements1D(const VariableArrayType *Vla) { llvm::Value *VlaSize = VLASizeMap[Vla->getSizeExpr()]; assert(VlaSize && "no size for VLA!"); assert(VlaSize->getType() == SizeTy); - return { VlaSize, Vla->getElementType() }; + return {VlaSize, Vla->getElementType()}; } void CodeGenFunction::EmitVariablyModifiedType(QualType type) { @@ -2370,7 +2398,7 @@ void CodeGenFunction::EmitVariablyModifiedType(QualType type) { } while (type->isVariablyModifiedType()); } -Address CodeGenFunction::EmitVAListRef(const Expr* E) { +Address CodeGenFunction::EmitVAListRef(const Expr *E) { if (getContext().getBuiltinVaListType()->isArrayType()) return EmitPointerWithAlignment(E); return EmitLValue(E).getAddress(*this); @@ -2394,9 +2422,11 @@ CodeGenFunction::protectFromPeepholes(RValue rvalue) { // is trunc(zext) folding, but if we add more, we can easily // extend this protection. - if (!rvalue.isScalar()) return PeepholeProtection(); + if (!rvalue.isScalar()) + return PeepholeProtection(); llvm::Value *value = rvalue.getScalarVal(); - if (!isa(value)) return PeepholeProtection(); + if (!isa(value)) + return PeepholeProtection(); // Just make an extra bitcast. assert(HaveInsertPoint()); @@ -2409,7 +2439,8 @@ CodeGenFunction::protectFromPeepholes(RValue rvalue) { } void CodeGenFunction::unprotectFromPeepholes(PeepholeProtection protection) { - if (!protection.Inst) return; + if (!protection.Inst) + return; // In theory, we could try to duplicate the peepholes now, but whatever. protection.Inst->eraseFromParent(); @@ -2520,8 +2551,8 @@ Address CodeGenFunction::EmitFieldAnnotations(const FieldDecl *D, return Address(V, Addr.getAlignment()); } - llvm::Function *F = CGM.getIntrinsic(llvm::Intrinsic::ptr_annotation, - CGM.Int8PtrTy); + llvm::Function *F = + CGM.getIntrinsic(llvm::Intrinsic::ptr_annotation, CGM.Int8PtrTy); for (const auto *I : D->specific_attrs()) { V = Builder.CreateBitCast(V, CGM.Int8PtrTy); @@ -2546,8 +2577,7 @@ Address CodeGenFunction::EmitIntelFPGAFieldAnnotations(SourceLocation Location, // llvm.ptr.annotation intrinsic accepts a pointer to integer of any width - // don't perform bitcasts if value is integer if (VTy->getPointerElementType()->isIntegerTy()) { - llvm::Function *F = - CGM.getIntrinsic(llvm::Intrinsic::ptr_annotation, VTy); + llvm::Function *F = CGM.getIntrinsic(llvm::Intrinsic::ptr_annotation, VTy); V = EmitAnnotationCall(F, V, AnnotStr, Location); return Address(V, Addr.getAlignment()); @@ -2564,7 +2594,7 @@ Address CodeGenFunction::EmitIntelFPGAFieldAnnotations(SourceLocation Location, return Address(V, Addr.getAlignment()); } -CodeGenFunction::CGCapturedStmtInfo::~CGCapturedStmtInfo() { } +CodeGenFunction::CGCapturedStmtInfo::~CGCapturedStmtInfo() {} CodeGenFunction::SanitizerScope::SanitizerScope(CodeGenFunction *CGF) : CGF(CGF) { @@ -2627,8 +2657,7 @@ void CodeGenFunction::checkTargetFeatures(SourceLocation Loc, // Return if the builtin doesn't have any required features. if (FeatureList.empty()) return; - assert(FeatureList.find(' ') == StringRef::npos && - "Space in feature list"); + assert(FeatureList.find(' ') == StringRef::npos && "Space in feature list"); TargetFeatures TF(CallerFeatureMap); if (!TF.hasRequiredFeatures(FeatureList)) CGM.getDiags().Report(Loc, diag::err_builtin_needs_feature) @@ -2656,12 +2685,12 @@ void CodeGenFunction::checkTargetFeatures(SourceLocation Loc, ReqFeatures.push_back(F.getKey()); } if (!llvm::all_of(ReqFeatures, [&](StringRef Feature) { - if (!CallerFeatureMap.lookup(Feature)) { - MissingFeature = Feature.str(); - return false; - } - return true; - })) + if (!CallerFeatureMap.lookup(Feature)) { + MissingFeature = Feature.str(); + return false; + } + return true; + })) CGM.getDiags().Report(Loc, diag::err_function_needs_feature) << FD->getDeclName() << TargetDecl->getDeclName() << MissingFeature; } diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index 9ffa0ce900698..d93c6e30b497d 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -197,8 +197,7 @@ bool Sema::isKnownGoodSYCLDecl(const Decl *D) { if (const FunctionDecl *FD = dyn_cast(D)) { const IdentifierInfo *II = FD->getIdentifier(); const DeclContext *DC = FD->getDeclContext(); - if (II && II->isStr("__spirv_ocl_printf") && - !FD->isDefined() && + if (II && II->isStr("__spirv_ocl_printf") && !FD->isDefined() && FD->getLanguageLinkage() == CXXLanguageLinkage && DC->getEnclosingNamespaceContext()->isTranslationUnit()) return true; @@ -1666,6 +1665,25 @@ class SyclKernelDeclCreator : public SyclKernelFieldHandler { void addParam(const FieldDecl *FD, QualType FieldTy) { ParamDesc newParamDesc = makeParamDesc(FD, FieldTy); + + // zahira +#if 0 + // zahira + StringRef FuncName = KernelDecl->getName(); + StringRef ArgName = FD->getName(); + std::string ArgType = FieldTy.getAsString(); + SourceLocation ArgLoc = FD->getLocation(); + + printf("AddParam FD\n"); + printf("FuncName: %s\n", FuncName.data()); + printf("ArgName: %s \n", ArgName.data()); + printf("ArgType: %s \n", ArgType.data()); +#endif +#if 1 + SemaRef.getDiagnostics().OptReportHandler.AddKernelArgs( + KernelDecl, FD->getName(), FieldTy.getAsString(), + FD->getLocation()); +#endif addParam(newParamDesc, FieldTy); } @@ -1676,6 +1694,22 @@ class SyclKernelDeclCreator : public SyclKernelFieldHandler { StringRef Name = "_arg__base"; ParamDesc newParamDesc = makeParamDesc(SemaRef.getASTContext(), Name, FieldTy); + // zahira +#if 0 + StringRef FuncName = KernelDecl->getName(); + StringRef ArgName = "_arg__base"; + std::string ArgType = FieldTy.getAsString(); + SourceLocation KernelArgLoc = BS.getBaseTypeLoc(); + printf("AddParam BS\n"); + printf("FuncName: %s\n", FuncName.data()); + printf("ArgName: %s \n", Name.data()); + printf("ArgType: %s \n", ArgType.data()); +#endif +#if 1 + SemaRef.getDiagnostics().OptReportHandler.AddKernelArgs( + KernelDecl, "_arg__base", FieldTy.getAsString(), + BS.getBaseTypeLoc()); +#endif addParam(newParamDesc, FieldTy); } // Add a parameter with specified name and type @@ -3588,8 +3622,8 @@ bool Sema::checkSYCLDeviceFunction(SourceLocation Loc, FunctionDecl *Callee) { SemaDiagnosticBuilder(DiagKind, Loc, diag::err_sycl_restrict, Caller, *this) << Sema::KernelCallUndefinedFunction; - SemaDiagnosticBuilder(DiagKind, Callee->getLocation(), diag::note_previous_decl, - Caller, *this) + SemaDiagnosticBuilder(DiagKind, Callee->getLocation(), + diag::note_previous_decl, Caller, *this) << Callee; return DiagKind != SemaDiagnosticBuilder::K_Immediate && @@ -3612,8 +3646,7 @@ void Sema::finalizeSYCLDelayedAnalysis(const FunctionDecl *Caller, bool NotDefinedNoAttr = !Callee->isDefined() && !HasAttr; if (NotDefinedNoAttr && !Callee->getBuiltinID()) { - Diag(Loc, diag::err_sycl_restrict) - << Sema::KernelCallUndefinedFunction; + Diag(Loc, diag::err_sycl_restrict) << Sema::KernelCallUndefinedFunction; Diag(Callee->getLocation(), diag::note_previous_decl) << Callee; Diag(Caller->getLocation(), diag::note_called_by) << Caller; } From 8b5d07cee94922ab58aa6ebe477ccd4ce8066d0a Mon Sep 17 00:00:00 2001 From: Zahira Ammarguellat Date: Fri, 2 Apr 2021 08:36:11 -0700 Subject: [PATCH 02/19] OptReport for kernel arguments. --- .../clang/Basic/SyclOptReportHandler.h | 1 - clang/lib/CodeGen/CodeGenAction.cpp | 56 +- clang/lib/CodeGen/CodeGenAction.cpp.mine | 1248 ----------------- clang/lib/CodeGen/CodeGenFunction.cpp | 39 +- 4 files changed, 85 insertions(+), 1259 deletions(-) delete mode 100644 clang/lib/CodeGen/CodeGenAction.cpp.mine diff --git a/clang/include/clang/Basic/SyclOptReportHandler.h b/clang/include/clang/Basic/SyclOptReportHandler.h index 9d9bdbc8ba713..f57451c38d163 100644 --- a/clang/include/clang/Basic/SyclOptReportHandler.h +++ b/clang/include/clang/Basic/SyclOptReportHandler.h @@ -41,7 +41,6 @@ class SyclOptReportHandler { void AddKernelArgs(const FunctionDecl *FD, StringRef ArgName, std::string ArgType, SourceLocation ArgLoc) { Map[FD].emplace_back(ArgName, ArgType, ArgLoc); - auto It = Map.find(FD); } SmallVector &getInfo(const FunctionDecl *FD) { auto It = Map.find(FD); diff --git a/clang/lib/CodeGen/CodeGenAction.cpp b/clang/lib/CodeGen/CodeGenAction.cpp index 73e1ee7013dc4..6e644f140f322 100644 --- a/clang/lib/CodeGen/CodeGenAction.cpp +++ b/clang/lib/CodeGen/CodeGenAction.cpp @@ -304,11 +304,16 @@ namespace clang { return; LLVMContext &Ctx = getModule()->getContext(); + // zahira + //LLVMContext::InlineAsmDiagHandlerTy OldHandler = + // Ctx.getInlineAsmDiagnosticHandler(); + // end zahira + std::unique_ptr OldDiagnosticHandler = Ctx.getDiagnosticHandler(); Ctx.setDiagnosticHandler(std::make_unique( CodeGenOpts, this)); - +#if 0 Expected> OptRecordFileOrErr = setupLLVMOptimizationRemarks( Ctx, CodeGenOpts.OptRecordFile, CodeGenOpts.OptRecordPasses, @@ -326,7 +331,7 @@ namespace clang { if (OptRecordFile && CodeGenOpts.getProfileUse() != CodeGenOptions::ProfileNone) Ctx.setDiagnosticsHotnessRequested(true); - +#endif // The parallel_for_work_group legalization pass can emit calls to // builtins function. Definitions of those builtins can be provided in // LinkModule. We force the pass to legalize the code before the link @@ -349,9 +354,10 @@ namespace clang { getModule(), Action, std::move(AsmOutStream)); Ctx.setDiagnosticHandler(std::move(OldDiagnosticHandler)); - +#if 0 if (OptRecordFile) OptRecordFile->keep(); +#endif } void HandleTagDeclDefinition(TagDecl *D) override { @@ -1046,8 +1052,52 @@ CodeGenAction::loadModule(MemoryBufferRef MBRef) { return {}; } +namespace { + // Handles the initialization and cleanup of the OptRecordFile. This + // customization allows initialization before the clang codegen runs + // so it can also emit to the opt report. + struct OptRecordFileRAII { + std::unique_ptr OptRecordFile; + std::unique_ptr OldDiagnosticHandler; + llvm::LLVMContext &Ctx; + + OptRecordFileRAII(CodeGenAction &CGA, llvm::LLVMContext &Ctx, + BackendConsumer &BC) + : Ctx(Ctx) { + CompilerInstance &CI = CGA.getCompilerInstance(); + auto &CodeGenOpts = CI.getCodeGenOpts(); + + OldDiagnosticHandler = Ctx.getDiagnosticHandler(); + + Ctx.setDiagnosticHandler( + std::make_unique(CodeGenOpts, &BC)); + + Expected> OptRecordFileOrErr = + setupLLVMOptimizationRemarks( + Ctx, CodeGenOpts.OptRecordFile, CodeGenOpts.OptRecordPasses, + CodeGenOpts.OptRecordFormat, CodeGenOpts.DiagnosticsWithHotness, + CodeGenOpts.DiagnosticsHotnessThreshold); + + if (Error E = OptRecordFileOrErr.takeError()) + reportOptRecordError(std::move(E), CI.getDiagnostics(), CodeGenOpts); + else + OptRecordFile = std::move(*OptRecordFileOrErr); + + if (OptRecordFile && + CodeGenOpts.getProfileUse() != CodeGenOptions::ProfileNone) + Ctx.setDiagnosticsHotnessRequested(true); + } + ~OptRecordFileRAII() { + Ctx.setDiagnosticHandler(std::move(OldDiagnosticHandler)); + if (OptRecordFile) + OptRecordFile->keep(); + } + }; +} // namespace + void CodeGenAction::ExecuteAction() { if (getCurrentFileKind().getLanguage() != Language::LLVM_IR) { + OptRecordFileRAII ORF(*this, *VMContext, *BEConsumer); this->ASTFrontendAction::ExecuteAction(); return; } diff --git a/clang/lib/CodeGen/CodeGenAction.cpp.mine b/clang/lib/CodeGen/CodeGenAction.cpp.mine deleted file mode 100644 index 7948c1244737d..0000000000000 --- a/clang/lib/CodeGen/CodeGenAction.cpp.mine +++ /dev/null @@ -1,1248 +0,0 @@ -//===--- CodeGenAction.cpp - LLVM Code Generation Frontend Action ---------===// -// -// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. -// See https://llvm.org/LICENSE.txt for license information. -// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception -// -//===----------------------------------------------------------------------===// - -#include "clang/CodeGen/CodeGenAction.h" -#include "CodeGenModule.h" -#include "CoverageMappingGen.h" -#include "MacroPPCallbacks.h" -#include "clang/AST/ASTConsumer.h" -#include "clang/AST/ASTContext.h" -#include "clang/AST/DeclCXX.h" -#include "clang/AST/DeclGroup.h" -#include "clang/Basic/DiagnosticFrontend.h" -#include "clang/Basic/FileManager.h" -#include "clang/Basic/LangStandard.h" -#include "clang/Basic/SourceManager.h" -#include "clang/Basic/TargetInfo.h" -#include "clang/CodeGen/BackendUtil.h" -#include "clang/CodeGen/ModuleBuilder.h" -#include "clang/Driver/DriverDiagnostic.h" -#include "clang/Frontend/CompilerInstance.h" -#include "clang/Frontend/FrontendDiagnostic.h" -#include "clang/Lex/Preprocessor.h" -#include "llvm/Bitcode/BitcodeReader.h" -#include "llvm/CodeGen/MachineOptimizationRemarkEmitter.h" -#include "llvm/IR/DebugInfo.h" -#include "llvm/IR/DiagnosticInfo.h" -#include "llvm/IR/DiagnosticPrinter.h" -#include "llvm/IR/GlobalValue.h" -#include "llvm/IR/LLVMContext.h" -#include "llvm/IR/LLVMRemarkStreamer.h" -#include "llvm/IR/LegacyPassManager.h" -#include "llvm/IR/Module.h" -#include "llvm/IRReader/IRReader.h" -#include "llvm/LTO/LTOBackend.h" -#include "llvm/Linker/Linker.h" -#include "llvm/Pass.h" -#include "llvm/SYCLLowerIR/LowerWGScope.h" -#include "llvm/Support/MemoryBuffer.h" -#include "llvm/Support/SourceMgr.h" -#include "llvm/Support/TimeProfiler.h" -#include "llvm/Support/Timer.h" -#include "llvm/Support/ToolOutputFile.h" -#include "llvm/Support/YAMLTraits.h" -#include "llvm/Transforms/IPO/Internalize.h" - -#include -using namespace clang; -using namespace llvm; - -namespace clang { - class BackendConsumer; - class ClangDiagnosticHandler final : public DiagnosticHandler { - public: - ClangDiagnosticHandler(const CodeGenOptions &CGOpts, BackendConsumer *BCon) - : CodeGenOpts(CGOpts), BackendCon(BCon) {} - - bool handleDiagnostics(const DiagnosticInfo &DI) override; - - bool isAnalysisRemarkEnabled(StringRef PassName) const override { - return (CodeGenOpts.OptimizationRemarkAnalysisPattern && - CodeGenOpts.OptimizationRemarkAnalysisPattern->match(PassName)); - } - bool isMissedOptRemarkEnabled(StringRef PassName) const override { - return (CodeGenOpts.OptimizationRemarkMissedPattern && - CodeGenOpts.OptimizationRemarkMissedPattern->match(PassName)); - } - bool isPassedOptRemarkEnabled(StringRef PassName) const override { - return (CodeGenOpts.OptimizationRemarkPattern && - CodeGenOpts.OptimizationRemarkPattern->match(PassName)); - } - - bool isAnyRemarkEnabled() const override { - return (CodeGenOpts.OptimizationRemarkAnalysisPattern || - CodeGenOpts.OptimizationRemarkMissedPattern || - CodeGenOpts.OptimizationRemarkPattern); - } - - private: - const CodeGenOptions &CodeGenOpts; - BackendConsumer *BackendCon; - }; - - static void reportOptRecordError(Error E, DiagnosticsEngine &Diags, - const CodeGenOptions CodeGenOpts) { - handleAllErrors( - std::move(E), - [&](const LLVMRemarkSetupFileError &E) { - Diags.Report(diag::err_cannot_open_file) - << CodeGenOpts.OptRecordFile << E.message(); - }, - [&](const LLVMRemarkSetupPatternError &E) { - Diags.Report(diag::err_drv_optimization_remark_pattern) - << E.message() << CodeGenOpts.OptRecordPasses; - }, - [&](const LLVMRemarkSetupFormatError &E) { - Diags.Report(diag::err_drv_optimization_remark_format) - << CodeGenOpts.OptRecordFormat; - }); - } - - class BackendConsumer : public ASTConsumer { - using LinkModule = CodeGenAction::LinkModule; - - virtual void anchor(); - DiagnosticsEngine &Diags; - BackendAction Action; - const HeaderSearchOptions &HeaderSearchOpts; - const CodeGenOptions &CodeGenOpts; - const TargetOptions &TargetOpts; - const LangOptions &LangOpts; - std::unique_ptr AsmOutStream; - ASTContext *Context; - - Timer LLVMIRGeneration; - unsigned LLVMIRGenerationRefCount; - - /// True if we've finished generating IR. This prevents us from generating - /// additional LLVM IR after emitting output in HandleTranslationUnit. This - /// can happen when Clang plugins trigger additional AST deserialization. - bool IRGenFinished = false; - - bool TimerIsEnabled = false; - - std::unique_ptr Gen; - - SmallVector LinkModules; - - // This is here so that the diagnostic printer knows the module a diagnostic - // refers to. - llvm::Module *CurLinkModule = nullptr; - - public: - BackendConsumer(BackendAction Action, DiagnosticsEngine &Diags, - const HeaderSearchOptions &HeaderSearchOpts, - const PreprocessorOptions &PPOpts, - const CodeGenOptions &CodeGenOpts, - const TargetOptions &TargetOpts, - const LangOptions &LangOpts, const std::string &InFile, - SmallVector LinkModules, - std::unique_ptr OS, LLVMContext &C, - CoverageSourceInfo *CoverageInfo = nullptr) - : Diags(Diags), Action(Action), HeaderSearchOpts(HeaderSearchOpts), - CodeGenOpts(CodeGenOpts), TargetOpts(TargetOpts), LangOpts(LangOpts), - AsmOutStream(std::move(OS)), Context(nullptr), - LLVMIRGeneration("irgen", "LLVM IR Generation Time"), - LLVMIRGenerationRefCount(0), - Gen(CreateLLVMCodeGen(Diags, InFile, HeaderSearchOpts, PPOpts, - CodeGenOpts, C, CoverageInfo)), - LinkModules(std::move(LinkModules)) { - TimerIsEnabled = CodeGenOpts.TimePasses; - llvm::TimePassesIsEnabled = CodeGenOpts.TimePasses; - llvm::TimePassesPerRun = CodeGenOpts.TimePassesPerRun; - } - - // This constructor is used in installing an empty BackendConsumer - // to use the clang diagnostic handler for IR input files. It avoids - // initializing the OS field. - BackendConsumer(BackendAction Action, DiagnosticsEngine &Diags, - const HeaderSearchOptions &HeaderSearchOpts, - const PreprocessorOptions &PPOpts, - const CodeGenOptions &CodeGenOpts, - const TargetOptions &TargetOpts, - const LangOptions &LangOpts, - SmallVector LinkModules, LLVMContext &C, - CoverageSourceInfo *CoverageInfo = nullptr) - : Diags(Diags), Action(Action), HeaderSearchOpts(HeaderSearchOpts), - CodeGenOpts(CodeGenOpts), TargetOpts(TargetOpts), LangOpts(LangOpts), - Context(nullptr), - LLVMIRGeneration("irgen", "LLVM IR Generation Time"), - LLVMIRGenerationRefCount(0), - Gen(CreateLLVMCodeGen(Diags, "", HeaderSearchOpts, PPOpts, - CodeGenOpts, C, CoverageInfo)), - LinkModules(std::move(LinkModules)) { - TimerIsEnabled = CodeGenOpts.TimePasses; - llvm::TimePassesIsEnabled = CodeGenOpts.TimePasses; - llvm::TimePassesPerRun = CodeGenOpts.TimePassesPerRun; - } - llvm::Module *getModule() const { return Gen->GetModule(); } - std::unique_ptr takeModule() { - return std::unique_ptr(Gen->ReleaseModule()); - } - - CodeGenerator *getCodeGenerator() { return Gen.get(); } - - void HandleCXXStaticMemberVarInstantiation(VarDecl *VD) override { - Gen->HandleCXXStaticMemberVarInstantiation(VD); - } - - void Initialize(ASTContext &Ctx) override { - assert(!Context && "initialized multiple times"); - - Context = &Ctx; - - if (TimerIsEnabled) - LLVMIRGeneration.startTimer(); - - Gen->Initialize(Ctx); - - if (TimerIsEnabled) - LLVMIRGeneration.stopTimer(); - } - - bool HandleTopLevelDecl(DeclGroupRef D) override { - PrettyStackTraceDecl CrashInfo(*D.begin(), SourceLocation(), - Context->getSourceManager(), - "LLVM IR generation of declaration"); - - // Recurse. - if (TimerIsEnabled) { - LLVMIRGenerationRefCount += 1; - if (LLVMIRGenerationRefCount == 1) - LLVMIRGeneration.startTimer(); - } - - Gen->HandleTopLevelDecl(D); - - if (TimerIsEnabled) { - LLVMIRGenerationRefCount -= 1; - if (LLVMIRGenerationRefCount == 0) - LLVMIRGeneration.stopTimer(); - } - - return true; - } - - void HandleInlineFunctionDefinition(FunctionDecl *D) override { - PrettyStackTraceDecl CrashInfo(D, SourceLocation(), - Context->getSourceManager(), - "LLVM IR generation of inline function"); - if (TimerIsEnabled) - LLVMIRGeneration.startTimer(); - - Gen->HandleInlineFunctionDefinition(D); - - if (TimerIsEnabled) - LLVMIRGeneration.stopTimer(); - } - - void HandleInterestingDecl(DeclGroupRef D) override { - // Ignore interesting decls from the AST reader after IRGen is finished. - if (!IRGenFinished) - HandleTopLevelDecl(D); - } - - // Links each entry in LinkModules into our module. Returns true on error. - bool LinkInModules() { - for (auto &LM : LinkModules) { - if (LM.PropagateAttrs) - for (Function &F : *LM.Module) { - // Skip intrinsics. Keep consistent with how intrinsics are created - // in LLVM IR. - if (F.isIntrinsic()) - continue; - Gen->CGM().addDefaultFunctionDefinitionAttributes(F); - } - - CurLinkModule = LM.Module.get(); - - bool Err; - if (LM.Internalize) { - Err = Linker::linkModules( - *getModule(), std::move(LM.Module), LM.LinkFlags, - [](llvm::Module &M, const llvm::StringSet<> &GVS) { - internalizeModule(M, [&GVS](const llvm::GlobalValue &GV) { - return !GV.hasName() || (GVS.count(GV.getName()) == 0); - }); - }); - } else { - Err = Linker::linkModules(*getModule(), std::move(LM.Module), - LM.LinkFlags); - } - - if (Err) - return true; - } - return false; // success - } - - void HandleTranslationUnit(ASTContext &C) override { - { - llvm::TimeTraceScope TimeScope("Frontend"); - PrettyStackTraceString CrashInfo("Per-file LLVM IR generation"); - if (TimerIsEnabled) { - LLVMIRGenerationRefCount += 1; - if (LLVMIRGenerationRefCount == 1) - LLVMIRGeneration.startTimer(); - } - - Gen->HandleTranslationUnit(C); - - if (TimerIsEnabled) { - LLVMIRGenerationRefCount -= 1; - if (LLVMIRGenerationRefCount == 0) - LLVMIRGeneration.stopTimer(); - } - - IRGenFinished = true; - } - - // Silently ignore if we weren't initialized for some reason. - if (!getModule()) - return; - - // Install an inline asm handler so that diagnostics get printed through - // our diagnostics hooks. - LLVMContext &Ctx = getModule()->getContext(); - LLVMContext::InlineAsmDiagHandlerTy OldHandler = - Ctx.getInlineAsmDiagnosticHandler(); - void *OldContext = Ctx.getInlineAsmDiagnosticContext(); - Ctx.setInlineAsmDiagnosticHandler(InlineAsmDiagHandler, this); -#if !INTEL_CUSTOMIZATION - std::unique_ptr \ - OldDiagnosticHandler = - Ctx.getDiagnosticHandler(); - Ctx.setDiagnosticHandler(std::make_unique( - CodeGenOpts, this)); - - Expected> OptRecordFileOrErr = - setupLLVMOptimizationRemarks( - Ctx, CodeGenOpts.OptRecordFile, CodeGenOpts.OptRecordPasses, - CodeGenOpts.OptRecordFormat, CodeGenOpts.DiagnosticsWithHotness, - CodeGenOpts.DiagnosticsHotnessThreshold); - - if (Error E = OptRecordFileOrErr.takeError()) { - reportOptRecordError(std::move(E), Diags, CodeGenOpts); - return; - } - - std::unique_ptr OptRecordFile = - std::move(*OptRecordFileOrErr); - - if (OptRecordFile && - CodeGenOpts.getProfileUse() != CodeGenOptions::ProfileNone) - Ctx.setDiagnosticsHotnessRequested(true); -#endif // INTEL_CUSTOMIZATION - - // The parallel_for_work_group legalization pass can emit calls to - // builtins function. Definitions of those builtins can be provided in - // LinkModule. We force the pass to legalize the code before the link - // happens. - if (LangOpts.SYCLIsDevice) { - PrettyStackTraceString CrashInfo("Pre-linking SYCL passes"); - legacy::PassManager PreLinkingSyclPasses; - PreLinkingSyclPasses.add(llvm::createSYCLLowerWGScopePass()); - PreLinkingSyclPasses.run(*getModule()); - } - - // Link each LinkModule into our module. - if (LinkInModules()) - return; - - EmbedBitcode(getModule(), CodeGenOpts, llvm::MemoryBufferRef()); - - EmitBackendOutput(Diags, HeaderSearchOpts, CodeGenOpts, TargetOpts, - LangOpts, C.getTargetInfo().getDataLayout(), - getModule(), Action, std::move(AsmOutStream)); - - Ctx.setInlineAsmDiagnosticHandler(OldHandler, OldContext); -#if !INTEL_CUSTOMIZATION - Ctx.setDiagnosticHandler(std::move(OldDiagnosticHandler)); - - if (OptRecordFile) - OptRecordFile->keep(); -#endif - } - - void HandleTagDeclDefinition(TagDecl *D) override { - PrettyStackTraceDecl CrashInfo(D, SourceLocation(), - Context->getSourceManager(), - "LLVM IR generation of declaration"); - Gen->HandleTagDeclDefinition(D); - } - - void HandleTagDeclRequiredDefinition(const TagDecl *D) override { - Gen->HandleTagDeclRequiredDefinition(D); - } - - void CompleteTentativeDefinition(VarDecl *D) override { - Gen->CompleteTentativeDefinition(D); - } - - void CompleteExternalDeclaration(VarDecl *D) override { - Gen->CompleteExternalDeclaration(D); - } - - void AssignInheritanceModel(CXXRecordDecl *RD) override { - Gen->AssignInheritanceModel(RD); - } - - void HandleVTable(CXXRecordDecl *RD) override { - Gen->HandleVTable(RD); - } - - static void InlineAsmDiagHandler(const llvm::SMDiagnostic &SM,void *Context, - unsigned LocCookie) { - SourceLocation Loc = SourceLocation::getFromRawEncoding(LocCookie); - ((BackendConsumer*)Context)->InlineAsmDiagHandler2(SM, Loc); - } - - /// Get the best possible source location to represent a diagnostic that - /// may have associated debug info. - const FullSourceLoc - getBestLocationFromDebugLoc(const llvm::DiagnosticInfoWithLocationBase &D, - bool &BadDebugInfo, StringRef &Filename, - unsigned &Line, unsigned &Column) const; - - void InlineAsmDiagHandler2(const llvm::SMDiagnostic &, - SourceLocation LocCookie); - - void DiagnosticHandlerImpl(const llvm::DiagnosticInfo &DI); - /// Specialized handler for InlineAsm diagnostic. - /// \return True if the diagnostic has been successfully reported, false - /// otherwise. - bool InlineAsmDiagHandler(const llvm::DiagnosticInfoInlineAsm &D); - /// Specialized handler for StackSize diagnostic. - /// \return True if the diagnostic has been successfully reported, false - /// otherwise. - bool StackSizeDiagHandler(const llvm::DiagnosticInfoStackSize &D); - /// Specialized handler for unsupported backend feature diagnostic. - void UnsupportedDiagHandler(const llvm::DiagnosticInfoUnsupported &D); - /// Specialized handlers for optimization remarks. - /// Note that these handlers only accept remarks and they always handle - /// them. - void EmitOptimizationMessage(const llvm::DiagnosticInfoOptimizationBase &D, - unsigned DiagID); - void - OptimizationRemarkHandler(const llvm::DiagnosticInfoOptimizationBase &D); - void OptimizationRemarkHandler( - const llvm::OptimizationRemarkAnalysisFPCommute &D); - void OptimizationRemarkHandler( - const llvm::OptimizationRemarkAnalysisAliasing &D); - void OptimizationFailureHandler( - const llvm::DiagnosticInfoOptimizationFailure &D); - }; - - void BackendConsumer::anchor() {} -} - -bool ClangDiagnosticHandler::handleDiagnostics(const DiagnosticInfo &DI) { - BackendCon->DiagnosticHandlerImpl(DI); - return true; -} - -/// ConvertBackendLocation - Convert a location in a temporary llvm::SourceMgr -/// buffer to be a valid FullSourceLoc. -static FullSourceLoc ConvertBackendLocation(const llvm::SMDiagnostic &D, - SourceManager &CSM) { - // Get both the clang and llvm source managers. The location is relative to - // a memory buffer that the LLVM Source Manager is handling, we need to add - // a copy to the Clang source manager. - const llvm::SourceMgr &LSM = *D.getSourceMgr(); - - // We need to copy the underlying LLVM memory buffer because llvm::SourceMgr - // already owns its one and clang::SourceManager wants to own its one. - const MemoryBuffer *LBuf = - LSM.getMemoryBuffer(LSM.FindBufferContainingLoc(D.getLoc())); - - // Create the copy and transfer ownership to clang::SourceManager. - // TODO: Avoid copying files into memory. - std::unique_ptr CBuf = - llvm::MemoryBuffer::getMemBufferCopy(LBuf->getBuffer(), - LBuf->getBufferIdentifier()); - // FIXME: Keep a file ID map instead of creating new IDs for each location. - FileID FID = CSM.createFileID(std::move(CBuf)); - - // Translate the offset into the file. - unsigned Offset = D.getLoc().getPointer() - LBuf->getBufferStart(); - SourceLocation NewLoc = - CSM.getLocForStartOfFile(FID).getLocWithOffset(Offset); - return FullSourceLoc(NewLoc, CSM); -} - - -/// InlineAsmDiagHandler2 - This function is invoked when the backend hits an -/// error parsing inline asm. The SMDiagnostic indicates the error relative to -/// the temporary memory buffer that the inline asm parser has set up. -void BackendConsumer::InlineAsmDiagHandler2(const llvm::SMDiagnostic &D, - SourceLocation LocCookie) { - // There are a couple of different kinds of errors we could get here. First, - // we re-format the SMDiagnostic in terms of a clang diagnostic. - - // Strip "error: " off the start of the message string. - StringRef Message = D.getMessage(); - if (Message.startswith("error: ")) - Message = Message.substr(7); - - // If the SMDiagnostic has an inline asm source location, translate it. - FullSourceLoc Loc; - if (D.getLoc() != SMLoc()) - Loc = ConvertBackendLocation(D, Context->getSourceManager()); - - unsigned DiagID; - switch (D.getKind()) { - case llvm::SourceMgr::DK_Error: - DiagID = diag::err_fe_inline_asm; - break; - case llvm::SourceMgr::DK_Warning: - DiagID = diag::warn_fe_inline_asm; - break; - case llvm::SourceMgr::DK_Note: - DiagID = diag::note_fe_inline_asm; - break; - case llvm::SourceMgr::DK_Remark: - llvm_unreachable("remarks unexpected"); - } - // If this problem has clang-level source location information, report the - // issue in the source with a note showing the instantiated - // code. - if (LocCookie.isValid()) { - Diags.Report(LocCookie, DiagID).AddString(Message); - - if (D.getLoc().isValid()) { - DiagnosticBuilder B = Diags.Report(Loc, diag::note_fe_inline_asm_here); - // Convert the SMDiagnostic ranges into SourceRange and attach them - // to the diagnostic. - for (const std::pair &Range : D.getRanges()) { - unsigned Column = D.getColumnNo(); - B << SourceRange(Loc.getLocWithOffset(Range.first - Column), - Loc.getLocWithOffset(Range.second - Column)); - } - } - return; - } - - // Otherwise, report the backend issue as occurring in the generated .s file. - // If Loc is invalid, we still need to report the issue, it just gets no - // location info. - Diags.Report(Loc, DiagID).AddString(Message); -} - -#define ComputeDiagID(Severity, GroupName, DiagID) \ - do { \ - switch (Severity) { \ - case llvm::DS_Error: \ - DiagID = diag::err_fe_##GroupName; \ - break; \ - case llvm::DS_Warning: \ - DiagID = diag::warn_fe_##GroupName; \ - break; \ - case llvm::DS_Remark: \ - llvm_unreachable("'remark' severity not expected"); \ - break; \ - case llvm::DS_Note: \ - DiagID = diag::note_fe_##GroupName; \ - break; \ - } \ - } while (false) - -#define ComputeDiagRemarkID(Severity, GroupName, DiagID) \ - do { \ - switch (Severity) { \ - case llvm::DS_Error: \ - DiagID = diag::err_fe_##GroupName; \ - break; \ - case llvm::DS_Warning: \ - DiagID = diag::warn_fe_##GroupName; \ - break; \ - case llvm::DS_Remark: \ - DiagID = diag::remark_fe_##GroupName; \ - break; \ - case llvm::DS_Note: \ - DiagID = diag::note_fe_##GroupName; \ - break; \ - } \ - } while (false) - -bool -BackendConsumer::InlineAsmDiagHandler(const llvm::DiagnosticInfoInlineAsm &D) { - unsigned DiagID; - ComputeDiagID(D.getSeverity(), inline_asm, DiagID); - std::string Message = D.getMsgStr().str(); - - // If this problem has clang-level source location information, report the - // issue as being a problem in the source with a note showing the instantiated - // code. - SourceLocation LocCookie = - SourceLocation::getFromRawEncoding(D.getLocCookie()); - if (LocCookie.isValid()) - Diags.Report(LocCookie, DiagID).AddString(Message); - else { - // Otherwise, report the backend diagnostic as occurring in the generated - // .s file. - // If Loc is invalid, we still need to report the diagnostic, it just gets - // no location info. - FullSourceLoc Loc; - Diags.Report(Loc, DiagID).AddString(Message); - } - // We handled all the possible severities. - return true; -} - -bool -BackendConsumer::StackSizeDiagHandler(const llvm::DiagnosticInfoStackSize &D) { - if (D.getSeverity() != llvm::DS_Warning) - // For now, the only support we have for StackSize diagnostic is warning. - // We do not know how to format other severities. - return false; - - if (const Decl *ND = Gen->GetDeclForMangledName(D.getFunction().getName())) { - // FIXME: Shouldn't need to truncate to uint32_t - Diags.Report(ND->getASTContext().getFullLoc(ND->getLocation()), - diag::warn_fe_frame_larger_than) - << static_cast(D.getStackSize()) << Decl::castToDeclContext(ND); - return true; - } - - return false; -} - -const FullSourceLoc BackendConsumer::getBestLocationFromDebugLoc( - const llvm::DiagnosticInfoWithLocationBase &D, bool &BadDebugInfo, - StringRef &Filename, unsigned &Line, unsigned &Column) const { - SourceManager &SourceMgr = Context->getSourceManager(); - FileManager &FileMgr = SourceMgr.getFileManager(); - SourceLocation DILoc; - - if (D.isLocationAvailable()) { - D.getLocation(Filename, Line, Column); - if (Line > 0) { - auto FE = FileMgr.getFile(Filename); - if (!FE) - FE = FileMgr.getFile(D.getAbsolutePath()); - if (FE) { - // If -gcolumn-info was not used, Column will be 0. This upsets the - // source manager, so pass 1 if Column is not set. - DILoc = SourceMgr.translateFileLineCol(*FE, Line, Column ? Column : 1); - } - } - BadDebugInfo = DILoc.isInvalid(); - } - - // If a location isn't available, try to approximate it using the associated - // function definition. We use the definition's right brace to differentiate - // from diagnostics that genuinely relate to the function itself. - FullSourceLoc Loc(DILoc, SourceMgr); - if (Loc.isInvalid()) - if (const Decl *FD = Gen->GetDeclForMangledName(D.getFunction().getName())) - Loc = FD->getASTContext().getFullLoc(FD->getLocation()); - - if (DILoc.isInvalid() && D.isLocationAvailable()) - // If we were not able to translate the file:line:col information - // back to a SourceLocation, at least emit a note stating that - // we could not translate this location. This can happen in the - // case of #line directives. - Diags.Report(Loc, diag::note_fe_backend_invalid_loc) - << Filename << Line << Column; - - return Loc; -} - -void BackendConsumer::UnsupportedDiagHandler( - const llvm::DiagnosticInfoUnsupported &D) { - // We only support warnings or errors. - assert(D.getSeverity() == llvm::DS_Error || - D.getSeverity() == llvm::DS_Warning); - - StringRef Filename; - unsigned Line, Column; - bool BadDebugInfo = false; - FullSourceLoc Loc; - std::string Msg; - raw_string_ostream MsgStream(Msg); - - // Context will be nullptr for IR input files, we will construct the diag - // message from llvm::DiagnosticInfoUnsupported. - if (Context != nullptr) { - Loc = getBestLocationFromDebugLoc(D, BadDebugInfo, Filename, Line, Column); - MsgStream << D.getMessage(); - } else { - DiagnosticPrinterRawOStream DP(MsgStream); - D.print(DP); - } - - auto DiagType = D.getSeverity() == llvm::DS_Error - ? diag::err_fe_backend_unsupported - : diag::warn_fe_backend_unsupported; - Diags.Report(Loc, DiagType) << MsgStream.str(); - - if (BadDebugInfo) - // If we were not able to translate the file:line:col information - // back to a SourceLocation, at least emit a note stating that - // we could not translate this location. This can happen in the - // case of #line directives. - Diags.Report(Loc, diag::note_fe_backend_invalid_loc) - << Filename << Line << Column; -} - -void BackendConsumer::EmitOptimizationMessage( - const llvm::DiagnosticInfoOptimizationBase &D, unsigned DiagID) { - // We only support warnings and remarks. - assert(D.getSeverity() == llvm::DS_Remark || - D.getSeverity() == llvm::DS_Warning); - - StringRef Filename; - unsigned Line, Column; - bool BadDebugInfo = false; - FullSourceLoc Loc; - std::string Msg; - raw_string_ostream MsgStream(Msg); - - // Context will be nullptr for IR input files, we will construct the remark - // message from llvm::DiagnosticInfoOptimizationBase. - if (Context != nullptr) { - Loc = getBestLocationFromDebugLoc(D, BadDebugInfo, Filename, Line, Column); - MsgStream << D.getMsg(); - } else { - DiagnosticPrinterRawOStream DP(MsgStream); - D.print(DP); - } - - if (D.getHotness()) - MsgStream << " (hotness: " << *D.getHotness() << ")"; - - Diags.Report(Loc, DiagID) - << AddFlagValue(D.getPassName()) - << MsgStream.str(); - - if (BadDebugInfo) - // If we were not able to translate the file:line:col information - // back to a SourceLocation, at least emit a note stating that - // we could not translate this location. This can happen in the - // case of #line directives. - Diags.Report(Loc, diag::note_fe_backend_invalid_loc) - << Filename << Line << Column; -} - -void BackendConsumer::OptimizationRemarkHandler( - const llvm::DiagnosticInfoOptimizationBase &D) { - // Without hotness information, don't show noisy remarks. - if (D.isVerbose() && !D.getHotness()) - return; - - if (D.isPassed()) { - // Optimization remarks are active only if the -Rpass flag has a regular - // expression that matches the name of the pass name in \p D. - if (CodeGenOpts.OptimizationRemarkPattern && - CodeGenOpts.OptimizationRemarkPattern->match(D.getPassName())) - EmitOptimizationMessage(D, diag::remark_fe_backend_optimization_remark); - } else if (D.isMissed()) { - // Missed optimization remarks are active only if the -Rpass-missed - // flag has a regular expression that matches the name of the pass - // name in \p D. - if (CodeGenOpts.OptimizationRemarkMissedPattern && - CodeGenOpts.OptimizationRemarkMissedPattern->match(D.getPassName())) - EmitOptimizationMessage( - D, diag::remark_fe_backend_optimization_remark_missed); - } else { - assert(D.isAnalysis() && "Unknown remark type"); - - bool ShouldAlwaysPrint = false; - if (auto *ORA = dyn_cast(&D)) - ShouldAlwaysPrint = ORA->shouldAlwaysPrint(); - - if (ShouldAlwaysPrint || - (CodeGenOpts.OptimizationRemarkAnalysisPattern && - CodeGenOpts.OptimizationRemarkAnalysisPattern->match(D.getPassName()))) - EmitOptimizationMessage( - D, diag::remark_fe_backend_optimization_remark_analysis); - } -} - -void BackendConsumer::OptimizationRemarkHandler( - const llvm::OptimizationRemarkAnalysisFPCommute &D) { - // Optimization analysis remarks are active if the pass name is set to - // llvm::DiagnosticInfo::AlwasyPrint or if the -Rpass-analysis flag has a - // regular expression that matches the name of the pass name in \p D. - - if (D.shouldAlwaysPrint() || - (CodeGenOpts.OptimizationRemarkAnalysisPattern && - CodeGenOpts.OptimizationRemarkAnalysisPattern->match(D.getPassName()))) - EmitOptimizationMessage( - D, diag::remark_fe_backend_optimization_remark_analysis_fpcommute); -} - -void BackendConsumer::OptimizationRemarkHandler( - const llvm::OptimizationRemarkAnalysisAliasing &D) { - // Optimization analysis remarks are active if the pass name is set to - // llvm::DiagnosticInfo::AlwasyPrint or if the -Rpass-analysis flag has a - // regular expression that matches the name of the pass name in \p D. - - if (D.shouldAlwaysPrint() || - (CodeGenOpts.OptimizationRemarkAnalysisPattern && - CodeGenOpts.OptimizationRemarkAnalysisPattern->match(D.getPassName()))) - EmitOptimizationMessage( - D, diag::remark_fe_backend_optimization_remark_analysis_aliasing); -} - -void BackendConsumer::OptimizationFailureHandler( - const llvm::DiagnosticInfoOptimizationFailure &D) { - EmitOptimizationMessage(D, diag::warn_fe_backend_optimization_failure); -} - -/// This function is invoked when the backend needs -/// to report something to the user. -void BackendConsumer::DiagnosticHandlerImpl(const DiagnosticInfo &DI) { - unsigned DiagID = diag::err_fe_inline_asm; - llvm::DiagnosticSeverity Severity = DI.getSeverity(); - // Get the diagnostic ID based. - switch (DI.getKind()) { - case llvm::DK_InlineAsm: - if (InlineAsmDiagHandler(cast(DI))) - return; - ComputeDiagID(Severity, inline_asm, DiagID); - break; - case llvm::DK_StackSize: - if (StackSizeDiagHandler(cast(DI))) - return; - ComputeDiagID(Severity, backend_frame_larger_than, DiagID); - break; - case DK_Linker: - assert(CurLinkModule); - // FIXME: stop eating the warnings and notes. - if (Severity != DS_Error) - return; - DiagID = diag::err_fe_cannot_link_module; - break; - case llvm::DK_OptimizationRemark: - // Optimization remarks are always handled completely by this - // handler. There is no generic way of emitting them. - OptimizationRemarkHandler(cast(DI)); - return; - case llvm::DK_OptimizationRemarkMissed: - // Optimization remarks are always handled completely by this - // handler. There is no generic way of emitting them. - OptimizationRemarkHandler(cast(DI)); - return; - case llvm::DK_OptimizationRemarkAnalysis: - // Optimization remarks are always handled completely by this - // handler. There is no generic way of emitting them. - OptimizationRemarkHandler(cast(DI)); - return; - case llvm::DK_OptimizationRemarkAnalysisFPCommute: - // Optimization remarks are always handled completely by this - // handler. There is no generic way of emitting them. - OptimizationRemarkHandler(cast(DI)); - return; - case llvm::DK_OptimizationRemarkAnalysisAliasing: - // Optimization remarks are always handled completely by this - // handler. There is no generic way of emitting them. - OptimizationRemarkHandler(cast(DI)); - return; - case llvm::DK_MachineOptimizationRemark: - // Optimization remarks are always handled completely by this - // handler. There is no generic way of emitting them. - OptimizationRemarkHandler(cast(DI)); - return; - case llvm::DK_MachineOptimizationRemarkMissed: - // Optimization remarks are always handled completely by this - // handler. There is no generic way of emitting them. - OptimizationRemarkHandler(cast(DI)); - return; - case llvm::DK_MachineOptimizationRemarkAnalysis: - // Optimization remarks are always handled completely by this - // handler. There is no generic way of emitting them. - OptimizationRemarkHandler(cast(DI)); - return; - case llvm::DK_OptimizationFailure: - // Optimization failures are always handled completely by this - // handler. - OptimizationFailureHandler(cast(DI)); - return; - case llvm::DK_Unsupported: - UnsupportedDiagHandler(cast(DI)); - return; - default: - // Plugin IDs are not bound to any value as they are set dynamically. - ComputeDiagRemarkID(Severity, backend_plugin, DiagID); - break; - } - std::string MsgStorage; - { - raw_string_ostream Stream(MsgStorage); - DiagnosticPrinterRawOStream DP(Stream); - DI.print(DP); - } - - if (DiagID == diag::err_fe_cannot_link_module) { - Diags.Report(diag::err_fe_cannot_link_module) - << CurLinkModule->getModuleIdentifier() << MsgStorage; - return; - } - - // Report the backend message using the usual diagnostic mechanism. - FullSourceLoc Loc; - Diags.Report(Loc, DiagID).AddString(MsgStorage); -} -#undef ComputeDiagID - -CodeGenAction::CodeGenAction(unsigned _Act, LLVMContext *_VMContext) - : Act(_Act), VMContext(_VMContext ? _VMContext : new LLVMContext), - OwnsVMContext(!_VMContext) {} - -CodeGenAction::~CodeGenAction() { - TheModule.reset(); - if (OwnsVMContext) - delete VMContext; -} - -bool CodeGenAction::hasIRSupport() const { return true; } - -void CodeGenAction::EndSourceFileAction() { - // If the consumer creation failed, do nothing. - if (!getCompilerInstance().hasASTConsumer()) - return; - - // Steal the module from the consumer. - TheModule = BEConsumer->takeModule(); -} - -std::unique_ptr CodeGenAction::takeModule() { - return std::move(TheModule); -} - -llvm::LLVMContext *CodeGenAction::takeLLVMContext() { - OwnsVMContext = false; - return VMContext; -} - -static std::unique_ptr -GetOutputStream(CompilerInstance &CI, StringRef InFile, BackendAction Action) { - switch (Action) { - case Backend_EmitAssembly: - return CI.createDefaultOutputFile(false, InFile, "s"); - case Backend_EmitLL: - return CI.createDefaultOutputFile(false, InFile, "ll"); - case Backend_EmitBC: - return CI.createDefaultOutputFile(true, InFile, "bc"); - case Backend_EmitNothing: - return nullptr; - case Backend_EmitMCNull: - return CI.createNullOutputFile(); - case Backend_EmitObj: - return CI.createDefaultOutputFile(true, InFile, "o"); - } - - llvm_unreachable("Invalid action!"); -} - -std::unique_ptr -CodeGenAction::CreateASTConsumer(CompilerInstance &CI, StringRef InFile) { - BackendAction BA = static_cast(Act); - std::unique_ptr OS = CI.takeOutputStream(); - if (!OS) - OS = GetOutputStream(CI, InFile, BA); - - if (BA != Backend_EmitNothing && !OS) - return nullptr; - - // Load bitcode modules to link with, if we need to. - if (LinkModules.empty()) - for (const CodeGenOptions::BitcodeFileToLink &F : - CI.getCodeGenOpts().LinkBitcodeFiles) { - auto BCBuf = CI.getFileManager().getBufferForFile(F.Filename); - if (!BCBuf) { - CI.getDiagnostics().Report(diag::err_cannot_open_file) - << F.Filename << BCBuf.getError().message(); - LinkModules.clear(); - return nullptr; - } - - Expected> ModuleOrErr = - getOwningLazyBitcodeModule(std::move(*BCBuf), *VMContext); - if (!ModuleOrErr) { - handleAllErrors(ModuleOrErr.takeError(), [&](ErrorInfoBase &EIB) { - CI.getDiagnostics().Report(diag::err_cannot_open_file) - << F.Filename << EIB.message(); - }); - LinkModules.clear(); - return nullptr; - } - LinkModules.push_back({std::move(ModuleOrErr.get()), F.PropagateAttrs, - F.Internalize, F.LinkFlags}); - } - - CoverageSourceInfo *CoverageInfo = nullptr; - // Add the preprocessor callback only when the coverage mapping is generated. - if (CI.getCodeGenOpts().CoverageMapping) - CoverageInfo = CodeGen::CoverageMappingModuleGen::setUpCoverageCallbacks( - CI.getPreprocessor()); - - std::unique_ptr Result(new BackendConsumer( - BA, CI.getDiagnostics(), CI.getHeaderSearchOpts(), - CI.getPreprocessorOpts(), CI.getCodeGenOpts(), CI.getTargetOpts(), - CI.getLangOpts(), std::string(InFile), std::move(LinkModules), - std::move(OS), *VMContext, CoverageInfo)); - BEConsumer = Result.get(); - - // Enable generating macro debug info only when debug info is not disabled and - // also macro debug info is enabled. - if (CI.getCodeGenOpts().getDebugInfo() != codegenoptions::NoDebugInfo && - CI.getCodeGenOpts().MacroDebugInfo) { - std::unique_ptr Callbacks = - std::make_unique(BEConsumer->getCodeGenerator(), - CI.getPreprocessor()); - CI.getPreprocessor().addPPCallbacks(std::move(Callbacks)); - } - - return std::move(Result); -} - -static void BitcodeInlineAsmDiagHandler(const llvm::SMDiagnostic &SM, - void *Context, - unsigned LocCookie) { - SM.print(nullptr, llvm::errs()); - - auto Diags = static_cast(Context); - unsigned DiagID; - switch (SM.getKind()) { - case llvm::SourceMgr::DK_Error: - DiagID = diag::err_fe_inline_asm; - break; - case llvm::SourceMgr::DK_Warning: - DiagID = diag::warn_fe_inline_asm; - break; - case llvm::SourceMgr::DK_Note: - DiagID = diag::note_fe_inline_asm; - break; - case llvm::SourceMgr::DK_Remark: - llvm_unreachable("remarks unexpected"); - } - - Diags->Report(DiagID).AddString("cannot compile inline asm"); -} - -std::unique_ptr -CodeGenAction::loadModule(MemoryBufferRef MBRef) { - CompilerInstance &CI = getCompilerInstance(); - SourceManager &SM = CI.getSourceManager(); - - // For ThinLTO backend invocations, ensure that the context - // merges types based on ODR identifiers. We also need to read - // the correct module out of a multi-module bitcode file. - if (!CI.getCodeGenOpts().ThinLTOIndexFile.empty()) { - VMContext->enableDebugTypeODRUniquing(); - - auto DiagErrors = [&](Error E) -> std::unique_ptr { - unsigned DiagID = - CI.getDiagnostics().getCustomDiagID(DiagnosticsEngine::Error, "%0"); - handleAllErrors(std::move(E), [&](ErrorInfoBase &EIB) { - CI.getDiagnostics().Report(DiagID) << EIB.message(); - }); - return {}; - }; - - Expected> BMsOrErr = getBitcodeModuleList(MBRef); - if (!BMsOrErr) - return DiagErrors(BMsOrErr.takeError()); - BitcodeModule *Bm = llvm::lto::findThinLTOModule(*BMsOrErr); - // We have nothing to do if the file contains no ThinLTO module. This is - // possible if ThinLTO compilation was not able to split module. Content of - // the file was already processed by indexing and will be passed to the - // linker using merged object file. - if (!Bm) { - auto M = std::make_unique("empty", *VMContext); - M->setTargetTriple(CI.getTargetOpts().Triple); - return M; - } - Expected> MOrErr = - Bm->parseModule(*VMContext); - if (!MOrErr) - return DiagErrors(MOrErr.takeError()); - return std::move(*MOrErr); - } - - llvm::SMDiagnostic Err; - if (std::unique_ptr M = parseIR(MBRef, Err, *VMContext)) - return M; - - // Translate from the diagnostic info to the SourceManager location if - // available. - // TODO: Unify this with ConvertBackendLocation() - SourceLocation Loc; - if (Err.getLineNo() > 0) { - assert(Err.getColumnNo() >= 0); - Loc = SM.translateFileLineCol(SM.getFileEntryForID(SM.getMainFileID()), - Err.getLineNo(), Err.getColumnNo() + 1); - } - - // Strip off a leading diagnostic code if there is one. - StringRef Msg = Err.getMessage(); - if (Msg.startswith("error: ")) - Msg = Msg.substr(7); - - unsigned DiagID = - CI.getDiagnostics().getCustomDiagID(DiagnosticsEngine::Error, "%0"); - - CI.getDiagnostics().Report(Loc, DiagID) << Msg; - return {}; -} - -// zahira -namespace { -// Handles the initialization and cleanup of the OptRecordFile. This -// customization allows initialization before the clang codegen runs -// so it can also emit to the opt report. -struct OptRecordFileRAII { - std::unique_ptr OptRecordFile; - std::unique_ptr OldDiagnosticHandler; - llvm::LLVMContext &Ctx; - - OptRecordFileRAII(CodeGenAction &CGA, llvm::LLVMContext &Ctx, - BackendConsumer &BC) - : Ctx(Ctx) { - CompilerInstance &CI = CGA.getCompilerInstance(); - auto &CodeGenOpts = CI.getCodeGenOpts(); - - OldDiagnosticHandler = Ctx.getDiagnosticHandler(); - - Ctx.setDiagnosticHandler( - std::make_unique(CodeGenOpts, &BC)); - - Expected> OptRecordFileOrErr = - setupLLVMOptimizationRemarks( - Ctx, CodeGenOpts.OptRecordFile, CodeGenOpts.OptRecordPasses, - CodeGenOpts.OptRecordFormat, CodeGenOpts.DiagnosticsWithHotness, - CodeGenOpts.DiagnosticsHotnessThreshold); - - if (Error E = OptRecordFileOrErr.takeError()) - reportOptRecordError(std::move(E), CI.getDiagnostics(), CodeGenOpts); - else - OptRecordFile = std::move(*OptRecordFileOrErr); - - if (OptRecordFile && - CodeGenOpts.getProfileUse() != CodeGenOptions::ProfileNone) - Ctx.setDiagnosticsHotnessRequested(true); - } - ~OptRecordFileRAII() { - Ctx.setDiagnosticHandler(std::move(OldDiagnosticHandler)); - if (OptRecordFile) - OptRecordFile->keep(); - } -}; -} // namespace - -void CodeGenAction::ExecuteAction() { - if (getCurrentFileKind().getLanguage() != Language::LLVM_IR) { - OptRecordFileRAII ORF(*this, *VMContext, *BEConsumer); - this->ASTFrontendAction::ExecuteAction(); - return; - } - - // If this is an IR file, we have to treat it specially. - BackendAction BA = static_cast(Act); - CompilerInstance &CI = getCompilerInstance(); - auto &CodeGenOpts = CI.getCodeGenOpts(); - auto &Diagnostics = CI.getDiagnostics(); - std::unique_ptr OS = - GetOutputStream(CI, getCurrentFile(), BA); - if (BA != Backend_EmitNothing && !OS) - return; - - SourceManager &SM = CI.getSourceManager(); - FileID FID = SM.getMainFileID(); - Optional MainFile = SM.getBufferOrNone(FID); - if (!MainFile) - return; - - TheModule = loadModule(*MainFile); - if (!TheModule) - return; - - const TargetOptions &TargetOpts = CI.getTargetOpts(); - if (TheModule->getTargetTriple() != TargetOpts.Triple) { - Diagnostics.Report(SourceLocation(), diag::warn_fe_override_module) - << TargetOpts.Triple; - TheModule->setTargetTriple(TargetOpts.Triple); - } - - // EmbedBitcode(TheModule.get(), CodeGenOpts, *MainFile); - - LLVMContext &Ctx = TheModule->getContext(); - Ctx.setInlineAsmDiagnosticHandler(BitcodeInlineAsmDiagHandler, &Diagnostics); - - // Restore any diagnostic handler previously set before returning from this - // function. - struct RAII { - LLVMContext &Ctx; - std::unique_ptr PrevHandler = Ctx.getDiagnosticHandler(); - ~RAII() { Ctx.setDiagnosticHandler(std::move(PrevHandler)); } - } _{Ctx}; - - // Set clang diagnostic handler. To do this we need to create a fake - // BackendConsumer. - BackendConsumer Result(BA, CI.getDiagnostics(), CI.getHeaderSearchOpts(), - CI.getPreprocessorOpts(), CI.getCodeGenOpts(), - CI.getTargetOpts(), CI.getLangOpts(), - std::move(LinkModules), *VMContext, nullptr); - // PR44896: Force DiscardValueNames as false. DiscardValueNames cannot be - // true here because the valued names are needed for reading textual IR. - Ctx.setDiscardValueNames(false); - OptRecordFileRAII ORF(*this, Ctx, Result); -#if 0 - Ctx.setDiagnosticHandler( - std::make_unique(CodeGenOpts, &Result)); - - Expected> OptRecordFileOrErr = - setupLLVMOptimizationRemarks( - Ctx, CodeGenOpts.OptRecordFile, CodeGenOpts.OptRecordPasses, - CodeGenOpts.OptRecordFormat, CodeGenOpts.DiagnosticsWithHotness, - CodeGenOpts.DiagnosticsHotnessThreshold); - - if (Error E = OptRecordFileOrErr.takeError()) { - reportOptRecordError(std::move(E), Diagnostics, CodeGenOpts); - return; - } - std::unique_ptr OptRecordFile = - std::move(*OptRecordFileOrErr); -#endif - EmitBackendOutput(Diagnostics, CI.getHeaderSearchOpts(), CodeGenOpts, - TargetOpts, CI.getLangOpts(), - CI.getTarget().getDataLayout(), TheModule.get(), BA, - std::move(OS)); -#if 0 - if (OptRecordFile) - OptRecordFile->keep(); -#endif -} - -// - -void EmitAssemblyAction::anchor() { } -EmitAssemblyAction::EmitAssemblyAction(llvm::LLVMContext *_VMContext) - : CodeGenAction(Backend_EmitAssembly, _VMContext) {} - -void EmitBCAction::anchor() { } -EmitBCAction::EmitBCAction(llvm::LLVMContext *_VMContext) - : CodeGenAction(Backend_EmitBC, _VMContext) {} - -void EmitLLVMAction::anchor() { } -EmitLLVMAction::EmitLLVMAction(llvm::LLVMContext *_VMContext) - : CodeGenAction(Backend_EmitLL, _VMContext) {} - -void EmitLLVMOnlyAction::anchor() { } -EmitLLVMOnlyAction::EmitLLVMOnlyAction(llvm::LLVMContext *_VMContext) - : CodeGenAction(Backend_EmitNothing, _VMContext) {} - -void EmitCodeGenOnlyAction::anchor() { } -EmitCodeGenOnlyAction::EmitCodeGenOnlyAction(llvm::LLVMContext *_VMContext) - : CodeGenAction(Backend_EmitMCNull, _VMContext) {} - -void EmitObjAction::anchor() { } -EmitObjAction::EmitObjAction(llvm::LLVMContext *_VMContext) - : CodeGenAction(Backend_EmitObj, _VMContext) {} diff --git a/clang/lib/CodeGen/CodeGenFunction.cpp b/clang/lib/CodeGen/CodeGenFunction.cpp index 6ccd5e771a5cc..b983cb010d8fa 100644 --- a/clang/lib/CodeGen/CodeGenFunction.cpp +++ b/clang/lib/CodeGen/CodeGenFunction.cpp @@ -757,6 +757,30 @@ void CodeGenFunction::EmitOpenCLKernelMetadata(const FunctionDecl *FD, Fn->setMetadata("initiation_interval", llvm::MDNode::get(Context, AttrMDArgs)); } + // zahira + llvm::DiagnosticLocation DL = SourceLocToDebugLoc(FD->getBody()->getBeginLoc()); +#if 0 + llvm::DiagnosticLocation DL = SourceLocToDebugLoc(Fn->getLoca + SyclOptReportHandler &OptReportHandler = CGM.getDiags().OptReportHandler; + if (OptReportHandler.HasOptReportInfo(FD)) { + llvm::OptimizationRemarkEmitter ORE(Fn); + for (auto &ORI : OptReportHandler.getInfo(FD)) { + llvm::DiagnosticLocation DL = SourceLocToDebugLoc(ORI.KernelArgLoc); + llvm::OptimizationRemark Remark("sycl", "Region", DL, + &Fn->getEntryBlock()); +#if 1 + printf("Name of FD: %s \n", FD->getName().data()); + printf("ArgName: %s \n", ORI.KernelArgName.data()); + printf("ArgType: %s \n", ORI.KernelArgType.data()); + printf("ArgLoc: %d \n", ORI.KernelArgLoc.getRawEncoding()); +#endif + Remark << llvm::ore::NV("KernelArg", ORI.KernelArgName); + Remark << llvm::ore::NV("ArgType", ORI.KernelArgType); + ORE.emit(Remark); + } + } +#endif + // end zahira } /// Determine whether the function F ends with a return stmt. @@ -1137,6 +1161,7 @@ void CodeGenFunction::StartFunction(GlobalDecl GD, QualType RetTy, ArgTypes.push_back(VD->getType()); QualType FnType = getContext().getFunctionType( RetTy, ArgTypes, FunctionProtoType::ExtProtoInfo(CC)); + DI->emitFunctionStart(GD, Loc, StartLoc, FnType, CurFn, CurFuncIsThunk); } @@ -1484,29 +1509,29 @@ void CodeGenFunction::GenerateCode(GlobalDecl GD, llvm::Function *Fn, // Emit the standard function prologue. StartFunction(GD, ResTy, Fn, FnInfo, Args, Loc, BodyRange.getBegin()); + // zahira +#if 1 SyclOptReportHandler &OptReportHandler = CGM.getDiags().OptReportHandler; - // zahira if (OptReportHandler.HasOptReportInfo(FD)) { llvm::OptimizationRemarkEmitter ORE(Fn); - //llvm::DiagnosticLocation DL = SourceLocToDebugLoc(Body->getBeginLoc()); for (auto &ORI : OptReportHandler.getInfo(FD)) { llvm::DiagnosticLocation DL = SourceLocToDebugLoc(ORI.KernelArgLoc); - llvm::OptimizationRemarkMissed Remark("SYCL", "ZAHIRA_Region", DL, - &Fn->getEntryBlock()); + llvm::OptimizationRemark Remark("sycl", "Region", DL, + &Fn->getEntryBlock()); #if 1 printf("Name of FD: %s \n", FD->getName().data()); printf("ArgName: %s \n", ORI.KernelArgName.data()); printf("ArgType: %s \n", ORI.KernelArgType.data()); printf("ArgLoc: %d \n", ORI.KernelArgLoc.getRawEncoding()); #endif - //llvm::OptimizationRemark Remark("sycl", "Kernelrgs", Fn); Remark << llvm::ore::NV("KernelArg", ORI.KernelArgName); - Remark << llvm::ore::NV("ArgType", ORI.KernelArgType); + Remark << llvm::ore::NV("ArgType", ORI.KernelArgType); + Remark << llvm::ore::NV("ArgLoc", ORI.KernelArgLoc.getRawEncoding()); ORE.emit(Remark); } } - +#endif // Generate the body of the function. PGO.assignRegionCounters(GD, CurFn); if (isa(FD)) From e69616ed803e09665ed32ec7d4c18bc97fbd16eb Mon Sep 17 00:00:00 2001 From: Zahira Ammarguellat Date: Mon, 5 Apr 2021 08:50:01 -0700 Subject: [PATCH 03/19] Improve the output. --- clang/lib/CodeGen/CodeGenFunction.cpp | 29 ++++++++++++++++++++------- clang/lib/Sema/SemaSYCL.cpp | 8 ++++---- 2 files changed, 26 insertions(+), 11 deletions(-) diff --git a/clang/lib/CodeGen/CodeGenFunction.cpp b/clang/lib/CodeGen/CodeGenFunction.cpp index b983cb010d8fa..3d944bb34fce6 100644 --- a/clang/lib/CodeGen/CodeGenFunction.cpp +++ b/clang/lib/CodeGen/CodeGenFunction.cpp @@ -1515,19 +1515,34 @@ void CodeGenFunction::GenerateCode(GlobalDecl GD, llvm::Function *Fn, if (OptReportHandler.HasOptReportInfo(FD)) { llvm::OptimizationRemarkEmitter ORE(Fn); + int count = 0; + for (auto &ORI : OptReportHandler.getInfo(FD)) { - llvm::DiagnosticLocation DL = SourceLocToDebugLoc(ORI.KernelArgLoc); - llvm::OptimizationRemark Remark("sycl", "Region", DL, - &Fn->getEntryBlock()); -#if 1 +#if 0 printf("Name of FD: %s \n", FD->getName().data()); + printf("Name of Fn: %s\n", Fn->getFunction().getName().data()); printf("ArgName: %s \n", ORI.KernelArgName.data()); printf("ArgType: %s \n", ORI.KernelArgType.data()); printf("ArgLoc: %d \n", ORI.KernelArgLoc.getRawEncoding()); #endif - Remark << llvm::ore::NV("KernelArg", ORI.KernelArgName); - Remark << llvm::ore::NV("ArgType", ORI.KernelArgType); - Remark << llvm::ore::NV("ArgLoc", ORI.KernelArgLoc.getRawEncoding()); + std::string prefix; + if (ORI.KernelArgName.empty()) + prefix = "&"; + llvm::DiagnosticLocation DL = SourceLocToDebugLoc(ORI.KernelArgLoc); + llvm::OptimizationRemark Remark("sycl", "Region", DL, + &Fn->getEntryBlock()); + Remark << "Argument " << llvm::ore::NV("Argument", count++) + << " for function kernel: " + << llvm::ore::NV(ORI.KernelArgName.empty() + ? "&" : "") << " " + << Fn->getName() << "." + << llvm::ore::NV(ORI.KernelArgName.empty() ? " " : ORI.KernelArgName) + << "(" << ORI.KernelArgType << ")"; +#if 0 + Remark << llvm::ore::NV("KernelArg ", ORI.KernelArgName) + << llvm::ore::NV("ArgType ", ORI.KernelArgType) + << llvm::ore::NV("ArgLoc ", ORI.KernelArgLoc.getRawEncoding()); + #endif ORE.emit(Remark); } } diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index 82296bc58e2da..0ce4802f18e38 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -1684,10 +1684,10 @@ class SyclKernelDeclCreator : public SyclKernelFieldHandler { void addParam(const FieldDecl *FD, QualType FieldTy) { ParamDesc newParamDesc = makeParamDesc(FD, FieldTy); - // zahira #if 0 // zahira StringRef FuncName = KernelDecl->getName(); + StringRef ArgName = FD->getName(); std::string ArgType = FieldTy.getAsString(); SourceLocation ArgLoc = FD->getLocation(); @@ -1704,7 +1704,7 @@ class SyclKernelDeclCreator : public SyclKernelFieldHandler { #endif addParam(newParamDesc, FieldTy); } - + void addParam(const CXXBaseSpecifier &BS, QualType FieldTy) { // TODO: There is no name for the base available, but duplicate names are // seemingly already possible, so we'll give them all the same name for now. @@ -1715,7 +1715,7 @@ class SyclKernelDeclCreator : public SyclKernelFieldHandler { // zahira #if 0 StringRef FuncName = KernelDecl->getName(); - StringRef ArgName = "_arg__base"; + StringRef ArgName = " "; std::string ArgType = FieldTy.getAsString(); SourceLocation KernelArgLoc = BS.getBaseTypeLoc(); printf("AddParam BS\n"); @@ -1725,7 +1725,7 @@ class SyclKernelDeclCreator : public SyclKernelFieldHandler { #endif #if 1 SemaRef.getDiagnostics().OptReportHandler.AddKernelArgs( - KernelDecl, "_arg__base", FieldTy.getAsString(), + KernelDecl, "", FieldTy.getAsString(), BS.getBaseTypeLoc()); #endif addParam(newParamDesc, FieldTy); From c44d670771929d7afd35ff80b008d7e215583052 Mon Sep 17 00:00:00 2001 From: Zahira Ammarguellat Date: Mon, 5 Apr 2021 12:16:33 -0700 Subject: [PATCH 04/19] cleanup Signed-off-by: Zahira Ammarguellat --- clang/lib/CodeGen/CodeGenAction.cpp | 27 +--------------- clang/lib/CodeGen/CodeGenFunction.cpp | 46 ++------------------------- clang/lib/Sema/SemaSYCL.cpp | 29 ----------------- 3 files changed, 3 insertions(+), 99 deletions(-) diff --git a/clang/lib/CodeGen/CodeGenAction.cpp b/clang/lib/CodeGen/CodeGenAction.cpp index 6e644f140f322..f02344ffe4086 100644 --- a/clang/lib/CodeGen/CodeGenAction.cpp +++ b/clang/lib/CodeGen/CodeGenAction.cpp @@ -304,34 +304,13 @@ namespace clang { return; LLVMContext &Ctx = getModule()->getContext(); - // zahira - //LLVMContext::InlineAsmDiagHandlerTy OldHandler = - // Ctx.getInlineAsmDiagnosticHandler(); - // end zahira std::unique_ptr OldDiagnosticHandler = Ctx.getDiagnosticHandler(); Ctx.setDiagnosticHandler(std::make_unique( CodeGenOpts, this)); -#if 0 - Expected> OptRecordFileOrErr = - setupLLVMOptimizationRemarks( - Ctx, CodeGenOpts.OptRecordFile, CodeGenOpts.OptRecordPasses, - CodeGenOpts.OptRecordFormat, CodeGenOpts.DiagnosticsWithHotness, - CodeGenOpts.DiagnosticsHotnessThreshold); - - if (Error E = OptRecordFileOrErr.takeError()) { - reportOptRecordError(std::move(E), Diags, CodeGenOpts); - return; - } - - std::unique_ptr OptRecordFile = - std::move(*OptRecordFileOrErr); + // The diagnostic handler is now processed in OptRecordFileRAII. - if (OptRecordFile && - CodeGenOpts.getProfileUse() != CodeGenOptions::ProfileNone) - Ctx.setDiagnosticsHotnessRequested(true); -#endif // The parallel_for_work_group legalization pass can emit calls to // builtins function. Definitions of those builtins can be provided in // LinkModule. We force the pass to legalize the code before the link @@ -354,10 +333,6 @@ namespace clang { getModule(), Action, std::move(AsmOutStream)); Ctx.setDiagnosticHandler(std::move(OldDiagnosticHandler)); -#if 0 - if (OptRecordFile) - OptRecordFile->keep(); -#endif } void HandleTagDeclDefinition(TagDecl *D) override { diff --git a/clang/lib/CodeGen/CodeGenFunction.cpp b/clang/lib/CodeGen/CodeGenFunction.cpp index bc98e7bf9d3a0..b9a672860428d 100644 --- a/clang/lib/CodeGen/CodeGenFunction.cpp +++ b/clang/lib/CodeGen/CodeGenFunction.cpp @@ -765,30 +765,6 @@ void CodeGenFunction::EmitOpenCLKernelMetadata(const FunctionDecl *FD, Fn->setMetadata("initiation_interval", llvm::MDNode::get(Context, AttrMDArgs)); } - // zahira - llvm::DiagnosticLocation DL = SourceLocToDebugLoc(FD->getBody()->getBeginLoc()); -#if 0 - llvm::DiagnosticLocation DL = SourceLocToDebugLoc(Fn->getLoca - SyclOptReportHandler &OptReportHandler = CGM.getDiags().OptReportHandler; - if (OptReportHandler.HasOptReportInfo(FD)) { - llvm::OptimizationRemarkEmitter ORE(Fn); - for (auto &ORI : OptReportHandler.getInfo(FD)) { - llvm::DiagnosticLocation DL = SourceLocToDebugLoc(ORI.KernelArgLoc); - llvm::OptimizationRemark Remark("sycl", "Region", DL, - &Fn->getEntryBlock()); -#if 1 - printf("Name of FD: %s \n", FD->getName().data()); - printf("ArgName: %s \n", ORI.KernelArgName.data()); - printf("ArgType: %s \n", ORI.KernelArgType.data()); - printf("ArgLoc: %d \n", ORI.KernelArgLoc.getRawEncoding()); -#endif - Remark << llvm::ore::NV("KernelArg", ORI.KernelArgName); - Remark << llvm::ore::NV("ArgType", ORI.KernelArgType); - ORE.emit(Remark); - } - } -#endif - // end zahira } /// Determine whether the function F ends with a return stmt. @@ -1517,25 +1493,12 @@ void CodeGenFunction::GenerateCode(GlobalDecl GD, llvm::Function *Fn, // Emit the standard function prologue. StartFunction(GD, ResTy, Fn, FnInfo, Args, Loc, BodyRange.getBegin()); - // zahira -#if 1 - SyclOptReportHandler &OptReportHandler = CGM.getDiags().OptReportHandler; + SyclOptReportHandler &OptReportHandler = CGM.getDiags().OptReportHandler; if (OptReportHandler.HasOptReportInfo(FD)) { llvm::OptimizationRemarkEmitter ORE(Fn); int count = 0; - for (auto &ORI : OptReportHandler.getInfo(FD)) { -#if 0 - printf("Name of FD: %s \n", FD->getName().data()); - printf("Name of Fn: %s\n", Fn->getFunction().getName().data()); - printf("ArgName: %s \n", ORI.KernelArgName.data()); - printf("ArgType: %s \n", ORI.KernelArgType.data()); - printf("ArgLoc: %d \n", ORI.KernelArgLoc.getRawEncoding()); -#endif - std::string prefix; - if (ORI.KernelArgName.empty()) - prefix = "&"; llvm::DiagnosticLocation DL = SourceLocToDebugLoc(ORI.KernelArgLoc); llvm::OptimizationRemark Remark("sycl", "Region", DL, &Fn->getEntryBlock()); @@ -1546,15 +1509,10 @@ void CodeGenFunction::GenerateCode(GlobalDecl GD, llvm::Function *Fn, << Fn->getName() << "." << llvm::ore::NV(ORI.KernelArgName.empty() ? " " : ORI.KernelArgName) << "(" << ORI.KernelArgType << ")"; -#if 0 - Remark << llvm::ore::NV("KernelArg ", ORI.KernelArgName) - << llvm::ore::NV("ArgType ", ORI.KernelArgType) - << llvm::ore::NV("ArgLoc ", ORI.KernelArgLoc.getRawEncoding()); - #endif ORE.emit(Remark); } } -#endif + // Generate the body of the function. PGO.assignRegionCounters(GD, CurFn); if (isa(FD)) diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index d1d58407a8ca7..4632aff6125ac 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -1690,25 +1690,9 @@ class SyclKernelDeclCreator : public SyclKernelFieldHandler { void addParam(const FieldDecl *FD, QualType FieldTy) { ParamDesc newParamDesc = makeParamDesc(FD, FieldTy); - -#if 0 - // zahira - StringRef FuncName = KernelDecl->getName(); - - StringRef ArgName = FD->getName(); - std::string ArgType = FieldTy.getAsString(); - SourceLocation ArgLoc = FD->getLocation(); - - printf("AddParam FD\n"); - printf("FuncName: %s\n", FuncName.data()); - printf("ArgName: %s \n", ArgName.data()); - printf("ArgType: %s \n", ArgType.data()); -#endif -#if 1 SemaRef.getDiagnostics().OptReportHandler.AddKernelArgs( KernelDecl, FD->getName(), FieldTy.getAsString(), FD->getLocation()); -#endif addParam(newParamDesc, FieldTy); } @@ -1719,22 +1703,9 @@ class SyclKernelDeclCreator : public SyclKernelFieldHandler { StringRef Name = "_arg__base"; ParamDesc newParamDesc = makeParamDesc(SemaRef.getASTContext(), Name, FieldTy); - // zahira -#if 0 - StringRef FuncName = KernelDecl->getName(); - StringRef ArgName = " "; - std::string ArgType = FieldTy.getAsString(); - SourceLocation KernelArgLoc = BS.getBaseTypeLoc(); - printf("AddParam BS\n"); - printf("FuncName: %s\n", FuncName.data()); - printf("ArgName: %s \n", Name.data()); - printf("ArgType: %s \n", ArgType.data()); -#endif -#if 1 SemaRef.getDiagnostics().OptReportHandler.AddKernelArgs( KernelDecl, "", FieldTy.getAsString(), BS.getBaseTypeLoc()); -#endif addParam(newParamDesc, FieldTy); } // Add a parameter with specified name and type From 87ce263b917ed05b5e391f4c85631278b129a221 Mon Sep 17 00:00:00 2001 From: Zahira Ammarguellat Date: Mon, 5 Apr 2021 13:01:53 -0700 Subject: [PATCH 05/19] Put back the indentation as it was Signed-off-by: Zahira Ammarguellat --- clang/lib/CodeGen/CodeGenFunction.cpp | 203 +++++++++++++------------- clang/lib/Sema/SemaSYCL.cpp | 6 +- 2 files changed, 102 insertions(+), 107 deletions(-) diff --git a/clang/lib/CodeGen/CodeGenFunction.cpp b/clang/lib/CodeGen/CodeGenFunction.cpp index b9a672860428d..a557f2fdb695b 100644 --- a/clang/lib/CodeGen/CodeGenFunction.cpp +++ b/clang/lib/CodeGen/CodeGenFunction.cpp @@ -103,20 +103,17 @@ llvm::fp::ExceptionBehavior clang::ToConstrainedExceptMD(LangOptions::FPExceptionModeKind Kind) { switch (Kind) { - case LangOptions::FPE_Ignore: - return llvm::fp::ebIgnore; - case LangOptions::FPE_MayTrap: - return llvm::fp::ebMayTrap; - case LangOptions::FPE_Strict: - return llvm::fp::ebStrict; + case LangOptions::FPE_Ignore: return llvm::fp::ebIgnore; + case LangOptions::FPE_MayTrap: return llvm::fp::ebMayTrap; + case LangOptions::FPE_Strict: return llvm::fp::ebStrict; } llvm_unreachable("Unsupported FP Exception Behavior"); } void CodeGenFunction::SetFPModel() { llvm::RoundingMode RM = getLangOpts().getFPRoundingMode(); - auto fpExceptionBehavior = - ToConstrainedExceptMD(getLangOpts().getFPExceptionMode()); + auto fpExceptionBehavior = ToConstrainedExceptMD(getLangOpts(). + getFPExceptionMode()); Builder.setDefaultConstrainedRounding(RM); Builder.setDefaultConstrainedExcept(fpExceptionBehavior); @@ -209,8 +206,7 @@ LValue CodeGenFunction::MakeNaturalAlignAddrLValue(llvm::Value *V, QualType T) { /// Given a value of type T* that may not be to a complete object, /// construct an l-value with the natural pointee alignment of T. -LValue CodeGenFunction::MakeNaturalAlignPointeeAddrLValue(llvm::Value *V, - QualType T) { +LValue CodeGenFunction::MakeNaturalAlignPointeeAddrLValue(llvm::Value *V, QualType T) { LValueBaseInfo BaseInfo; TBAAAccessInfo TBAAInfo; CharUnits Align = CGM.getNaturalTypeAlignment(T, &BaseInfo, &TBAAInfo, @@ -218,6 +214,7 @@ LValue CodeGenFunction::MakeNaturalAlignPointeeAddrLValue(llvm::Value *V, return MakeAddrLValue(Address(V, Align), T, BaseInfo, TBAAInfo); } + llvm::Type *CodeGenFunction::ConvertTypeForMem(QualType T) { return CGM.getTypes().ConvertTypeForMem(T); } @@ -306,7 +303,7 @@ llvm::DebugLoc CodeGenFunction::EmitReturnBlock() { // cleans up functions which started with a unified return block. if (ReturnBlock.getBlock()->hasOneUse()) { llvm::BranchInst *BI = - dyn_cast(*ReturnBlock.getBlock()->user_begin()); + dyn_cast(*ReturnBlock.getBlock()->user_begin()); if (BI && BI->isUnconditional() && BI->getSuccessor(0) == ReturnBlock.getBlock()) { // Record/return the DebugLoc of the simple 'return' expression to be used @@ -329,8 +326,7 @@ llvm::DebugLoc CodeGenFunction::EmitReturnBlock() { } static void EmitIfUsed(CodeGenFunction &CGF, llvm::BasicBlock *BB) { - if (!BB) - return; + if (!BB) return; if (!BB->use_empty()) return CGF.CurFn->getBasicBlockList().push_back(BB); delete BB; @@ -340,9 +336,9 @@ void CodeGenFunction::FinishFunction(SourceLocation EndLoc) { assert(BreakContinueStack.empty() && "mismatched push/pop in break/continue stack!"); - bool OnlySimpleReturnStmts = NumSimpleReturnExprs > 0 && - NumSimpleReturnExprs == NumReturnExprs && - ReturnBlock.getBlock()->use_empty(); + bool OnlySimpleReturnStmts = NumSimpleReturnExprs > 0 + && NumSimpleReturnExprs == NumReturnExprs + && ReturnBlock.getBlock()->use_empty(); // Usually the return expression is evaluated before the cleanup // code. If the function contains only a simple return statement, // such as a constant, the location before the cleanup code becomes @@ -407,7 +403,8 @@ void CodeGenFunction::FinishFunction(SourceLocation EndLoc) { EmitFunctionEpilog(*CurFnInfo, EmitRetDbgLoc, EndLoc); EmitEndEHSpec(CurCodeDecl); - assert(EHStack.empty() && "did not remove all scopes from cleanup stack!"); + assert(EHStack.empty() && + "did not remove all scopes from cleanup stack!"); // If someone did an indirect goto, emit the indirect goto block at the end of // the function. @@ -593,7 +590,8 @@ CodeGenFunction::DecodeAddrUsedInPrologue(llvm::Value *F, } void CodeGenFunction::EmitOpenCLKernelMetadata(const FunctionDecl *FD, - llvm::Function *Fn) { + llvm::Function *Fn) +{ if (!FD->hasAttr()) return; @@ -626,8 +624,7 @@ void CodeGenFunction::EmitOpenCLKernelMetadata(const FunctionDecl *FD, llvm::ConstantAsMetadata::get(Builder.getInt32(A->getXDim())), llvm::ConstantAsMetadata::get(Builder.getInt32(A->getYDim())), llvm::ConstantAsMetadata::get(Builder.getInt32(A->getZDim()))}; - Fn->setMetadata("work_group_size_hint", - llvm::MDNode::get(Context, AttrMDArgs)); + Fn->setMetadata("work_group_size_hint", llvm::MDNode::get(Context, AttrMDArgs)); } if (const ReqdWorkGroupSizeAttr *A = FD->getAttr()) { @@ -671,7 +668,7 @@ void CodeGenFunction::EmitOpenCLKernelMetadata(const FunctionDecl *FD, } if (const SYCLIntelNumSimdWorkItemsAttr *A = - FD->getAttr()) { + FD->getAttr()) { const auto *CE = dyn_cast(A->getValue()); assert(CE && "Not an integer constant expression"); Optional ArgVal = CE->getResultAsAPSInt(); @@ -693,7 +690,7 @@ void CodeGenFunction::EmitOpenCLKernelMetadata(const FunctionDecl *FD, } if (const SYCLIntelMaxGlobalWorkDimAttr *A = - FD->getAttr()) { + FD->getAttr()) { const auto *CE = dyn_cast(A->getValue()); assert(CE && "Not an integer constant expression"); Optional ArgVal = CE->getResultAsAPSInt(); @@ -768,7 +765,7 @@ void CodeGenFunction::EmitOpenCLKernelMetadata(const FunctionDecl *FD, } /// Determine whether the function F ends with a return stmt. -static bool endsWithReturn(const Decl *F) { +static bool endsWithReturn(const Decl* F) { const Stmt *Body = nullptr; if (auto *FD = dyn_cast_or_null(F)) Body = FD->getBody(); @@ -881,8 +878,7 @@ void CodeGenFunction::StartFunction(GlobalDecl GD, QualType RetTy, // Apply sanitizer attributes to the function. if (SanOpts.hasOneOf(SanitizerKind::Address | SanitizerKind::KernelAddress)) Fn->addFnAttr(llvm::Attribute::SanitizeAddress); - if (SanOpts.hasOneOf(SanitizerKind::HWAddress | - SanitizerKind::KernelHWAddress)) + if (SanOpts.hasOneOf(SanitizerKind::HWAddress | SanitizerKind::KernelHWAddress)) Fn->addFnAttr(llvm::Attribute::SanitizeHWAddress); if (SanOpts.has(SanitizerKind::MemTag)) Fn->addFnAttr(llvm::Attribute::SanitizeMemTag); @@ -906,8 +902,7 @@ void CodeGenFunction::StartFunction(GlobalDecl GD, QualType RetTy, IdentifierInfo *II = OMD->getSelector().getIdentifierInfoForSlot(0); if (OMD->getMethodFamily() == OMF_dealloc || OMD->getMethodFamily() == OMF_initialize || - (OMD->getSelector().isUnarySelector() && - II->isStr(".cxx_destruct"))) { + (OMD->getSelector().isUnarySelector() && II->isStr(".cxx_destruct"))) { markAsIgnoreThreadCheckingAtRuntime(Fn); } } @@ -1049,8 +1044,9 @@ void CodeGenFunction::StartFunction(GlobalDecl GD, QualType RetTy, if (llvm::Constant *PrologueSig = getPrologueSignature(CGM, FD)) { // Remove any (C++17) exception specifications, to allow calling e.g. a // noexcept function through a non-noexcept pointer. - auto ProtoTy = getContext().getFunctionTypeWithExceptionSpec( - FD->getType(), EST_None); + auto ProtoTy = + getContext().getFunctionTypeWithExceptionSpec(FD->getType(), + EST_None); llvm::Constant *FTRTTIConst = CGM.GetAddrOfRTTIDescriptor(ProtoTy, /*ForEH=*/true); llvm::Constant *FTRTTIConstEncoded = @@ -1145,7 +1141,6 @@ void CodeGenFunction::StartFunction(GlobalDecl GD, QualType RetTy, ArgTypes.push_back(VD->getType()); QualType FnType = getContext().getFunctionType( RetTy, ArgTypes, FunctionProtoType::ExtProtoInfo(CC)); - DI->emitFunctionStart(GD, Loc, StartLoc, FnType, CurFn, CurFuncIsThunk); } @@ -1177,16 +1172,14 @@ void CodeGenFunction::StartFunction(GlobalDecl GD, QualType RetTy, if (CGM.getCodeGenOpts().MNopMCount) { if (!CGM.getCodeGenOpts().CallFEntry) CGM.getDiags().Report(diag::err_opt_not_valid_without_opt) - << "-mnop-mcount" - << "-mfentry"; + << "-mnop-mcount" << "-mfentry"; Fn->addFnAttr("mnop-mcount"); } if (CGM.getCodeGenOpts().RecordMCount) { if (!CGM.getCodeGenOpts().CallFEntry) CGM.getDiags().Report(diag::err_opt_not_valid_without_opt) - << "-mrecord-mcount" - << "-mfentry"; + << "-mrecord-mcount" << "-mfentry"; Fn->addFnAttr("mrecord-mcount"); } } @@ -1196,7 +1189,7 @@ void CodeGenFunction::StartFunction(GlobalDecl GD, QualType RetTy, if (getContext().getTargetInfo().getTriple().getArch() != llvm::Triple::systemz) CGM.getDiags().Report(diag::err_opt_not_valid_on_target) - << "-mpacked-stack"; + << "-mpacked-stack"; Fn->addFnAttr("packed-stack"); } @@ -1239,7 +1232,8 @@ void CodeGenFunction::StartFunction(GlobalDecl GD, QualType RetTy, // Tell the epilog emitter to autorelease the result. We do this // now so that various specialized functions can suppress it // during their IR-generation. - if (getLangOpts().ObjCAutoRefCount && !CurFnInfo->isReturnsRetained() && + if (getLangOpts().ObjCAutoRefCount && + !CurFnInfo->isReturnsRetained() && RetTy->isObjCRetainableType()) AutoreleaseResult = true; } @@ -1257,7 +1251,8 @@ void CodeGenFunction::StartFunction(GlobalDecl GD, QualType RetTy, if (D && isa(D) && cast(D)->isInstance()) { CGM.getCXXABI().EmitInstanceFunctionProlog(*this); const CXXMethodDecl *MD = cast(D); - if (MD->getParent()->isLambda() && MD->getOverloadedOperator() == OO_Call) { + if (MD->getParent()->isLambda() && + MD->getOverloadedOperator() == OO_Call) { // We're in a lambda; figure out the captures. MD->getParent()->getCaptureFields(LambdaCaptureFields, LambdaThisCaptureField); @@ -1268,24 +1263,21 @@ void CodeGenFunction::StartFunction(GlobalDecl GD, QualType RetTy, // Get the lvalue for the field (which is a copy of the enclosing object // or contains the address of the enclosing object). - LValue ThisFieldLValue = - EmitLValueForLambdaField(LambdaThisCaptureField); + LValue ThisFieldLValue = EmitLValueForLambdaField(LambdaThisCaptureField); if (!LambdaThisCaptureField->getType()->isPointerType()) { - // If the enclosing object was captured by value, just use its - // address. + // If the enclosing object was captured by value, just use its address. CXXThisValue = ThisFieldLValue.getAddress(*this).getPointer(); } else { // Load the lvalue pointed to by the field, since '*this' was captured // by reference. - CXXThisValue = EmitLoadOfLValue(ThisFieldLValue, SourceLocation()) - .getScalarVal(); + CXXThisValue = + EmitLoadOfLValue(ThisFieldLValue, SourceLocation()).getScalarVal(); } } for (auto *FD : MD->getParent()->fields()) { if (FD->hasCapturedVLAType()) { - auto *ExprArg = - EmitLoadOfLValue(EmitLValueForLambdaField(FD), SourceLocation()) - .getScalarVal(); + auto *ExprArg = EmitLoadOfLValue(EmitLValueForLambdaField(FD), + SourceLocation()).getScalarVal(); auto VAT = FD->getCapturedVLAType(); VLASizeMap[VAT->getSizeExpr()] = ExprArg; } @@ -1318,8 +1310,8 @@ void CodeGenFunction::StartFunction(GlobalDecl GD, QualType RetTy, // If any of the arguments have a variably modified type, make sure to // emit the type size. - for (FunctionArgList::const_iterator i = Args.begin(), e = Args.end(); i != e; - ++i) { + for (FunctionArgList::const_iterator i = Args.begin(), e = Args.end(); + i != e; ++i) { const VarDecl *VD = *i; // Dig out the type as written from ParmVarDecls; it's unclear whether @@ -1389,8 +1381,7 @@ void CodeGenFunction::EmitBlockWithFallThrough(llvm::BasicBlock *BB, static void TryMarkNoThrow(llvm::Function *F) { // LLVM treats 'nounwind' on a function as part of the type, so we // can't do this on functions that can be overwritten. - if (F->isInterposable()) - return; + if (F->isInterposable()) return; for (llvm::BasicBlock &BB : *F) for (llvm::Instruction &I : BB) @@ -1519,7 +1510,8 @@ void CodeGenFunction::GenerateCode(GlobalDecl GD, llvm::Function *Fn, EmitDestructorBody(Args); else if (isa(FD)) EmitConstructorBody(Args); - else if (getLangOpts().CUDA && !getLangOpts().CUDAIsDevice && + else if (getLangOpts().CUDA && + !getLangOpts().CUDAIsDevice && FD->hasAttr()) CGM.getCUDARuntime().emitDeviceStub(*this, Args); else if (isa(FD) && @@ -1579,8 +1571,7 @@ void CodeGenFunction::GenerateCode(GlobalDecl GD, llvm::Function *Fn, /// that we can just remove the code. bool CodeGenFunction::ContainsLabel(const Stmt *S, bool IgnoreCaseStmts) { // Null statement, not a label! - if (!S) - return false; + if (!S) return false; // If this is a label, we have to emit the code, consider something like: // if (0) { ... foo: bar(); } goto foo; @@ -1612,8 +1603,7 @@ bool CodeGenFunction::ContainsLabel(const Stmt *S, bool IgnoreCaseStmts) { /// inside of it, this is fine. bool CodeGenFunction::containsBreak(const Stmt *S) { // Null statement, not a label! - if (!S) - return false; + if (!S) return false; // If this is a switch or loop that defines its own break scope, then we can // include it and anything inside of it. @@ -1633,8 +1623,7 @@ bool CodeGenFunction::containsBreak(const Stmt *S) { } bool CodeGenFunction::mightAddDeclToScope(const Stmt *S) { - if (!S) - return false; + if (!S) return false; // Some statement kinds add a scope and thus never add a decl to the current // scope. Note, this list is longer than the list of statements that might @@ -1680,11 +1669,11 @@ bool CodeGenFunction::ConstantFoldsToSimpleInteger(const Expr *Cond, // to bool. Expr::EvalResult Result; if (!Cond->EvaluateAsInt(Result, getContext())) - return false; // Not foldable, not integer or not fully evaluatable. + return false; // Not foldable, not integer or not fully evaluatable. llvm::APSInt Int = Result.Val.getInt(); if (!AllowLabels && CodeGenFunction::ContainsLabel(Cond)) - return false; // Contains a label. + return false; // Contains a label. ResultInt = Int; return true; @@ -1956,7 +1945,7 @@ void CodeGenFunction::EmitBranchOnBoolExpr(const Expr *Cond, // br(c ? throw x : y, t, f) -> br(c, br(throw x, t, f), br(y, t, f) // Fold this to: // br(c, throw x, br(y, t, f)) - EmitCXXThrowExpr(Throw, /*KeepInsertionPoint*/ false); + EmitCXXThrowExpr(Throw, /*KeepInsertionPoint*/false); return; } @@ -2015,8 +2004,8 @@ static void emitNonZeroVLAInit(CodeGenFunction &CGF, QualType baseType, CGBuilderTy &Builder = CGF.Builder; CharUnits baseSize = CGF.getContext().getTypeSizeInChars(baseType); - llvm::Value *baseSizeInChars = - llvm::ConstantInt::get(CGF.IntPtrTy, baseSize.getQuantity()); + llvm::Value *baseSizeInChars + = llvm::ConstantInt::get(CGF.IntPtrTy, baseSize.getQuantity()); Address begin = Builder.CreateElementBitCast(dest, CGF.Int8Ty, "vla.begin"); @@ -2034,7 +2023,8 @@ static void emitNonZeroVLAInit(CodeGenFunction &CGF, QualType baseType, llvm::PHINode *cur = Builder.CreatePHI(begin.getType(), 2, "vla.cur"); cur->addIncoming(begin.getPointer(), originBB); - CharUnits curAlign = dest.getAlignment().alignmentOfArrayElement(baseSize); + CharUnits curAlign = + dest.getAlignment().alignmentOfArrayElement(baseSize); // memcpy the individual element bit-pattern. Builder.CreateMemCpy(Address(cur, curAlign), src, baseSizeInChars, @@ -2042,7 +2032,7 @@ static void emitNonZeroVLAInit(CodeGenFunction &CGF, QualType baseType, // Go to the next element. llvm::Value *next = - Builder.CreateInBoundsGEP(CGF.Int8Ty, cur, baseSizeInChars, "vla.next"); + Builder.CreateInBoundsGEP(CGF.Int8Ty, cur, baseSizeInChars, "vla.next"); // Leave if that's the end of the VLA. llvm::Value *done = Builder.CreateICmpEQ(next, end, "vla-init.isdone"); @@ -2052,7 +2042,8 @@ static void emitNonZeroVLAInit(CodeGenFunction &CGF, QualType baseType, CGF.EmitBlock(contBB); } -void CodeGenFunction::EmitNullInitialization(Address DestPtr, QualType Ty) { +void +CodeGenFunction::EmitNullInitialization(Address DestPtr, QualType Ty) { // Ignore empty classes in C++. if (getLangOpts().CPlusPlus) { if (const RecordType *RT = Ty->getAs()) { @@ -2068,14 +2059,15 @@ void CodeGenFunction::EmitNullInitialization(Address DestPtr, QualType Ty) { // Get size and alignment info for this aggregate. CharUnits size = getContext().getTypeSizeInChars(Ty); - llvm::Value *SizeVal; + llvm::Value *SizeVal;() const VariableArrayType *vla; // Don't bother emitting a zero-byte memset. if (size.isZero()) { // But note that getTypeInfo returns 0 for a VLA. - if (const VariableArrayType *vlaType = dyn_cast_or_null( - getContext().getAsArrayType(Ty))) { + if (const VariableArrayType *vlaType = + dyn_cast_or_null( + getContext().getAsArrayType(Ty))) { auto VlaSize = getVLASize(vlaType); SizeVal = VlaSize.NumElts; CharUnits eltSize = getContext().getTypeSizeInChars(VlaSize.Type); @@ -2096,22 +2088,21 @@ void CodeGenFunction::EmitNullInitialization(Address DestPtr, QualType Ty) { // like -1, which happens to be the pattern used by member-pointers. if (!CGM.getTypes().isZeroInitializable(Ty)) { // For a VLA, emit a single element, then splat that over the VLA. - if (vla) - Ty = getContext().getBaseElementType(vla); + if (vla) Ty = getContext().getBaseElementType(vla); llvm::Constant *NullConstant = CGM.EmitNullConstant(Ty); - llvm::GlobalVariable *NullVariable = new llvm::GlobalVariable( - CGM.getModule(), NullConstant->getType(), - /*isConstant=*/true, llvm::GlobalVariable::PrivateLinkage, NullConstant, - Twine()); + llvm::GlobalVariable *NullVariable = + new llvm::GlobalVariable(CGM.getModule(), NullConstant->getType(), + /*isConstant=*/true, + llvm::GlobalVariable::PrivateLinkage, + NullConstant, Twine()); CharUnits NullAlign = DestPtr.getAlignment(); NullVariable->setAlignment(NullAlign.getAsAlign()); Address SrcPtr(Builder.CreateBitCast(NullVariable, Builder.getInt8PtrTy()), NullAlign); - if (vla) - return emitNonZeroVLAInit(*this, Ty, DestPtr, SrcPtr, SizeVal); + if (vla) return emitNonZeroVLAInit(*this, Ty, DestPtr, SrcPtr, SizeVal); // Get and call the appropriate llvm.memcpy overload. Builder.CreateMemCpy(DestPtr, SrcPtr, SizeVal, false); @@ -2144,8 +2135,8 @@ llvm::BasicBlock *CodeGenFunction::GetIndirectGotoBlock() { CGBuilderTy TmpBuilder(*this, createBasicBlock("indirectgoto")); // Create the PHI node that indirect gotos will add entries to. - llvm::Value *DestVal = - TmpBuilder.CreatePHI(Int8PtrTy, 0, "indirect.goto.dest"); + llvm::Value *DestVal = TmpBuilder.CreatePHI(Int8PtrTy, 0, + "indirect.goto.dest"); // Create the indirect branch instruction. IndirectBranch = TmpBuilder.CreateIndirectBr(DestVal); @@ -2185,7 +2176,7 @@ llvm::Value *CodeGenFunction::emitArrayLength(const ArrayType *origArrayType, // We have some number of constant-length arrays, so addr should // have LLVM type [M x [N x [...]]]*. Build a GEP that walks // down to the first element of addr. - SmallVector gepIndices; + SmallVector gepIndices; // GEP down to the array type. llvm::ConstantInt *zero = Builder.getInt32(0); @@ -2195,17 +2186,18 @@ llvm::Value *CodeGenFunction::emitArrayLength(const ArrayType *origArrayType, QualType eltType; llvm::ArrayType *llvmArrayType = - dyn_cast(addr.getElementType()); + dyn_cast(addr.getElementType()); while (llvmArrayType) { assert(isa(arrayType)); - assert(cast(arrayType)->getSize().getZExtValue() == - llvmArrayType->getNumElements()); + assert(cast(arrayType)->getSize().getZExtValue() + == llvmArrayType->getNumElements()); gepIndices.push_back(zero); countFromCLAs *= llvmArrayType->getNumElements(); eltType = arrayType->getElementType(); - llvmArrayType = dyn_cast(llvmArrayType->getElementType()); + llvmArrayType = + dyn_cast(llvmArrayType->getElementType()); arrayType = getContext().getAsArrayType(arrayType->getElementType()); assert((!llvmArrayType || arrayType) && "LLVM and Clang types are out-of-synch"); @@ -2233,7 +2225,8 @@ llvm::Value *CodeGenFunction::emitArrayLength(const ArrayType *origArrayType, baseType = eltType; - llvm::Value *numElements = llvm::ConstantInt::get(SizeTy, countFromCLAs); + llvm::Value *numElements = + llvm::ConstantInt::get(SizeTy, countFromCLAs); // If we had any VLA dimensions, factor them in. if (numVLAElements) @@ -2269,10 +2262,11 @@ CodeGenFunction::getVLASize(const VariableArrayType *type) { } } while ((type = getContext().getAsVariableArrayType(elementType))); - return {numElements, elementType}; + return { numElements, elementType }; } -CodeGenFunction::VlaSizePair CodeGenFunction::getVLAElements1D(QualType type) { +CodeGenFunction::VlaSizePair +CodeGenFunction::getVLAElements1D(QualType type) { const VariableArrayType *vla = getContext().getAsVariableArrayType(type); assert(vla && "type was not a variable array type!"); return getVLAElements1D(vla); @@ -2283,7 +2277,7 @@ CodeGenFunction::getVLAElements1D(const VariableArrayType *Vla) { llvm::Value *VlaSize = VLASizeMap[Vla->getSizeExpr()]; assert(VlaSize && "no size for VLA!"); assert(VlaSize->getType() == SizeTy); - return {VlaSize, Vla->getElementType()}; + return { VlaSize, Vla->getElementType() }; } void CodeGenFunction::EmitVariablyModifiedType(QualType type) { @@ -2432,7 +2426,7 @@ void CodeGenFunction::EmitVariablyModifiedType(QualType type) { } while (type->isVariablyModifiedType()); } -Address CodeGenFunction::EmitVAListRef(const Expr *E) { +Address CodeGenFunction::EmitVAListRef(const Expr* E) { if (getContext().getBuiltinVaListType()->isArrayType()) return EmitPointerWithAlignment(E); return EmitLValue(E).getAddress(*this); @@ -2456,11 +2450,9 @@ CodeGenFunction::protectFromPeepholes(RValue rvalue) { // is trunc(zext) folding, but if we add more, we can easily // extend this protection. - if (!rvalue.isScalar()) - return PeepholeProtection(); + if (!rvalue.isScalar()) return PeepholeProtection(); llvm::Value *value = rvalue.getScalarVal(); - if (!isa(value)) - return PeepholeProtection(); + if (!isa(value)) return PeepholeProtection(); // Just make an extra bitcast. assert(HaveInsertPoint()); @@ -2473,8 +2465,7 @@ CodeGenFunction::protectFromPeepholes(RValue rvalue) { } void CodeGenFunction::unprotectFromPeepholes(PeepholeProtection protection) { - if (!protection.Inst) - return; + if (!protection.Inst) return; // In theory, we could try to duplicate the peepholes now, but whatever. protection.Inst->eraseFromParent(); @@ -2577,7 +2568,8 @@ Address CodeGenFunction::EmitFieldAnnotations(const FieldDecl *D, // llvm.ptr.annotation intrinsic accepts a pointer to integer of any width - // don't perform bitcasts if value is integer if (VTy->getPointerElementType()->isIntegerTy()) { - llvm::Function *F = CGM.getIntrinsic(llvm::Intrinsic::ptr_annotation, VTy); + llvm::Function *F = + CGM.getIntrinsic(llvm::Intrinsic::ptr_annotation, VTy); for (const auto *I : D->specific_attrs()) V = EmitAnnotationCall(F, V, I->getAnnotation(), D->getLocation(), I); @@ -2585,8 +2577,8 @@ Address CodeGenFunction::EmitFieldAnnotations(const FieldDecl *D, return Address(V, Addr.getAlignment()); } - llvm::Function *F = - CGM.getIntrinsic(llvm::Intrinsic::ptr_annotation, CGM.Int8PtrTy); + llvm::Function *F = CGM.getIntrinsic(llvm::Intrinsic::ptr_annotation, + CGM.Int8PtrTy); for (const auto *I : D->specific_attrs()) { V = Builder.CreateBitCast(V, CGM.Int8PtrTy); @@ -2628,7 +2620,7 @@ Address CodeGenFunction::EmitIntelFPGAFieldAnnotations(SourceLocation Location, return Address(V, Addr.getAlignment()); } -CodeGenFunction::CGCapturedStmtInfo::~CGCapturedStmtInfo() {} +CodeGenFunction::CGCapturedStmtInfo::~CGCapturedStmtInfo() { } CodeGenFunction::SanitizerScope::SanitizerScope(CodeGenFunction *CGF) : CGF(CGF) { @@ -2691,7 +2683,8 @@ void CodeGenFunction::checkTargetFeatures(SourceLocation Loc, // Return if the builtin doesn't have any required features. if (FeatureList.empty()) return; - assert(FeatureList.find(' ') == StringRef::npos && "Space in feature list"); + assert(FeatureList.find(' ') == StringRef::npos && + "Space in feature list"); TargetFeatures TF(CallerFeatureMap); if (!TF.hasRequiredFeatures(FeatureList)) CGM.getDiags().Report(Loc, diag::err_builtin_needs_feature) @@ -2719,12 +2712,12 @@ void CodeGenFunction::checkTargetFeatures(SourceLocation Loc, ReqFeatures.push_back(F.getKey()); } if (!llvm::all_of(ReqFeatures, [&](StringRef Feature) { - if (!CallerFeatureMap.lookup(Feature)) { - MissingFeature = Feature.str(); - return false; - } - return true; - })) + if (!CallerFeatureMap.lookup(Feature)) { + MissingFeature = Feature.str(); + return false; + } + return true; + })) CGM.getDiags().Report(Loc, diag::err_function_needs_feature) << FD->getDeclName() << TargetDecl->getDeclName() << MissingFeature; } diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index 4632aff6125ac..3cd4bd07ffa85 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -197,7 +197,8 @@ bool Sema::isKnownGoodSYCLDecl(const Decl *D) { if (const FunctionDecl *FD = dyn_cast(D)) { const IdentifierInfo *II = FD->getIdentifier(); const DeclContext *DC = FD->getDeclContext(); - if (II && II->isStr("__spirv_ocl_printf") && !FD->isDefined() && + if (II && II->isStr("__spirv_ocl_printf") && + !FD->isDefined() && FD->getLanguageLinkage() == CXXLanguageLinkage && DC->getEnclosingNamespaceContext()->isTranslationUnit()) return true; @@ -3665,7 +3666,8 @@ void Sema::finalizeSYCLDelayedAnalysis(const FunctionDecl *Caller, bool NotDefinedNoAttr = !Callee->isDefined() && !HasAttr; if (NotDefinedNoAttr && !Callee->getBuiltinID()) { - Diag(Loc, diag::err_sycl_restrict) << Sema::KernelCallUndefinedFunction; + Diag(Loc, diag::err_sycl_restrict) + << Sema::KernelCallUndefinedFunction; Diag(Callee->getLocation(), diag::note_previous_decl) << Callee; Diag(Caller->getLocation(), diag::note_called_by) << Caller; } From 232a8a9cb655270f35e59a774a8b9208019c7e6f Mon Sep 17 00:00:00 2001 From: Zahira Ammarguellat Date: Mon, 5 Apr 2021 13:16:38 -0700 Subject: [PATCH 06/19] Put back the indentation as it was Signed-off-by: Zahira Ammarguellat --- clang/lib/CodeGen/CodeGenFunction.cpp | 30 +++++++++++++-------------- clang/lib/Sema/SemaSYCL.cpp | 2 +- 2 files changed, 16 insertions(+), 16 deletions(-) diff --git a/clang/lib/CodeGen/CodeGenFunction.cpp b/clang/lib/CodeGen/CodeGenFunction.cpp index a557f2fdb695b..505e6283477b6 100644 --- a/clang/lib/CodeGen/CodeGenFunction.cpp +++ b/clang/lib/CodeGen/CodeGenFunction.cpp @@ -112,8 +112,8 @@ clang::ToConstrainedExceptMD(LangOptions::FPExceptionModeKind Kind) { void CodeGenFunction::SetFPModel() { llvm::RoundingMode RM = getLangOpts().getFPRoundingMode(); - auto fpExceptionBehavior = ToConstrainedExceptMD(getLangOpts(). - getFPExceptionMode()); + auto fpExceptionBehavior = ToConstrainedExceptMD( + getLangOpts().getFPExceptionMode()); Builder.setDefaultConstrainedRounding(RM); Builder.setDefaultConstrainedExcept(fpExceptionBehavior); @@ -206,7 +206,8 @@ LValue CodeGenFunction::MakeNaturalAlignAddrLValue(llvm::Value *V, QualType T) { /// Given a value of type T* that may not be to a complete object, /// construct an l-value with the natural pointee alignment of T. -LValue CodeGenFunction::MakeNaturalAlignPointeeAddrLValue(llvm::Value *V, QualType T) { +LValue +CodeGenFunction::MakeNaturalAlignPointeeAddrLValue(llvm::Value *V, QualType T) { LValueBaseInfo BaseInfo; TBAAAccessInfo TBAAInfo; CharUnits Align = CGM.getNaturalTypeAlignment(T, &BaseInfo, &TBAAInfo, @@ -1172,14 +1173,14 @@ void CodeGenFunction::StartFunction(GlobalDecl GD, QualType RetTy, if (CGM.getCodeGenOpts().MNopMCount) { if (!CGM.getCodeGenOpts().CallFEntry) CGM.getDiags().Report(diag::err_opt_not_valid_without_opt) - << "-mnop-mcount" << "-mfentry"; + << "-mnop-mcount" << "-mfentry"; Fn->addFnAttr("mnop-mcount"); } if (CGM.getCodeGenOpts().RecordMCount) { if (!CGM.getCodeGenOpts().CallFEntry) CGM.getDiags().Report(diag::err_opt_not_valid_without_opt) - << "-mrecord-mcount" << "-mfentry"; + << "-mrecord-mcount" << "-mfentry"; Fn->addFnAttr("mrecord-mcount"); } } @@ -2004,7 +2005,7 @@ static void emitNonZeroVLAInit(CodeGenFunction &CGF, QualType baseType, CGBuilderTy &Builder = CGF.Builder; CharUnits baseSize = CGF.getContext().getTypeSizeInChars(baseType); - llvm::Value *baseSizeInChars + llvm::Value *baseSizeInChars = llvm::ConstantInt::get(CGF.IntPtrTy, baseSize.getQuantity()); Address begin = @@ -2059,7 +2060,7 @@ CodeGenFunction::EmitNullInitialization(Address DestPtr, QualType Ty) { // Get size and alignment info for this aggregate. CharUnits size = getContext().getTypeSizeInChars(Ty); - llvm::Value *SizeVal;() + llvm::Value *SizeVal; const VariableArrayType *vla; // Don't bother emitting a zero-byte memset. @@ -2094,8 +2095,8 @@ CodeGenFunction::EmitNullInitialization(Address DestPtr, QualType Ty) { llvm::GlobalVariable *NullVariable = new llvm::GlobalVariable(CGM.getModule(), NullConstant->getType(), - /*isConstant=*/true, - llvm::GlobalVariable::PrivateLinkage, + /*isConstant=*/true, + llvm::GlobalVariable::PrivateLinkage, NullConstant, Twine()); CharUnits NullAlign = DestPtr.getAlignment(); NullVariable->setAlignment(NullAlign.getAsAlign()); @@ -2129,13 +2130,12 @@ llvm::BlockAddress *CodeGenFunction::GetAddrOfLabel(const LabelDecl *L) { llvm::BasicBlock *CodeGenFunction::GetIndirectGotoBlock() { // If we already made the indirect branch for indirect goto, return its block. - if (IndirectBranch) - return IndirectBranch->getParent(); + if (IndirectBranch) return IndirectBranch->getParent(); CGBuilderTy TmpBuilder(*this, createBasicBlock("indirectgoto")); // Create the PHI node that indirect gotos will add entries to. - llvm::Value *DestVal = TmpBuilder.CreatePHI(Int8PtrTy, 0, + llvm::Value *DestVal = TmpBuilder.CreatePHI(Int8PtrTy, 0, "indirect.goto.dest"); // Create the indirect branch instruction. @@ -2225,8 +2225,8 @@ llvm::Value *CodeGenFunction::emitArrayLength(const ArrayType *origArrayType, baseType = eltType; - llvm::Value *numElements = - llvm::ConstantInt::get(SizeTy, countFromCLAs); + llvm::Value *numElements + = llvm::ConstantInt::get(SizeTy, countFromCLAs); // If we had any VLA dimensions, factor them in. if (numVLAElements) @@ -2569,7 +2569,7 @@ Address CodeGenFunction::EmitFieldAnnotations(const FieldDecl *D, // don't perform bitcasts if value is integer if (VTy->getPointerElementType()->isIntegerTy()) { llvm::Function *F = - CGM.getIntrinsic(llvm::Intrinsic::ptr_annotation, VTy); + CGM.getIntrinsic(llvm::Intrinsic::ptr_annotation, VTy); for (const auto *I : D->specific_attrs()) V = EmitAnnotationCall(F, V, I->getAnnotation(), D->getLocation(), I); diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index 3cd4bd07ffa85..9ee38a44bc608 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -197,7 +197,7 @@ bool Sema::isKnownGoodSYCLDecl(const Decl *D) { if (const FunctionDecl *FD = dyn_cast(D)) { const IdentifierInfo *II = FD->getIdentifier(); const DeclContext *DC = FD->getDeclContext(); - if (II && II->isStr("__spirv_ocl_printf") && + if (II && II->isStr("__spirv_ocl_printf") && !FD->isDefined() && FD->getLanguageLinkage() == CXXLanguageLinkage && DC->getEnclosingNamespaceContext()->isTranslationUnit()) From 4be5deb9bcf2d72b10eea0bd6e47485ce2b158fa Mon Sep 17 00:00:00 2001 From: Zahira Ammarguellat Date: Wed, 7 Apr 2021 18:14:51 -0700 Subject: [PATCH 07/19] Addding test case Signed-off-by: Zahira Ammarguellat --- clang/test/SemaSYCL/kernel-arg-opt-report.cpp | 42 +++++++++++++++++++ 1 file changed, 42 insertions(+) create mode 100644 clang/test/SemaSYCL/kernel-arg-opt-report.cpp diff --git a/clang/test/SemaSYCL/kernel-arg-opt-report.cpp b/clang/test/SemaSYCL/kernel-arg-opt-report.cpp new file mode 100644 index 0000000000000..aa5955a0b1b33 --- /dev/null +++ b/clang/test/SemaSYCL/kernel-arg-opt-report.cpp @@ -0,0 +1,42 @@ +// RUN: %clang_cc1 -triple spir64-unknown-unknown-sycldevice -fsycl-is-device \ +// RUN: -Wno-sycl-2017-compat -emit-llvm-bc %s -o %t-host.bc -opt-record-file %t-host.yaml + +#include "Inputs/sycl.hpp" + +class second_base { +public: + int *e; +}; + +class InnerFieldBase { +public: + int d; +}; +class InnerField : public InnerFieldBase { + int c; +}; + +struct base { +public: + int b; + InnerField obj; +}; + +struct derived : base, second_base { + int a; + + void operator()() const { + } +}; + +int main() { + cl::sycl::queue q; + + q.submit([&](cl::sycl::handler &cgh) { + derived f{}; + cgh.single_task(f); + }); + + return 0; +} + From 978d6b8fd9ae64a83ff1d6d59e9d35f9aaf42e97 Mon Sep 17 00:00:00 2001 From: Zahira Ammarguellat Date: Wed, 7 Apr 2021 18:34:06 -0700 Subject: [PATCH 08/19] Addding test case Signed-off-by: Zahira Ammarguellat --- clang/include/clang/Basic/SyclOptReportHandler.h | 7 +++---- clang/lib/CodeGen/CGSYCLRuntime.cpp | 5 +++-- clang/lib/CodeGen/CodeGenAction.cpp | 11 ++++++----- clang/lib/CodeGen/CodeGenFunction.cpp | 10 ++++------ clang/lib/Sema/SemaSYCL.cpp | 8 +++----- 5 files changed, 19 insertions(+), 22 deletions(-) diff --git a/clang/include/clang/Basic/SyclOptReportHandler.h b/clang/include/clang/Basic/SyclOptReportHandler.h index f57451c38d163..936953ba430aa 100644 --- a/clang/include/clang/Basic/SyclOptReportHandler.h +++ b/clang/include/clang/Basic/SyclOptReportHandler.h @@ -30,10 +30,9 @@ class SyclOptReportHandler { SourceLocation KernelArgLoc; OptReportInfo(llvm::StringRef ArgName, std::string ArgType, - SourceLocation ArgLoc) : - KernelArgName(ArgName), - KernelArgType(ArgType), - KernelArgLoc(ArgLoc) {} + SourceLocation ArgLoc) + : KernelArgName(ArgName), KernelArgType(ArgType), KernelArgLoc(ArgLoc) { + } }; llvm::DenseMap> Map; diff --git a/clang/lib/CodeGen/CGSYCLRuntime.cpp b/clang/lib/CodeGen/CGSYCLRuntime.cpp index 03aebb500c52c..e0914b38c4eb7 100644 --- a/clang/lib/CodeGen/CGSYCLRuntime.cpp +++ b/clang/lib/CodeGen/CGSYCLRuntime.cpp @@ -14,9 +14,10 @@ #include "CodeGenFunction.h" #include "clang/AST/Attr.h" #include "clang/AST/Decl.h" -#include "llvm/IR/Instructions.h" -#include "llvm/Analysis/OptimizationRemarkEmitter.h" #include "clang/Basic/SourceLocation.h" +#include "llvm/Analysis/OptimizationRemarkEmitter.h" +#include "llvm/IR/Instructions.h" + #include using namespace clang; diff --git a/clang/lib/CodeGen/CodeGenAction.cpp b/clang/lib/CodeGen/CodeGenAction.cpp index f02344ffe4086..e69757aa1e74c 100644 --- a/clang/lib/CodeGen/CodeGenAction.cpp +++ b/clang/lib/CodeGen/CodeGenAction.cpp @@ -305,6 +305,7 @@ namespace clang { LLVMContext &Ctx = getModule()->getContext(); + std::unique_ptr OldDiagnosticHandler = Ctx.getDiagnosticHandler(); Ctx.setDiagnosticHandler(std::make_unique( @@ -1028,10 +1029,10 @@ CodeGenAction::loadModule(MemoryBufferRef MBRef) { } namespace { - // Handles the initialization and cleanup of the OptRecordFile. This - // customization allows initialization before the clang codegen runs - // so it can also emit to the opt report. - struct OptRecordFileRAII { + // Handles the initialization and cleanup of the OptRecordFile. This + // customization allows initialization before the clang codegen runs + // so it can also emit to the opt report. + struct OptRecordFileRAII { std::unique_ptr OptRecordFile; std::unique_ptr OldDiagnosticHandler; llvm::LLVMContext &Ctx; @@ -1067,7 +1068,7 @@ namespace { if (OptRecordFile) OptRecordFile->keep(); } - }; +}; } // namespace void CodeGenAction::ExecuteAction() { diff --git a/clang/lib/CodeGen/CodeGenFunction.cpp b/clang/lib/CodeGen/CodeGenFunction.cpp index 505e6283477b6..76b79572a89ec 100644 --- a/clang/lib/CodeGen/CodeGenFunction.cpp +++ b/clang/lib/CodeGen/CodeGenFunction.cpp @@ -1485,7 +1485,6 @@ void CodeGenFunction::GenerateCode(GlobalDecl GD, llvm::Function *Fn, // Emit the standard function prologue. StartFunction(GD, ResTy, Fn, FnInfo, Args, Loc, BodyRange.getBegin()); - SyclOptReportHandler &OptReportHandler = CGM.getDiags().OptReportHandler; if (OptReportHandler.HasOptReportInfo(FD)) { llvm::OptimizationRemarkEmitter ORE(Fn); @@ -1496,10 +1495,10 @@ void CodeGenFunction::GenerateCode(GlobalDecl GD, llvm::Function *Fn, &Fn->getEntryBlock()); Remark << "Argument " << llvm::ore::NV("Argument", count++) << " for function kernel: " - << llvm::ore::NV(ORI.KernelArgName.empty() - ? "&" : "") << " " + << llvm::ore::NV(ORI.KernelArgName.empty() ? "&" : "") << " " << Fn->getName() << "." - << llvm::ore::NV(ORI.KernelArgName.empty() ? " " : ORI.KernelArgName) + << llvm::ore::NV(ORI.KernelArgName.empty() ? " " + : ORI.KernelArgName) << "(" << ORI.KernelArgType << ")"; ORE.emit(Remark); } @@ -2568,8 +2567,7 @@ Address CodeGenFunction::EmitFieldAnnotations(const FieldDecl *D, // llvm.ptr.annotation intrinsic accepts a pointer to integer of any width - // don't perform bitcasts if value is integer if (VTy->getPointerElementType()->isIntegerTy()) { - llvm::Function *F = - CGM.getIntrinsic(llvm::Intrinsic::ptr_annotation, VTy); + llvm::Function *F = CGM.getIntrinsic(llvm::Intrinsic::ptr_annotation, VTy); for (const auto *I : D->specific_attrs()) V = EmitAnnotationCall(F, V, I->getAnnotation(), D->getLocation(), I); diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index 9ee38a44bc608..b23a394a226a9 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -1692,11 +1692,10 @@ class SyclKernelDeclCreator : public SyclKernelFieldHandler { void addParam(const FieldDecl *FD, QualType FieldTy) { ParamDesc newParamDesc = makeParamDesc(FD, FieldTy); SemaRef.getDiagnostics().OptReportHandler.AddKernelArgs( - KernelDecl, FD->getName(), FieldTy.getAsString(), - FD->getLocation()); + KernelDecl, FD->getName(), FieldTy.getAsString(), FD->getLocation()); addParam(newParamDesc, FieldTy); } - + void addParam(const CXXBaseSpecifier &BS, QualType FieldTy) { // TODO: There is no name for the base available, but duplicate names are // seemingly already possible, so we'll give them all the same name for now. @@ -1705,8 +1704,7 @@ class SyclKernelDeclCreator : public SyclKernelFieldHandler { ParamDesc newParamDesc = makeParamDesc(SemaRef.getASTContext(), Name, FieldTy); SemaRef.getDiagnostics().OptReportHandler.AddKernelArgs( - KernelDecl, "", FieldTy.getAsString(), - BS.getBaseTypeLoc()); + KernelDecl, "", FieldTy.getAsString(), BS.getBaseTypeLoc()); addParam(newParamDesc, FieldTy); } // Add a parameter with specified name and type From 32454acc0f06228766da512210c0ab7daac80a81 Mon Sep 17 00:00:00 2001 From: Zahira Ammarguellat Date: Wed, 14 Apr 2021 17:50:17 -0700 Subject: [PATCH 09/19] Formatting Signed-off-by: Zahira Ammarguellat --- clang/lib/CodeGen/CodeGenFunction.cpp | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) diff --git a/clang/lib/CodeGen/CodeGenFunction.cpp b/clang/lib/CodeGen/CodeGenFunction.cpp index 76b79572a89ec..74c49d4ffd744 100644 --- a/clang/lib/CodeGen/CodeGenFunction.cpp +++ b/clang/lib/CodeGen/CodeGenFunction.cpp @@ -1494,11 +1494,11 @@ void CodeGenFunction::GenerateCode(GlobalDecl GD, llvm::Function *Fn, llvm::OptimizationRemark Remark("sycl", "Region", DL, &Fn->getEntryBlock()); Remark << "Argument " << llvm::ore::NV("Argument", count++) - << " for function kernel: " + << " for function kernel: " << llvm::ore::NV(ORI.KernelArgName.empty() ? "&" : "") << " " - << Fn->getName() << "." + << Fn->getName() << "." << llvm::ore::NV(ORI.KernelArgName.empty() ? " " - : ORI.KernelArgName) + : ORI.KernelArgName) << "(" << ORI.KernelArgType << ")"; ORE.emit(Remark); } From d31606265dedd62d28e0b50d6819a14e72b20594 Mon Sep 17 00:00:00 2001 From: Zahira Ammarguellat Date: Wed, 14 Apr 2021 17:58:51 -0700 Subject: [PATCH 10/19] Formatting Signed-off-by: Zahira Ammarguellat --- clang/include/clang/Basic/SyclOptReportHandler.h | 2 +- clang/lib/CodeGen/CodeGenAction.cpp | 9 ++++----- 2 files changed, 5 insertions(+), 6 deletions(-) diff --git a/clang/include/clang/Basic/SyclOptReportHandler.h b/clang/include/clang/Basic/SyclOptReportHandler.h index 936953ba430aa..e5855def66286 100644 --- a/clang/include/clang/Basic/SyclOptReportHandler.h +++ b/clang/include/clang/Basic/SyclOptReportHandler.h @@ -31,7 +31,7 @@ class SyclOptReportHandler { OptReportInfo(llvm::StringRef ArgName, std::string ArgType, SourceLocation ArgLoc) - : KernelArgName(ArgName), KernelArgType(ArgType), KernelArgLoc(ArgLoc) { + : KernelArgName(ArgName), KernelArgType(ArgType), KernelArgLoc(ArgLoc) { } }; llvm::DenseMap> Map; diff --git a/clang/lib/CodeGen/CodeGenAction.cpp b/clang/lib/CodeGen/CodeGenAction.cpp index e69757aa1e74c..41b6d82035eec 100644 --- a/clang/lib/CodeGen/CodeGenAction.cpp +++ b/clang/lib/CodeGen/CodeGenAction.cpp @@ -304,7 +304,6 @@ namespace clang { return; LLVMContext &Ctx = getModule()->getContext(); - std::unique_ptr OldDiagnosticHandler = Ctx.getDiagnosticHandler(); @@ -1029,10 +1028,10 @@ CodeGenAction::loadModule(MemoryBufferRef MBRef) { } namespace { - // Handles the initialization and cleanup of the OptRecordFile. This - // customization allows initialization before the clang codegen runs - // so it can also emit to the opt report. - struct OptRecordFileRAII { +// Handles the initialization and cleanup of the OptRecordFile. This +// customization allows initialization before the clang codegen runs +// so it can also emit to the opt report. +struct OptRecordFileRAII { std::unique_ptr OptRecordFile; std::unique_ptr OldDiagnosticHandler; llvm::LLVMContext &Ctx; From aeed8072dc7999d0fa098d09f605f662faff14aa Mon Sep 17 00:00:00 2001 From: Zahira Ammarguellat Date: Thu, 15 Apr 2021 12:58:05 -0700 Subject: [PATCH 11/19] After review fixes Signed-off-by: Zahira Ammarguellat --- clang/test/SemaSYCL/kernel-arg-opt-report.cpp | 63 ++++++++++++++++++- 1 file changed, 61 insertions(+), 2 deletions(-) diff --git a/clang/test/SemaSYCL/kernel-arg-opt-report.cpp b/clang/test/SemaSYCL/kernel-arg-opt-report.cpp index aa5955a0b1b33..fe71db3e833e3 100644 --- a/clang/test/SemaSYCL/kernel-arg-opt-report.cpp +++ b/clang/test/SemaSYCL/kernel-arg-opt-report.cpp @@ -1,5 +1,8 @@ // RUN: %clang_cc1 -triple spir64-unknown-unknown-sycldevice -fsycl-is-device \ // RUN: -Wno-sycl-2017-compat -emit-llvm-bc %s -o %t-host.bc -opt-record-file %t-host.yaml +// RUN: FileCheck -check-prefix=CHECK --input-file %t-host.yaml %s +// The test generates remarks about the kernel argument, their location and type +// in the resulting yaml file. #include "Inputs/sycl.hpp" @@ -22,6 +25,63 @@ struct base { InnerField obj; }; +//CHECK: --- !Passed +//CHECK: Pass:{{.*}}sycl +//CHECK: Name:{{.*}}Region +//CHECK: DebugLoc:{{.*}} { File: '{{.*}}kernel-arg-opt-report.cpp', +//CHECK: Line: 85, Column: 18 } +//CHECK: Function: _ZTS7derived +//CHECK: Args: +//CHECK-NEXT: String: 'Argument ' +//CHECK-NEXT: Argument: '0' +//CHECK-NEXT: String: ' for function kernel: ' +//CHECK-NEXT: String: '&' +//CHECK-NEXT: String: ' ' +//CHECK-NEXT: String: _ZTS7derived +//CHECK-NEXT: String: . +//CHECK-NEXT: String: ' ' +//CHECK-NEXT: String: '(' +//CHECK-NEXT: String: struct base +//CHECK-NEXT: String: ')' + +//CHECK: --- !Passed +//CHECK: Pass:{{.*}}sycl +//CHECK: Name:{{.*}}Region +//CHECK: DebugLoc:{{.*}} { File: '{{.*}}kernel-arg-opt-report.cpp', +//CHECK: Line: 11, Column: 8 } +//CHECK: Function: _ZTS7derived +//CHECK: Args: +//CHECK-NEXT: String: 'Argument ' +//CHECK-NEXT: Argument: '1' +//CHECK-NEXT: String: ' for function kernel: ' +//CHECK-NEXT: String: '' +//CHECK-NEXT: String: ' ' +//CHECK-NEXT: String: _ZTS7derived +//CHECK-NEXT: String: . +//CHECK-NEXT: String: e +//CHECK-NEXT: String: '(' +//CHECK-NEXT: String: struct __wrapper_class +//CHECK-NEXT: String: ')' + +//CHECK: --- !Passed +//CHECK: Pass:{{.*}}sycl +//CHECK: Name:{{.*}}Region +//CHECK: DebugLoc:{{.*}} { File: '{{.*}}kernel-arg-opt-report.cpp', +//CHECK: Line: 86, Column: 7 } +//CHECK: Function: _ZTS7derived +//CHECK: Args: +//CHECK-NEXT: String: 'Argument ' +//CHECK-NEXT: Argument: '2' +//CHECK-NEXT: String: ' for function kernel: ' +//CHECK-NEXT: String: '' +//CHECK-NEXT: String: ' ' +//CHECK-NEXT: String: _ZTS7derived +//CHECK-NEXT: String: . +//CHECK-NEXT: String: a +//CHECK-NEXT: String: '(' +//CHECK-NEXT: String: int +//CHECK-NEXT: String: ')' + struct derived : base, second_base { int a; @@ -30,7 +90,7 @@ struct derived : base, second_base { }; int main() { - cl::sycl::queue q; + sycl::queue q; q.submit([&](cl::sycl::handler &cgh) { derived f{}; @@ -39,4 +99,3 @@ int main() { return 0; } - From 6ce72687567f8b2b092fb3b603e152b9ab6a6481 Mon Sep 17 00:00:00 2001 From: Zahira Ammarguellat Date: Tue, 20 Apr 2021 12:27:15 -0700 Subject: [PATCH 12/19] Review comments fixes Signed-off-by: Zahira Ammarguellat --- clang/include/clang/Basic/SyclOptReportHandler.h | 4 ++-- clang/lib/CodeGen/CodeGenAction.cpp | 2 +- clang/lib/CodeGen/CodeGenFunction.cpp | 6 +++--- 3 files changed, 6 insertions(+), 6 deletions(-) diff --git a/clang/include/clang/Basic/SyclOptReportHandler.h b/clang/include/clang/Basic/SyclOptReportHandler.h index e5855def66286..ba246f7958970 100644 --- a/clang/include/clang/Basic/SyclOptReportHandler.h +++ b/clang/include/clang/Basic/SyclOptReportHandler.h @@ -41,13 +41,13 @@ class SyclOptReportHandler { std::string ArgType, SourceLocation ArgLoc) { Map[FD].emplace_back(ArgName, ArgType, ArgLoc); } - SmallVector &getInfo(const FunctionDecl *FD) { + SmallVector &GetInfo(const FunctionDecl *FD) { auto It = Map.find(FD); assert(It != Map.end()); return It->second; } - bool HasOptReportInfo(const FunctionDecl *FD) { + bool HasOptReportInfo(const FunctionDecl *FD) const { return Map.find(FD) != Map.end(); } }; diff --git a/clang/lib/CodeGen/CodeGenAction.cpp b/clang/lib/CodeGen/CodeGenAction.cpp index 41b6d82035eec..deae8ed0e640c 100644 --- a/clang/lib/CodeGen/CodeGenAction.cpp +++ b/clang/lib/CodeGen/CodeGenAction.cpp @@ -1040,7 +1040,7 @@ struct OptRecordFileRAII { BackendConsumer &BC) : Ctx(Ctx) { CompilerInstance &CI = CGA.getCompilerInstance(); - auto &CodeGenOpts = CI.getCodeGenOpts(); + CodeGenOptions &CodeGenOpts = CI.getCodeGenOpts(); OldDiagnosticHandler = Ctx.getDiagnosticHandler(); diff --git a/clang/lib/CodeGen/CodeGenFunction.cpp b/clang/lib/CodeGen/CodeGenFunction.cpp index a90f15d0a2d1d..2c69e4e1f4977 100644 --- a/clang/lib/CodeGen/CodeGenFunction.cpp +++ b/clang/lib/CodeGen/CodeGenFunction.cpp @@ -1525,12 +1525,12 @@ void CodeGenFunction::GenerateCode(GlobalDecl GD, llvm::Function *Fn, SyclOptReportHandler &OptReportHandler = CGM.getDiags().OptReportHandler; if (OptReportHandler.HasOptReportInfo(FD)) { llvm::OptimizationRemarkEmitter ORE(Fn); - int count = 0; - for (auto &ORI : OptReportHandler.getInfo(FD)) { + int Count = 0; + for (auto &ORI : OptReportHandler.GetInfo(FD)) { llvm::DiagnosticLocation DL = SourceLocToDebugLoc(ORI.KernelArgLoc); llvm::OptimizationRemark Remark("sycl", "Region", DL, &Fn->getEntryBlock()); - Remark << "Argument " << llvm::ore::NV("Argument", count++) + Remark << "Argument " << llvm::ore::NV("Argument", Count++) << " for function kernel: " << llvm::ore::NV(ORI.KernelArgName.empty() ? "&" : "") << " " << Fn->getName() << "." From ddad6f91d056f6b133429c0c8ee7716b97c5fca1 Mon Sep 17 00:00:00 2001 From: Zahira Ammarguellat Date: Wed, 21 Apr 2021 08:06:18 -0700 Subject: [PATCH 13/19] Review comments fixes Signed-off-by: Zahira Ammarguellat --- clang/include/clang/Basic/SyclOptReportHandler.h | 6 +++--- clang/lib/Sema/SemaSYCL.cpp | 3 ++- 2 files changed, 5 insertions(+), 4 deletions(-) diff --git a/clang/include/clang/Basic/SyclOptReportHandler.h b/clang/include/clang/Basic/SyclOptReportHandler.h index ba246f7958970..4795803340fcb 100644 --- a/clang/include/clang/Basic/SyclOptReportHandler.h +++ b/clang/include/clang/Basic/SyclOptReportHandler.h @@ -25,11 +25,11 @@ class FunctionDecl; class SyclOptReportHandler { private: struct OptReportInfo { - StringRef KernelArgName; + std::string KernelArgName; std::string KernelArgType; SourceLocation KernelArgLoc; - OptReportInfo(llvm::StringRef ArgName, std::string ArgType, + OptReportInfo(std::string ArgName, std::string ArgType, SourceLocation ArgLoc) : KernelArgName(ArgName), KernelArgType(ArgType), KernelArgLoc(ArgLoc) { } @@ -37,7 +37,7 @@ class SyclOptReportHandler { llvm::DenseMap> Map; public: - void AddKernelArgs(const FunctionDecl *FD, StringRef ArgName, + void AddKernelArgs(const FunctionDecl *FD, std::string ArgName, std::string ArgType, SourceLocation ArgLoc) { Map[FD].emplace_back(ArgName, ArgType, ArgLoc); } diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index 0d13ff8857558..de843112b6ec3 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -1769,7 +1769,8 @@ class SyclKernelDeclCreator : public SyclKernelFieldHandler { void addParam(const FieldDecl *FD, QualType FieldTy) { ParamDesc newParamDesc = makeParamDesc(FD, FieldTy); SemaRef.getDiagnostics().OptReportHandler.AddKernelArgs( - KernelDecl, FD->getName(), FieldTy.getAsString(), FD->getLocation()); + KernelDecl, FD->getName().data(), FieldTy.getAsString(), + FD->getLocation()); addParam(newParamDesc, FieldTy); } From 0af442253fe8c21b0d5d7925d857877d705f68bd Mon Sep 17 00:00:00 2001 From: Zahira Ammarguellat Date: Wed, 21 Apr 2021 12:24:09 -0700 Subject: [PATCH 14/19] Review comments fixes Signed-off-by: Zahira Ammarguellat --- clang/lib/CodeGen/CGSYCLRuntime.cpp | 1 - clang/lib/CodeGen/CodeGenFunction.cpp | 19 ++++++++++--------- 2 files changed, 10 insertions(+), 10 deletions(-) diff --git a/clang/lib/CodeGen/CGSYCLRuntime.cpp b/clang/lib/CodeGen/CGSYCLRuntime.cpp index e0914b38c4eb7..eb6e512c05a44 100644 --- a/clang/lib/CodeGen/CGSYCLRuntime.cpp +++ b/clang/lib/CodeGen/CGSYCLRuntime.cpp @@ -17,7 +17,6 @@ #include "clang/Basic/SourceLocation.h" #include "llvm/Analysis/OptimizationRemarkEmitter.h" #include "llvm/IR/Instructions.h" - #include using namespace clang; diff --git a/clang/lib/CodeGen/CodeGenFunction.cpp b/clang/lib/CodeGen/CodeGenFunction.cpp index 2c69e4e1f4977..255bc0dd475a2 100644 --- a/clang/lib/CodeGen/CodeGenFunction.cpp +++ b/clang/lib/CodeGen/CodeGenFunction.cpp @@ -1525,18 +1525,19 @@ void CodeGenFunction::GenerateCode(GlobalDecl GD, llvm::Function *Fn, SyclOptReportHandler &OptReportHandler = CGM.getDiags().OptReportHandler; if (OptReportHandler.HasOptReportInfo(FD)) { llvm::OptimizationRemarkEmitter ORE(Fn); - int Count = 0; - for (auto &ORI : OptReportHandler.GetInfo(FD)) { - llvm::DiagnosticLocation DL = SourceLocToDebugLoc(ORI.KernelArgLoc); + for (auto ORI : llvm::enumerate(OptReportHandler.GetInfo(FD))) { + llvm::DiagnosticLocation DL = + SourceLocToDebugLoc(ORI.value().KernelArgLoc); llvm::OptimizationRemark Remark("sycl", "Region", DL, &Fn->getEntryBlock()); - Remark << "Argument " << llvm::ore::NV("Argument", Count++) + Remark << "Argument " << llvm::ore::NV("Argument", ORI.index()) << " for function kernel: " - << llvm::ore::NV(ORI.KernelArgName.empty() ? "&" : "") << " " - << Fn->getName() << "." - << llvm::ore::NV(ORI.KernelArgName.empty() ? " " - : ORI.KernelArgName) - << "(" << ORI.KernelArgType << ")"; + << llvm::ore::NV(ORI.value().KernelArgName.empty() ? "&" : "") + << " " << Fn->getName() << "." + << llvm::ore::NV(ORI.value().KernelArgName.empty() + ? " " + : ORI.value().KernelArgName) + << "(" << ORI.value().KernelArgType << ")"; ORE.emit(Remark); } } From f99e200249e14920c4a10da922dd32aea4c7b85e Mon Sep 17 00:00:00 2001 From: Zahira Ammarguellat Date: Thu, 22 Apr 2021 06:46:07 -0700 Subject: [PATCH 15/19] Review comments fixes Signed-off-by: Zahira Ammarguellat --- clang/include/clang/Basic/Diagnostic.h | 7 ++++++- clang/lib/CodeGen/CodeGenFunction.cpp | 3 ++- clang/lib/Sema/SemaSYCL.cpp | 4 ++-- 3 files changed, 10 insertions(+), 4 deletions(-) diff --git a/clang/include/clang/Basic/Diagnostic.h b/clang/include/clang/Basic/Diagnostic.h index 6d674dc1995ba..f389a39e84a54 100644 --- a/clang/include/clang/Basic/Diagnostic.h +++ b/clang/include/clang/Basic/Diagnostic.h @@ -295,6 +295,7 @@ class DiagnosticsEngine : public RefCountedBase { DiagnosticConsumer *Client = nullptr; std::unique_ptr Owner; SourceManager *SourceMgr = nullptr; + SyclOptReportHandler OptReportHandler; /// Mapping information for diagnostics. /// @@ -548,7 +549,11 @@ class DiagnosticsEngine : public RefCountedBase { LLVM_DUMP_METHOD void dump() const; LLVM_DUMP_METHOD void dump(StringRef DiagName) const; - SyclOptReportHandler OptReportHandler; + /// Retrieve the report SyclOptReport info. + SyclOptReportHandler &getSYCLOptReportHandler() { return OptReportHandler; } + const SyclOptReportHandler &getSYCLOptReportHandler() const { + return OptReportHandler; + } const IntrusiveRefCntPtr &getDiagnosticIDs() const { return Diags; diff --git a/clang/lib/CodeGen/CodeGenFunction.cpp b/clang/lib/CodeGen/CodeGenFunction.cpp index 255bc0dd475a2..fe2cd5cae1d0d 100644 --- a/clang/lib/CodeGen/CodeGenFunction.cpp +++ b/clang/lib/CodeGen/CodeGenFunction.cpp @@ -1522,7 +1522,8 @@ void CodeGenFunction::GenerateCode(GlobalDecl GD, llvm::Function *Fn, // Emit the standard function prologue. StartFunction(GD, ResTy, Fn, FnInfo, Args, Loc, BodyRange.getBegin()); - SyclOptReportHandler &OptReportHandler = CGM.getDiags().OptReportHandler; + SyclOptReportHandler &OptReportHandler = + CGM.getDiags().getSYCLOptReportHandler(); if (OptReportHandler.HasOptReportInfo(FD)) { llvm::OptimizationRemarkEmitter ORE(Fn); for (auto ORI : llvm::enumerate(OptReportHandler.GetInfo(FD))) { diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index de843112b6ec3..a1f2852cdb9dd 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -1768,7 +1768,7 @@ class SyclKernelDeclCreator : public SyclKernelFieldHandler { void addParam(const FieldDecl *FD, QualType FieldTy) { ParamDesc newParamDesc = makeParamDesc(FD, FieldTy); - SemaRef.getDiagnostics().OptReportHandler.AddKernelArgs( + SemaRef.getDiagnostics().getSYCLOptReportHandler().AddKernelArgs( KernelDecl, FD->getName().data(), FieldTy.getAsString(), FD->getLocation()); addParam(newParamDesc, FieldTy); @@ -1781,7 +1781,7 @@ class SyclKernelDeclCreator : public SyclKernelFieldHandler { StringRef Name = "_arg__base"; ParamDesc newParamDesc = makeParamDesc(SemaRef.getASTContext(), Name, FieldTy); - SemaRef.getDiagnostics().OptReportHandler.AddKernelArgs( + SemaRef.getDiagnostics().getSYCLOptReportHandler().AddKernelArgs( KernelDecl, "", FieldTy.getAsString(), BS.getBaseTypeLoc()); addParam(newParamDesc, FieldTy); } From 6a87209592c05ca11a6fd3fa2857f20b24a26881 Mon Sep 17 00:00:00 2001 From: Zahira Ammarguellat Date: Thu, 22 Apr 2021 13:06:22 -0700 Subject: [PATCH 16/19] Review comments fixes Signed-off-by: Zahira Ammarguellat --- clang/include/clang/Basic/SyclOptReportHandler.h | 8 ++++---- clang/lib/CodeGen/CodeGenAction.cpp | 5 ++--- 2 files changed, 6 insertions(+), 7 deletions(-) diff --git a/clang/include/clang/Basic/SyclOptReportHandler.h b/clang/include/clang/Basic/SyclOptReportHandler.h index 4795803340fcb..107acc872d583 100644 --- a/clang/include/clang/Basic/SyclOptReportHandler.h +++ b/clang/include/clang/Basic/SyclOptReportHandler.h @@ -31,17 +31,17 @@ class SyclOptReportHandler { OptReportInfo(std::string ArgName, std::string ArgType, SourceLocation ArgLoc) - : KernelArgName(ArgName), KernelArgType(ArgType), KernelArgLoc(ArgLoc) { - } + : KernelArgName(std::move(ArgName)), KernelArgType(std::move(ArgType)), + KernelArgLoc(ArgLoc) {} }; - llvm::DenseMap> Map; + llvm::DenseMap> Map; public: void AddKernelArgs(const FunctionDecl *FD, std::string ArgName, std::string ArgType, SourceLocation ArgLoc) { Map[FD].emplace_back(ArgName, ArgType, ArgLoc); } - SmallVector &GetInfo(const FunctionDecl *FD) { + SmallVector &GetInfo(const FunctionDecl *FD) { auto It = Map.find(FD); assert(It != Map.end()); return It->second; diff --git a/clang/lib/CodeGen/CodeGenAction.cpp b/clang/lib/CodeGen/CodeGenAction.cpp index deae8ed0e640c..8175574ff6a52 100644 --- a/clang/lib/CodeGen/CodeGenAction.cpp +++ b/clang/lib/CodeGen/CodeGenAction.cpp @@ -1038,12 +1038,11 @@ struct OptRecordFileRAII { OptRecordFileRAII(CodeGenAction &CGA, llvm::LLVMContext &Ctx, BackendConsumer &BC) - : Ctx(Ctx) { + : Ctx(Ctx), OldDiagnosticHandler(Ctx.getDiagnosticHandler()) { + CompilerInstance &CI = CGA.getCompilerInstance(); CodeGenOptions &CodeGenOpts = CI.getCodeGenOpts(); - OldDiagnosticHandler = Ctx.getDiagnosticHandler(); - Ctx.setDiagnosticHandler( std::make_unique(CodeGenOpts, &BC)); From eac6f25ce28c1c4f140399eb131d10b608ae1af6 Mon Sep 17 00:00:00 2001 From: Zahira Ammarguellat Date: Thu, 22 Apr 2021 14:01:03 -0700 Subject: [PATCH 17/19] Review comments fixes Signed-off-by: Zahira Ammarguellat --- clang/include/clang/Basic/SyclOptReportHandler.h | 2 +- clang/lib/CodeGen/CodeGenFunction.cpp | 3 ++- 2 files changed, 3 insertions(+), 2 deletions(-) diff --git a/clang/include/clang/Basic/SyclOptReportHandler.h b/clang/include/clang/Basic/SyclOptReportHandler.h index 107acc872d583..2cb6650ba7a29 100644 --- a/clang/include/clang/Basic/SyclOptReportHandler.h +++ b/clang/include/clang/Basic/SyclOptReportHandler.h @@ -7,7 +7,7 @@ //===----------------------------------------------------------------------===// /// /// \file -/// Defines clang::SyclOptReport class. +/// Defines clang::SyclOptReportHandler class. /// //===----------------------------------------------------------------------===// diff --git a/clang/lib/CodeGen/CodeGenFunction.cpp b/clang/lib/CodeGen/CodeGenFunction.cpp index fe2cd5cae1d0d..f61b1b97b427d 100644 --- a/clang/lib/CodeGen/CodeGenFunction.cpp +++ b/clang/lib/CodeGen/CodeGenFunction.cpp @@ -2645,7 +2645,8 @@ Address CodeGenFunction::EmitIntelFPGAFieldAnnotations(SourceLocation Location, // llvm.ptr.annotation intrinsic accepts a pointer to integer of any width - // don't perform bitcasts if value is integer if (VTy->getPointerElementType()->isIntegerTy()) { - llvm::Function *F = CGM.getIntrinsic(llvm::Intrinsic::ptr_annotation, VTy); + llvm::Function *F = + CGM.getIntrinsic(llvm::Intrinsic::ptr_annotation, VTy); V = EmitAnnotationCall(F, V, AnnotStr, Location); return Address(V, Addr.getAlignment()); From 84a2ceea579aa885d21cdcecc9d35abd27cbd8b3 Mon Sep 17 00:00:00 2001 From: Zahira Ammarguellat Date: Fri, 23 Apr 2021 06:36:32 -0700 Subject: [PATCH 18/19] Review comments fixes Signed-off-by: Zahira Ammarguellat --- clang/include/clang/Basic/SyclOptReportHandler.h | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/clang/include/clang/Basic/SyclOptReportHandler.h b/clang/include/clang/Basic/SyclOptReportHandler.h index 2cb6650ba7a29..af3ba5ae1c0a8 100644 --- a/clang/include/clang/Basic/SyclOptReportHandler.h +++ b/clang/include/clang/Basic/SyclOptReportHandler.h @@ -1,4 +1,4 @@ -//===---------------------- SyclOptReport.h ---------------------*- C++ -*-===// +//===---------------------- SyclOptReportHandler.h --------------*- C++ -*-===// // // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. // See https://llvm.org/LICENSE.txt for license information. From 8efb79ac73897bb6ad5022e398a82a1ff7f6eae8 Mon Sep 17 00:00:00 2001 From: Zahira Ammarguellat Date: Fri, 23 Apr 2021 07:18:56 -0700 Subject: [PATCH 19/19] Review comments fixes Signed-off-by: Zahira Ammarguellat --- clang/lib/CodeGen/CodeGenFunction.cpp | 10 ++++------ 1 file changed, 4 insertions(+), 6 deletions(-) diff --git a/clang/lib/CodeGen/CodeGenFunction.cpp b/clang/lib/CodeGen/CodeGenFunction.cpp index f61b1b97b427d..aabfb649174b1 100644 --- a/clang/lib/CodeGen/CodeGenFunction.cpp +++ b/clang/lib/CodeGen/CodeGenFunction.cpp @@ -1529,16 +1529,14 @@ void CodeGenFunction::GenerateCode(GlobalDecl GD, llvm::Function *Fn, for (auto ORI : llvm::enumerate(OptReportHandler.GetInfo(FD))) { llvm::DiagnosticLocation DL = SourceLocToDebugLoc(ORI.value().KernelArgLoc); + std::string KAN = ORI.value().KernelArgName; llvm::OptimizationRemark Remark("sycl", "Region", DL, &Fn->getEntryBlock()); Remark << "Argument " << llvm::ore::NV("Argument", ORI.index()) << " for function kernel: " - << llvm::ore::NV(ORI.value().KernelArgName.empty() ? "&" : "") - << " " << Fn->getName() << "." - << llvm::ore::NV(ORI.value().KernelArgName.empty() - ? " " - : ORI.value().KernelArgName) - << "(" << ORI.value().KernelArgType << ")"; + << llvm::ore::NV(KAN.empty() ? "&" : "") << " " << Fn->getName() + << "." << llvm::ore::NV(KAN.empty() ? " " : KAN) << "(" + << ORI.value().KernelArgType << ")"; ORE.emit(Remark); } }