Merge from 'sycl' to 'sycl-web' (20 commits)

Author: iclsrc

.github/workflows/sycl_precommit.yml

@@ -27,11 +27,13 @@ jobs:
     container:
       image: ghcr.io/intel/llvm/sycl_ubuntu2004_nightly:no-drivers
     steps:
+    - name: 'PR commits + 1'
+      run: echo "PR_FETCH_DEPTH=$(( ${{ github.event.pull_request.commits }} + 1 ))" >> "${GITHUB_ENV}"
     - uses: actions/checkout@v2
       with:
         ref: ${{ github.event.pull_request.head.sha }}
         persist-credentials: false
-        fetch-depth: 2
+        fetch-depth: ${{ env.PR_FETCH_DEPTH }}
     - name: Run clang-format
       uses: ./devops/actions/clang-format
 

clang/include/clang/Basic/Attr.td

@@ -2388,6 +2388,18 @@ def SYCLIntelFPGANofusion : StmtAttr {
   let Documentation = [SYCLIntelFPGANofusionAttrDocs];
 }
 
+def SYCLIntelFPGAMaxReinvocationDelay : StmtAttr {
+  let Spellings = [CXX11<"intel", "max_reinvocation_delay">];
+  let Subjects = SubjectList<[ForStmt, CXXForRangeStmt, WhileStmt, DoStmt],
+                             ErrorDiag, "'for', 'while', and 'do' statements">;
+  let Args = [ExprArgument<"NExpr">];
+  let LangOpts = [SYCLIsDevice, SilentlyIgnoreSYCLIsHost];
+  let IsStmtDependent = 1;
+  let Documentation = [SYCLIntelFPGAMaxReinvocationDelayAttrDocs];
+}
+def : MutualExclusions<[SYCLIntelFPGADisableLoopPipelining,
+                        SYCLIntelFPGAMaxReinvocationDelay]>;
+
 def IntelFPGALocalNonConstVar : SubsetSubject<Var,
                               [{S->hasLocalStorage() &&
                                 S->getKind() != Decl::ImplicitParam &&

clang/include/clang/Basic/AttrDocs.td

@@ -3318,7 +3318,7 @@ disables pipelining of the loop or function data path, causing the loop
 or function to be executed serially. Cannot be used on the same loop or
 function, or in conjunction with ``max_interleaving``,
 ``speculated_iterations``, ``max_concurrency``, ``initiation_interval``,
-or ``ivdep``.
+``ivdep``, or ``max_reinvocation_delay``.
 
 .. code-block:: c++
 
@@ -3451,6 +3451,31 @@ loop should not be fused with any adjacent loop.
   }];
 }
 
+def SYCLIntelFPGAMaxReinvocationDelayAttrDocs : Documentation {
+  let Category = DocCatVariable;
+  let Heading = "intel::max_reinvocation_delay";
+  let Content = [{
+This attribute applies to a loop. Specifies the maximum number of cycles allowed
+on the delay between the launch of the last iteration of a loop invocation and
+the launch of the first iteration of a new loop invocation. Parameter N is
+mandatory, and is a positive integer. Cannot be used on the same loop in
+conjunction with disable_loop_pipelining.
+
+.. code-block:: c++
+
+  void foo() {
+    int var = 0;
+    [[intel::max_reinvocation_delay(1)]]
+    for (int i = 0; sycl::log10((float)(x)) < 10; i++) var++;
+  }
+
+  template<int N>
+  void bar() {
+    [[intel::max_reinvocation_delay(N)]] for(;;) { }
+  }
+  }];
+}
+
 def SYCLIntelLoopFuseDocs : Documentation {
   let Category = DocCatFunction;
   let Heading = "loop_fuse, loop_fuse_independent";

clang/include/clang/Basic/DiagnosticDriverKinds.td

@@ -345,6 +345,8 @@ def warn_drv_sycl_offload_target_duplicate : Warning<
 def warn_drv_sycl_target_missing : Warning<
   "linked binaries do not contain expected '%0' target; found targets: '%1'">,
   InGroup<SyclTarget>;
+def err_drv_multiple_target_with_forced_target : Error<
+  "multiple target usage with '%0' is not supported with '%1'">;
 def err_drv_failed_to_deduce_target_from_arch : Error<
   "failed to deduce triple for target architecture '%0'; specify the triple "
   "using '-fopenmp-targets' and '-Xopenmp-target' instead.">;

clang/include/clang/Basic/DiagnosticSemaKinds.td

@@ -7221,9 +7221,11 @@ def warn_format_nonliteral : Warning<
   InGroup<FormatNonLiteral>, DefaultIgnore;
 
 def err_sycl_device_global_incorrect_scope : Error<
-  "'device_global' variables must be static or declared at namespace scope">;
+  "'device_global' variable must be a static data member or declared in global or namespace scope">;
 def err_sycl_device_global_not_publicly_accessible: Error<
-  "'device_global' member variable %0 is not publicly accessible from namespace scope">;
+  "'device_global' member variable %0 should be publicly accessible from namespace scope">;
+def err_sycl_device_global_array : Error<
+  "'device_global' array is not allowed">;
 
 def err_unexpected_interface : Error<
   "unexpected interface name %0: expected expression">;

clang/include/clang/Driver/Options.td

@@ -2860,6 +2860,10 @@ def fsycl_link_targets_EQ : CommaJoined<["-"], "fsycl-link-targets=">,
   Flags<[NoXarchOption, CC1Option, CoreOption, Deprecated]>,
   HelpText<"Specify comma-separated list of triples SYCL offloading targets "
            "to produce linked device images (deprecated)">;
+def fsycl_force_target_EQ : Joined<["-"], "fsycl-force-target=">,
+  Flags<[NoXarchOption, CoreOption]>,
+  HelpText<"Force the usage of the given triple when extracting device code "
+           "from any given objects on the command line">;
 def fsycl_device_code_split_EQ : Joined<["-"], "fsycl-device-code-split=">,
    Flags<[CC1Option, CoreOption]>, HelpText<"Perform SYCL device code split: per_kernel (device code module is "
   "created for each SYCL kernel) | per_source (device code module is created for each source (translation unit)) | off (no device code split). | auto (use heuristic to select the best way of splitting device code). "

clang/include/clang/Sema/Sema.h

@@ -2296,6 +2296,9 @@ class Sema final {
                                              Expr *E);
   SYCLIntelFPGALoopCoalesceAttr *
   BuildSYCLIntelFPGALoopCoalesceAttr(const AttributeCommonInfo &CI, Expr *E);
+  SYCLIntelFPGAMaxReinvocationDelayAttr *
+  BuildSYCLIntelFPGAMaxReinvocationDelayAttr(const AttributeCommonInfo &CI, 
+                                             Expr *E);
 
   bool CheckQualifiedFunctionForTypeId(QualType T, SourceLocation Loc);
 

clang/lib/CodeGen/CGLoopInfo.cpp

@@ -611,6 +611,15 @@ MDNode *LoopInfo::createMetadata(
                             llvm::Type::getInt32Ty(Ctx), VC.second))};
     LoopProperties.push_back(MDNode::get(Ctx, Vals));
   }
+  
+  if (Attrs.SYCLMaxReinvocationDelayNCycles) {
+    Metadata *Vals[] = {
+        MDString::get(Ctx, "llvm.loop.intel.max_reinvocation_delay.count"),
+        ConstantAsMetadata::get(
+          ConstantInt::get(llvm::Type::getInt32Ty(Ctx),
+                           *Attrs.SYCLMaxReinvocationDelayNCycles))};
+    LoopProperties.push_back(MDNode::get(Ctx, Vals));
+  }
 
   LoopProperties.insert(LoopProperties.end(), AdditionalLoopProperties.begin(),
                         AdditionalLoopProperties.end());
@@ -645,6 +654,7 @@ void LoopAttributes::clear() {
   SYCLMaxInterleavingNInvocations.reset();
   SYCLSpeculatedIterationsNIterations.reset();
   SYCLIntelFPGAVariantCount.clear();
+  SYCLMaxReinvocationDelayNCycles.reset();
   UnrollCount = 0;
   UnrollAndJamCount = 0;
   VectorizeEnable = LoopAttributes::Unspecified;
@@ -681,6 +691,7 @@ LoopInfo::LoopInfo(BasicBlock *Header, const LoopAttributes &Attrs,
       !Attrs.SYCLMaxInterleavingNInvocations &&
       !Attrs.SYCLSpeculatedIterationsNIterations &&
       Attrs.SYCLIntelFPGAVariantCount.empty() && Attrs.UnrollCount == 0 &&
+      !Attrs.SYCLMaxReinvocationDelayNCycles &&
       Attrs.UnrollAndJamCount == 0 && !Attrs.PipelineDisabled &&
       Attrs.PipelineInitiationInterval == 0 &&
       Attrs.VectorizePredicateEnable == LoopAttributes::Unspecified &&
@@ -1012,6 +1023,9 @@ void LoopInfoStack::push(BasicBlock *Header, clang::ASTContext &Ctx,
   // emitted
   // For attribute nofusion:
   // 'llvm.loop.fusion.disable' metadata will be emitted
+  // For attribute max_reinvocation_delay:
+  // n - 'llvm.loop.intel.max_reinvocation_delay.count, i32 n' metadata will be
+  // emitted
   for (const auto *A : Attrs) {
     if (const auto *IntelFPGAIVDep = dyn_cast<SYCLIntelFPGAIVDepAttr>(A))
       addSYCLIVDepInfo(Header->getContext(), IntelFPGAIVDep->getSafelenValue(),
@@ -1076,6 +1090,14 @@ void LoopInfoStack::push(BasicBlock *Header, clang::ASTContext &Ctx,
 
     if (isa<SYCLIntelFPGANofusionAttr>(A))
       setSYCLNofusionEnable();
+
+    if (const auto *IntelFPGAMaxReinvocationDelay =
+            dyn_cast<SYCLIntelFPGAMaxReinvocationDelayAttr>(A)) {
+      const auto *CE = cast<ConstantExpr>(
+          IntelFPGAMaxReinvocationDelay->getNExpr());
+      llvm::APSInt ArgVal = CE->getResultAsAPSInt();
+      setSYCLMaxReinvocationDelayNCycles(ArgVal.getSExtValue());
+    }
   }
 
   setMustProgress(MustProgress);

clang/lib/CodeGen/CGLoopInfo.h

@@ -134,6 +134,9 @@ struct LoopAttributes {
   /// Value for llvm.loop.intel.speculated.iterations.count metadata.
   llvm::Optional<unsigned> SYCLSpeculatedIterationsNIterations;
 
+  // Value for llvm.loop.intel.max_reinvocation_delay metadata.
+  llvm::Optional<unsigned> SYCLMaxReinvocationDelayNCycles;
+
   /// llvm.unroll.
   unsigned UnrollCount;
 
@@ -410,6 +413,11 @@ class LoopInfoStack {
   /// Set no progress for the next loop pushed.
   void setMustProgress(bool P) { StagedAttrs.MustProgress = P; }
 
+  /// Set value of max reinvocation delay for the next loop pushed.
+  void setSYCLMaxReinvocationDelayNCycles(unsigned C) {
+    StagedAttrs.SYCLMaxReinvocationDelayNCycles = C;
+  }
+
 private:
   /// Returns true if there is LoopInfo on the stack.
   bool hasInfo() const { return !Active.empty(); }

clang/lib/Driver/Driver.cpp

@@ -785,6 +785,8 @@ static bool addSYCLDefaultTriple(Compilation &C,
   /// Returns true if a triple is added to SYCLTriples, false otherwise
   if (!C.getDriver().isSYCLDefaultTripleImplied())
     return false;
+  if (C.getInputArgs().hasArg(options::OPT_fsycl_force_target_EQ))
+    return false;
   for (const auto &SYCLTriple : SYCLTriples) {
     if (SYCLTriple.getSubArch() == llvm::Triple::NoSubArch &&
         SYCLTriple.isSPIR())
@@ -1057,6 +1059,14 @@ void Driver::CreateOffloadingDeviceToolChains(Compilation &C,
       C.getInputArgs().getLastArg(options::OPT_fsycl_device_code_split_EQ),
       {"per_kernel", "per_source", "auto", "off"});
 
+  Arg *SYCLForceTarget =
+      getArgRequiringSYCLRuntime(options::OPT_fsycl_force_target_EQ);
+  if (SYCLForceTarget) {
+    StringRef Val(SYCLForceTarget->getValue());
+    llvm::Triple TT(MakeSYCLDeviceTriple(Val));
+    if (!isValidSYCLTriple(TT))
+      Diag(clang::diag::err_drv_invalid_sycl_target) << Val;
+  }
   bool HasSYCLTargetsOption = SYCLTargets || SYCLLinkTargets || SYCLAddTargets;
   llvm::StringMap<StringRef> FoundNormalizedTriples;
   llvm::SmallVector<llvm::Triple, 4> UniqueSYCLTriplesVec;
@@ -1066,6 +1076,15 @@ void Driver::CreateOffloadingDeviceToolChains(Compilation &C,
     Arg *SYCLTargetsValues = SYCLTargets ? SYCLTargets : SYCLLinkTargets;
     if (SYCLTargetsValues) {
       if (SYCLTargetsValues->getNumValues()) {
+
+        // Multiple targets are currently not supported when using
+        // -fsycl-force-target as the bundler does not allow for multiple
+        // outputs of the same target.
+        if (SYCLForceTarget && SYCLTargetsValues->getNumValues() > 1)
+          Diag(clang::diag::err_drv_multiple_target_with_forced_target)
+              << SYCLTargetsValues->getAsString(C.getInputArgs())
+              << SYCLForceTarget->getAsString(C.getInputArgs());
+
         for (StringRef Val : SYCLTargetsValues->getValues()) {
           llvm::Triple TT(MakeSYCLDeviceTriple(Val));
           if (!isValidSYCLTriple(TT)) {