[SYCL][Graph] Implementation of whole graph update

.github/workflows/sycl-aws.yml

@@ -1,5 +1,7 @@
 name: Start/Stop AWS instance
 
+permissions: read-all
+
 on:
   workflow_call:
     inputs:

.github/workflows/sycl-containers.yaml

@@ -21,6 +21,8 @@ on:
       - 'devops/scripts/install_build_tools.sh'
       - '.github/workflows/sycl-containers.yaml'
 
+permissions: read-all
+
 jobs:
   base_image_ubuntu2204:
     if: github.repository == 'intel/llvm'

.github/workflows/sycl-detect-changes.yml

@@ -7,6 +7,8 @@ on:
         description: Matched filters
         value: ${{ jobs.need_check.outputs.filters }}
 
+permissions: read-all
+
 jobs:
   need_check:
     name: Decide which tests could be affected by the changes

.github/workflows/sycl-docs.yml

@@ -11,8 +11,12 @@ on:
     - 'clang/docs/**'
     - 'sycl/doc/**'
 
+permissions: read-all
+
 jobs:
   build:
+    permissions:
+      contents: write  # for Git to git push
     runs-on: ubuntu-latest
     if: github.repository == 'intel/llvm'
     steps:

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

@@ -103,6 +103,8 @@ on:
         options:
           - 3
 
+permissions: read-all
+
 jobs:
   build:
     name: Build + LIT

.github/workflows/sycl-linux-matrix-e2e-on-nightly.yml

@@ -10,6 +10,8 @@ on:
           Format: '{"VAR1":"VAL1","VAR2":"VAL2",...}'
         default: '{"LIT_FILTER":""}'
 
+permissions: read-all
+
 jobs:
   linux_e2e_on_nightly:
     name: E2E on Nightly

.github/workflows/sycl-linux-precommit-aws.yml

@@ -14,6 +14,9 @@ on:
     types:
       - completed
 
+permissions:
+  contents: read
+
 jobs:
   create-check:
     runs-on: [Linux, build]

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

@@ -29,6 +29,8 @@ concurrency:
   group: ${{ github.workflow }}-${{ github.event.pull_request.number || github.ref }}
   cancel-in-progress: true
 
+permissions: read-all
+
 jobs:
   detect_changes:
     uses: ./.github/workflows/sycl-detect-changes.yml

.github/workflows/sycl-linux-run-tests.yml

@@ -132,6 +132,9 @@ on:
           - false
           - true
 
+permissions:
+  contents: read
+
 jobs:
   run:
     name: ${{ inputs.name }}

.github/workflows/sycl-macos-build-and-test.yml

@@ -19,6 +19,8 @@ on:
         required: false
         default: ""
 
+permissions: read-all
+
 jobs:
   build:
     name: Build

.github/workflows/sycl-nightly.yml

@@ -5,6 +5,8 @@ on:
   schedule:
     - cron: '0 3 * * *'
 
+permissions: read-all
+
 jobs:
   ubuntu2204_build:
     if: github.repository == 'intel/llvm'

.github/workflows/sycl-post-commit.yml

@@ -19,6 +19,8 @@ on:
     - ./devops/actions/cleanup
     - ./devops/actions/cached_checkout
 
+permissions: read-all
+
 jobs:
   build-lin:
     name: Linux (Self build + shared libraries + no-assertions)

.github/workflows/sycl-stale-issues.yml

@@ -4,8 +4,13 @@ on:
   schedule:
     - cron: '30 1 * * *'
 
+permissions: read-all
+
 jobs:
   close-issues:
+    permissions:
+      issues: write  # for actions/stale to close stale issues
+      pull-requests: write  # for actions/stale to close stale PRs
     runs-on: ubuntu-20.04
     steps:
     - uses: actions/stale@v9

.github/workflows/sycl-sync-main.yml

@@ -3,8 +3,12 @@ name: main branch sync
 on:
   workflow_dispatch:
 
+permissions: read-all
+
 jobs:
   sync:
+    permissions:
+      contents: write  # for Git to git push
     runs-on: ubuntu-20.04
     if: github.repository == 'intel/llvm'
     steps:

.github/workflows/sycl-update-gpu-driver.yml

@@ -5,8 +5,12 @@ on:
     - cron:  '0 3 * * 2'
   workflow_dispatch:
 
+permissions: read-all
+
 jobs:
   update_driver_linux:
+    permissions:
+      contents: write  # for Git to git push
     runs-on: ubuntu-20.04
     if: github.repository == 'intel/llvm'
     steps:
@@ -26,5 +30,6 @@ jobs:
         git checkout -B $BRANCH
         git add -u
         git commit -m "[GHA] Uplift Linux GPU RT version to $NEW_DRIVER_VERSION" || exit 0   # exit if commit is empty
+        git show
         git push https://$GITHUB_TOKEN@github.com/${{ github.repository }} ${BRANCH}
         gh pr create --head $BRANCH --title "[GHA] Uplift Linux GPU RT version to $NEW_DRIVER_VERSION" --body "Scheduled drivers uplift"

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

@@ -50,6 +50,7 @@ on:
         type: choice
         options:
           - 3
+permissions: read-all
 
 jobs:
   build:

.github/workflows/sycl-windows-precommit.yml

@@ -23,8 +23,7 @@ on:
     - 'devops/containers/**'
     - 'devops/actions/build_container/**'
 
-permissions:
-  contents: read
+permissions: read-all
 
 concurrency:
   #  Cancel a currently running workflow from the same PR, branch or tag.

.github/workflows/sycl-windows-run-tests.yml

@@ -32,6 +32,9 @@ on:
         type: string
         default: '{}'
         required: False
+
+permissions: read-all
+
 jobs:
   run:
     name: ${{ inputs.name }}

clang/include/clang/Basic/DiagnosticSemaKinds.td

@@ -191,7 +191,7 @@ def err_intel_sycl_alloca_wrong_arg
             "'sycl::kernel_handler &'. Got %0">;
 def err_intel_sycl_alloca_wrong_type
     : Error<"__builtin_intel_sycl_alloca can only return 'sycl::private_ptr' "
-            "to a cv-unqualified object type. Got %0">;
+            "to a cv-unqualified trivial type. Got %0">;
 def err_intel_sycl_alloca_wrong_size
     : Error<"__builtin_intel_sycl_alloca must be passed a specialization "
             "constant of integral value type as a template argument. Got %1 (%0)">;

clang/lib/CodeGen/BackendUtil.cpp

@@ -1138,10 +1138,10 @@ void EmitAssemblyHelper::RunOptimizationPipeline(
 
       // Add SPIRITTAnnotations pass to the pass manager if
       // -fsycl-instrument-device-code option was passed. This option can be
-      // used only with spir triple.
+      // used only with spir or spirv triple.
       if (CodeGenOpts.SPIRITTAnnotations) {
         assert(
-            TargetTriple.isSPIR() &&
+            TargetTriple.isSPIROrSPIRV() &&
             "ITT annotations can only be added to a module with spir target");
         MPM.addPass(SPIRITTAnnotationsPass());
       }

clang/lib/Driver/Driver.cpp

@@ -806,10 +806,10 @@ static bool isValidSYCLTriple(llvm::Triple T) {
     return true;
 
   // Check for invalid SYCL device triple values.
-  // Non-SPIR arch.
-  if (!T.isSPIR())
+  // Non-SPIR/SPIRV arch.
+  if (!T.isSPIROrSPIRV())
     return false;
-  // SPIR arch, but has invalid SubArch for AOT.
+  // SPIR/SPIRV arch, but has invalid SubArch for AOT.
   StringRef A(T.getArchName());
   if (T.getSubArch() == llvm::Triple::NoSubArch &&
       ((T.getArch() == llvm::Triple::spir && !A.equals("spir")) ||
@@ -833,7 +833,7 @@ static bool addSYCLDefaultTriple(Compilation &C,
     return false;
   for (const auto &SYCLTriple : SYCLTriples) {
     if (SYCLTriple.getSubArch() == llvm::Triple::NoSubArch &&
-        SYCLTriple.isSPIR())
+        SYCLTriple.isSPIROrSPIRV())
       return false;
     // If we encounter a known non-spir* target, do not add the default triple.
     if (SYCLTriple.isNVPTX() || SYCLTriple.isAMDGCN())
@@ -2392,8 +2392,9 @@ void Driver::PrintHelp(bool ShowHidden) const {
 }
 
 llvm::Triple Driver::MakeSYCLDeviceTriple(StringRef TargetArch) const {
-  SmallVector<StringRef, 5> SYCLAlias = {"spir", "spir64", "spir64_fpga",
-                                         "spir64_x86_64", "spir64_gen"};
+  SmallVector<StringRef, 5> SYCLAlias = {
+      "spir",       "spir64",  "spir64_fpga", "spir64_x86_64",
+      "spir64_gen", "spirv32", "spirv64"};
   if (std::find(SYCLAlias.begin(), SYCLAlias.end(), TargetArch) !=
       SYCLAlias.end()) {
     llvm::Triple TT;
@@ -3699,7 +3700,7 @@ bool Driver::checkForSYCLDefaultDevice(Compilation &C,
   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)
+      if ((TT.isSPIROrSPIRV()) && TT.getSubArch() == llvm::Triple::NoSubArch)
         // Default triple found
         return false;
     }
@@ -4177,7 +4178,7 @@ class OffloadingActionBuilder final {
       // supported GPUs.  sm_20 code should work correctly, if
       // suboptimally, on all newer GPUs.
       if (GpuArchList.empty()) {
-        if (ToolChains.front()->getTriple().isSPIRV())
+        if (ToolChains.front()->getTriple().isSPIROrSPIRV())
           GpuArchList.push_back(CudaArch::Generic);
         else
           GpuArchList.push_back(DefaultCudaArch);
@@ -4449,7 +4450,7 @@ class OffloadingActionBuilder final {
             // compiler phases, including backend and assemble phases.
             ActionList AL;
             Action *BackendAction = nullptr;
-            if (ToolChains.front()->getTriple().isSPIRV()) {
+            if (ToolChains.front()->getTriple().isSPIROrSPIRV()) {
               // Emit LLVM bitcode for SPIR-V targets. SPIR-V device tool chain
               // (HIPSPVToolChain) runs post-link LLVM IR passes.
               types::ID Output = Args.hasArg(options::OPT_S)
@@ -4770,7 +4771,7 @@ class OffloadingActionBuilder final {
         DA->registerDependentActionInfo(
             ToolChains[I], /*BoundArch=*/StringRef(), Action::OFK_OpenMP);
 
-        if (!ToolChains[I]->getTriple().isSPIR()) {
+        if (!ToolChains[I]->getTriple().isSPIROrSPIRV()) {
           // Create object from the deps bitcode.
           auto *BA = C.MakeAction<BackendJobAction>(DA, types::TY_PP_Asm);
           auto *AA = C.MakeAction<AssembleJobAction>(BA, types::TY_Object);
@@ -5311,7 +5312,7 @@ class OffloadingActionBuilder final {
       auto TargetTriple = TC->getTriple();
       auto IsNVPTX = TargetTriple.isNVPTX();
       auto IsAMDGCN = TargetTriple.isAMDGCN();
-      auto IsSPIR = TargetTriple.isSPIR();
+      auto IsSPIR = TargetTriple.isSPIROrSPIRV();
       bool IsSpirvAOT = TargetTriple.isSPIRAOT();
       const bool IsSYCLNativeCPU =
           TC->getAuxTriple() &&
@@ -6167,7 +6168,7 @@ class OffloadingActionBuilder final {
         // If any section found is an 'image' based object that was created
         // with the intention of not requiring the matching SYCL target, do
         // not emit the diagnostic.
-        if (SyclTarget.TC->getTriple().isSPIR()) {
+        if (SyclTarget.TC->getTriple().isSPIROrSPIRV()) {
           bool SectionFound = false;
           for (auto Section : UniqueSections) {
             if (SectionFound)
@@ -6749,7 +6750,7 @@ class OffloadingActionBuilder final {
            ++TI) {
         HasFPGATarget |= TI->second->getTriple().getSubArch() ==
                          llvm::Triple::SPIRSubArch_fpga;
-        HasSPIRTarget |= TI->second->getTriple().isSPIR();
+        HasSPIRTarget |= TI->second->getTriple().isSPIROrSPIRV();
       }
       bool isArchive = !(HostAction->getType() == types::TY_Object &&
                          isObjectFile(InputArg->getAsString(Args)));
@@ -9035,8 +9036,8 @@ InputInfoList Driver::BuildJobsForActionNoCache(
       } else if (types::isFPGA(JA->getType())) {
         std::string Ext(types::getTypeTempSuffix(JA->getType()));
         types::ID TI = types::TY_Object;
-        if (EffectiveTriple.isSPIR()) {
-          if (!UI.DependentToolChain->getTriple().isSPIR())
+        if (EffectiveTriple.isSPIROrSPIRV()) {
+          if (!UI.DependentToolChain->getTriple().isSPIROrSPIRV())
             continue;
           // Output file from unbundle is FPGA device. Name the file
           // accordingly.
@@ -10120,6 +10121,8 @@ const ToolChain &Driver::getOffloadingDeviceToolChain(
       switch (Target.getArch()) {
       case llvm::Triple::spir:
       case llvm::Triple::spir64:
+      case llvm::Triple::spirv32:
+      case llvm::Triple::spirv64:
         TC = std::make_unique<toolchains::SYCLToolChain>(*this, Target, HostTC,
                                                          Args);
         break;

clang/lib/Driver/OffloadBundler.cpp

@@ -614,11 +614,11 @@ class ObjectFileHandler final : public FileHandler {
       // input is a bitcode for spir target we need to remove module-level
       // inline asm from it, if there is one, and recreate the buffer with new
       // contents.
-      // TODO: remove this workaround once spir target gets asm parser.
+      // TODO: remove this workaround once spir/spirv target gets asm parser.
       if (isBitcode((const unsigned char *)Buf->getBufferStart(),
                     (const unsigned char *)Buf->getBufferEnd()))
         if (getTargetTriple(BundlerConfig.TargetNames[I], BundlerConfig)
-                .isSPIR()) {
+                .isSPIROrSPIRV()) {
           SMDiagnostic Err;
           std::unique_ptr<Module> Mod = parseIR(*Buf, Err, Context);
           if (!Mod)

clang/lib/Driver/SanitizerArgs.cpp

@@ -1139,8 +1139,9 @@ void SanitizerArgs::addArgs(const ToolChain &TC, const llvm::opt::ArgList &Args,
       return;
     GPUSanitize = true;
   }
-  // SPIR sanitizer support is experimental and will pass a fixed set of flags
-  if (TC.getTriple().isSPIR()) {
+  // SPIR/SPIRV sanitizer support is experimental and will pass a fixed set of
+  // flags
+  if (TC.getTriple().isSPIROrSPIRV()) {
     if (Sanitizers.has(SanitizerKind::Address)) {
       CmdArgs.push_back("-fsanitize=address");
       CmdArgs.push_back("-fsanitize-address-use-after-return=never");

clang/lib/Driver/ToolChain.cpp

@@ -1537,8 +1537,8 @@ llvm::opt::DerivedArgList *ToolChain::TranslateOffloadTargetArgs(
         DAL->append(A);
         continue;
       }
-      // SPIR-V special case for -mlong-double
-      if (getTriple().isSPIR() &&
+      // SPIR/SPIR-V special case for -mlong-double
+      if (getTriple().isSPIROrSPIRV() &&
           A->getOption().matches(options::OPT_LongDouble_Group)) {
         DAL->append(A);
         continue;