Merge remote-tracking branch 'intel/sycl' into Alcpz/impl-SYCLcompat

Author: Alcpz

.github/workflows/sycl_exp_precommit.yml

@@ -31,18 +31,6 @@ jobs:
       # prefer widespread gzip compression.
       artifact_archive_name: sycl_linux.tar.gz
 
-  ubuntu2204_opaque_pointers_build_test:
-    if: github.repository == 'intel/llvm'
-    uses: ./.github/workflows/sycl_linux_build_and_test.yml
-    needs: test_matrix
-    secrets: inherit
-    with:
-      build_cache_root: "/__w/"
-      build_cache_suffix: opaque_pointers
-      build_artifact_suffix: opaque_pointers
-      build_configure_extra_args: "--hip --cuda --enable-esimd-emulator --cmake-opt=-DSPIRV_ENABLE_OPAQUE_POINTERS=TRUE"
-      merge_ref: ''
-
   windows_default:
     name: Windows
     if: github.repository == 'intel/llvm'

.github/workflows/sycl_nightly.yml

@@ -1,8 +1,16 @@
 name: E2E on AWS CUDA
+run-name: E2E on AWS CUDA - ${{ github.event.workflow_run.display_title }}
+# We have to keep pre-commit AWS CUDA testing in a separate workflow because we
+# need access to AWS secret and that isn't available on pull_request jobs for
+# PRs from forks. And github's "require approval for all outside collaborators"
+# is bypassed on pull_request_target.
+#
+# Also, we use commit status and not check suite/run (which, in theory, is more
+# powerful) due to https://github.com/orgs/community/discussions/24616.
 
 on:
   workflow_run:
-    workflows: [SYCL Experimental Pre-Commit]
+    workflows: [SYCL Pre Commit on Linux]
     types:
       - completed
 

.github/workflows/sycl_precommit_aws.yml

@@ -1,11 +1,13 @@
 name: SYCL Pre Commit on Linux
 
 on:
-  pull_request_target:
+  # We rely on "Fork pull request workflows from outside collaborators" -
+  # "Require approval for all outside collaborators" at
+  # https://github.com/intel/llvm/settings/actions for security.
+  pull_request:
     branches:
     - sycl
     - sycl-devops-pr/**
-    - llvmspirv_pulldown
     # Do not run builds if changes are only in the following locations
     paths-ignore:
     - '.github/ISSUE_TEMPLATE/**'
@@ -15,32 +17,19 @@ on:
     - 'clang/docs/**'
     - '**.md'
     - '**.rst'
-    # Changes in CI won't have any effect with pull_request_target
-    - '.github/workflows'
-    # For CI-related files we explicitly skip all the jobs below even if there
-    # were other (non-ignored) files modified in this PR.
-    - 'devops/*/**'
-
-permissions:
-  contents: read
 
 jobs:
   detect_changes:
     uses: ./.github/workflows/sycl_detect_changes.yml
 
   lint:
-    needs: [detect_changes]
-    if: |
-      github.event.pull_request.head.repo.full_name == 'intel/llvm' ||
-      !contains(needs.detect_changes.outputs.filters, 'ci')
     runs-on: [Linux, build]
     container:
       image: ghcr.io/intel/llvm/sycl_ubuntu2204_nightly:no-drivers
       options: -u 1001:1001
     steps:
     - uses: actions/checkout@v3
       with:
-        ref: ${{ github.base_ref }}
         sparse-checkout: |
           devops/actions/cached_checkout
           devops/actions/clang-format
@@ -61,36 +50,55 @@ jobs:
       with:
         path: src
 
-  # This job generates matrix of tests for SYCL End-to-End tests
-  test_matrix:
-    needs: [detect_changes]
-    if: |
-      github.event.pull_request.head.repo.full_name == 'intel/llvm' ||
-      !contains(needs.detect_changes.outputs.filters, 'ci')
-    name: Generate Test Matrix
-    uses: ./.github/workflows/sycl_gen_test_matrix.yml
-    with:
-      ref: ${{ github.event.pull_request.head.sha }}
-      lts_config: "hip_amdgpu;lin_intel;esimd_emu;cuda_aws"
-
-  linux_default:
-    name: Linux
-    # Only build and test patches, that have passed all linter checks, because
-    # the next commit is likely to be a follow-up on that job.
-    needs: [lint, test_matrix, detect_changes]
+  build:
+    needs: [lint, detect_changes]
     if: |
       always()
       && (success() || contains(github.event.pull_request.labels.*.name, 'ignore-lint'))
-      && (github.event.pull_request.head.repo.full_name == 'intel/llvm'
-          || !contains(needs.detect_changes.outputs.filters, 'ci'))
-    uses: ./.github/workflows/sycl_linux_build_and_test.yml
-    secrets: inherit
+    uses: ./.github/workflows/sycl_linux_build.yml
     with:
-      build_ref: ${{ github.event.pull_request.head.sha }}
-      merge_ref: ${{ github.event.pull_request.base.sha }}
+      build_ref: ${{ github.sha }}
+      merge_ref: ''
       build_cache_root: "/__w/"
       build_artifact_suffix: "default"
       build_cache_suffix: "default"
-      lts_matrix: ${{ needs.test_matrix.outputs.lts_lx_matrix }}
-      lts_aws_matrix: ${{ needs.test_matrix.outputs.lts_aws_matrix }}
       changes: ${{ needs.detect_changes.outputs.filters }}
+
+  test:
+    needs: [build, detect_changes]
+    strategy:
+      fail-fast: false
+      matrix:
+        include:
+          - name: ESIMD Emu
+            runner: '["Linux", "x86-cpu"]'
+            image: ghcr.io/intel/llvm/ubuntu2204_build:latest
+            image_options: -u 1001
+            target_devices: ext_intel_esimd_emulator:gpu
+          - name: AMD/HIP
+            runner: '["Linux", "amdgpu"]'
+            image: ghcr.io/intel/llvm/ubuntu2204_build:latest
+            image_options: -u 1001 --device=/dev/dri --device=/dev/kfd
+            target_devices: ext_oneapi_hip:gpu
+          - name: Intel
+            runner: '["Linux", "gen12"]'
+            image: ghcr.io/intel/llvm/ubuntu2204_intel_drivers:latest
+            image_options: -u 1001 --device=/dev/dri --privileged --cap-add SYS_ADMIN
+            target_devices: ext_oneapi_level_zero:gpu;opencl:gpu;opencl:cpu
+            reset_gpu: ${{ contains(needs.detect_changes.outputs.filters, 'drivers') }}
+    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 }}
+      reset_gpu: ${{ matrix.reset_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 }}
+

.github/workflows/sycl_precommit_linux.yml

@@ -4031,6 +4031,7 @@ def VecTypeHint : InheritableAttr {
   let Spellings = [GNU<"vec_type_hint">, CXX11<"sycl", "vec_type_hint">];
   let Args = [TypeArgument<"TypeHint">];
   let Subjects = SubjectList<[Function], ErrorDiag>;
+  let SupportsNonconformingLambdaSyntax = 1;
   let Documentation = [Undocumented];
 }
 

clang/include/clang/Basic/Attr.td

@@ -11979,6 +11979,9 @@ def warn_ivdep_attribute_argument : Warning<
 def warn_attribute_spelling_deprecated : Warning<
   "attribute %0 is deprecated">,
   InGroup<DeprecatedAttributes>;
+def warn_attribute_deprecated_ignored : Warning<
+  "attribute %0 is deprecated; attribute ignored">,
+  InGroup<DeprecatedAttributes>;
 def note_spelling_suggestion : Note<
   "did you mean to use %0 instead?">;
 def warn_attribute_requires_non_negative_integer_argument :

clang/include/clang/Basic/DiagnosticSemaKinds.td

@@ -4550,8 +4550,11 @@ static void handleSYCLIntelLoopFuseAttr(Sema &S, Decl *D, const ParsedAttr &A) {
 
 static void handleVecTypeHint(Sema &S, Decl *D, const ParsedAttr &AL) {
   // This attribute is deprecated without replacement in SYCL 2020 mode.
-  if (S.LangOpts.getSYCLVersion() > LangOptions::SYCL_2017)
-    S.Diag(AL.getLoc(), diag::warn_attribute_spelling_deprecated) << AL;
+  // Ignore the attribute in SYCL 2020.
+  if (S.LangOpts.getSYCLVersion() > LangOptions::SYCL_2017) {
+    S.Diag(AL.getLoc(), diag::warn_attribute_deprecated_ignored) << AL;
+    return;
+  }
 
   // If the attribute is used with the [[sycl::vec_type_hint]] spelling in SYCL
   // 2017 mode, we want to warn about using the newer name in the older

clang/lib/Sema/SemaDeclAttr.cpp

@@ -0,0 +1,26 @@
+// RUN: %clang_cc1 -fsycl-is-device -sycl-std=2020 -internal-isystem %S/Inputs -fsyntax-only -verify %s
+
+// Test which verifies [[sycl::vec_type_hint()]] is accepted
+// with non-conforming lambda syntax.
+
+// NOTE: This attribute is not supported in the SYCL backends.
+// To be minimally conformant with SYCL2020, attribute is
+// accepted by the Clang FE with a warning. No additional
+// semantic handling or IR generation is done for this
+// attribute.
+
+#include "sycl.hpp"
+
+struct test {};
+
+using namespace sycl;
+queue q;
+
+void bar() {
+  q.submit([&](handler &h) {
+    h.single_task<class kernelname>(
+        // expected-warning@+1 {{attribute 'vec_type_hint' is deprecated; attribute ignored}}
+        []() [[sycl::vec_type_hint(test)]] {});
+  });
+}
+

clang/test/SemaSYCL/vec-type-hint-2.cpp

@@ -7,4 +7,4 @@
 
 // __attribute__((vec_type_hint)) is deprecated without replacement in SYCL
 // 2020 mode, but is allowed in SYCL 2017 and OpenCL modes.
-KERNEL __attribute__((vec_type_hint(int))) void foo() {} // sycl-2020-warning {{attribute 'vec_type_hint' is deprecated}}
+KERNEL __attribute__((vec_type_hint(int))) void foo() {} // sycl-2020-warning {{attribute 'vec_type_hint' is deprecated; attribute ignored}}

clang/test/SemaSYCL/vec-type-hint.cpp

@@ -63,6 +63,7 @@ static const char *LegalSYCLFunctions[] = {
     "^sycl::_V1::ext::oneapi::sub_group::.+",
     "^sycl::_V1::ext::oneapi::experimental::spec_constant<.+>::.+",
     "^sycl::_V1::ext::oneapi::experimental::this_sub_group",
+    "^sycl::_V1::ext::oneapi::experimental::uniform<.+>::.+",
     "^sycl::_V1::ext::oneapi::bfloat16::.+",
     "^sycl::_V1::ext::oneapi::experimental::if_architecture_is"};
 

llvm/lib/SYCLLowerIR/ESIMD/ESIMDVerifier.cpp

@@ -734,8 +734,8 @@ std::enable_if_t<__FAST_MATH_GENFLOAT(T), T> sin(T x) __NOEXC {
 
 // svgenfloat sincos (svgenfloat x, genfloatptr cosval)
 template <typename T, typename T2>
-std::enable_if_t<
-    detail::is_svgenfloat<T>::value && detail::is_genfloatptr<T2>::value, T>
+std::enable_if_t<__FAST_MATH_GENFLOAT(T) && detail::is_genfloatptr<T2>::value,
+                 T>
 sincos(T x, T2 cosval) __NOEXC {
   detail::check_vector_size<T, T2>();
   return __sycl_std::__invoke_sincos<T>(x, cosval);
@@ -2500,6 +2500,23 @@ std::enable_if_t<detail::is_svgenfloatf<T>::value, T> cos(T x) __NOEXC {
   return native::cos(x);
 }
 
+// svgenfloat sincos (svgenfloat x, genfloatptr cosval)
+// This is a performance optimization to ensure that sincos isn't slower than a
+// pair of sin/cos executed separately. Theoretically, calling non-native sincos
+// might be faster than calling native::sin plus native::cos separately and we'd
+// need some kind of cost model to make the right decision (and move this
+// entirely to the JIT/AOT compilers). However, in practice, this simpler
+// solution seems to work just fine and matches how sin/cos above are optimized
+// for the fast math path.
+template <typename T, typename T2>
+std::enable_if_t<
+    detail::is_svgenfloatf<T>::value && detail::is_genfloatptr<T2>::value, T>
+sincos(T x, T2 cosval) __NOEXC {
+  detail::check_vector_size<T, T2>();
+  *cosval = native::cos(x);
+  return native::sin(x);
+}
+
 // svgenfloatf exp (svgenfloatf x)
 template <typename T>
 std::enable_if_t<detail::is_svgenfloatf<T>::value, T> exp(T x) __NOEXC {

sycl/include/sycl/builtins.hpp

@@ -371,7 +371,8 @@ constexpr bool has_struct_arg(Ret (*)(Args...)) {
 
 template <typename Ret, typename... Args>
 constexpr bool has_struct_ret(Ret (*)(Args...)) {
-  return std::is_class_v<Ret> && !is_simd_or_mask_type<Ret>::value;
+  return std::is_class_v<Ret> && !is_simd_or_mask_type<Ret>::value &&
+         !is_uniform_type<Ret>::value;
 }
 
 template <typename Ret, typename... Args>