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

  CONFLICT (content): Merge conflict in .github/workflows/email-check.yaml

Author: t-demchuk

.github/workflows/docs.yml

@@ -65,7 +65,7 @@ jobs:
           fetch-depth: 1
       - name: Get subprojects that have doc changes
         id: docs-changed-subprojects
-        uses: tj-actions/changed-files@v42
+        uses: tj-actions/changed-files@v44
         with:
           files_yaml: |
             llvm:

.github/workflows/email-check.yaml

@@ -28,6 +28,7 @@ jobs:
 
       - name: Validate author email
         if: ${{ endsWith(steps.author.outputs.EMAIL, 'noreply.github.com')  }}
+        uses: actions/github-script@v7
         env:
           COMMENT: >-
             ⚠️ We detected that you are using a GitHub private e-mail address to contribute to the repo.<br/>

.github/workflows/issue-write.yml

@@ -27,7 +27,7 @@ jobs:
           name: workflow-args
 
       - name: 'Comment on PR'
-        uses: actions/github-script@v3
+        uses: actions/github-script@v7
         with:
           github-token: ${{ secrets.GITHUB_TOKEN }}
           script: |

.github/workflows/pr-code-format.yml

@@ -26,7 +26,7 @@ jobs:
 
       - name: Get changed files
         id: changed-files
-        uses: tj-actions/changed-files@v42
+        uses: tj-actions/changed-files@v44
         with:
           separator: ","
           skip_initial_fetch: true

.github/workflows/scorecard.yml

@@ -57,6 +57,6 @@ jobs:
 
       # Upload the results to GitHub's code scanning dashboard.
       - name: "Upload to code-scanning"
-        uses: github/codeql-action/upload-sarif@b7bf0a3ed3ecfa44160715d7c442788f65f0f923 # v3.23.2
+        uses: github/codeql-action/upload-sarif@1b1aada464948af03b950897e5eb522f92603cc2 # v3.24.9
         with:
           sarif_file: results.sarif

.github/workflows/sycl-containers.yaml

@@ -28,6 +28,8 @@ jobs:
     if: github.repository == 'intel/llvm'
     name: Base Ubuntu 22.04 Docker image
     runs-on: ubuntu-22.04
+    permissions:
+      packages: write
     steps:
       - name: Checkout
         uses: actions/checkout@v4
@@ -47,6 +49,8 @@ jobs:
     if: github.repository == 'intel/llvm'
     name: Build Ubuntu Docker image
     runs-on: ubuntu-22.04
+    permissions:
+      packages: write
     steps:
       - name: Checkout
         uses: actions/checkout@v4
@@ -69,6 +73,8 @@ jobs:
     if: github.repository == 'intel/llvm'
     name: Intel Drivers Ubuntu 22.04 Docker image
     runs-on: ubuntu-22.04
+    permissions:
+      packages: write
     needs: base_image_ubuntu2204
     steps:
       - name: Checkout
@@ -107,6 +113,8 @@ jobs:
     if: github.repository == 'intel/llvm'
     name: Intel Drivers (unstable) Ubuntu 22.04 Docker image
     runs-on: ubuntu-22.04
+    permissions:
+      packages: write
     needs: base_image_ubuntu2204
     steps:
       - name: Checkout

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

@@ -20,7 +20,7 @@ jobs:
       filters: ${{ steps.result.outputs.result }}
     steps:
       - name: Check file changes
-        uses: dorny/paths-filter@0bc4621a3135347011ad047f9ecf449bf72ce2bd
+        uses: dorny/paths-filter@de90cc6fb38fc0963ad72b210f1f284cd68cea36
         id: changes
         with:
           filters: |

.github/workflows/sycl-nightly.yml

@@ -146,7 +146,7 @@ jobs:
           echo "TAG=$(date +'%Y-%m-%d')-${GITHUB_SHA::7}" >> "$GITHUB_OUTPUT"
         fi
     - name: Upload binaries
-      uses: softprops/action-gh-release@de2c0eb89ae2a093876385947365aca7b0e5f844
+      uses: softprops/action-gh-release@9d7c94cfd0a1f3ed45544c887983e9fa900f0564
       with:
         files: |
           sycl_linux.tar.gz
@@ -160,6 +160,8 @@ jobs:
   ubuntu2204_docker_build_push:
     if: github.repository == 'intel/llvm'
     runs-on: [Linux, build]
+    permissions:
+      packages: write
     needs: ubuntu2204_build
     steps:
     - uses: actions/checkout@v4

clang/lib/CodeGen/CGBuiltin.cpp

@@ -23929,9 +23929,25 @@ CodeGenFunction::EmitIntelSYCLAllocaBuiltin(const CallExpr *E,
     // an `alloca` or an equivalent construct in later compilation stages.
     IRBuilderBase::InsertPointGuard IPG(Builder);
     Builder.SetInsertPoint(AllocaInsertPt);
-    return Builder.CreateIntrinsic(
+    llvm::CallInst *CI = Builder.CreateIntrinsic(
         AllocaTy, Intrinsic::sycl_alloca,
         {UID, SpecConstPtr, RTBufferPtr, EltTyConst, Align}, nullptr, "alloca");
+
+    // Propagate function used aspects.
+    llvm::Function *F = CI->getCalledFunction();
+    constexpr llvm::StringLiteral MDName = "sycl_used_aspects";
+    if (!F->getMetadata(MDName)) {
+      auto *AspectAttr = FD->getAttr<SYCLUsesAspectsAttr>();
+      assert(AspectAttr && AspectAttr->aspects_size() == 1 &&
+             "Expecting a single aspect");
+      llvm::APSInt AspectInt =
+          (*AspectAttr->aspects_begin())->EvaluateKnownConstInt(getContext());
+      llvm::Type *I32Ty = Builder.getInt32Ty();
+      llvm::Constant *C = llvm::Constant::getIntegerValue(I32Ty, AspectInt);
+      llvm::Metadata *AspectMD = llvm::ConstantAsMetadata::get(C);
+      F->setMetadata(MDName, llvm::MDNode::get(Builder.getContext(), AspectMD));
+    }
+    return CI;
   }();
 
   // Perform AS cast if needed.

clang/lib/Driver/Driver.cpp

@@ -7393,19 +7393,6 @@ 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)) {
-        Action *PreprocessAction =
-            C.MakeAction<PreprocessJobAction>(Current, types::TY_Dependencies);
-        PreprocessAction->propagateHostOffloadInfo(Action::OFK_SYCL,
-                                                   /*BoundArch=*/nullptr);
-        Actions.push_back(PreprocessAction);
-      }
-
       if (Phase == phases::Precompile && ExtractAPIAction) {
         ExtractAPIAction->addHeaderInput(Current);
         Current = nullptr;
@@ -8162,9 +8149,6 @@ 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)
-  //
   // OffloadClass of type TY_Nothing: device-only output will place many outputs
   // into a single offloading action. We should count all inputs to the action
   // as outputs. Also ignore device-only outputs if we're compiling with
@@ -8180,10 +8164,7 @@ 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->getKind() == Action::PreprocessJobClass &&
-             A->getType() == types::TY_Dependencies &&
-             C.getArgs().hasArg(options::OPT_fsycl))))
+             A->getInputs().front()->getKind() == Action::IfsMergeJobClass)))
         ++NumOutputs;
       else if (A->getKind() == Action::OffloadClass &&
                A->getType() == types::TY_Nothing &&

clang/lib/Driver/ToolChains/Clang.cpp

@@ -1076,6 +1076,8 @@ void Clang::AddPreprocessingOptions(Compilation &C, const JobAction &JA,
   // 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.
+  // TODO: remove this when/if we can improve the host compilation situation
+  // when dealing with the temporary file generated for the footer.
   if (ContainsAppendFooterAction(&JA))
     ArgM = nullptr;
 
@@ -1092,12 +1094,6 @@ 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)) {

clang/test/CodeGenSYCL/Inputs/private_alloca.hpp

@@ -10,6 +10,7 @@ namespace experimental {
 
 template <typename ElementType, auto &Size, access::decorated DecorateAddress>
 __SYCL_BUILTIN_ALIAS(__builtin_intel_sycl_alloca)
+[[__sycl_detail__::__uses_aspects__(sycl::aspect::ext_oneapi_private_alloca)]]
 multi_ptr<ElementType, access::address_space::private_space,
           DecorateAddress> private_alloca(kernel_handler &h);
 

clang/test/CodeGenSYCL/Inputs/sycl.hpp

@@ -101,6 +101,7 @@ enum class __SYCL_TYPE(aspect) aspect { // #AspectEnum
   custom = 4,
   fp16 = 5,
   fp64 = 6,
+  ext_oneapi_private_alloca = 7,
 };
 
 using access::target;

clang/test/CodeGenSYCL/aspect_enum.cpp

@@ -3,11 +3,12 @@
 // Tests for IR of [[__sycl_detail__::sycl_type(aspect)]] enum.
 #include "sycl.hpp"
 
-// CHECK: !sycl_aspects = !{![[HOST:[0-9]+]], ![[CPU:[0-9]+]], ![[GPU:[0-9]+]], ![[ACC:[0-9]+]], ![[CUSTOM:[0-9]+]], ![[FP16:[0-9]+]], ![[FP64:[0-9]+]]}
+// CHECK: !sycl_aspects = !{![[HOST:[0-9]+]], ![[CPU:[0-9]+]], ![[GPU:[0-9]+]], ![[ACC:[0-9]+]], ![[CUSTOM:[0-9]+]], ![[FP16:[0-9]+]], ![[FP64:[0-9]+]], ![[PRIVATE_ALLOCA:[0-9]+]]}
 // CHECK: ![[HOST]] = !{!"host", i32 0}
 // CHECK: ![[CPU]] = !{!"cpu", i32 1}
 // CHECK: ![[GPU]] = !{!"gpu", i32 2}
 // CHECK: ![[ACC]] = !{!"accelerator", i32 3}
 // CHECK: ![[CUSTOM]] = !{!"custom", i32 4}
 // CHECK: ![[FP16]] = !{!"fp16", i32 5}
 // CHECK: ![[FP64]] = !{!"fp64", i32 6}
+// CHECK: ![[PRIVATE_ALLOCA]] = !{!"ext_oneapi_private_alloca", i32 7}

clang/test/CodeGenSYCL/builtin-alloca.cpp

@@ -46,3 +46,12 @@ SYCL_EXTERNAL void test(sycl::kernel_handler &kh) {
   auto ptr1 = sycl::ext::oneapi::experimental::private_alloca<int, intSize, sycl::access::decorated::legacy>(kh);
   auto ptr2 = sycl::ext::oneapi::experimental::private_alloca<myStruct, intSize, sycl::access::decorated::no>(kh);
 }
+
+// CHECK: declare !sycl_used_aspects ![[#USED_ASPECTS:]] ptr @llvm.sycl.alloca.p0.p4.p4.p4.f64
+
+// CHECK: declare !sycl_used_aspects ![[#USED_ASPECTS]] ptr @llvm.sycl.alloca.p0.p4.p4.p4.i32
+
+// CHECK: declare !sycl_used_aspects ![[#USED_ASPECTS]] ptr @llvm.sycl.alloca.p0.p4.p4.p4.s_struct.myStructs
+
+// CHECK-DAG: ![[#USED_ASPECTS]] = !{i32 [[#PRIVATE_ALLOCA_ASPECT:]]}
+// CHECK-DAG: !{!"ext_oneapi_private_alloca", i32 [[#PRIVATE_ALLOCA_ASPECT]]}

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

@@ -76,8 +76,7 @@
 /// Check behaviors for dependency generation
 // RUN:  %clangxx -fsycl -MD -c %s -### 2>&1 \
 // RUN:   | FileCheck -check-prefix DEP_GEN %s
-// DEP_GEN:  clang{{.*}} "-fsycl-is-host"
-// DEP_GEN-SAME: "-Eonly"
+// DEP_GEN:  clang{{.*}} "-fsycl-is-device"
 // DEP_GEN-SAME: "-dependency-file"
 // DEP_GEN-SAME: "-MT"
 // DEP_GEN-SAME: "-internal-isystem" "{{.*}}{{[/\\]+}}include{{[/\\]+}}sycl"
@@ -88,19 +87,18 @@
 /// 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)" {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)" {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)
+// DEP_GEN_PHASES: 0: input, "[[INPUTFILE:.+\.cpp]]", c++, (device-sycl)
+// DEP_GEN_PHASES: 1: preprocessor, {0}, c++-cpp-output, (device-sycl)
+// DEP_GEN_PHASES: 2: compiler, {1}, ir, (device-sycl)
+// DEP_GEN_PHASES: 3: offload, "device-sycl (spir64-unknown-unknown)" {2}, ir
+// DEP_GEN_PHASES: 4: input, "[[INPUTFILE]]", c++, (host-sycl)
+// DEP_GEN_PHASES: 5: append-footer, {4}, c++, (host-sycl)
+// DEP_GEN_PHASES: 6: preprocessor, {5}, c++-cpp-output, (host-sycl)
+// DEP_GEN_PHASES: 7: offload, "host-sycl (x86_64-unknown-linux-gnu)" {6}, "device-sycl (spir64-unknown-unknown)" {2}, c++-cpp-output
+// DEP_GEN_PHASES: 8: compiler, {7}, ir, (host-sycl)
+// DEP_GEN_PHASES: 9: backend, {8}, assembler, (host-sycl)
+// DEP_GEN_PHASES: 10: assembler, {9}, object, (host-sycl)
+// DEP_GEN_PHASES: 11: clang-offload-bundler, {3, 10}, object, (host-sycl)
 
 /// Allow for -o and preprocessing
 // RUN:  %clangxx -fsycl -MD -c %s -o dummy -### 2>&1 \

llvm/include/llvm/SYCLLowerIR/DeviceConfigFile.td

@@ -72,6 +72,7 @@ def AspectExt_oneapi_is_component : Aspect<"ext_oneapi_is_component">;
 def AspectExt_oneapi_graph : Aspect<"ext_oneapi_graph">;
 def AspectExt_intel_fpga_task_sequence : Aspect<"ext_intel_fpga_task_sequence">;
 def AspectExt_oneapi_limited_graph : Aspect<"ext_oneapi_limited_graph">;
+def AspectExt_oneapi_private_alloca : Aspect<"ext_oneapi_private_alloca">;
 // Deprecated aspects
 def AspectInt64_base_atomics : Aspect<"int64_base_atomics">;
 def AspectInt64_extended_atomics : Aspect<"int64_extended_atomics">;
@@ -123,7 +124,8 @@ def : TargetInfo<"__TestAspectList",
     AspectExt_oneapi_mipmap, AspectExt_oneapi_mipmap_anisotropy, AspectExt_oneapi_mipmap_level_reference, AspectExt_intel_esimd,
     AspectExt_oneapi_ballot_group, AspectExt_oneapi_fixed_size_group, AspectExt_oneapi_opportunistic_group,
     AspectExt_oneapi_tangle_group, AspectExt_intel_matrix, AspectExt_oneapi_is_composite, AspectExt_oneapi_is_component,
-    AspectExt_oneapi_graph, AspectExt_intel_fpga_task_sequence, AspectExt_oneapi_limited_graph],
+    AspectExt_oneapi_graph, AspectExt_intel_fpga_task_sequence, AspectExt_oneapi_limited_graph,
+    AspectExt_oneapi_private_alloca],
     []>;
 // This definition serves the only purpose of testing whether the deprecated aspect list defined in here and in SYCL RT
 // match.

llvm/test/tools/sycl-post-link/spec-constants/SYCL-alloca-error.ll

@@ -0,0 +1,24 @@
+; RUN: sycl-post-link -spec-const=emulation %s 2>&1 | FileCheck %s
+
+; This test checks the `-spec-const` pass on SPIR-V targets and emulation mode,
+; i.e., on AOT SPIR-V targets. In this scenario, 'llvm.sycl.alloca' intrinsics
+; must be left unmodified.
+
+; Note that coming from clang this case should never be reached.
+
+; CHECK: sycl-post-link NOTE: no modifications to the input LLVM IR have been made
+
+target triple = "spir64_x86_64"
+
+%"class.sycl::_V1::specialization_id" = type { i64 }
+
+@size_i64 = addrspace(1) constant %"class.sycl::_V1::specialization_id" { i64 10 }, align 8
+
+@size_i64_stable_name = private unnamed_addr constant [36 x i8] c"_ZTS14name_generatorIL_Z8size_i64EE\00", align 1
+
+define dso_local void @private_alloca() {
+  call ptr @llvm.sycl.alloca.p0.p4.p4.p4.f64(ptr addrspace(4) addrspacecast (ptr @size_i64_stable_name to ptr addrspace(4)), ptr addrspace(4) addrspacecast (ptr addrspace(1) @size_i64 to ptr addrspace(4)), ptr addrspace(4) null, double 0.000000e+00, i64 8)
+  ret void
+}
+
+declare ptr @llvm.sycl.alloca.p0.p4.p4.p4.f32(ptr addrspace(4), ptr addrspace(4), ptr addrspace(4), float, i64)

llvm/test/tools/sycl-post-link/spec-constants/SYCL-alloca.ll

@@ -1,7 +1,10 @@
 ; RUN: sycl-post-link -spec-const=native < %s -S -o %t.table
-; RUN: FileCheck %s -check-prefixes=CHECK-RT < %t_0.ll
+; RUN: FileCheck %s -check-prefixes=CHECK,CHECK-RT < %t_0.ll
 ; RUN: FileCheck %s --check-prefixes=CHECK-PROPS < %t_0.prop
 
+; RUN: sycl-post-link -spec-const=emulation < %s -S -o %t.table
+; RUN: FileCheck %s -check-prefixes=CHECK,CHECK-EMULATION < %t_0.ll
+
 ; This test checks that the post link tool is able to correctly transform
 ; SYCL alloca intrinsics in SPIR-V devices.
 
@@ -10,9 +13,9 @@
 %"class.sycl::_V1::specialization_id.1" = type { i16 }
 %my_range = type { ptr addrspace(4), ptr addrspace(4) }
 
-@size_i64 = internal addrspace(1) constant %"class.sycl::_V1::specialization_id" { i64 10 }, align 8
-@size_i32 = internal addrspace(1) constant %"class.sycl::_V1::specialization_id.0" { i32 120 }, align 4
-@size_i16 = internal addrspace(1) constant %"class.sycl::_V1::specialization_id.1" { i16 1 }, align 2
+@size_i64 = addrspace(1) constant %"class.sycl::_V1::specialization_id" { i64 10 }, align 8
+@size_i32 = addrspace(1) constant %"class.sycl::_V1::specialization_id.0" { i32 120 }, align 4
+@size_i16 = addrspace(1) constant %"class.sycl::_V1::specialization_id.1" { i16 1 }, align 2
 
 ; Check that the following globals are preserved: even though they are not used
 ; in the module anymore, they could still be referenced by debug info metadata
@@ -30,10 +33,19 @@
 define dso_local void @private_alloca() {
 ; CHECK-RT: [[LENGTH:%.*]] = call i32 @_Z20__spirv_SpecConstantii(i32 1, i32 120)
 ; CHECK-RT: {{.*}} = alloca double, i32 [[LENGTH]], align 8
+
+; CHECK-EMULATION:  alloca double, i32 120, align 8
   call ptr @llvm.sycl.alloca.p0.p4.p4.p4.f64(ptr addrspace(4) addrspacecast (ptr @size_i32_stable_name to ptr addrspace(4)), ptr addrspace(4) addrspacecast (ptr addrspace(1) @size_i32 to ptr addrspace(4)), ptr addrspace(4) null, double 0.000000e+00, i64 8)
 ; CHECK-RT: [[LENGTH:%.*]] = call i64 @_Z20__spirv_SpecConstantix(i32 0, i64 10)
 ; CHECK-RT: {{.*}} = alloca float, i64 [[LENGTH]], align 8
+
+; CHECK-EMULATION:  alloca float, i64 10, align 8
   call ptr @llvm.sycl.alloca.p0.p4.p4.p4.f32(ptr addrspace(4) addrspacecast (ptr @size_i64_stable_name to ptr addrspace(4)), ptr addrspace(4) addrspacecast (ptr addrspace(1) @size_i64 to ptr addrspace(4)), ptr addrspace(4) null, float 0.000000e+00, i64 8)
+
+; CHECK-RT:  %[[LENGTH:.*]] = call i16 @_Z20__spirv_SpecConstantis(i32 2, i16 1)
+; CHECK-RT:  {{.*}} = alloca %my_range, i16 %[[LENGTH]], align 64
+
+; CHECK-EMULATION:  alloca %my_range, i16 1, align 64
   call ptr @llvm.sycl.alloca.p0.p4.p4.p4.s_my_range(ptr addrspace(4) addrspacecast (ptr @size_i16_stable_name to ptr addrspace(4)), ptr addrspace(4) addrspacecast (ptr addrspace(1) @size_i16 to ptr addrspace(4)), ptr addrspace(4) null, %my_range zeroinitializer, i64 64)
   ret void
 }