Merge branch 'sycl' into review/yang/msan_private_mem_2

Author: AllanZyne

.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@v45
+        uses: tj-actions/changed-files@v46
         with:
           files_yaml: |
             llvm:

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

@@ -34,7 +34,7 @@ jobs:
 
       - name: Get changed files
         id: changed-files
-        uses: tj-actions/changed-files@v45
+        uses: tj-actions/changed-files@v46
         with:
           separator: ","
           skip_initial_fetch: true

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

@@ -167,8 +167,7 @@ jobs:
           -DCMAKE_CXX_COMPILER_LAUNCHER=ccache \
           -DLLVM_INSTALL_UTILS=ON \
           -DNATIVECPU_USE_OCK=Off \
-          -DLLVM_EXPERIMENTAL_TARGETS_TO_BUILD=SPIRV \
-          --level_zero_v1_and_v2
+          -DLLVM_EXPERIMENTAL_TARGETS_TO_BUILD=SPIRV
     - name: Compile
       id: build
       run: cmake --build $GITHUB_WORKSPACE/build --target sycl-toolchain

.github/workflows/sycl-nightly.yml

@@ -202,12 +202,12 @@ jobs:
     with:
       mode: stop
 
-  build-sycl-cts:
+  build-sycl-cts-linux:
     needs: ubuntu2204_build
     if: ${{ always() && !cancelled() && needs.ubuntu2204_build.outputs.build_conclusion == 'success' }}
     uses: ./.github/workflows/sycl-linux-run-tests.yml
     with:
-      name: Build SYCL-CTS
+      name: Build SYCL-CTS for Linux
       runner: '["Linux", "build"]'
       cts_testing_mode: 'build-only'
       image_options: -u 1001 --device=/dev/dri --privileged --cap-add SYS_ADMIN
@@ -216,9 +216,10 @@ jobs:
       sycl_toolchain_artifact: sycl_linux_default
       sycl_toolchain_archive: ${{ needs.ubuntu2204_build.outputs.artifact_archive_name }}
       sycl_toolchain_decompress_command: ${{ needs.ubuntu2204_build.outputs.artifact_decompress_command }}
+      sycl_cts_artifact: sycl_cts_bin_linux
 
-  run-sycl-cts:
-    needs: [ubuntu2204_build, build-sycl-cts]
+  run-sycl-cts-linux:
+    needs: [ubuntu2204_build, build-sycl-cts-linux]
     if: ${{ always() && !cancelled() && needs.ubuntu2204_build.outputs.build_conclusion == 'success' }}
     strategy:
       fail-fast: false
@@ -245,7 +246,41 @@ jobs:
       sycl_toolchain_artifact: sycl_linux_default
       sycl_toolchain_archive: ${{ needs.ubuntu2204_build.outputs.artifact_archive_name }}
       sycl_toolchain_decompress_command: ${{ needs.ubuntu2204_build.outputs.artifact_decompress_command }}
-      sycl_cts_artifact: sycl_cts_bin
+      sycl_cts_artifact: sycl_cts_bin_linux
+
+  build-sycl-cts-win:
+    needs: build-win
+    if: ${{ always() && !cancelled() && needs.build-win.outputs.build_conclusion == 'success' }}
+    uses: ./.github/workflows/sycl-windows-run-tests.yml
+    with:
+      name: Build SYCL-CTS for Windows
+      runner: '["Windows", "build-e2e"]'
+      cts_testing_mode: 'build-only'
+      tests_selector: cts
+      ref: ${{ github.sha }}
+      sycl_toolchain_archive: ${{ needs.build-win.outputs.artifact_archive_name }}
+      sycl_cts_artifact: sycl_cts_bin_win
+
+  run-sycl-cts-win:
+    needs: [build-win, build-sycl-cts-win]
+    if: ${{ always() && !cancelled() && needs.build-win.outputs.build_conclusion == 'success' }}
+    strategy:
+      fail-fast: false
+      matrix:
+        include:
+          - name: SYCL-CTS on L0 gen12
+            runner: '["Windows", "gen12"]'
+            target_devices: level_zero:gpu
+    uses: ./.github/workflows/sycl-windows-run-tests.yml
+    with:
+      name: ${{ matrix.name }}
+      runner: ${{ matrix.runner }}
+      cts_testing_mode: 'run-only'
+      target_devices: ${{ matrix.target_devices }}
+      tests_selector: cts
+      ref: ${{ github.sha }}
+      sycl_toolchain_archive: ${{ needs.build-win.outputs.artifact_archive_name }}
+      sycl_cts_artifact: sycl_cts_bin_win
 
   aggregate_benchmark_results:
     if: github.repository == 'intel/llvm' && !cancelled()

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

@@ -29,6 +29,10 @@ concurrency:
 permissions: read-all
 
 jobs:
+  detect_changes:
+    if: ${{ github.event_name == 'pull_request' }}
+    uses: ./.github/workflows/sycl-detect-changes.yml
+
   build-lin:
     name: Linux (Self build + no-assertions)
     if: github.repository == 'intel/llvm'
@@ -40,7 +44,7 @@ jobs:
       build_configure_extra_args: --no-assertions --hip --cuda --native_cpu -DSYCL_ENABLE_STACK_PRINTING=ON -DSYCL_LIB_WITH_DEBUG_SYMBOL=ON
 
   e2e-lin:
-    needs: [build-lin]
+    needs: [detect_changes, build-lin]
     if: ${{ always() && !cancelled() && needs.build-lin.outputs.build_conclusion == 'success' }}
     strategy:
       fail-fast: false
@@ -94,6 +98,13 @@ jobs:
       sycl_toolchain_archive: ${{ needs.build-lin.outputs.artifact_archive_name }}
       sycl_toolchain_decompress_command: ${{ needs.build-lin.outputs.artifact_decompress_command }}
 
+      # Do not install drivers on AMD and CUDA runners.
+      install_igc_driver: >-
+        ${{ github.event_name == 'pull_request' }} &&
+        ${{ !contains(matrix.target_devices, 'cuda') &&
+        !contains(matrix.target_devices, 'hip') &&
+        contains(needs.detect_changes.outputs.filters, 'drivers') }}
+
   build-win:
     if: |
       always()

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

@@ -10,6 +10,18 @@ on:
         type: string
         required: True
 
+      target_devices:
+        type: string
+        required: False
+      extra_cmake_args:
+        type: string
+        required: False
+      tests_selector:
+        description: |
+          Two possible options: "e2e" and "cts".
+        type: string
+        default: "e2e"
+
       extra_lit_opts:
         description: |
           Extra options to be added to LIT_OPTS.
@@ -28,6 +40,10 @@ on:
           Note: it doesn't affect ./devops/actions/run-tests/* as these actions
           call checkout again and therefore override the devops directory, so
           configs/dependecies from input.ref are used.
+      tests_ref:
+        type: string
+        required: False
+        description: Commit SHA or branch to checkout e2e/cts tests.
 
       sycl_toolchain_artifact:
         type: string
@@ -48,6 +64,23 @@ on:
         required: false
         default: "cl"
 
+      cts_testing_mode:
+        description: |
+          Testing mode to run SYCL-CTS in, can be either `full`, `build-only`
+          or `run-only`. In `build-only` mode an artifact of the CTS binaries
+          will be uploaded.
+        type: string
+        default: 'full'
+
+      sycl_cts_artifact:
+        type: string
+        default: ''
+        required: False
+      artifact_retention_days:
+        description: 'E2E/SYCL-CTS binaries artifact retention period.'
+        type: string
+        default: 1
+
 permissions: read-all
 
 jobs:
@@ -86,6 +119,7 @@ jobs:
     - name: Register cleanup after job is finished
       uses: ./devops/actions/cleanup
     - uses: ./devops/actions/cached_checkout
+      if: inputs.tests_selector == 'e2e'
       with:
         path: llvm
         ref: ${{ inputs.ref || github.sha }}
@@ -103,16 +137,19 @@ jobs:
     - name: Setup SYCL toolchain
       run: |
         echo "PATH=$env:GITHUB_WORKSPACE\\install\\bin;$env:PATH" | Out-File -FilePath $env:GITHUB_ENV -Encoding utf8 -Append
+        echo "LIB=$env:GITHUB_WORKSPACE\\install\\lib;$env:LIB" | Out-File -FilePath $env:GITHUB_ENV -Encoding utf8 -Append
     - run: |
         sycl-ls
     - run: |
         sycl-ls --verbose
     - name: Configure E2E with Level Zero target
+      if: inputs.tests_selector == 'e2e'
       shell: cmd
       run: |
         mkdir build-e2e
         cmake -GNinja -B build-e2e -S.\llvm\sycl\test-e2e -DSYCL_TEST_E2E_TARGETS="level_zero:gpu" -DCMAKE_CXX_COMPILER="clang++" -DLEVEL_ZERO_LIBS_DIR="D:\\github\\level-zero_win-sdk\\lib" -DLEVEL_ZERO_INCLUDE="D:\\github\\level-zero_win-sdk\\include" -DLLVM_LIT="..\\llvm\\llvm\\utils\\lit\\lit.py"
     - name: Run End-to-End tests
+      if: inputs.tests_selector == 'e2e'
       shell: bash
       run: |
         # Run E2E tests.
@@ -121,6 +158,18 @@ jobs:
         fi
         export LIT_OPTS="-v --no-progress-bar --show-unsupported --show-pass --show-xfail --max-time 3600 --time-tests ${{ inputs.extra_lit_opts }}"
         cmake --build build-e2e --target check-sycl-e2e
+
+    - name: Run SYCL CTS Tests
+      if: inputs.tests_selector == 'cts'
+      uses: ./devops/actions/run-tests/windows/cts
+      with:
+        ref: ${{ inputs.tests_ref || 'main' }}
+        extra_cmake_args: ${{ inputs.extra_cmake_args }}
+        cts_testing_mode: ${{ inputs.cts_testing_mode }}
+        sycl_cts_artifact: ${{ inputs.sycl_cts_artifact }}
+        target_devices: ${{ inputs.target_devices }}
+        retention-days: ${{ inputs.artifact_retention_days }}
+
     - name:  Detect hung tests
       if: always()
       shell: powershell
@@ -135,4 +184,4 @@ jobs:
       if: always()
       run: |
         rmdir /q /s install
-        rmdir /q /s build-e2e
+        if exist build-e2e rmdir /q /s build-e2e

buildbot/configure.py

@@ -70,13 +70,13 @@ def do_configure(args, passthrough_args):
     if sys.platform != "darwin":
         # For more info on the enablement of level_zero_v2 refer to this document:
         # https://github.com/intel/llvm/blob/sycl/unified-runtime/source/adapters/level_zero/v2/README.md
-        if args.level_zero_v2:
-            sycl_enabled_backends.append("level_zero_v2")
-        elif args.level_zero_v1_and_v2:
+        if args.level_zero_adapter_version == "V1":
             sycl_enabled_backends.append("level_zero")
+        if args.level_zero_adapter_version == "V2":
             sycl_enabled_backends.append("level_zero_v2")
-        else:
+        if args.level_zero_adapter_version == "ALL":
             sycl_enabled_backends.append("level_zero")
+            sycl_enabled_backends.append("level_zero_v2")
 
     # lld is needed on Windows or for the HIP adapter on AMD
     if platform.system() == "Windows" or (args.hip and args.hip_platform == "AMD"):
@@ -340,12 +340,11 @@ def main():
         help="choose hardware platform for HIP backend",
     )
     parser.add_argument(
-        "--level_zero_v2", action="store_true", help="Enable SYCL Level Zero V2"
-    )
-    parser.add_argument(
-        "--level_zero_v1_and_v2",
-        action="store_true",
-        help="Enable SYCL Level Zero Legacy and V2",
+        "--level_zero_adapter_version",
+        type=str,
+        choices=["V1", "V2", "ALL"],
+        default="ALL",
+        help="Choose version of Level Zero adapter to build",
     )
     parser.add_argument(
         "--host-target",

clang/lib/CodeGen/Targets/SPIR.cpp

@@ -38,9 +38,20 @@ class CommonSPIRABIInfo : public DefaultABIInfo {
 ABIArgInfo CommonSPIRABIInfo::classifyKernelArgumentType(QualType Ty) const {
   Ty = useFirstFieldIfTransparentUnion(Ty);
 
-  if (getContext().getLangOpts().SYCLIsDevice && isAggregateTypeForABI(Ty)) {
+  if (getContext().getLangOpts().SYCLIsDevice) {
+    if (const BuiltinType *BT = Ty->getAs<BuiltinType>()) {
+      switch (BT->getKind()) {
+      case BuiltinType::Bool:
+        // Bool / i1 isn't a legal kernel argument in SPIR-V.
+        // Coerce the type to follow the host representation of bool.
+        return ABIArgInfo::getDirect(CGT.ConvertTypeForMem(Ty));
+      default:
+        break;
+      }
+    }
     // Pass all aggregate types allowed by Sema by value.
-    return getNaturalAlignIndirect(Ty);
+    if (isAggregateTypeForABI(Ty))
+      return getNaturalAlignIndirect(Ty);
   }
 
   return DefaultABIInfo::classifyArgumentType(Ty);

clang/lib/Sema/SPIRVBuiltins.td

@@ -897,15 +897,15 @@ foreach name = ["BitCount"] in {
 // 3.32.20. Barrier Instructions
 
 foreach name = ["ControlBarrier"] in {
-  // TODO: Allow enum flags instead of UInt ?
-  // TODO: We should enforce that the UInt must be a literal.
-  def : SPVBuiltin<name, [Void, UInt, UInt, UInt], Attr.Convergent>;
+  // TODO: Allow enum flags instead of Int ?
+  // TODO: We should enforce that the Int must be a literal.
+  def : SPVBuiltin<name, [Void, Int, Int, Int], Attr.Convergent>;
 }
 
 foreach name = ["MemoryBarrier"] in {
-  // TODO: Allow enum flags instead of UInt ?
-  // TODO: We should enforce that the UInt must be a literal.
-  def : SPVBuiltin<name, [Void, UInt, UInt]>;
+  // TODO: Allow enum flags instead of Int ?
+  // TODO: We should enforce that the Int must be a literal.
+  def : SPVBuiltin<name, [Void, Int, Int]>;
 }
 
 // 3.32.21. Group and Subgroup Instructions

clang/test/CodeGenSPIRV/spirv-builtin-lookup.cpp

@@ -15,3 +15,19 @@ double acos(double val) {
 }
 
 // CHECK: declare noundef double @_Z16__spirv_ocl_acosd(double noundef)
+
+void control_barrier() {
+  // CHECK-LABEL: @_Z15control_barrierv
+  // CHECK: call void @_Z22__spirv_ControlBarrieriii
+  __spirv_ControlBarrier(2, 2, 912);
+}
+
+// CHECK: declare void @_Z22__spirv_ControlBarrieriii(i32 noundef, i32 noundef, i32 noundef)
+
+void memory_barrier() {
+  // CHECK-LABEL: @_Z14memory_barrierv
+  // CHECK: call void @_Z21__spirv_MemoryBarrierii(
+  __spirv_MemoryBarrier(0, 2);
+}
+
+// CHECK: declare void @_Z21__spirv_MemoryBarrierii(i32 noundef, i32 noundef)

clang/test/CodeGenSYCL/Inputs/sycl.hpp

@@ -19,7 +19,7 @@ extern "C" int printf(const char* fmt, ...);
 
 #ifdef __SYCL_DEVICE_ONLY__
 __attribute__((convergent)) extern __attribute__((sycl_device)) void
-__spirv_ControlBarrier(int, int, int) noexcept;
+__spirv_ControlBarrier(int, int, int);
 #endif
 
 // Dummy runtime classes to model SYCL API.

clang/test/CodeGenSYCL/bool-kernel-argument.cpp

@@ -0,0 +1,26 @@
+// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -triple spir64-unknown-unknown -disable-llvm-passes -emit-llvm %s -o - | FileCheck %s
+
+#include "sycl.hpp"
+
+void take_bool(bool) {}
+
+int main() {
+  bool test = false;
+  sycl::queue q;
+
+  // CHECK: @_ZTSZZ4mainENKUlRN4sycl3_V17handlerEE_clES2_E11test_kernel(i8 {{.*}} [[ARG:%[A-Za-z_0-9]*]]
+  // CHECK: %__SYCLKernel = alloca
+  // CHECK: %test = getelementptr inbounds nuw %class.anon, ptr addrspace(4) %__SYCLKernel.ascast
+  // CHECK: store i8 %{{.*}}, ptr addrspace(4) %test
+  // CHECK: call spir_func void @_ZZZ4mainENKUlRN4sycl3_V17handlerEE_clES2_ENKUlvE_clEv
+  //
+  // CHECK: define {{.*}} @_Z9take_boolb(i1
+  q.submit([&](sycl::handler &h) {
+    h.single_task<class test_kernel>([=]() {
+      (void)test;
+      take_bool(test);
+    });
+  });
+
+  return 0;
+}