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

Author: iclsrc

.github/CODEOWNERS

@@ -125,6 +125,13 @@ sycl/test-e2e/InvokeSimd/ @intel/dpcpp-esimd-reviewers
 buildbot/ @intel/dpcpp-devops-reviewers
 devops/ @intel/dpcpp-devops-reviewers
 sycl/cts_exclude_filter/ @intel/dpcpp-devops-reviewers
+.ci/ @intel/dpcpp-devops-reviewers
+
+# Third-party dependencies updated by dependabot
+**/requirements.txt @intel/dpcpp-devops-reviewers
+**/requirements-hashed.txt @intel/dpcpp-devops-reviewers
+**/package-lock.json @intel/dpcpp-devops-reviewers
+llvm/utils/git/requirements_formatting.txt @intel/dpcpp-devops-reviewers
 
 # dev-igc driver update
 devops/dependencies-igc-dev.json @intel/sycl-matrix-reviewers @intel/dpcpp-esimd-reviewers @intel/dpcpp-devops-reviewers

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

@@ -79,18 +79,35 @@ jobs:
 
       toolchain_artifact: sycl_linux_ubuntu22
 
+  compat_read_exclude:
+    name: Read compatibility testing exclude list
+    runs-on: [Linux, build]
+    outputs:
+      FILTER: ${{ steps.result.outputs.FILTER }}
+    steps:
+    - uses: actions/checkout@v4
+      with:
+        sparse-checkout: |
+          devops/
+    - name: Register cleanup after job is finished
+      uses: ./devops/actions/cleanup
+    - id: result
+      shell: bash
+      run: |
+        # Transform to format expected by `llvm-lit --filter-out "pattern1|pattern2|..."`.
+        # First, remove comments/empty lines, then join lines with "|" as separator.
+        echo FILTER="$(grep -v '^#\|^\W*$' devops/compat_ci_exclude.sycl-rel-6_2 | paste -sd '|')" >> $GITHUB_OUTPUT
+
   run_prebuilt_e2e_tests:
-    needs: [build, detect_changes]
+    needs: [build, detect_changes, compat_read_exclude]
     if: ${{ always() && !cancelled() && needs.build.outputs.build_conclusion == 'success' }}
     strategy:
       fail-fast: false
       matrix:
         include:
           - name: GEN 12 Integrated
             runner: '["Linux", "gen12"]'
-            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
-            extra_lit_opts: --param gpu-intel-gen12=True
           - name: NVIDIA/CUDA
             runner: '["Linux", "cuda"]'
             image_options: -u 1001 --gpus all --cap-add SYS_ADMIN
@@ -102,45 +119,43 @@ jobs:
             extra_lit_opts: -j 1
           - name: Intel Arc A-Series Graphics
             runner: '["Linux", "arc"]'
-            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;level_zero_v2:gpu
-            extra_lit_opts: --param matrix-xmx8=True
-          - name: E2E tests with dev igc on Intel Arc A-Series Graphics
+          - name: Dev IGC / Intel Arc A-Series Graphics
             runner: '["Linux", "arc"]'
             image: ghcr.io/intel/llvm/ubuntu2404_intel_drivers:devigc
-            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
-            extra_lit_opts: --param matrix-xmx8=True
             use_igc_dev: true
             env: '{"LIT_FILTER":"Matrix/"}'
-          - name: E2E tests on Intel Ponte Vecchio GPU
+          - name: Intel Ponte Vecchio GPU
             runner: '["Linux", "pvc"]'
-            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;level_zero_v2:gpu
-          - name: Dev IGC on Intel Ponte Vecchio GPU
+          - name: Dev IGC / Intel Ponte Vecchio GPU
             runner: '["Linux", "pvc"]'
             image: ghcr.io/intel/llvm/ubuntu2404_intel_drivers:devigc
-            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
             use_igc_dev: true
             env: '{"LIT_FILTER":"Matrix/"}'
           - name: Intel Battlemage Graphics
             runner: '["Linux", "bmg"]'
-            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_v1:gpu;level_zero_v2:gpu
           - name: Preview Mode
             runner: '["Linux", "gen12"]'
-            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
             extra_lit_opts: --param test-preview-mode=True
             e2e_binaries_artifact: e2e_bin_preview
+          - name: ABI compatibility against sycl-rel-6_2
+            runner: '["Linux", "pvc"]'
+            image: ghcr.io/intel/llvm/sycl_prebuilt_tests:sycl-rel-6_2
+            target_devices: level_zero:gpu
+            extra_lit_opts: '--param test-preview-mode=False --filter-out "${{ needs.compat_read_exclude.outputs.FILTER }}"'
+            e2e_binaries_artifact: 'in-container'
 
     uses: ./.github/workflows/sycl-linux-run-tests.yml
     with:
       name: ${{ matrix.name }}
       runner: ${{ matrix.runner }}
       image: ${{ matrix.image }}
-      image_options: ${{ matrix.image_options }}
+      image_options: ${{ matrix.image_options || '-u 1001 --device=/dev/dri -v /dev/dri/by-path:/dev/dri/by-path --privileged --cap-add SYS_ADMIN' }}
       target_devices: ${{ matrix.target_devices }}
       extra_lit_opts: ${{ matrix.extra_lit_opts }}
       repo_ref: ${{ github.sha }}

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

@@ -62,8 +62,9 @@ on:
         description: |
           When set in modes other than `run-only` results in artifact upload.
           For `run-only` mode, if specified, means downloading pre-built
-          binaries from the artifact, otherwise they are expected to be part of
-          the container.
+          binaries from the artifact. If set to special value `in-container`
+          then the binaries are expected to be stored in the container image
+          instead of coming from an artifact.
         type: string
         default: ''
         required: False

.github/workflows/sycl-nightly.yml

@@ -77,34 +77,46 @@ jobs:
             image_options: -u 1001 --device=/dev/dri --device=/dev/kfd
             target_devices: hip:gpu
 
+          - name: Preview mode on AMD/HIP
+            runner: '["Linux", "amdgpu"]'
+            image_options: -u 1001 --device=/dev/dri --device=/dev/kfd
+            target_devices: hip:gpu
+            extra_lit_opts: --param test-preview-mode=True
+
           - name: NVIDIA/CUDA
             runner: '["Linux", "cuda"]'
             image_options: -u 1001 --gpus all --cap-add SYS_ADMIN
             target_devices: cuda:gpu
 
+          - name: Preview mode on NVIDIA/CUDA
+            runner: '["Linux", "cuda"]'
+            image_options: -u 1001 --gpus all --cap-add SYS_ADMIN
+            target_devices: cuda:gpu
+            extra_lit_opts: --param test-preview-mode=True
+
           - name: Intel L0 Gen12 GPU
             runner: '["Linux", "gen12"]'
-            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
 
           - name: Intel L0 Ponte Vecchio GPU
             runner: '["Linux", "pvc"]'
-            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
 
           - name: Intel L0 Battlemage GPU
             runner: '["Linux", "bmg"]'
-            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
 
+          - name: Preview mode on Intel L0 Battlemage GPU
+            runner: '["Linux", "bmg"]'
+            target_devices: level_zero:gpu
+            extra_lit_opts: --param test-preview-mode=True
+
           - name: Intel L0 Arc A-Series GPU
             runner: '["Linux", "arc"]'
-            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
 
           - name: Intel OCL Gen12 GPU
             runner: '["Linux", "gen12"]'
-            image_options: -u 1001 --device=/dev/dri -v /dev/dri/by-path:/dev/dri/by-path --privileged --cap-add SYS_ADMIN
             target_devices: opencl:gpu
 
           - name: OCL CPU (AMD)
@@ -124,15 +136,14 @@ jobs:
 
           - name: Preview mode on SPR/PVC
             runner: '["Linux", "pvc"]'
-            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
             extra_lit_opts: --param test-preview-mode=True
 
     uses: ./.github/workflows/sycl-linux-run-tests.yml
     with:
       name: ${{ matrix.name }}
       runner: ${{ matrix.runner }}
-      image_options: ${{ matrix.image_options }}
+      image_options: ${{ matrix.image_options || '-u 1001 --device=/dev/dri -v /dev/dri/by-path:/dev/dri/by-path --privileged --cap-add SYS_ADMIN' }}
       target_devices: ${{ matrix.target_devices }}
       tests_selector: e2e
       extra_lit_opts: "--param 'cxx_flags=-D_GLIBCXX_USE_CXX11_ABI=0' ${{ matrix.extra_lit_opts }}"

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

@@ -57,7 +57,6 @@ jobs:
             target_devices: level_zero:gpu
           - name: Intel Arc A-Series Graphics with Level Zero
             runner: '["Linux", "arc"]'
-            extra_lit_opts: --param matrix-xmx8=True
           # Performance tests below. Specifics:
           #  - only run performance tests (use LIT_FILTER env)
           #  - ask llvm-lit to show all the output, even for PASS (-a)

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

@@ -25,6 +25,7 @@ on:
     - '.github/workflows/trivy.yml'
     - 'devops/containers/**'
     - 'devops/actions/build_container/**'
+    - 'devops/compat_ci_exclude.sycl-rel-6_2'
     - 'unified-runtime/examples/**'
     - 'unified-runtime/scripts/**'
     - 'unified-runtime/test/**'

clang/lib/Sema/SemaSYCL.cpp

@@ -18,6 +18,7 @@
 #include "clang/AST/SYCLKernelInfo.h"
 #include "clang/AST/StmtSYCL.h"
 #include "clang/AST/TemplateArgumentVisitor.h"
+#include "clang/AST/Type.h"
 #include "clang/AST/TypeOrdering.h"
 #include "clang/AST/TypeVisitor.h"
 #include "clang/Analysis/CallGraph.h"
@@ -6737,23 +6738,75 @@ class FreeFunctionPrinter {
   /// returned string Example:
   /// \code
   ///  template <typename T1, typename T2>
-  ///  void foo(T1 a, T2 b);
+  ///  void foo(T1 a, int b, T2 c);
   /// \endcode
-  /// returns string "T1 a, T2 b"
+  /// returns string "T1, int, T2"
   std::string
   getTemplatedParamList(const llvm::ArrayRef<clang::ParmVarDecl *> Parameters,
                         PrintingPolicy Policy) {
     bool FirstParam = true;
     llvm::SmallString<128> ParamList;
     llvm::raw_svector_ostream ParmListOstream{ParamList};
     Policy.SuppressTagKeyword = true;
+
     for (ParmVarDecl *Param : Parameters) {
       if (FirstParam)
         FirstParam = false;
       else
         ParmListOstream << ", ";
-      ParmListOstream << Param->getType().getAsString(Policy);
-      ParmListOstream << " " << Param->getNameAsString();
+
+      // There are cases when we can't directly use neither the original
+      // argument type, nor its canonical version. An example would be:
+      // template<typename T>
+      // void kernel(sycl::accessor<T, 1>);
+      // template void kernel(sycl::accessor<int, 1>);
+      // Accessor has multiple non-type template arguments with default values
+      // and non-qualified type will not include necessary namespaces for all
+      // of them. Qualified type will have that information, but all references
+      // to T will be replaced to something like type-argument-0
+      // What we do instead is we iterate template arguments of both versions
+      // of a type in sync and take elements from one or another to get the best
+      // of both: proper references to template arguments of a kernel itself and
+      // fully-qualified names for enumerations.
+      //
+      // Moral of the story: drop integration header ASAP (but that is blocked
+      // by support for 3rd-party host compilers, which is important).
+      QualType T = Param->getType();
+      QualType CT = T.getCanonicalType();
+
+      auto *ET = dyn_cast<ElaboratedType>(T.getTypePtr());
+      if (!ET) {
+        ParmListOstream << T.getAsString(Policy);
+        continue;
+      }
+
+      auto *TST =
+          dyn_cast<TemplateSpecializationType>(ET->getNamedType().getTypePtr());
+      auto *CTST = dyn_cast<TemplateSpecializationType>(CT.getTypePtr());
+      if (!TST || !CTST) {
+        ParmListOstream << T.getAsString(Policy);
+        continue;
+      }
+
+      TemplateName TN = TST->getTemplateName();
+      auto SpecArgs = TST->template_arguments();
+      auto DeclArgs = CTST->template_arguments();
+
+      TN.getAsTemplateDecl()->printQualifiedName(ParmListOstream);
+      ParmListOstream << "<";
+
+      for (size_t I = 0, E = std::max(DeclArgs.size(), SpecArgs.size()),
+                  SE = SpecArgs.size();
+           I < E; ++I) {
+        if (I != 0)
+          ParmListOstream << ", ";
+        if (I < SE) // A specialized argument exists, use it
+          SpecArgs[I].print(Policy, ParmListOstream, false /* IncludeType */);
+        else // Print a canonical form of a default argument
+          DeclArgs[I].print(Policy, ParmListOstream, false /* IncludeType */);
+      }
+
+      ParmListOstream << ">";
     }
     return ParamList.str().str();
   }

clang/test/CodeGenSYCL/free-function-kernel-templated-arg-with-enum.cpp

@@ -0,0 +1,54 @@
+// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -triple spir64-unknown-unknown -sycl-std=2020 -fsycl-int-header=%t.h %s
+// RUN: FileCheck -input-file=%t.h %s
+//
+// The purpose of this test is to ensure that forward declarations of free
+// function kernels are emitted properly.
+// However, this test checks a specific scenario:
+// - free function kernel is a function template
+// - its argument is templated and has non-type template parameter (with default
+//   value) that is an enumeration defined within a namespace
+
+namespace ns {
+
+enum class enum_A { A, B, C };
+
+template<typename T, enum_A V = enum_A::B>
+class feature_A {};
+
+namespace nested {
+enum class enum_B { A, B, C };
+
+template<typename T, int V, enum_B V2 = enum_B::A, enum_A V3 = enum_A::C>
+struct feature_B {};
+}
+
+inline namespace nested_inline {
+namespace nested2 {
+enum class enum_C { A, B, C };
+
+template<int V = 42, enum_C V2 = enum_C::B>
+struct feature_C {};
+}
+} // namespace nested_inline
+} // namespace ns
+
+template<typename T>
+[[__sycl_detail__::add_ir_attributes_function("sycl-nd-range-kernel", 2)]]
+void templated_on_A(ns::feature_A<T> Arg) {}
+template void templated_on_A(ns::feature_A<int>);
+
+// CHECK: template <typename T> void templated_on_A(ns::feature_A<T, ns::enum_A::B>);
+
+template<typename T, int V = 42>
+[[__sycl_detail__::add_ir_attributes_function("sycl-nd-range-kernel", 2)]]
+void templated_on_B(ns::nested::feature_B<T, V> Arg) {}
+template void templated_on_B(ns::nested::feature_B<int, 12>);
+
+// CHECK: template <typename T, int V> void templated_on_B(ns::nested::feature_B<T, V, ns::nested::enum_B::A, ns::enum_A::C>);
+
+template<int V>
+[[__sycl_detail__::add_ir_attributes_function("sycl-nd-range-kernel", 2)]]
+void templated_on_C(ns::nested2::feature_C<V> Arg) {}
+template void templated_on_C(ns::nested2::feature_C<42>);
+
+// CHECK: template <int V> void templated_on_C(ns::nested2::feature_C<V, ns::nested2::enum_C::B>);