Merge branch 'sycl' into aaron/stopReportingFPExtensions

Author: aarongreig

.github/workflows/sycl-linux-build.yml

@@ -212,6 +212,12 @@ jobs:
       if: always() && !cancelled() && contains(inputs.changes, 'libdevice')
       run: |
         cmake --build $GITHUB_WORKSPACE/build --target check-libdevice
+    - name: Check E2E test requirements
+      if: always() && !cancelled() && !contains(inputs.changes, 'sycl')
+      run: |
+        # TODO consider moving this to Dockerfile.
+        export LD_LIBRARY_PATH=/usr/local/cuda/compat/:/usr/local/cuda/lib64:$LD_LIBRARY_PATH
+        LIT_OPTS="--allow-empty-runs" LIT_FILTER="e2e_test_requirements" cmake --build $GITHUB_WORKSPACE/build --target check-sycl
     - name: Install
       if: ${{ always() && !cancelled() && steps.build.conclusion == 'success' }}
       # TODO replace utility installation with a single CMake target

buildbot/configure.py

@@ -64,7 +64,7 @@ def do_configure(args):
 
     sycl_enable_xpti_tracing = "ON"
     xpti_enable_werror = "OFF"
-    llvm_enable_zstd = "OFF"
+    llvm_enable_zstd = "ON"
 
     if sys.platform != "darwin":
         sycl_enabled_backends.append("level_zero")
@@ -134,8 +134,6 @@ def do_configure(args):
 
         # For clang-format, clang-tidy and code coverage
         llvm_enable_projects += ";clang-tools-extra;compiler-rt"
-        # Build with zstd enabled on CI.
-        llvm_enable_zstd = "ON"
         if sys.platform != "darwin":
             # libclc is required for CI validation
             libclc_enabled = True

clang/include/clang/Driver/Action.h

@@ -698,6 +698,14 @@ class OffloadWrapperJobAction : public JobAction {
   // Get the compilation step setting.
   bool getCompileStep() const { return CompileStep; }
 
+  // Set the individual wrapping setting.  This is used to tell the wrapper job
+  // action that the wrapping (and subsequent compile step) should be done
+  // with for-each instead of using -batch.
+  void setWrapIndividualFiles() { WrapIndividualFiles = true; }
+
+  // Get the individual wrapping setting.
+  bool getWrapIndividualFiles() const { return WrapIndividualFiles; }
+
   // Set the offload kind for the current wrapping job action.  Default usage
   // is to use the kind of the current toolchain.
   void setOffloadKind(OffloadKind SetKind) { Kind = SetKind; }
@@ -707,6 +715,7 @@ class OffloadWrapperJobAction : public JobAction {
 
 private:
   bool CompileStep = true;
+  bool WrapIndividualFiles = false;
   OffloadKind Kind = OFK_None;
 };
 

clang/lib/Driver/Driver.cpp

@@ -5536,9 +5536,58 @@ class OffloadingActionBuilder final {
                   BundlingActions, types::TY_Object);
               if (auto *OWA = dyn_cast<OffloadWrapperJobAction>(DeviceAction))
                 OWA->setOffloadKind(Action::OFK_Host);
+              // The Backend compilation step performed here is being done for
+              // creating FPGA archives.  The possible split binaries after
+              // sycl-post-link need to be individually wrapped as opposed to
+              // being passed into the clang-offload-wrapper via a table and
+              // using the -batch option - effectively creating a single
+              // binary.  The resulting archive created from -fsycl-link should
+              // not contain the singular binary, but should be individual
+              // binaries to be consumed later by either the -fsycl-link=image
+              // device compilation step or being linked into the final exe.
+              //
+              // Typical compile flow:
+              //               .bc
+              //                |
+              //     sycl-post-link -split=kernel
+              //                |
+              //       +--------+--------+
+              //       |        |        |
+              //     split1   split2   split3
+              //       |        |        |
+              // llvm-spirv llvm-spirv llvm-spirv
+              //       |        |        |
+              //     ocloc    ocloc    ocloc
+              //       |        |        |
+              //       +--------+--------+
+              //                |
+              //     clang-offload-wrapper -batch
+              //                |
+              //               .o
+              //
+              // Individual wrap compile flow:
+              //               .bc
+              //                |
+              //     sycl-post-link -split=kernel
+              //                |
+              //       +--------+--------+
+              //       |        |        |
+              //     split1   split2   split3
+              //       |        |        |
+              // llvm-spirv llvm-spirv llvm-spirv
+              //       |        |        |
+              //     ocloc    ocloc    ocloc
+              //       |        |        |
+              //     wrap     wrap     wrap
+              //       |        |        |
+              //      .o       .o       .o
+              //
               Action *CompiledDeviceAction =
-                  C.MakeAction<OffloadWrapperJobAction>(WrapperItems,
-                                                        types::TY_Object);
+                  C.MakeAction<OffloadWrapperJobAction>(FPGAAOTAction,
+                                                        types::TY_Tempfilelist);
+              if (auto *OWA =
+                      dyn_cast<OffloadWrapperJobAction>(CompiledDeviceAction))
+                OWA->setWrapIndividualFiles();
               addDeps(CompiledDeviceAction, TC, BoundArch);
             }
             addDeps(DeviceAction, TC, BoundArch);
@@ -5812,6 +5861,9 @@ class OffloadingActionBuilder final {
           };
 
           Action *ExtractIRFilesAction = createExtractIRFilesAction();
+          // Device binaries that are individually wrapped when creating an
+          // FPGA Archive.
+          ActionList FPGAArchiveWrapperInputs;
 
           if (IsNVPTX || IsAMDGCN) {
             JobAction *FinAction =
@@ -5897,6 +5949,7 @@ class OffloadingActionBuilder final {
                 FileTableTformJobAction::COL_CODE,
                 FileTableTformJobAction::COL_CODE);
             WrapperInputs.push_back(ReplaceFilesAction);
+            FPGAArchiveWrapperInputs.push_back(BuildCodeAction);
           }
           if (SkipWrapper) {
             // Wrapper step not requested.
@@ -5931,8 +5984,11 @@ class OffloadingActionBuilder final {
                 if (auto *OWA = dyn_cast<OffloadWrapperJobAction>(DeviceAction))
                   OWA->setOffloadKind(Action::OFK_Host);
                 Action *CompiledDeviceAction =
-                    C.MakeAction<OffloadWrapperJobAction>(WrapperInputs,
-                                                          types::TY_Object);
+                    C.MakeAction<OffloadWrapperJobAction>(
+                        FPGAArchiveWrapperInputs, types::TY_Tempfilelist);
+                if (auto *OWA =
+                        dyn_cast<OffloadWrapperJobAction>(CompiledDeviceAction))
+                  OWA->setWrapIndividualFiles();
                 addDeps(CompiledDeviceAction, TC, nullptr);
               }
               addDeps(DeviceAction, TC, nullptr);
@@ -6461,7 +6517,7 @@ class OffloadingActionBuilder final {
           if (GpuInitHasErrors)
             return true;
 
-          int GenIndex = 0;
+          size_t GenIndex = 0;
           // Fill SYCLTargetInfoList
           for (auto &TT : SYCLTripleList) {
             auto TCIt = llvm::find_if(

clang/lib/Driver/ToolChains/Clang.cpp

@@ -10263,8 +10263,18 @@ void OffloadWrapper::ConstructJob(Compilation &C, const JobAction &JA,
     const InputInfo &I = Inputs[0];
     assert(I.isFilename() && "Invalid input.");
 
-    if (I.getType() == types::TY_Tempfiletable ||
-        I.getType() == types::TY_Tempfilelist || IsEmbeddedIR)
+    // TODO: The embedded compilation step after the wrapping step restricts
+    // the ability to control the 'for each' methodology used when performing
+    // device code splitting.  We set the individual wrap behavior when we know
+    // the wrapping and compile step should be done individually.  Ideally this
+    // would be controlled at the JobAction creation, but we cannot do that
+    // until the compilation of the wrap is it's own JobAction.
+    bool IndividualWrapCompile = WrapperJob.getWrapIndividualFiles();
+    const InputInfo TempOutput(types::TY_LLVM_BC, WrapperFileName,
+                               WrapperFileName);
+    if (!IndividualWrapCompile &&
+        (I.getType() == types::TY_Tempfiletable ||
+         I.getType() == types::TY_Tempfilelist || IsEmbeddedIR))
       // Input files are passed via the batch job file table.
       WrapperArgs.push_back(C.getArgs().MakeArgString("-batch"));
     WrapperArgs.push_back(C.getArgs().MakeArgString(I.getFilename()));
@@ -10273,7 +10283,17 @@ void OffloadWrapper::ConstructJob(Compilation &C, const JobAction &JA,
         JA, *this, ResponseFileSupport::None(),
         TCArgs.MakeArgString(getToolChain().GetProgramPath(getShortName())),
         WrapperArgs, std::nullopt);
-    C.addCommand(std::move(Cmd));
+
+    if (IndividualWrapCompile) {
+      // When wrapping FPGA device binaries for FPGA archives, create individual
+      // wrapped and compiled entries for the archive.
+      StringRef ParallelJobs =
+          C.getArgs().getLastArgValue(options::OPT_fsycl_max_parallel_jobs_EQ);
+      clang::driver::tools::SYCL::constructLLVMForeachCommand(
+          C, JA, std::move(Cmd), Inputs, TempOutput, this, "", "bc",
+          ParallelJobs);
+    } else
+      C.addCommand(std::move(Cmd));
 
     if (WrapperCompileEnabled) {
       // TODO Use TC.SelectTool().
@@ -10296,9 +10316,19 @@ void OffloadWrapper::ConstructJob(Compilation &C, const JobAction &JA,
       SmallString<128> ClangPath(C.getDriver().Dir);
       llvm::sys::path::append(ClangPath, "clang");
       const char *Clang = C.getArgs().MakeArgString(ClangPath);
-      C.addCommand(std::make_unique<Command>(JA, *this,
-                                             ResponseFileSupport::None(), Clang,
-                                             ClangArgs, std::nullopt));
+      auto PostWrapCompileCmd =
+          std::make_unique<Command>(JA, *this, ResponseFileSupport::None(),
+                                    Clang, ClangArgs, std::nullopt);
+      if (IndividualWrapCompile) {
+        StringRef ParallelJobs = C.getArgs().getLastArgValue(
+            options::OPT_fsycl_max_parallel_jobs_EQ);
+        InputInfoList Inputs;
+        Inputs.push_back(TempOutput);
+        clang::driver::tools::SYCL::constructLLVMForeachCommand(
+            C, JA, std::move(PostWrapCompileCmd), Inputs, Output, this, "",
+            "bc", ParallelJobs);
+      } else
+        C.addCommand(std::move(PostWrapCompileCmd));
     }
     return;
   } // end of SYCL flavor of offload wrapper command creation

clang/lib/Driver/ToolChains/HIPUtility.cpp

@@ -324,7 +324,10 @@ void HIP::constructHIPFatbinCommand(Compilation &C, const JobAction &JA,
       Args.MakeArgString(std::string("-output=").append(Output));
   BundlerArgs.push_back(BundlerOutputArg);
 
-  addOffloadCompressArgs(Args, BundlerArgs);
+  // For SYCL, the compression is occurring during the wrapping step, so we do
+  // not want to do additional compression here.
+  if (!JA.isDeviceOffloading(Action::OFK_SYCL))
+    addOffloadCompressArgs(Args, BundlerArgs);
 
   const char *Bundler = Args.MakeArgString(
       T.getToolChain().GetProgramPath("clang-offload-bundler"));

clang/lib/Driver/ToolChains/SYCL.cpp

@@ -1780,6 +1780,16 @@ void SYCLToolChain::AddImpliedTargetArgs(const llvm::Triple &Triple,
     ArgStringList TargArgs;
     Args.AddAllArgValues(TargArgs, options::OPT_Xs, options::OPT_Xs_separate);
     Args.AddAllArgValues(TargArgs, options::OPT_Xsycl_backend);
+    // For -Xsycl-target-backend=<triple> the triple value is used to push
+    // specific options to the matching device compilation using that triple.
+    // Scrutinize this to make sure we are only checking the values needed
+    // for the current device compilation.
+    for (auto *A : Args) {
+      if (!A->getOption().matches(options::OPT_Xsycl_backend_EQ))
+        continue;
+      if (getDriver().MakeSYCLDeviceTriple(A->getValue()) == Triple)
+        TargArgs.push_back(A->getValue(1));
+    }
     // Check for any -device settings.
     std::string DevArg;
     if (IsJIT || Device == "pvc" || hasPVCDevice(TargArgs, DevArg)) {

clang/lib/Sema/SemaSYCL.cpp

@@ -3167,7 +3167,7 @@ class SyclKernelDeclCreator : public SyclKernelFieldHandler {
 //        // code
 //      }
 //
-//      [[intel::reqd_sub_group_size(4)]] void operator()(sycl::id<1> id) const
+//      [[sycl::reqd_sub_group_size(4)]] void operator()(sycl::id<1> id) const
 //      {
 //        // code
 //      }

clang/lib/Sema/SemaSYCLDeclAttr.cpp

@@ -132,6 +132,13 @@ void SemaSYCL::checkDeprecatedSYCLAttributeSpelling(const ParsedAttr &A,
     return;
   }
 
+  // Additionally, diagnose deprecated [[intel::reqd_sub_group_size]] spelling
+  if (A.getKind() == ParsedAttr::AT_IntelReqdSubGroupSize && A.getScopeName() &&
+      A.getScopeName()->isStr("intel")) {
+    diagnoseDeprecatedAttribute(A, "sycl", "reqd_sub_group_size");
+    return;
+  }
+
   // Diagnose SYCL 2020 spellings in later SYCL modes.
   if (getLangOpts().getSYCLVersion() >= LangOptions::SYCL_2020) {
     // All attributes in the cl vendor namespace are deprecated in favor of a

clang/test/CodeGenSYCL/kernel-op-calls.cpp

@@ -11,7 +11,7 @@ class Functor1 {
 public:
  Functor1(){}
 
- [[intel::reqd_sub_group_size(4)]] void operator()(sycl::id<1> id) const {}
+ [[sycl::reqd_sub_group_size(4)]] void operator()(sycl::id<1> id) const {}
 
  [[sycl::work_group_size_hint(1, 2, 3)]] void operator()(sycl::id<2> id) const {}