Merge remote-tracking branch 'public/sycl' into private/s-kanaev/assert-abort-impl

Author: Sergey Kanaev

CONTRIBUTING.md

@@ -66,11 +66,6 @@ There are 3 types of tests which are used for DPC++ toolchain validation:
 should not have hardware (e.g. GPU, FPGA, etc.) or external software
 dependencies (e.g. OpenCL, Level Zero, CUDA runtimes). All tests not following
 this approach should be moved to DPC++ end-to-end or SYCL-CTS tests.
-However, the tests for a feature under active development requiring atomic
-change for tests and product can be put to
-[sycl/test/on-device](../../sycl/test/on-device) temporarily. It is developer
-responsibility to move the tests to DPC++ E2E test suite or SYCL-CTS once
-the feature is stabilized.
 
     **Guidelines for adding DPC++ in-tree LIT tests (DPC++ Clang FE tests)**:
     - Use `sycl::` namespace instead of `cl::sycl::`

README.md

@@ -1,37 +1,30 @@
 # Intel Project for LLVM\* technology
 
-## Introduction
+Intel staging area for llvm.org contribution. Home for Intel LLVM-based projects:
 
-Intel staging area for llvm.org contribution.
-Home for Intel LLVM-based projects:
+## oneAPI Data Parallel C++ compiler
 
-* oneAPI Data Parallel C++ compiler - see
-  [sycl](https://github.com/intel/llvm/tree/sycl) branch. More information on
-  oneAPI and DPC++ is available at
-  [https://www.oneapi.com/](https://www.oneapi.com/)
-  * [![Linux Post Commit Checks](https://github.com/intel/llvm/workflows/Linux%20Post%20Commit%20Checks/badge.svg)](https://github.com/intel/llvm/actions?query=workflow%3A%22Linux+Post+Commit+Checks%22)
-    [![Generate Doxygen documentation](https://github.com/intel/llvm/workflows/Generate%20Doxygen%20documentation/badge.svg)](https://github.com/intel/llvm/actions?query=workflow%3A%22Generate+Doxygen+documentation%22)
-* Late-outline OpenMP\* and OpenMP\* Offload - see
-  [openmp](https://github.com/intel/llvm/tree/openmp) branch.
+[![](https://spec.oneapi.io/oneapi-logo-white-scaled.jpg)](https://www.oneapi.io/)
 
-## License
+See [sycl](https://github.com/intel/llvm/tree/sycl) branch and
+[DPC++ Documentation](https://intel.github.io/llvm-docs/).
 
-See [LICENSE.txt](sycl/LICENSE.TXT) for details.
+[![Linux Post Commit Checks](https://github.com/intel/llvm/workflows/Linux%20Post%20Commit%20Checks/badge.svg)](https://github.com/intel/llvm/actions?query=workflow%3A%22Linux+Post+Commit+Checks%22)
+[![Generate Doxygen documentation](https://github.com/intel/llvm/workflows/Generate%20Doxygen%20documentation/badge.svg)](https://github.com/intel/llvm/actions?query=workflow%3A%22Generate+Doxygen+documentation%22)
 
-## Contributing
+DPC++ is an open, cross-architecture language built upon the ISO C++ and Khronos
+SYCL\* standards. DPC++ extends these standards with a number of extensions,
+which can be found in [sycl/doc/extensions](sycl/doc/extensions) directory.
 
-See [CONTRIBUTING.md](CONTRIBUTING.md) for details.
+## Late-outline OpenMP\* and OpenMP\* Offload
+See [openmp](https://github.com/intel/llvm/tree/openmp) branch.
 
-## Sub-projects Documentation
+# License
 
-* oneAPI Data Parallel C++ compiler - See
-  [DPC++ Documentation](https://intel.github.io/llvm-docs/)
-* Late-outline OpenMP\* and OpenMP\* Offload - TBD
+See [LICENSE.txt](sycl/LICENSE.TXT) for details.
 
-## DPC++ extensions
+# Contributing
 
-DPC++ is an open, cross-architecture language built upon the ISO C++ and Khronos
-SYCL\* standards. DPC++ extends these standards with a number of extensions,
-which can be found in [sycl/doc/extensions](sycl/doc/extensions) directory.
+See [CONTRIBUTING.md](CONTRIBUTING.md) for details.
 
 *\*Other names and brands may be claimed as the property of others.*

buildbot/configure.py

@@ -61,6 +61,9 @@ def do_configure(args):
         if args.rocm_platform == 'AMD':
             llvm_targets_to_build += ';AMDGPU'
             libclc_targets_to_build += ';amdgcn--;amdgcn--amdhsa'
+
+            # The ROCm plugin for AMD uses lld for linking
+            llvm_enable_projects += ';lld'
         elif args.rocm_platform == 'NVIDIA' and not args.cuda:
             llvm_targets_to_build += ';NVPTX'
             libclc_targets_to_build += ';nvptx64--;nvptx64--nvidiacl'

clang/include/clang/Driver/Options.td

@@ -2628,14 +2628,6 @@ def fsycl_device_code_split_EQ : Joined<["-"], "fsycl-device-code-split=">,
 def fsycl_device_code_split : Flag<["-"], "fsycl-device-code-split">, Alias<fsycl_device_code_split_EQ>,
   AliasArgs<["auto"]>, Flags<[CC1Option, CoreOption]>,
   HelpText<"Perform SYCL device code split in the 'auto' mode, i.e. use heuristic to distribute device code across modules">;
-def fsycl_device_code_split_esimd : Flag<["-"], "fsycl-device-code-split-esimd">,
-  Flags<[CC1Option, CoreOption]>, HelpText<"Split SYCL and ESIMD kernels into separate modules">;
-def fno_sycl_device_code_split_esimd : Flag<["-"], "fno-sycl-device-code-split-esimd">,
-  Flags<[CC1Option, CoreOption]>, HelpText<"Don't split SYCL and ESIMD kernels into separate modules">;
-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">,

clang/lib/Driver/Driver.cpp

@@ -2932,6 +2932,18 @@ bool Driver::checkForSYCLDefaultDevice(Compilation &C,
   if (Args.hasArg(options::OPT_fno_sycl_link_spirv))
     return false;
 
+  // Do not do the check if the default device is passed in -fsycl-targets
+  // or if -fsycl-targets isn't passed (that implies default device)
+  if (const Arg *A = Args.getLastArg(options::OPT_fsycl_targets_EQ)) {
+    for (const char *Val : A->getValues()) {
+      llvm::Triple TT(C.getDriver().MakeSYCLDeviceTriple(Val));
+      if (TT.isSPIR() && TT.getSubArch() == llvm::Triple::NoSubArch)
+        // Default triple found
+        return false;
+    }
+  } else if (!Args.hasArg(options::OPT_fintelfpga))
+    return false;
+
   SmallVector<const char *, 16> AllArgs(getLinkerArgs(C, Args, true));
   for (StringRef Arg : AllArgs) {
     if (hasSYCLDefaultSection(C, Arg))
@@ -5492,6 +5504,16 @@ void Driver::BuildActions(Compilation &C, DerivedArgList &Args,
         break;
       }
 
+      // When performing -fsycl based compilations and generating dependency
+      // information, perform a specific dependency generation compilation which
+      // is not based on the source + footer compilation.
+      if (Phase == phases::Preprocess && Args.hasArg(options::OPT_fsycl) &&
+          Args.hasArg(options::OPT_M_Group) &&
+          !Args.hasArg(options::OPT_fno_sycl_use_footer)) {
+        Actions.push_back(
+            C.MakeAction<PreprocessJobAction>(Current, types::TY_Dependencies));
+      }
+
       // FIXME: Should we include any prior module file outputs as inputs of
       // later actions in the same command line?
 
@@ -5856,6 +5878,9 @@ void Driver::BuildJobs(Compilation &C) const {
   // we are also generating .o files. So we allow more than one output file in
   // this case as well.
   //
+  // Preprocessing job performed for -fsycl enabled compilation specifically
+  // for dependency generation (TY_Dependencies)
+  //
   if (FinalOutput) {
     unsigned NumOutputs = 0;
     unsigned NumIfsOutputs = 0;
@@ -5866,7 +5891,10 @@ void Driver::BuildJobs(Compilation &C) const {
              A->getKind() == clang::driver::Action::CompileJobClass &&
              0 == NumIfsOutputs++) ||
             (A->getKind() == Action::BindArchClass && A->getInputs().size() &&
-             A->getInputs().front()->getKind() == Action::IfsMergeJobClass)))
+             A->getInputs().front()->getKind() == Action::IfsMergeJobClass) ||
+            (A->getKind() == Action::PreprocessJobClass &&
+             A->getType() == types::TY_Dependencies &&
+             C.getArgs().hasArg(options::OPT_fsycl))))
         ++NumOutputs;
 
     if (NumOutputs > 1) {

clang/lib/Driver/ToolChains/Clang.cpp

@@ -1203,6 +1203,12 @@ void Clang::AddPreprocessingOptions(Compilation &C, const JobAction &JA,
     C.getDriver().addFPGATempDepFile(DepFile, BaseName);
   };
 
+  // Do not add dependency generation information when compiling the source +
+  // footer combination.  The dependency generation is done in a separate
+  // compile step so we can retain original source information.
+  if (ContainsAppendFooterAction(&JA))
+    ArgM = nullptr;
+
   if (ArgM) {
     // Determine the output location.
     const char *DepFile;
@@ -1216,6 +1222,12 @@ void Clang::AddPreprocessingOptions(Compilation &C, const JobAction &JA,
             DepFile, Clang::getBaseInputName(Args, Inputs[0]));
     } else if (Output.getType() == types::TY_Dependencies) {
       DepFile = Output.getFilename();
+      if (!ContainsAppendFooterAction(&JA) && Args.hasArg(options::OPT_fsycl) &&
+          !Args.hasArg(options::OPT_fno_sycl_use_footer) &&
+          !JA.isDeviceOffloading(Action::OFK_SYCL))
+        // Name the dependency file for the specific dependency generation
+        // step created for the integration footer enabled compilation.
+        DepFile = getDependencyFileName(Args, Inputs);
     } else if (!ArgMD) {
       DepFile = "-";
     } else if (IsIntelFPGA && JA.isDeviceOffloading(Action::OFK_SYCL)) {
@@ -8933,13 +8945,8 @@ void SYCLPostLink::ConstructJob(Compilation &C, const JobAction &JA,
     // Symbol file and specialization constant info generation is mandatory -
     // add options unconditionally
     addArgs(CmdArgs, TCArgs, {"-symbols"});
-    // By default we split SYCL and ESIMD kernels into separate modules
-    if (TCArgs.hasFlag(options::OPT_fsycl_device_code_split_esimd,
-                       options::OPT_fno_sycl_device_code_split_esimd, true))
-      addArgs(CmdArgs, TCArgs, {"-split-esimd"});
-    if (TCArgs.hasFlag(options::OPT_fsycl_device_code_lower_esimd,
-                       options::OPT_fno_sycl_device_code_lower_esimd, true))
-      addArgs(CmdArgs, TCArgs, {"-lower-esimd"});
+    addArgs(CmdArgs, TCArgs, {"-split-esimd"});
+    addArgs(CmdArgs, TCArgs, {"-lower-esimd"});
   }
   addArgs(CmdArgs, TCArgs,
           {StringRef(getSYCLPostLinkOptimizationLevel(TCArgs))});

clang/test/Driver/sycl-int-footer.cpp

@@ -67,3 +67,35 @@
 // FOOTER_PATH: append-file{{.*}} "--append=dummy_dir{{(/|\\\\)}}{{.*}}-footer-{{.*}}.h"
 // FOOTER_PATH-SAME: "--output=dummy_dir{{(/|\\\\)}}[[APPENDEDSRC:.+\.cpp]]"
 // FOOTER_PATH: clang{{.*}} "-x" "c++" "dummy_dir{{(/|\\\\)}}[[APPENDEDSRC]]"
+
+/// Check behaviors for dependency generation
+// RUN:  %clangxx -fsycl -MD -c %s -### 2>&1 \
+// RUN:   | FileCheck -check-prefix DEP_GEN %s
+// DEP_GEN:  clang{{.*}} "-Eonly"
+// DEP_GEN-SAME: "-dependency-file"
+// DEP_GEN-SAME: "-MT"
+// DEP_GEN-SAME: "-x" "c++" "[[INPUTFILE:.+\.cpp]]"
+// DEP_GEN: append-file{{.*}} "[[INPUTFILE]]"
+// DEP_GEN-NOT: clang{{.*}} "-dependency-file"
+
+/// Dependency generation phases
+// RUN:  %clangxx -target x86_64-unknown-linux-gnu -fsycl -MD -c %s -ccc-print-phases 2>&1 \
+// RUN:   | FileCheck -check-prefix DEP_GEN_PHASES %s
+// DEP_GEN_PHASES: 0: input, "[[INPUTFILE:.+\.cpp]]", c++, (host-sycl)
+// DEP_GEN_PHASES: 1: preprocessor, {0}, dependencies
+// DEP_GEN_PHASES: 2: input, "[[INPUTFILE]]", c++, (device-sycl)
+// DEP_GEN_PHASES: 3: preprocessor, {2}, c++-cpp-output, (device-sycl)
+// DEP_GEN_PHASES: 4: compiler, {3}, ir, (device-sycl)
+// DEP_GEN_PHASES: 5: offload, "device-sycl (spir64-unknown-unknown-sycldevice)" {4}, ir
+// DEP_GEN_PHASES: 6: append-footer, {0}, c++, (host-sycl)
+// DEP_GEN_PHASES: 7: preprocessor, {6}, c++-cpp-output, (host-sycl)
+// DEP_GEN_PHASES: 8: offload, "host-sycl (x86_64-unknown-linux-gnu)" {7}, "device-sycl (spir64-unknown-unknown-sycldevice)" {4}, c++-cpp-output
+// DEP_GEN_PHASES: 9: compiler, {8}, ir, (host-sycl)
+// DEP_GEN_PHASES: 10: backend, {9}, assembler, (host-sycl)
+// DEP_GEN_PHASES: 11: assembler, {10}, object, (host-sycl)
+// DEP_GEN_PHASES: 12: clang-offload-bundler, {5, 11}, object, (host-sycl)
+
+/// Allow for -o and preprocessing
+// RUN:  %clangxx -fsycl -MD -c %s -o dummy -### 2>&1 \
+// RUN:   | FileCheck -check-prefix DEP_GEN_OUT_ERROR %s
+// DEP_GEN_OUT_ERROR-NOT: cannot specify -o when generating multiple output files

clang/test/Driver/sycl-offload-intelfpga.cpp

@@ -322,11 +322,11 @@
 /// -fintelfpga dependency file check with host .d enabled
 // RUN: %clangxx -### -MMD -fsycl -fintelfpga -Xshardware %t-1.cpp %t-2.cpp 2>&1 \
 // RUN:  | FileCheck -check-prefix=CHK-FPGA-DEP-FILES-HOST %s
-// CHK-FPGA-DEP-FILES-HOST: clang{{.*}} "-dependency-file" "[[INPUT1:.+\.d]]" "-MT" "{{.*}}.o"
-// CHK-FPGA-DEP-FILES-HOST: clang{{.*}} "-dependency-file" "[[INPUT2:.+\.d]]" "-MT" "{{.*}}.o"
+// CHK-FPGA-DEP-FILES-HOST: clang{{.*}} "-triple" "spir64_fpga-unknown-unknown-sycldevice"{{.*}} "-dependency-file" "[[INPUT1:.+\.d]]" "-MT" "{{.*}}.o"
+// CHK-FPGA-DEP-FILES-HOST: clang{{.*}} "-triple" "spir64_fpga-unknown-unknown-sycldevice"{{.*}} "-dependency-file" "[[INPUT2:.+\.d]]" "-MT" "{{.*}}.o"
 // CHK-FPGA-DEP-FILES-HOST: aoc{{.*}} "-dep-files={{.*}}[[INPUT1]],{{.*}}[[INPUT2]]"
-// CHK-FPGA-DEP-FILES-HOST: clang{{.*}} "-fsycl-is-host"{{.*}} "-dependency-file"
-// CHK-FPGA-DEP-FILES-HOST: clang{{.*}} "-fsycl-is-host"{{.*}} "-dependency-file"
+// CHK-FPGA-DEP-FILES-HOST: clang{{.*}} "-fsycl-is-host"{{.*}}
+// CHK-FPGA-DEP-FILES-HOST: clang{{.*}} "-fsycl-is-host"{{.*}}
 
 /// -fintelfpga dependency file generation test to object
 // RUN: %clangxx -### -fsycl -fintelfpga -target x86_64-unknown-linux-gnu %t-1.cpp %t-2.cpp -c 2>&1 \
@@ -348,7 +348,7 @@
 // RUN:  | FileCheck -check-prefixes=CHK-FPGA-DEP-FILES3,CHK-FPGA-DEP-FILES3-LIN %s
 // RUN: %clang_cl -### --target=x86_64-pc-windows-msvc -fsycl -fintelfpga %t-1.cpp -c -Fodummy.obj 2>&1 \
 // RUN:  | FileCheck -check-prefixes=CHK-FPGA-DEP-FILES3,CHK-FPGA-DEP-FILES3-WIN %s
-// CHK-FPGA-DEP-FILES3: clang{{.*}} "-dependency-file" "[[OUTPUT:.+\.d]]"
+// CHK-FPGA-DEP-FILES3: clang{{.*}} "-triple" "spir64_fpga-unknown-unknown-sycldevice"{{.*}} "-dependency-file" "[[OUTPUT:.+\.d]]"
 // CHK-FPGA-DEP-FILES3-LIN: clang-offload-bundler{{.*}} "-type=o" "-targets=sycl-spir64_fpga-unknown-unknown-sycldevice,host-x86_64-unknown-linux-gnu,sycl-fpga_dep" {{.*}} "-inputs={{.*}}.bc,{{.*}}.o,[[OUTPUT]]"
 // CHK-FPGA-DEP-FILES3-WIN: clang-offload-bundler{{.*}} "-type=o" "-targets=sycl-spir64_fpga-unknown-unknown-sycldevice,host-x86_64-pc-windows-msvc,sycl-fpga_dep" {{.*}} "-inputs={{.*}}.bc,{{.*}}.obj,[[OUTPUT]]"
 

clang/test/Driver/sycl-offload-with-split.c

@@ -299,25 +299,15 @@
 // Check ESIMD device code split.
 // RUN:   %clang    -### -fsycl %s 2>&1 | FileCheck %s -check-prefixes=CHK-ESIMD-SPLIT
 // RUN:   %clang_cl -### -fsycl %s 2>&1 | FileCheck %s -check-prefixes=CHK-ESIMD-SPLIT
-// RUN:   %clang    -### -fsycl -fsycl-device-code-split-esimd %s 2>&1 | FileCheck %s -check-prefixes=CHK-ESIMD-SPLIT
-// RUN:   %clang_cl -### -fsycl -fsycl-device-code-split-esimd %s 2>&1 | FileCheck %s -check-prefixes=CHK-ESIMD-SPLIT
-// RUN:   %clang    -### -fsycl -fno-sycl-device-code-split-esimd %s 2>&1 | FileCheck %s -check-prefixes=CHK-NO-ESIMD-SPLIT
-// RUN:   %clang_cl -### -fsycl -fno-sycl-device-code-split-esimd %s 2>&1 | FileCheck %s -check-prefixes=CHK-NO-ESIMD-SPLIT
 // RUN:   %clang    -### -fsycl -fintelfpga %s 2>&1 | FileCheck %s -check-prefixes=CHK-ESIMD-SPLIT
 // RUN:   %clang    -### -fsycl -fsycl-targets=spir64_fpga-unknown-unknown-sycldevice %s 2>&1 | FileCheck %s -check-prefixes=CHK-ESIMD-SPLIT
 // RUN:   %clang_cl -### -fsycl -fintelfpga %s 2>&1 | FileCheck %s -check-prefixes=CHK-ESIMD-SPLIT
 // CHK-ESIMD-SPLIT: sycl-post-link{{.*}} "-split-esimd"
-// CHK-NO-ESIMD-SPLIT-NOT: sycl-post-link{{.*}} "-split-esimd"
 
 // Check lowering of ESIMD device code.
 // RUN:   %clang    -### -fsycl %s 2>&1 | FileCheck %s -check-prefixes=CHK-ESIMD-LOWER
 // RUN:   %clang_cl -### -fsycl %s 2>&1 | FileCheck %s -check-prefixes=CHK-ESIMD-LOWER
-// RUN:   %clang    -### -fsycl -fsycl-device-code-lower-esimd %s 2>&1 | FileCheck %s -check-prefixes=CHK-ESIMD-LOWER
-// RUN:   %clang_cl -### -fsycl -fsycl-device-code-lower-esimd %s 2>&1 | FileCheck %s -check-prefixes=CHK-ESIMD-LOWER
-// RUN:   %clang    -### -fsycl -fno-sycl-device-code-lower-esimd %s 2>&1 | FileCheck %s -check-prefixes=CHK-NO-ESIMD-LOWER
-// RUN:   %clang_cl -### -fsycl -fno-sycl-device-code-lower-esimd %s 2>&1 | FileCheck %s -check-prefixes=CHK-NO-ESIMD-LOWER
 // RUN:   %clang    -### -fsycl -fintelfpga %s 2>&1 | FileCheck %s -check-prefixes=CHK-ESIMD-LOWER
 // RUN:   %clang    -### -fsycl -fsycl-targets=spir64_fpga-unknown-unknown-sycldevice %s 2>&1 | FileCheck %s -check-prefixes=CHK-ESIMD-LOWER
 // RUN:   %clang_cl -### -fsycl -fintelfpga %s 2>&1 | FileCheck %s -check-prefixes=CHK-ESIMD-LOWER
 // CHK-ESIMD-LOWER: sycl-post-link{{.*}} "-lower-esimd"
-// CHK-NO-ESIMD-LOWER-NOT: sycl-post-link{{.*}} "-lower-esimd"

clang/test/Driver/sycl-offload.cpp

@@ -96,3 +96,10 @@
 // RUN:    | FileCheck -check-prefix NO_IMPLIED_DEVICE %s
 // NO_IMPLIED_DEVICE: clang-offload-bundler{{.*}} "-type=o" "-targets=sycl-spir64-unknown-unknown-sycldevice"{{.*}} "-check-section"
 // NO_IMPLIED_DEVICE-NOT: clang-offload-bundler{{.*}} "-targets={{.*}}spir64-unknown-unknown-sycldevice{{.*}}" "-unbundle"
+
+/// Passing in the default triple should allow for -Xsycl-target options
+// RUN:  %clangxx -### -target x86_64-unknown-linux-gnu -fsycl -fsycl-targets=spir64 -Xsycl-target-backend=spir64 -DFOO -Xsycl-target-linker=spir64 -DFOO2 %S/Inputs/SYCL/objlin64.o 2>&1 \
+// RUN:    | FileCheck -check-prefixes=SYCL_TARGET_OPT %s
+// RUN:  %clangxx -### -target x86_64-unknown-linux-gnu -fsycl -Xsycl-target-backend=spir64 -DFOO -Xsycl-target-linker=spir64 -DFOO2 %S/Inputs/SYCL/objlin64.o 2>&1 \
+// RUN:    | FileCheck -check-prefixes=SYCL_TARGET_OPT %s
+// SYCL_TARGET_OPT: clang-offload-wrapper{{.*}} "-compile-opts=-DFOO" "-link-opts=-DFOO2"