Merge branch 'sycl' of https://github.com/intel/llvm into ianayl/2way-prefetch

Author: ianayl

.github/workflows/sycl-containers-igc-dev.yaml

@@ -33,7 +33,7 @@ jobs:
             imagefile: ubuntu2404_intel_drivers
             tag: devigc
             build_args: |
-                  "use_latest=false"
+                  "use_unstable_driver=false"
                   "use_igc_dev=true"
     steps:
       - name: Checkout

.github/workflows/sycl-containers.yaml

@@ -43,33 +43,37 @@ jobs:
             file: ubuntu2404_base
             tag: latest
             build_args: ""
-          - name: Build Ubuntu Docker image
+          - name: Build Ubuntu 22.04 Docker image
             file: ubuntu2204_build
             tag: latest
             build_args: ""
+          - name: Build Ubuntu 24.04 Docker image
+            file: ubuntu2404_build
+            tag: latest
+            build_args: ""
           - name: Build Ubuntu 24.04 oneAPI Docker image
             file: ubuntu2404_build_oneapi
             tag: latest
             build_args: ""
           - name: Intel Drivers Ubuntu 22.04 Docker image
             file: ubuntu2204_intel_drivers
             tag: latest
-            build_args: "use_latest=false"
+            build_args: "use_unstable_driver=false"
           - name: Intel Drivers Ubuntu 24.04 Docker image
             file: ubuntu2404_intel_drivers
             tag: latest
-            build_args: "use_latest=false"
+            build_args: "use_unstable_driver=false"
           - name: Intel Drivers (unstable) Ubuntu 24.04 Docker image
             file: ubuntu2404_intel_drivers
             tag: unstable
-            build_args: "use_latest=true"
+            build_args: "use_unstable_driver=true"
           - name: Build + Intel Drivers Ubuntu 22.04 Docker image
             file: ubuntu2204_intel_drivers
             tag: alldeps
             build_args: |
               base_image=ghcr.io/intel/llvm/ubuntu2204_build
               base_tag=latest
-              use_latest=false
+              use_unstable_driver=false
     steps:
       - name: Checkout
         uses: actions/checkout@v4

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

@@ -72,6 +72,53 @@ jobs:
           else
             echo 'arc_tests="Matrix/"' >> "$GITHUB_OUTPUT"
           fi
+
+  build_e2e_tests:
+    needs: [build]
+    if: ${{ always() && !cancelled() && needs.build.outputs.build_conclusion == 'success' }}
+    uses: ./.github/workflows/sycl-linux-run-tests.yml
+    with:
+      name: Build e2e tests
+      runner: '["Linux", "build"]'
+      image: ghcr.io/intel/llvm/ubuntu2404_intel_drivers:latest
+      image_options: -u 1001
+      ref: ${{ github.sha }}
+      merge_ref: ''
+      sycl_toolchain_artifact: sycl_linux_default
+      sycl_toolchain_archive: ${{ needs.build.outputs.artifact_archive_name }}
+      sycl_toolchain_decompress_command: ${{ needs.build.outputs.artifact_decompress_command }}
+      e2e_testing_mode: 'build-only'
+  run_prebuilt_e2e_tests:
+    needs: [build, build_e2e_tests]
+    if: ${{ always() && !cancelled() && needs.build.outputs.build_conclusion == 'success' }}
+    strategy:
+      fail-fast: false
+      matrix:
+        include:
+          - name: Intel
+            runner: '["Linux", "gen12"]'
+            image: ghcr.io/intel/llvm/ubuntu2404_intel_drivers:latest
+            image_options: -u 1001 --device=/dev/dri -v /dev/dri/by-path:/dev/dri/by-path --privileged --cap-add SYS_ADMIN
+            target_devices: level_zero:gpu;opencl:gpu;opencl:cpu
+            reset_intel_gpu: true
+            extra_lit_opts: --param gpu-intel-gen12=True
+    uses: ./.github/workflows/sycl-linux-run-tests.yml
+    with:
+      name: ${{ matrix.name }}
+      runner: ${{ matrix.runner }}
+      image: ${{ matrix.image }}
+      image_options: ${{ matrix.image_options }}
+      target_devices: ${{ matrix.target_devices }}
+      extra_lit_opts: --param fallback-to-build-if-requires-build-and-run=True ${{ matrix.extra_lit_opts }}
+      reset_intel_gpu: ${{ matrix.reset_intel_gpu }}
+      ref: ${{ github.sha }}
+      merge_ref: ''
+      sycl_toolchain_artifact: sycl_linux_default
+      sycl_toolchain_archive: ${{ needs.build.outputs.artifact_archive_name }}
+      sycl_toolchain_decompress_command: ${{ needs.build.outputs.artifact_decompress_command }}
+      e2e_binaries_artifact: sycl_e2e_bin_default
+      e2e_testing_mode: 'run-only'
+
   test:
     needs: [build, detect_changes, determine_arc_tests]
     if: ${{ always() && !cancelled() && needs.build.outputs.build_conclusion == 'success' }}
@@ -90,13 +137,6 @@ jobs:
             image_options: -u 1001 --device=/dev/dri --device=/dev/kfd
             target_devices: ext_oneapi_hip:gpu
             reset_intel_gpu: false
-          - name: Intel
-            runner: '["Linux", "gen12"]'
-            image: ghcr.io/intel/llvm/ubuntu2404_intel_drivers:latest
-            image_options: -u 1001 --device=/dev/dri -v /dev/dri/by-path:/dev/dri/by-path --privileged --cap-add SYS_ADMIN
-            target_devices: level_zero:gpu;opencl:gpu;opencl:cpu
-            reset_intel_gpu: true
-            extra_lit_opts: --param gpu-intel-gen12=True
           - name: E2E tests on Intel Arc A-Series Graphics
             runner: '["Linux", "arc"]'
             image: ghcr.io/intel/llvm/ubuntu2404_intel_drivers:latest

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

@@ -19,7 +19,7 @@ on:
 
       target_devices:
         type: string
-        required: True
+        required: False
       extra_cmake_args:
         type: string
         required: False
@@ -59,6 +59,30 @@ on:
         default: ''
         required: False
 
+      e2e_binaries_artifact:
+        description: |
+          By setting this the E2E binaries folder will not be created, rather it
+          will be downloaded and extracted from the specified artifact. When
+          running tests in `run-only` mode this must be provided.
+        type: string
+        default: ''
+        required: False
+      e2e_testing_mode:
+        description: |
+          Testing mode to run E2E tests in, can be either `full`, `build-only`
+          or `run-only`. In `build-only` mode an artifact of the E2E binaries
+          will be uploaded.
+        type: string
+        default: 'full'
+      artifact_suffix:
+        description: 'Suffix for E2E binaries artifact that is output when in `build-only`.'
+        type: string
+        default: 'default'
+      retention-days:
+        description: 'E2E binaries artifact retention period.'
+        type: string
+        default: 1
+
       reset_intel_gpu:
         type: string
         required: False
@@ -87,6 +111,7 @@ on:
           - '["amdgpu"]'
           - '["Linux", "arc"]'
           - '["cts-cpu"]'
+          - '["Linux", "build"]'
       image:
         description: |
           Use option ending with ":build" for AMDGPU, ":latest" for the rest.
@@ -142,6 +167,11 @@ on:
         options:
           - false
           - true
+      e2e_testing_mode:
+        type: choice
+        options:
+          - "full"
+          - "build-only"
 
 permissions:
   contents: read
@@ -270,8 +300,19 @@ jobs:
             cat /usr/local/lib/igc/IGCTAG.txt
           fi
 
+    - name: Download E2E Binaries
+      if: inputs.e2e_binaries_artifact != ''
+      uses: actions/download-artifact@v4
+      with:
+        name: ${{ inputs.e2e_binaries_artifact }}
+    - name: Extract E2E Binaries
+      if: inputs.e2e_binaries_artifact != ''
+      run: |
+        mkdir build-e2e
+        tar -I 'zstd' -xf e2e_binaries.tar.zst -C build-e2e
+
     - name: Deduce E2E CMake options
-      if: inputs.tests_selector == 'e2e'
+      if: inputs.tests_selector == 'e2e' && inputs.e2e_binaries_artifact == ''
       id: cmake_opts
       shell: bash
       env:
@@ -281,14 +322,14 @@ jobs:
           echo "opts=$CMAKE_EXTRA_ARGS" >> $GITHUB_OUTPUT
         fi
     - name: Configure E2E tests
-      if: inputs.tests_selector == 'e2e'
+      if: inputs.tests_selector == 'e2e' && inputs.e2e_binaries_artifact == ''
       run: |
-        cmake -GNinja -B./build-e2e -S./llvm/sycl/test-e2e -DSYCL_TEST_E2E_TARGETS="${{ inputs.target_devices }}" -DCMAKE_CXX_COMPILER="$(which clang++)" -DLLVM_LIT="$PWD/llvm/llvm/utils/lit/lit.py" ${{ steps.cmake_opts.outputs.opts }}
+        cmake -GNinja -B./build-e2e -S./llvm/sycl/test-e2e -DCMAKE_CXX_COMPILER="$(which clang++)" -DLLVM_LIT="$PWD/llvm/llvm/utils/lit/lit.py" ${{ steps.cmake_opts.outputs.opts }}
     - name: SYCL End-to-end tests
       shell: bash {0}
       if: inputs.tests_selector == 'e2e'
       env:
-        LIT_OPTS: -v --no-progress-bar --show-unsupported --show-pass --show-xfail --max-time 3600 --time-tests ${{ inputs.extra_lit_opts }}
+        LIT_OPTS: -v --no-progress-bar --show-unsupported --show-pass --show-xfail --max-time 3600 --time-tests --param test-mode=${{ inputs.e2e_testing_mode }} --param sycl_devices=${{ inputs.target_devices }} ${{ inputs.extra_lit_opts }}
       run: |
         ninja -C build-e2e check-sycl-e2e > e2e.log 2>&1
         exit_code=$?
@@ -375,3 +416,13 @@ jobs:
         grep 'exit code: [^0]' -r logs >> $GITHUB_STEP_SUMMARY
 
         exit $ret
+    - name: Pack E2E binaries
+      if: ${{ always() && !cancelled() && inputs.e2e_testing_mode == 'build-only'}}
+      run: tar -I 'zstd -9' -cf e2e_binaries.tar.zst -C ./build-e2e .
+    - name: Upload E2E binaries
+      if: ${{ always() && !cancelled() && inputs.e2e_testing_mode == 'build-only'}}
+      uses: actions/upload-artifact@v4
+      with:
+        name: sycl_e2e_bin_${{ inputs.artifact_suffix }}
+        path: e2e_binaries.tar.zst
+        retention-days: ${{ inputs.retention-days }}

devops/containers/ubuntu2204_intel_drivers.Dockerfile

@@ -5,7 +5,7 @@ FROM $base_image:$base_tag
 
 ENV DEBIAN_FRONTEND=noninteractive
 
-ARG use_latest=true
+ARG use_unstable_driver=true
 
 USER root
 
@@ -18,7 +18,7 @@ COPY dependencies.json /
 RUN mkdir /runtimes
 ENV INSTALL_LOCATION=/runtimes
 RUN --mount=type=secret,id=github_token \
-    if [ "$use_latest" = "true" ]; then \
+    if [ "$use_unstable_driver" = "true" ]; then \
       install_driver_opt=" --use-latest"; \
     else \
       install_driver_opt=" dependencies.json"; \

devops/containers/ubuntu2404_build.Dockerfile

@@ -0,0 +1,42 @@
+FROM nvidia/cuda:12.6.3-devel-ubuntu24.04
+
+ENV DEBIAN_FRONTEND=noninteractive
+
+USER root
+
+# Install SYCL prerequisites
+COPY scripts/install_build_tools.sh /install.sh
+RUN /install.sh
+
+SHELL ["/bin/bash", "-ec"]
+
+# Make the directory if it doesn't exist yet.
+# This location is recommended by the distribution maintainers.
+RUN mkdir --parents --mode=0755 /etc/apt/keyrings
+# Download the key, convert the signing-key to a full
+# keyring required by apt and store in the keyring directory
+RUN wget https://repo.radeon.com/rocm/rocm.gpg.key -O - | \
+gpg --dearmor | tee /etc/apt/keyrings/rocm.gpg > /dev/null && \
+# Add rocm repo
+echo "deb [arch=amd64 signed-by=/etc/apt/keyrings/rocm.gpg] https://repo.radeon.com/amdgpu/6.3/ubuntu noble main" \
+    | tee /etc/apt/sources.list.d/amdgpu.list && \
+echo "deb [arch=amd64 signed-by=/etc/apt/keyrings/rocm.gpg] https://repo.radeon.com/rocm/apt/6.3 noble main" \
+    |  tee --append /etc/apt/sources.list.d/rocm.list && \
+echo -e 'Package: *\nPin: release o=repo.radeon.com\nPin-Priority: 600' \
+    | tee /etc/apt/preferences.d/rocm-pin-600 && \
+echo -e 'Package: *\nPin: release o=repo.radeon.com\nPin-Priority: 600' \
+    | tee /etc/apt/preferences.d/rocm-pin-600
+# Install the ROCM kernel driver
+RUN apt update && apt install -yqq rocm-dev && \
+  apt-get clean && \
+  rm -rf /var/lib/apt/lists/*
+
+COPY scripts/create-sycl-user.sh /user-setup.sh
+RUN /user-setup.sh
+
+COPY scripts/docker_entrypoint.sh /docker_entrypoint.sh
+
+USER sycl
+
+ENTRYPOINT ["/docker_entrypoint.sh"]
+

devops/containers/ubuntu2404_intel_drivers.Dockerfile

@@ -5,7 +5,7 @@ FROM $base_image:$base_tag
 
 ENV DEBIAN_FRONTEND=noninteractive
 
-ARG use_latest=true
+ARG use_unstable_driver=true
 
 USER root
 
@@ -18,7 +18,7 @@ COPY dependencies.json /
 RUN mkdir /runtimes
 ENV INSTALL_LOCATION=/runtimes
 RUN --mount=type=secret,id=github_token \
-    if [ "$use_latest" = "true" ]; then \
+    if [ "$use_unstable_driver" = "true" ]; then \
       install_driver_opt=" --use-latest"; \
     else \
       install_driver_opt=" dependencies.json"; \

llvm/lib/SYCLLowerIR/SanitizerKernelMetadata.cpp

@@ -38,6 +38,29 @@ PreservedAnalyses SanitizerKernelMetadataPass::run(Module &M,
   auto &DL = M.getDataLayout();
   auto &Ctx = M.getContext();
 
+  // Fix device global type, by wrapping a structure type
+  {
+    assert(KernelMetadata->getValueType()->isArrayTy());
+
+    auto *KernelMetadataOld = KernelMetadata;
+
+    StructType *StructTypeWithArray = StructType::create(Ctx);
+    StructTypeWithArray->setBody(KernelMetadataOld->getValueType());
+
+    KernelMetadata = new GlobalVariable(
+        M, StructTypeWithArray, false, GlobalValue::ExternalLinkage,
+        ConstantStruct::get(StructTypeWithArray,
+                            KernelMetadataOld->getInitializer()),
+        "", nullptr, GlobalValue::NotThreadLocal, 1); // Global AddressSpace
+    KernelMetadata->takeName(KernelMetadataOld);
+    KernelMetadata->setUnnamedAddr(GlobalValue::UnnamedAddr::Local);
+    KernelMetadata->setDSOLocal(true);
+    KernelMetadata->copyAttributesFrom(KernelMetadataOld);
+    KernelMetadata->copyMetadata(KernelMetadataOld, 0);
+
+    KernelMetadataOld->eraseFromParent();
+  }
+
   // Fix attributes
   KernelMetadata->addAttribute(
       "sycl-device-global-size",

llvm/test/tools/sycl-post-link/device-sanitizer/asan.ll

@@ -18,7 +18,7 @@ $_ZTSZZ4mainENKUlRN4sycl3_V17handlerEE_clES2_E8MyKernel = comdat any
 
 @__asan_kernel = internal addrspace(1) constant [55 x i8] c"_ZTSZZ4mainENKUlRN4sycl3_V17handlerEE_clES2_E8MyKernel\00"
 @__AsanKernelMetadata = appending dso_local local_unnamed_addr addrspace(1) global [1 x { i64, i64 }] [{ i64, i64 } { i64 ptrtoint (ptr addrspace(1) @__asan_kernel to i64), i64 54 }] #2
-; CHECK-IR: @__AsanKernelMetadata {{.*}} !spirv.Decorations
+; CHECK-IR: @__AsanKernelMetadata = dso_local local_unnamed_addr addrspace(1) global %0 { {{.*}} }, !spirv.Decorations
 @__spirv_BuiltInGlobalInvocationId = external dso_local local_unnamed_addr addrspace(1) constant <3 x i64>, align 32
 @__asan_func = internal addrspace(2) constant [106 x i8] c"typeinfo name for main::'lambda'(sycl::_V1::handler&)::operator()(sycl::_V1::handler&) const::MyKernelR_4\00"
 

llvm/test/tools/sycl-post-link/device-sanitizer/msan.ll

@@ -18,7 +18,7 @@ $_ZTSZZ4mainENKUlRN4sycl3_V17handlerEE_clES2_E8MyKernel = comdat any
 
 @__msan_kernel = internal addrspace(1) constant [55 x i8] c"_ZTSZZ4mainENKUlRN4sycl3_V17handlerEE_clES2_E8MyKernel\00"
 @__MsanKernelMetadata = appending dso_local local_unnamed_addr addrspace(1) global [1 x { i64, i64 }] [{ i64, i64 } { i64 ptrtoint (ptr addrspace(1) @__msan_kernel to i64), i64 54 }] #0
-; CHECK-IR: @__MsanKernelMetadata {{.*}} !spirv.Decorations
+; CHECK-IR: @__MsanKernelMetadata = dso_local local_unnamed_addr addrspace(1) global %0 { {{.*}} }, !spirv.Decorations
 @__spirv_BuiltInGlobalInvocationId = external dso_local local_unnamed_addr addrspace(1) constant <3 x i64>, align 32
 @__asan_func = internal addrspace(2) constant [106 x i8] c"typeinfo name for main::'lambda'(sycl::_V1::handler&)::operator()(sycl::_V1::handler&) const::MyKernelR_4\00"