[SYCL][NATIVECPU] Enable source-based code coverage in Native CPU (#1…

…5073)

Supports [clang's source-based code
coverage](https://clang.llvm.org/docs/SourceBasedCodeCoverage.html) to
enable code coverage testing of SYCL applications via the Native CPU
SYCL target.

Clang's `-fprofile-instr-generate -fcoverage-mapping` options can now be
used with the `native_cpu` SYCL target to compile/instrument host and
device code, enabling 'llvm-cov' to render a coverage report after
running the SYCL application (see also updated documentation in this
PR).

Subsequent PRs will enable in NativeCPU more of the currently
unsupported options for device compilation, also for performance
profiling.

**Details and explanations for the changes in this PR:**
This PR tests coverage options on the existing NativeCPU vector-add test
by adding an additional invocation with previously disabled options
`-fprofile-instr-generate -fcoverage-mapping -mllvm
-system-headers-coverage`. Enabling these options on device code caused
an [assert in the upstream clang profiling code generation
tools](https://github.com/intel/llvm/blob/b023d407862bd853ba5881c34985f99d039d856c/clang/lib/CodeGen/CoverageMappingGen.cpp#L960)
due to the invalid source location on the AST for the implicitly
generated kernel body, specifically the compound statement containing
the kernel body. This PR honors this upstream clang assert by replacing
the invalid source location in the compound statement with the source
location of the kernel body. Using this now valid source location
maintains the location (of the kernel caller function) currently tested
by [non-upstream-llvm lit test
`CodeGenSYCL/debug-info-srcpos-kernel.cpp`](https://github.com/intel/llvm/blob/sycl/clang/test/CodeGenSYCL/debug-info-srcpos-kernel.cpp),
but exposed an issue that led to a change in behavior in
[non-llvm-upstream lit test
`SemaSYCL/kernel-arg-opt-report.cpp`](https://github.com/intel/llvm/blob/sycl/clang/test/SemaSYCL/kernel-arg-opt-report.cpp),
which was due to the previously invalid source location causing the
compiler to skip code to set the current location. To restore the
original behavior of this test (checking for the location of the kernel
functor, as opposed to the kernel caller function) this PR temporarily
(and only for the purpose of generating the report) sets the current
location to the location of the kernel argument using the upstream clang
utility
[clang::CodeGen::ApplyDebugLocation](https://github.com/intel/llvm/blob/b023d407862bd853ba5881c34985f99d039d856c/clang/lib/CodeGen/CGDebugInfo.h#L860).

---------

Co-authored-by: Michael Toguchi <[email protected]>

clang/lib/CodeGen/CodeGenFunction.cpp

@@ -1789,6 +1789,9 @@ void CodeGenFunction::GenerateCode(GlobalDecl GD, llvm::Function *Fn,
     if (SyclOptReport.HasOptReportInfo(FD)) {
       llvm::OptimizationRemarkEmitter ORE(Fn);
       for (auto ORI : llvm::enumerate(SyclOptReport.GetInfo(FD))) {
+        // Temporarily apply arg location to ensure SourceLocToDebugLoc
+        // picks up the expected file.
+        ApplyDebugLocation TempApplyLoc(*this, ORI.value().KernelArgLoc);
         llvm::DiagnosticLocation DL =
             SourceLocToDebugLoc(ORI.value().KernelArgLoc);
         StringRef NameInDesc = ORI.value().KernelArgDescName;

clang/lib/Driver/ToolChains/SYCL.cpp

@@ -1618,6 +1618,23 @@ static std::vector<OptSpecifier> getUnsupportedOpts(void) {
   return UnsupportedOpts;
 }
 
+// Currently supported options by SYCL NativeCPU device compilation
+static inline bool SupportedByNativeCPU(const SYCLToolChain &TC,
+                                        const OptSpecifier &Opt) {
+  if (!TC.IsSYCLNativeCPU)
+    return false;
+
+  switch (Opt.getID()) {
+  case options::OPT_fcoverage_mapping:
+  case options::OPT_fno_coverage_mapping:
+  case options::OPT_fprofile_instr_generate:
+  case options::OPT_fprofile_instr_generate_EQ:
+  case options::OPT_fno_profile_instr_generate:
+    return true;
+  }
+  return false;
+}
+
 SYCLToolChain::SYCLToolChain(const Driver &D, const llvm::Triple &Triple,
                              const ToolChain &HostTC, const ArgList &Args)
     : ToolChain(D, Triple, Args), HostTC(HostTC),
@@ -1629,6 +1646,9 @@ SYCLToolChain::SYCLToolChain(const Driver &D, const llvm::Triple &Triple,
   // Diagnose unsupported options only once.
   for (OptSpecifier Opt : getUnsupportedOpts()) {
     if (const Arg *A = Args.getLastArg(Opt)) {
+      // Native CPU can support options unsupported by other targets.
+      if (SupportedByNativeCPU(*this, Opt))
+        continue;
       // All sanitizer options are not currently supported, except
       // AddressSanitizer
       if (A->getOption().getID() == options::OPT_fsanitize_EQ &&
@@ -1669,6 +1689,9 @@ SYCLToolChain::TranslateArgs(const llvm::opt::DerivedArgList &Args,
     bool Unsupported = false;
     for (OptSpecifier UnsupportedOpt : getUnsupportedOpts()) {
       if (Opt.matches(UnsupportedOpt)) {
+        // NativeCPU should allow most normal cpu options.
+        if (SupportedByNativeCPU(*this, Opt.getID()))
+          continue;
         if (Opt.getID() == options::OPT_fsanitize_EQ &&
             A->getValues().size() == 1) {
           std::string SanitizeVal = A->getValue();

clang/lib/Sema/SemaSYCL.cpp

@@ -3625,8 +3625,11 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler {
     BodyStmts.insert(BodyStmts.end(), FinalizeStmts.begin(),
                      FinalizeStmts.end());
 
+    SourceLocation LL = NewBody ? NewBody->getBeginLoc() : SourceLocation();
+    SourceLocation LR = NewBody ? NewBody->getEndLoc() : SourceLocation();
+
     return CompoundStmt::Create(SemaSYCLRef.getASTContext(), BodyStmts,
-                                FPOptionsOverride(), {}, {});
+                                FPOptionsOverride(), LL, LR);
   }
 
   void annotateHierarchicalParallelismAPICalls() {

clang/test/Driver/sycl-native-cpu.cpp

@@ -24,3 +24,9 @@
 
 // RUN: %clangxx -fsycl -fsycl-targets=spir64 %s -### 2>&1 | FileCheck -check-prefix=CHECK-NONATIVECPU %s
 // CHECK-NONATIVECPU-NOT: "-D" "__SYCL_NATIVE_CPU__"
+
+// Checking that coverage testing options are accepted by native_cpu, and that device and host compilation invocations receive the same options
+// RUN: %clangxx -fsycl -fsycl-targets=native_cpu -Werror -fno-profile-instr-generate -fprofile-instr-generate -fno-coverage-mapping -fcoverage-mapping -### %s 2>&1 | FileCheck %s --check-prefix=CHECK_COV_INVO
+// CHECK_COV_INVO:{{.*}}clang{{.*}}-fsycl-is-device{{.*}}"-fsycl-is-native-cpu" "-D" "__SYCL_NATIVE_CPU__"{{.*}}"-fprofile-instrument=clang"{{.*}}"-fcoverage-mapping" "-fcoverage-compilation-dir={{.*}}"
+// CHECK_COV_INVO:{{.*}}clang{{.*}}"-fsycl-is-host"{{.*}}"-fprofile-instrument=clang"{{.*}}"-fcoverage-mapping" "-fcoverage-compilation-dir={{.*}}"
+

clang/test/SemaSYCL/kernel_functor_location.cpp

@@ -0,0 +1,22 @@
+// RUN: %clang_cc1 -fsycl-is-device -ast-dump %s | FileCheck %s
+//
+// Checks that the compound statement of the implicitly generated kernel body
+// has a valid source location (containing "line"). Previously this location
+// was invalid containing "<<invalid sloc>>" which causes asserts in the
+// llvm profiling tools.
+
+#include "Inputs/sycl.hpp"
+
+struct Functor {
+  void operator()() const {}
+};
+
+// CHECK: FunctionDecl {{.*}} _ZTS7Functor 'void ()'
+// CHECK-NEXT: |-CompoundStmt {{.*}} <{{.*}}line{{.*}}>
+
+int main() {
+
+  sycl::queue().submit([&](sycl::handler &cgh) {
+    cgh.single_task(Functor{});
+  });
+}

sycl/doc/design/SYCLNativeCPU.md

@@ -1,10 +1,10 @@
 # SYCL Native CPU
 
-The SYCL Native CPU flow aims at treating the host CPU as a "first class citizen", providing a SYCL implementation that targets CPUs of various different architectures, with no other dependencies than DPC++ itself, while bringing performances comparable to state-of-the-art CPU backends.
+The SYCL Native CPU flow aims at treating the host CPU as a "first class citizen", providing a SYCL implementation that targets CPUs of various different architectures, with no other dependencies than DPC++ itself, while bringing performances comparable to state-of-the-art CPU backends. SYCL Native CPU also provides some initial/experimental support for LLVM's [source-based code coverage tools](https://clang.llvm.org/docs/SourceBasedCodeCoverage.html) (see also section [Code coverage](#code-coverage)).
 
 # Compiler and runtime options
 
-The SYCL Native CPU flow is enabled by setting `native_cpu` as a `sycl-target` (please note that currently doing so overrides any other SYCL target specified in the compiler invocation):
+The SYCL Native CPU flow is enabled by setting `native_cpu` as a `sycl-target`:
 
 ```
 clang++ -fsycl -fsycl-targets=native_cpu <input> -o <output>
@@ -28,9 +28,16 @@ clang++ <device-ir> -o <device-o>
 clang++ -L<sycl-lib-path> -lsycl <device-o> <host-o> -o <output>
 ```
 
+Note that SYCL Native CPU co-exists alongside the other SYCL targets. For example, the following command line builds SYCL code simultaneously for SYCL Native CPU and for OpenCL.
+
+```
+clang++ -fsycl -fsycl-targets=native_cpu,spir64 <input> -o <output>
+```
+The application can then run on either SYCL target by setting the DPC++ `ONEAPI_DEVICE_SELECTOR` environment variable accordingly.
+
 ## Configuring DPC++ with SYCL Native CPU
 
-SYCL Native CPU needs to be enabled explictly when configuring DPC++, using `--native_cpu`, e.g.
+SYCL Native CPU needs to be enabled explicitly when configuring DPC++, using `--native_cpu`, e.g.
 
 ```
 python buildbot/configure.py \
@@ -86,7 +93,19 @@ Whole Function Vectorization is enabled by default, and can be controlled throug
 * `-mllvm -sycl-native-cpu-no-vecz`: disable Whole Function Vectorization.
 * `-mllvm -sycl-native-cpu-vecz-width`: sets the vector width to the specified value, defaults to 8.
 
-For more details on how the Whole Function Vectorizer is integrated for SYCL Native CPU, refer to the [Technical details[(#technical-details) section.
+For more details on how the Whole Function Vectorizer is integrated for SYCL Native CPU, refer to the [Technical details](#technical-details) section.
+
+# Code coverage
+
+SYCL Native CPU has experimental support for LLVM's source-based [code coverage](https://clang.llvm.org/docs/SourceBasedCodeCoverage.html). This enables coverage testing across device and host code.
+Example usage:
+
+```bash
+clang.exe -fsycl -fsycl-targets=native_cpu  -fprofile-instr-generate -fcoverage-mapping %fname% -o vector-add.exe
+.\vector-add.exe
+llvm-profdata merge -sparse default.profraw -o foo.profdata
+llvm-cov show .\vector-add.exe -instr-profile=foo.profdata
+```
 
 ## Ongoing work
 
@@ -95,7 +114,7 @@ For more details on how the Whole Function Vectorizer is integrated for SYCL Nat
 * Subgroup support
 * Performance optimizations
 
-### Please note that Windows support is temporarily disabled due to some implementation details, it will be reinstantiated soon.
+### Please note that Windows is partially supported but temporarily disabled due to some implementation details, it will be re-enabled soon.
 
 # Technical details
 
@@ -140,13 +159,13 @@ entry:
 }
 ```
 
-For the SYCL Native CPU target, the device compiler is in charge of materializing the SPIRV builtins (such as `@__spirv_BuiltInGlobalInvocationId`), so that they can be correctly updated by the runtime when executing the kernel. This is performed by the [PrepareSYCLNativeCPU pass](https://github.com/intel/llvm/blob/sycl/llvm/lib/SYCLLowerIR/PrepareSYCLNativeCPU.cpp).
+For the SYCL Native CPU target, the device compiler is in charge of materializing the SPIRV builtins (such as `@__spirv_BuiltInGlobalInvocationId`), so that they can be correctly updated by the runtime when executing the kernel. This is performed by the [PrepareSYCLNativeCPU pass](https://github.com/intel/llvm/blob/sycl/llvm/lib/SYCLNativeCPUUtils/PrepareSYCLNativeCPU.cpp).
 The PrepareSYCLNativeCPUPass also emits a `subhandler` function, which receives the kernel arguments from the SYCL runtime (packed in a vector), unpacks them, and forwards only the used ones to the actual kernel. 
 
 
 ## PrepareSYCLNativeCPU Pass
 
-This pass will add a pointer to a `nativecpu_state` struct as kernel argument to all the kernel functions, and it will replace all the uses of SPIRV builtins with the return value of appropriately defined functions, which will read the requested information from the `__nativecpu_state` struct. The `__nativecpu_state` struct and the builtin functions are defined in [native_cpu.hpp](https://github.com/intel/llvm/blob/sycl/sycl/include/sycl/detail/native_cpu.hpp).
+This pass will add a pointer to a `native_cpu::state` struct as kernel argument to all the kernel functions, and it will replace all the uses of SPIRV builtins with the return value of appropriately defined functions, which will read the requested information from the `native_cpu::state` struct. The `native_cpu::state` struct is defined in the [native_cpu UR adapter](https://github.com/oneapi-src/unified-runtime/blob/main/source/adapters/native_cpu/nativecpu_state.hpp) and the builtin functions are defined in the [native_cpu device library](https://github.com/intel/llvm/blob/sycl/libdevice/nativecpu_utils.cpp).
 
 
 The resulting IR is:
@@ -188,11 +207,11 @@ entry:
 }
 ```
 
-As you can see, the `subhandler` steals the kernel's function name, and receives two pointer arguments: the first one points to the kernel arguments from the SYCL runtime, and the second one to the `__nativecpu_state` struct.
+As you can see, the `subhandler` steals the kernel's function name, and receives two pointer arguments: the first one points to the kernel arguments from the SYCL runtime, and the second one to the `nativecpu::state` struct.
 
 ## Handling barriers 
 
-On SYCL Native CPU, calls to `__spirv_ControlBarrier` are handled using the `WorkItemLoopsPass` from the oneAPI Construction Kit. This pass handles barriers by splitting the kernel between calls calls to `__spirv_ControlBarrier`, and creating a wrapper that runs the subkernels over the local range. In order to correctly interface to the oneAPI Construction Kit pass pipeline, SPIRV builtins are converted to `mux` builtins (used by the OCK) by the `ConvertToMuxBuiltinsSYCLNativeCPUPass`.
+On SYCL Native CPU, calls to `__spirv_ControlBarrier` are handled using the `WorkItemLoopsPass` from the oneAPI Construction Kit. This pass handles barriers by splitting the kernel between calls to `__spirv_ControlBarrier`, and creating a wrapper that runs the subkernels over the local range. In order to correctly interface to the oneAPI Construction Kit pass pipeline, SPIRV builtins are defined in the device library to call the corresponding `mux` builtins (used by the OCK).
 
 ## Vectorization
 

sycl/test/native_cpu/vector-add.cpp

@@ -15,6 +15,13 @@
 // RUN: %clangxx -fsycl -fsycl-targets=native_cpu -mllvm -sycl-native-cpu-vecz-width=4 %s -g -o %t-vec
 // RUN: env ONEAPI_DEVICE_SELECTOR="native_cpu:cpu" %t-vec
 
+// Ensure coverage options work in the compiler invocations.
+// For builds with asserts enabled we also need to pass the option
+// -mllvm -system-headers-coverage
+// We need to also check if clang-rt is built and then run the executable and
+// verify the (profiling) outputs.
+// RUN: %clangxx -fsycl -fsycl-targets=native_cpu %s -fprofile-instr-generate -fcoverage-mapping -mllvm -system-headers-coverage -c -o %t
+
 #include <sycl/sycl.hpp>
 
 #include <array>