Generate an opt report of kernel arguments.

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"
@@ -547,6 +548,8 @@ class DiagnosticsEngine : public RefCountedBase<DiagnosticsEngine> {
   LLVM_DUMP_METHOD void dump() const;
   LLVM_DUMP_METHOD void dump(StringRef DiagName) const;
 
+  SyclOptReportHandler OptReportHandler;
+
   const IntrusiveRefCntPtr<DiagnosticIDs> &getDiagnosticIDs() const {
     return Diags;
   }

clang/include/clang/Basic/SyclOptReportHandler.h

@@ -0,0 +1,57 @@
+//===---------------------- 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<const FunctionDecl *, SmallVector<OptReportInfo, 4>> Map;
+
+public:
+  void AddKernelArgs(const FunctionDecl *FD, StringRef ArgName,
+                     std::string ArgType, SourceLocation ArgLoc) {
+    Map[FD].emplace_back(ArgName, ArgType, ArgLoc);
+  }
+  SmallVector<OptReportInfo, 4> &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

clang/lib/CodeGen/CGSYCLRuntime.cpp

@@ -14,7 +14,10 @@
 #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 <assert.h>
 
 using namespace clang;

clang/lib/CodeGen/CodeGenAction.cpp

@@ -304,28 +304,12 @@ namespace clang {
         return;
 
       LLVMContext &Ctx = getModule()->getContext();
+
       std::unique_ptr<DiagnosticHandler> OldDiagnosticHandler =
           Ctx.getDiagnosticHandler();
       Ctx.setDiagnosticHandler(std::make_unique<ClangDiagnosticHandler>(
         CodeGenOpts, this));
-
-      Expected<std::unique_ptr<llvm::ToolOutputFile>> 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<llvm::ToolOutputFile> 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,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<llvm::ToolOutputFile> OptRecordFile;
+  std::unique_ptr<DiagnosticHandler> 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<ClangDiagnosticHandler>(CodeGenOpts, &BC));
+
+    Expected<std::unique_ptr<llvm::ToolOutputFile>> 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;
   }

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().OptReportHandler;
+  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());
+      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 << ")";
+      ORE.emit(Remark);
+    }
+  }
+
   // Generate the body of the function.
   PGO.assignRegionCounters(GD, CurFn);
   if (isa<CXXDestructorDecl>(FD))
@@ -2618,8 +2638,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());

clang/lib/Sema/SemaSYCL.cpp

@@ -1768,6 +1768,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());
     addParam(newParamDesc, FieldTy);
   }
 
@@ -1778,6 +1780,8 @@ class SyclKernelDeclCreator : public SyclKernelFieldHandler {
     StringRef Name = "_arg__base";
     ParamDesc newParamDesc =
         makeParamDesc(SemaRef.getASTContext(), Name, FieldTy);
+    SemaRef.getDiagnostics().OptReportHandler.AddKernelArgs(
+        KernelDecl, "", FieldTy.getAsString(), BS.getBaseTypeLoc());
     addParam(newParamDesc, FieldTy);
   }
   // Add a parameter with specified name and type

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;
+}