[SYCL] Add ITT annotation instructions (#3299)

This patch introduces InstrumentalAnnotationsPass pass that
adds ITT instrumentations for barrier and barrier-like functions
and atomic instructions. The pass is being included to
compiler flow only when "-fsycl-instrument-device-code" option
passed to clang driver.

Following annotations were added:
- __itt_offload_wi_start_wrapper
  Notify tools work-item execution has started
- __itt_offload_wi_resume_wrapper
  Notify tools work-item execution resumed (e.g. after barrier)
- __itt_offload_wi_finish_wrapper
  Notify tools work-item execution has finished
- __itt_offload_wg_barrier
  Notify tools work-item has reached a barrier
- __itt_offload_atomic_op_start
   Atomic operation markup
- __itt_offload_atomic_op_finish
  Atomic operation markup

Signed-off-by: Dmitry Sidorov <dmitry.sidorov@intel.com>

Author: MrSidims

.github/CODEOWNERS

@@ -114,3 +114,8 @@ SYCLLowerIR/ @kbobrovs @DenisBakhvalov
 esimd/ @kbobrovs @DenisBakhvalov
 sycl/include/CL/sycl/INTEL/esimd.hpp @kbobrovs @DenisBakhvalov
 sycl/doc/extensions/ExplicitSIMD/ @kbobrovs
+
+# ITT annotations
+llvm/lib/Transforms/Instrumentation/SPIRITTAnnotations.cpp @MrSidims
+llvm/include/llvm/Transforms/Instrumentation/SPIRITTAnnotations.h @MrSidims
+

clang/include/clang/Basic/CodeGenOptions.def

@@ -417,6 +417,9 @@ CODEGENOPT(PassByValueIsNoAlias, 1, 0)
 /// according to the field declaring type width.
 CODEGENOPT(AAPCSBitfieldWidth, 1, 1)
 
+// Whether to instrument SPIR device code with ITT annotations
+CODEGENOPT(SPIRITTAnnotations, 1, 0)
+
 #undef CODEGENOPT
 #undef ENUM_CODEGENOPT
 #undef VALUE_CODEGENOPT

clang/include/clang/Driver/Options.td

@@ -2418,6 +2418,10 @@ def fsycl_device_code_lower_esimd : Flag<["-"], "fsycl-device-code-lower-esimd">
   Flags<[CC1Option, CoreOption]>, HelpText<"Lower ESIMD-specific constructs">;
 def fno_sycl_device_code_lower_esimd : Flag<["-"], "fno-sycl-device-code-lower-esimd">,
   Flags<[CC1Option, CoreOption]>, HelpText<"Do not lower ESIMD-specific constructs">;
+def fsycl_instrument_device_code : Flag<["-"], "fsycl-instrument-device-code">,
+  Group<sycl_Group>, Flags<[CC1Option, CoreOption]>,
+  HelpText<"Add ITT instrumentation intrinsics calls">,
+  MarshallingInfoFlag<CodeGenOpts<"SPIRITTAnnotations">>;
 defm sycl_id_queries_fit_in_int: OptInFFlag<"sycl-id-queries-fit-in-int", "Assume", "Do not assume", " that SYCL ID queries fit within MAX_INT.", [CC1Option,CoreOption], LangOpts<"SYCLValueFitInMaxInt">>;
 def fsycl_use_bitcode : Flag<["-"], "fsycl-use-bitcode">,
   Flags<[CC1Option, CoreOption]>, HelpText<"Use LLVM bitcode instead of SPIR-V in fat objects">;

clang/lib/CodeGen/BackendUtil.cpp

@@ -72,6 +72,7 @@
 #include "llvm/Transforms/Instrumentation/InstrProfiling.h"
 #include "llvm/Transforms/Instrumentation/MemProfiler.h"
 #include "llvm/Transforms/Instrumentation/MemorySanitizer.h"
+#include "llvm/Transforms/Instrumentation/SPIRITTAnnotations.h"
 #include "llvm/Transforms/Instrumentation/SanitizerCoverage.h"
 #include "llvm/Transforms/Instrumentation/ThreadSanitizer.h"
 #include "llvm/Transforms/ObjCARC.h"
@@ -948,6 +949,16 @@ void EmitAssemblyHelper::EmitAssembly(BackendAction Action,
       LangOpts.EnableDAEInSpirKernels)
     PerModulePasses.add(createDeadArgEliminationSYCLPass());
 
+  // Add SPIRITTAnnotations pass to the pass manager if
+  // -fsycl-instrument-device-code option was passed. This option can be
+  // used only with spir triple.
+  if (CodeGenOpts.SPIRITTAnnotations) {
+    if (!llvm::Triple(TheModule->getTargetTriple()).isSPIR())
+      llvm::report_fatal_error(
+          "ITT annotations can only by added to a module with spir target");
+    PerModulePasses.add(createSPIRITTAnnotationsPass());
+  }
+
   switch (Action) {
   case Backend_EmitNothing:
     break;

clang/lib/Driver/ToolChains/Clang.cpp

@@ -5881,6 +5881,15 @@ void Clang::ConstructJob(Compilation &C, const JobAction &JA,
   if (Args.hasFlag(options::OPT_fsycl, options::OPT_fno_sycl, false))
     Args.AddLastArg(CmdArgs, options::OPT_sycl_std_EQ);
 
+  // Forward -fsycl-instrument-device-code option to cc1. This option can only
+  // be used with spir triple.
+  if (Arg *A = Args.getLastArg(options::OPT_fsycl_instrument_device_code)) {
+    if (!Triple.isSPIR())
+      D.Diag(diag::err_drv_unsupported_opt_for_target)
+          << A->getAsString(Args) << TripleStr;
+    CmdArgs.push_back("-fsycl-instrument-device-code");
+  }
+
   if (IsHIP) {
     if (Args.hasFlag(options::OPT_fhip_new_launch_api,
                      options::OPT_fno_hip_new_launch_api, true))

clang/test/CodeGenSYCL/kernel-simple-instrumentation.cpp

@@ -0,0 +1,21 @@
+/// Check if start/finish ITT annotations are being added during compilation of
+/// SYCL device code
+
+// RUN: %clang_cc1 -fsycl-is-device -fsycl-instrument-device-code -triple spir64-unknown-unknown-sycldevice -emit-llvm %s -o - | FileCheck %s
+
+// CHECK: kernel_function
+// CHECK-NEXT: entry:
+// CHECK-NEXT: call void @__itt_offload_wi_start_wrapper()
+// CHECK: call void @__itt_offload_wi_finish_wrapper()
+// CHECK-NEXT: ret void
+
+#include "Inputs/sycl.hpp"
+
+int main() {
+  cl::sycl::accessor<int, 1, cl::sycl::access::mode::read_write> accessorA;
+  cl::sycl::kernel_single_task<class kernel_function>(
+      [=]() {
+        accessorA.use();
+      });
+  return 0;
+}

clang/test/Driver/sycl-instrumentation.c

@@ -0,0 +1,14 @@
+/// Check that SPIR ITT instrumentation is disabled by default:
+// RUN: %clang -### %s 2>&1 \
+// RUN: | FileCheck -check-prefix=CHECK-DEFAULT %s
+// CHECK-DEFAULT-NOT: "-fsycl-instrument-device-code"
+
+/// Check if "fsycl_instrument_device_code" is passed to -cc1:
+// RUN: %clang -### -fsycl-instrument-device-code %s 2>&1 \
+// RUN: | FileCheck -check-prefix=CHECK-ENABLED %s
+// CHECK-ENABLED: "-cc1"{{.*}} "-fsycl-instrument-device-code"
+
+/// Check if "fsycl_instrument_device_code" usage with a non-spirv target
+/// results in an error.
+// RUN: %clang -### -fsycl-instrument-device-code --target=x86 %s 2>&1
+// expected-error{{unsupported option '-fsycl-instrument-device-code' for target 'x86_64-unknown-linux-gnu'}}

llvm/include/llvm/InitializePasses.h

@@ -429,6 +429,7 @@ void initializeStripSymbolsPass(PassRegistry&);
 void initializeStructurizeCFGLegacyPassPass(PassRegistry &);
 void initializeSYCLLowerWGScopeLegacyPassPass(PassRegistry &);
 void initializeSYCLLowerESIMDLegacyPassPass(PassRegistry &);
+void initializeSPIRITTAnnotationsLegacyPassPass(PassRegistry &);
 void initializeESIMDLowerLoadStorePass(PassRegistry &);
 void initializeESIMDLowerVecArgLegacyPassPass(PassRegistry &);
 void initializeTailCallElimPass(PassRegistry&);

llvm/include/llvm/LinkAllPasses.h

@@ -48,6 +48,7 @@
 #include "llvm/Transforms/InstCombine/InstCombine.h"
 #include "llvm/Transforms/Instrumentation.h"
 #include "llvm/Transforms/Instrumentation/BoundsChecking.h"
+#include "llvm/Transforms/Instrumentation/SPIRITTAnnotations.h"
 #include "llvm/Transforms/ObjCARC.h"
 #include "llvm/Transforms/Scalar.h"
 #include "llvm/Transforms/Scalar/GVN.h"
@@ -204,6 +205,7 @@ namespace {
       (void)llvm::createSYCLLowerESIMDPass();
       (void)llvm::createESIMDLowerLoadStorePass();
       (void)llvm::createESIMDLowerVecArgPass();
+      (void)llvm::createSPIRITTAnnotationsPass();
       std::string buf;
       llvm::raw_string_ostream os(buf);
       (void) llvm::createPrintModulePass(os);

llvm/include/llvm/Transforms/Instrumentation/SPIRITTAnnotations.h

@@ -0,0 +1,27 @@
+//===----- SPIRITTAnnotations.h - SPIR Instrumental Annotations Pass ------===//
+//
+// 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
+//
+//===----------------------------------------------------------------------===//
+//
+// A transformation pass which adds instrumental calls to annotate SPIR
+// synchronization instructions. This can be used for kernel profiling.
+//===----------------------------------------------------------------------===//
+
+#pragma once
+
+#include "llvm/IR/Module.h"
+#include "llvm/IR/PassManager.h"
+
+namespace llvm {
+
+class SPIRITTAnnotationsPass : public PassInfoMixin<SPIRITTAnnotationsPass> {
+public:
+  PreservedAnalyses run(Module &M, ModuleAnalysisManager &MAM);
+};
+
+ModulePass *createSPIRITTAnnotationsPass();
+
+} // namespace llvm

llvm/lib/Transforms/Instrumentation/CMakeLists.txt

@@ -15,6 +15,7 @@ add_llvm_component_library(LLVMInstrumentation
   PGOMemOPSizeOpt.cpp
   PoisonChecking.cpp
   SanitizerCoverage.cpp
+  SPIRITTAnnotations.cpp
   ValueProfileCollector.cpp
   ThreadSanitizer.cpp
   HWAddressSanitizer.cpp