Merge remote-tracking branch 'intel_llvm/sycl' into syclTypeAttr

Author: elizabethandrews

.github/workflows/sycl_precommit.yml

@@ -8,6 +8,7 @@ on:
     paths-ignore:
     - '.github/ISSUE_TEMPLATE/**'
     - '.github/CODEOWNERS'
+    - '.github/workflows/sycl_update_gpu_driver.yml'
     - 'devops/containers/**'
     - 'devops/scripts/install_drivers.sh'
     - 'devops/scripts/install_build_tools.sh'
@@ -26,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
 

.github/workflows/sycl_update_gpu_driver.yml

@@ -18,41 +18,13 @@ jobs:
     - name: Create Pull Request
       env:
         BRANCH: ci/update_gpu_driver-linux-${{ env.NEW_DRIVER_VERSION }}
-        LLVMBOT_TOKEN: ${{ secrets.LLVM_MAIN_SYNC_BBSYCL_TOKEN }}
-        GITHUB_TOKEN: ${{ secrets.GITHUB_TOKEN }}
+        GITHUB_TOKEN: ${{ secrets.LLVM_MAIN_SYNC_BBSYCL_TOKEN }}
       run: |
         cd $GITHUB_WORKSPACE
-        # Set fake identity to fulfil git requirements
         git config --global user.name "GitHub Actions"
         git config --global user.email "[email protected]"
         git checkout -B $BRANCH
         git add -u
-        git commit -m "[GHA] Uplift GPU RT version for Linux CI" || exit 0   # exit if commit is empty
-        git push https://[email protected]/${{ github.repository }} ${BRANCH}
-        gh pr create --head $BRANCH --title "[GHA] Uplift GPU RT version for Linux CI" --body "Uplift GPU RT version for Linux to $NEW_DRIVER_VERSION"
-
-  update_driver_linux_staging:
-    runs-on: ubuntu-latest
-    if: github.repository == 'intel/llvm'
-    steps:
-    - uses: actions/checkout@v2
-    - name: Update dependencies file
-      run: |
-        version="$(python3 devops/scripts/update_drivers.py linux_staging)"
-        echo 'NEW_DRIVER_VERSION='$version >> $GITHUB_ENV
-    - name: Update sycl Branch
-      env:
-        BRANCH: ci/update_gpu_driver-linux_staging-${{ env.NEW_DRIVER_VERSION }}
-        LLVMBOT_TOKEN: ${{ secrets.LLVM_MAIN_SYNC_BBSYCL_TOKEN }}
-        GITHUB_TOKEN: ${{ secrets.GITHUB_TOKEN }}
-      run: |
-        cd $GITHUB_WORKSPACE
-        # Set fake identity to fulfil git requirements
-        git config --global user.name "GitHub Actions"
-        git config --global user.email "[email protected]"
-        git checkout -B $BRANCH
-        git add -u
-        git commit -m "[GHA] Uplift GPU RT version for Nightly Builds" || exit 0   # exit if commit is empty
-        git push https://[email protected]/${{ github.repository }} ${BRANCH}
-        gh pr create --head $BRANCH --title "[GHA] Uplift GPU RT version for Nightly Builds" --body "Uplift GPU RT version for Linux to $NEW_DRIVER_VERSION"
-
+        git commit -m "[GHA] Uplift Linux GPU RT version to $NEW_DRIVER_VERSION" || exit 0   # exit if commit is empty
+        git push https://[email protected]/${{ github.repository }} ${BRANCH}
+        gh pr create --head $BRANCH --title "[GHA] Uplift Linux GPU RT version to $NEW_DRIVER_VERSION" --body "Scheduled drivers uplift"

clang/include/clang/Basic/Attr.td

@@ -1247,18 +1247,18 @@ def CUDAShared : InheritableAttr {
 }
 def : MutualExclusions<[CUDAConstant, CUDAShared, HIPManaged]>;
 
+def GlobalStorageNonLocalVar : SubsetSubject<Var,
+                                             [{S->hasGlobalStorage() &&
+                                               !S->isLocalVarDeclOrParm()}],
+                                             "global variables">;
+
 def SYCLDevice : InheritableAttr {
   let Spellings = [GNU<"sycl_device">];
-  let Subjects = SubjectList<[Function]>;
+  let Subjects = SubjectList<[Function, GlobalStorageNonLocalVar]>;
   let LangOpts = [SYCLIsDevice, SilentlyIgnoreSYCLIsHost];
   let Documentation = [SYCLDeviceDocs];
 }
 
-def GlobalStorageNonLocalVar : SubsetSubject<Var,
-                                             [{S->hasGlobalStorage() &&
-                                               !S->isLocalVarDeclOrParm()}],
-                                             "global variables">;
-
 def SYCLGlobalVar : InheritableAttr {
   let Spellings = [GNU<"sycl_global_var">];
   let Subjects = SubjectList<[GlobalStorageNonLocalVar], ErrorDiag>;
@@ -2401,6 +2401,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

@@ -3314,7 +3314,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++
 
@@ -3447,6 +3447,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

@@ -342,6 +342,9 @@ def err_drv_sycl_missing_amdgpu_arch : Error<
 def warn_drv_sycl_offload_target_duplicate : Warning<
   "SYCL offloading target '%0' is similar to target '%1' already specified; "
   "will be ignored">, InGroup<SyclTarget>;
+def warn_drv_sycl_target_missing : Warning<
+  "linked binaries do not contain expected '%0' target; found targets: '%1'">,
+  InGroup<SyclTarget>;
 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

@@ -11727,6 +11727,9 @@ def err_sycl_restrict : Error<
   "|use a const static or global variable that is neither zero-initialized "
   "nor constant-initialized"
   "}0">;
+def err_sycl_external_global : Error<
+  "invalid reference to 'device_global' variable; external 'device_global'"
+  " variable must be marked with SYCL_EXTERNAL macro">;
 def warn_sycl_kernel_too_big_args : Warning<
   "size of kernel arguments (%0 bytes) may exceed the supported maximum "
   "of %1 bytes on some devices">, InGroup<SyclStrict>, ShowInSystemHeader;
@@ -11758,9 +11761,12 @@ def err_sycl_function_attribute_mismatch : Error<
   "SYCL kernel without %0 attribute can't call a function with this attribute">;
 def err_sycl_x_y_z_arguments_must_be_one : Error<
   "all %0 attribute arguments must be '1' when the %1 attribute argument is '0'">;
-def err_sycl_attribute_internal_function
+def err_sycl_attribute_internal_decl
     : Error<"%0 attribute cannot be applied to a "
-            "static function or function in an anonymous namespace">;
+            "static %select{function|variable}1 or %select{function|variable}1 "
+            "in an anonymous namespace">;
+def err_sycl_attribute_not_device_global
+    : Error<"%0 attribute can only be applied to 'device_global' variables">;
 def err_sycl_compiletime_property_duplication : Error<
   "can't apply %0 property twice to the same accessor">;
 def err_sycl_invalid_property_list_param_number : Error<

clang/include/clang/Driver/Driver.h

@@ -769,6 +769,10 @@ class Driver {
   bool checkForOffloadStaticLib(Compilation &C,
                                 llvm::opt::DerivedArgList &Args) const;
 
+  /// Checks for any mismatch of targets and provided input binaries.
+  void checkForOffloadMismatch(Compilation &C,
+                               llvm::opt::DerivedArgList &Args) const;
+
   /// Track filename used for the FPGA dependency info.
   mutable llvm::StringMap<const std::string> FPGATempDepFiles;
 

clang/include/clang/Driver/Options.td

@@ -2867,6 +2867,12 @@ defm sycl_instrument_device_code
           BothFlags<[CC1Option, CoreOption], " Instrumentation and Tracing "
                     "Technology (ITT) instrumentation intrinsics calls "
                     "(experimental)">>;
+def fsycl_link_huge_device_code : Flag<["-"], "fsycl-link-huge-device-code">,
+  Group<sycl_Group>, HelpText<"Generate and use a custom linker script for huge"
+  " device code sections">;
+def fno_sycl_link_huge_device_code : Flag<["-"], "fno-sycl-link-huge-device-code">,
+  Group<sycl_Group>, HelpText<"Do not generate or use a custom linker script"
+  " for huge device code sections (default)">;
 defm sycl_id_queries_fit_in_int: BoolFOption<"sycl-id-queries-fit-in-int",
   LangOpts<"SYCLValueFitInMaxInt">, DefaultTrue,
   PosFlag<SetTrue, [], "Assume">, NegFlag<SetFalse, [], "Do not assume">,

clang/include/clang/Sema/Sema.h

@@ -2287,6 +2287,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/CodeGen/CodeGenModule.cpp

@@ -840,13 +840,15 @@ void CodeGenModule::Release() {
         llvm::MDString::get(Ctx, CodeGenOpts.MemoryProfileOutput));
   }
 
-  if ((LangOpts.CUDAIsDevice || LangOpts.isSYCL()) && getTriple().isNVPTX()) {
+  if ((LangOpts.CUDAIsDevice || LangOpts.SYCLIsDevice) && getTriple().isNVPTX()) {
     // Indicate whether __nvvm_reflect should be configured to flush denormal
     // floating point values to 0.  (This corresponds to its "__CUDA_FTZ"
     // property.)
     getModule().addModuleFlag(llvm::Module::Override, "nvvm-reflect-ftz",
-                              CodeGenOpts.FP32DenormalMode.Output !=
-                                  llvm::DenormalMode::IEEE);
+                              (CodeGenOpts.FP32DenormalMode.Output !=
+                                  llvm::DenormalMode::IEEE) ||
+                              (CodeGenOpts.FPDenormalMode.Output !=
+                                  llvm::DenormalMode::IEEE));
     getModule().addModuleFlag(llvm::Module::Override, "nvvm-reflect-prec-sqrt",
                               getTarget().getTargetOpts().NVVMCudaPrecSqrt);
   }

clang/lib/CodeGen/CodeGenTypes.cpp

@@ -51,6 +51,54 @@ void CodeGenTypes::addRecordTypeName(const RecordDecl *RD,
                                      StringRef suffix) {
   SmallString<256> TypeName;
   llvm::raw_svector_ostream OS(TypeName);
+  // If RD is spirv_JointMatrixINTEL type, mangle differently.
+  if (CGM.getTriple().isSPIRV() || CGM.getTriple().isSPIR()) {
+    if (RD->getQualifiedNameAsString() == "__spv::__spirv_JointMatrixINTEL") {
+      if (auto TemplateDecl = dyn_cast<ClassTemplateSpecializationDecl>(RD)) {
+        ArrayRef<TemplateArgument> TemplateArgs =
+            TemplateDecl->getTemplateArgs().asArray();
+        OS << "spirv.JointMatrixINTEL.";
+        for (auto &TemplateArg : TemplateArgs) {
+          OS << "_";
+          if (TemplateArg.getKind() == TemplateArgument::Type) {
+            llvm::Type *TTy = ConvertType(TemplateArg.getAsType());
+            if (TTy->isIntegerTy()) {
+              switch (TTy->getIntegerBitWidth()) {
+              case 8:
+                OS << "char";
+                break;
+              case 16:
+                OS << "short";
+                break;
+              case 32:
+                OS << "int";
+                break;
+              case 64:
+                OS << "long";
+                break;
+              default:
+                OS << "i" << TTy->getIntegerBitWidth();
+                break;
+              }
+            } else if (TTy->isBFloatTy())
+              OS << "bfloat16";
+            else if (TTy->isStructTy()) {
+              StringRef LlvmTyName = TTy->getStructName();
+              // Emit half/bfloat16 for sycl[::*]::{half,bfloat16}
+              if (LlvmTyName.startswith("class.sycl::") ||
+                  LlvmTyName.startswith("class.__sycl_internal::"))
+                LlvmTyName = LlvmTyName.rsplit("::").second;
+              OS << LlvmTyName;
+            } else
+              TTy->print(OS, false, true);
+          } else if (TemplateArg.getKind() == TemplateArgument::Integral)
+            OS << TemplateArg.getAsIntegral();
+        }
+        Ty->setName(OS.str());
+        return;
+      }
+    }
+  }
   OS << RD->getKindName() << '.';
 
   // FIXME: We probably want to make more tweaks to the printing policy. For