diff --git a/clang/include/clang/Basic/Diagnostic.h b/clang/include/clang/Basic/Diagnostic.h index 3b915fb15a891..f389a39e84a54 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" @@ -294,6 +295,7 @@ class DiagnosticsEngine : public RefCountedBase { DiagnosticConsumer *Client = nullptr; std::unique_ptr Owner; SourceManager *SourceMgr = nullptr; + SyclOptReportHandler OptReportHandler; /// Mapping information for diagnostics. /// @@ -547,6 +549,12 @@ class DiagnosticsEngine : public RefCountedBase { LLVM_DUMP_METHOD void dump() const; LLVM_DUMP_METHOD void dump(StringRef DiagName) const; + /// 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/include/clang/Basic/SyclOptReportHandler.h b/clang/include/clang/Basic/SyclOptReportHandler.h new file mode 100644 index 0000000000000..af3ba5ae1c0a8 --- /dev/null +++ b/clang/include/clang/Basic/SyclOptReportHandler.h @@ -0,0 +1,57 @@ +//===---------------------- 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. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// +/// +/// \file +/// Defines clang::SyclOptReportHandler 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 { + std::string KernelArgName; + std::string KernelArgType; + SourceLocation KernelArgLoc; + + OptReportInfo(std::string ArgName, std::string ArgType, + SourceLocation ArgLoc) + : KernelArgName(std::move(ArgName)), KernelArgType(std::move(ArgType)), + KernelArgLoc(ArgLoc) {} + }; + 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) { + auto It = Map.find(FD); + assert(It != Map.end()); + return It->second; + } + + bool HasOptReportInfo(const FunctionDecl *FD) const { + 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..eb6e512c05a44 100644 --- a/clang/lib/CodeGen/CGSYCLRuntime.cpp +++ b/clang/lib/CodeGen/CGSYCLRuntime.cpp @@ -14,6 +14,8 @@ #include "CodeGenFunction.h" #include "clang/AST/Attr.h" #include "clang/AST/Decl.h" +#include "clang/Basic/SourceLocation.h" +#include "llvm/Analysis/OptimizationRemarkEmitter.h" #include "llvm/IR/Instructions.h" #include diff --git a/clang/lib/CodeGen/CodeGenAction.cpp b/clang/lib/CodeGen/CodeGenAction.cpp index 73e1ee7013dc4..8175574ff6a52 100644 --- a/clang/lib/CodeGen/CodeGenAction.cpp +++ b/clang/lib/CodeGen/CodeGenAction.cpp @@ -304,28 +304,12 @@ namespace clang { return; LLVMContext &Ctx = getModule()->getContext(); + 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); + // The diagnostic handler is now processed in OptRecordFileRAII. // The parallel_for_work_group legalization pass can emit calls to // builtins function. Definitions of those builtins can be provided in @@ -349,9 +333,6 @@ namespace clang { getModule(), Action, std::move(AsmOutStream)); Ctx.setDiagnosticHandler(std::move(OldDiagnosticHandler)); - - if (OptRecordFile) - OptRecordFile->keep(); } void HandleTagDeclDefinition(TagDecl *D) override { @@ -1046,8 +1027,51 @@ 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), OldDiagnosticHandler(Ctx.getDiagnosticHandler()) { + + CompilerInstance &CI = CGA.getCompilerInstance(); + CodeGenOptions &CodeGenOpts = CI.getCodeGenOpts(); + + 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/CodeGenFunction.cpp b/clang/lib/CodeGen/CodeGenFunction.cpp index 693743252f246..aabfb649174b1 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" @@ -1521,6 +1522,25 @@ 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().getSYCLOptReportHandler(); + if (OptReportHandler.HasOptReportInfo(FD)) { + llvm::OptimizationRemarkEmitter ORE(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(KAN.empty() ? "&" : "") << " " << Fn->getName() + << "." << llvm::ore::NV(KAN.empty() ? " " : KAN) << "(" + << ORI.value().KernelArgType << ")"; + ORE.emit(Remark); + } + } + // Save parameters for coroutine function. if (Body && isa_and_nonnull(Body)) for (const auto *ParamDecl : FD->parameters()) diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index c351fd343fe9b..b0459d97c01f0 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -1768,6 +1768,9 @@ class SyclKernelDeclCreator : public SyclKernelFieldHandler { void addParam(const FieldDecl *FD, QualType FieldTy) { ParamDesc newParamDesc = makeParamDesc(FD, FieldTy); + SemaRef.getDiagnostics().getSYCLOptReportHandler().AddKernelArgs( + KernelDecl, FD->getName().data(), FieldTy.getAsString(), + FD->getLocation()); addParam(newParamDesc, FieldTy); } @@ -1778,6 +1781,8 @@ class SyclKernelDeclCreator : public SyclKernelFieldHandler { StringRef Name = "_arg__base"; ParamDesc newParamDesc = makeParamDesc(SemaRef.getASTContext(), Name, FieldTy); + SemaRef.getDiagnostics().getSYCLOptReportHandler().AddKernelArgs( + KernelDecl, "", FieldTy.getAsString(), BS.getBaseTypeLoc()); addParam(newParamDesc, FieldTy); } // Add a parameter with specified name and type 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..fe71db3e833e3 --- /dev/null +++ b/clang/test/SemaSYCL/kernel-arg-opt-report.cpp @@ -0,0 +1,101 @@ +// 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" + +class second_base { +public: + int *e; +}; + +class InnerFieldBase { +public: + int d; +}; +class InnerField : public InnerFieldBase { + int c; +}; + +struct base { +public: + int b; + 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; + + void operator()() const { + } +}; + +int main() { + sycl::queue q; + + q.submit([&](cl::sycl::handler &cgh) { + derived f{}; + cgh.single_task(f); + }); + + return 0; +}