Merge remote-tracking branch 'intel/sycl' into steffen/refactor_kc

Author: steffenlarsen

.github/CODEOWNERS

@@ -63,12 +63,18 @@ sycl/test-e2e/Adapters/ @intel/unified-runtime-reviewers
 sycl/ur_win_proxy_loader @intel/llvm-reviewers-runtime
 sycl/test-e2e/Adapters/dll-detach-order.cpp @intel/llvm-reviewers-runtime
 
-# CUDA specific runtime implementations
+# CUDA and HIP
 sycl/include/sycl/ext/oneapi/experimental/cuda/ @intel/llvm-reviewers-cuda
-
-# CUDA and HIP device code tests
 sycl/test/check_device_code/cuda/ @intel/llvm-reviewers-cuda
 sycl/test/check_device_code/hip/ @intel/llvm-reviewers-cuda
+llvm/include/llvm/SYCLLowerIR/GlobalOffset.h @intel/llvm-reviewers-cuda
+llvm/lib/SYCLLowerIR/GlobalOffset.cpp @intel/llvm-reviewers-cuda
+llvm/include/llvm/SYCLLowerIR/LocalAccessorToSharedMemory.h @intel/llvm-reviewers-cuda
+llvm/lib/SYCLLowerIR/LocalAccessorToSharedMemory.cpp @intel/llvm-reviewers-cuda
+llvm/include/llvm/SYCLLowerIR/SYCLCreateNVVMAnnotations.h @intel/llvm-reviewers-cuda
+llvm/lib/SYCLLowerIR/SYCLCreateNVVMAnnotations.cpp @intel/llvm-reviewers-cuda
+llvm/lib/Target/NVPTX @intel/llvm-reviewers-cuda
+llvm/lib/Target/AMDGPU @intel/llvm-reviewers-cuda
 
 # XPTI instrumentation utilities
 xpti/ @intel/llvm-reviewers-runtime

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

@@ -74,6 +74,7 @@ jobs:
               - 'sycl/test-e2e/(ESIMD|InvokeSimd)/**'
             ur:
               - 'unified-runtime/**'
+              - .github/workflows/ur-*
 
       - name: Set output
         id: result

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

@@ -77,13 +77,7 @@ jobs:
           - 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
-            reset_intel_gpu: true
-            extra_lit_opts: --param matrix-xmx8=True
-          - name: Intel Arc A-Series Graphics for L0 V2 adapter
-            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_v2:gpu
+            target_devices: level_zero:gpu;opencl:gpu;level_zero_v2:gpu
             reset_intel_gpu: true
             extra_lit_opts: --param matrix-xmx8=True
           - name: E2E tests with dev igc on Intel Arc A-Series Graphics

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

@@ -322,6 +322,7 @@ jobs:
       uses: ./devops/actions/run-tests/cts
       with:
         ref: ${{ inputs.tests_ref || 'main' }}
+        cts_exclude_ref: ${{ inputs.repo_ref }}
         extra_cmake_args: ${{ inputs.extra_cmake_args }}
         cts_testing_mode: ${{ inputs.cts_testing_mode }}
         sycl_cts_artifact: ${{ inputs.sycl_cts_artifact }}

.github/workflows/ur-build-hw.yml

@@ -94,9 +94,14 @@ jobs:
     - name: Checkout LLVM
       uses: actions/checkout@b4ffde65f46336ab88eb53be808477a3936bae11 # v4.1.1
 
-    - name: Install UR python dependencies
+    # Latest distros do not allow global pip installation
+    - name: Install UR python dependencies in venv
       working-directory: ${{github.workspace}}/unified-runtime
-      run: pip install -r third_party/requirements.txt
+      run: |
+        python3 -m venv .venv
+        . .venv/bin/activate
+        echo "$PATH" >> $GITHUB_PATH
+        pip install -r third_party/requirements.txt
 
     - name: Download DPC++
       run: |

clang/lib/Sema/SemaDeclAttr.cpp

@@ -4725,7 +4725,8 @@ void Sema::AddModeAttr(Decl *D, const AttributeCommonInfo &CI,
 
   if (NewElemTy.isNull()) {
     // Only emit diagnostic on host for 128-bit mode attribute
-    if (!(DestWidth == 128 && getLangOpts().CUDAIsDevice))
+    if (!(DestWidth == 128 &&
+          (getLangOpts().CUDAIsDevice || getLangOpts().SYCLIsDevice)))
       Diag(AttrLoc, diag::err_machine_mode) << 1 /*Unsupported*/ << Name;
     return;
   }

clang/test/SemaSYCL/float128.cpp

@@ -5,6 +5,7 @@
 
 sycl::queue deviceQueue;
 
+typedef _Complex float __cfloat128 __attribute__ ((__mode__ (__TC__)));
 typedef __float128 BIGTY;
 
 template <class T>

devops/actions/run-tests/cts/action.yml

@@ -3,8 +3,10 @@ name: 'Run SYCL CTS tests'
 inputs:
   ref:
     description: "Commit SHA or branch to checkout tests"
-    required: false
-    default: "main"
+    required: true
+  cts_exclude_ref:
+    description: "Commit SHA or branch to checkout the cts_exclude_filter dir"
+    required: true
   extra_cmake_args:
     required: false
   cts_testing_mode:
@@ -19,6 +21,18 @@ inputs:
 runs:
   using: "composite"
   steps:
+  - name: Checkout cts_exclude_filter folder
+    uses: actions/checkout@v4
+    with:
+      ref: ${{ inputs.cts_exclude_ref }}
+      path: cts_exclude
+      sparse-checkout: |
+        sycl/cts_exclude_filter
+  - name: Move sycl to root
+    shell: bash
+    run: |
+      mv cts_exclude/sycl .
+      rm -rf cts_exclude
   - name: Checkout SYCL CTS tests
     if: inputs.cts_testing_mode != 'run-only'
     uses: ./devops/actions/cached_checkout
@@ -42,11 +56,11 @@ runs:
       # If CTS_TESTS_TO_BUILD is null - use filter
       if [ -z "$CTS_TESTS_TO_BUILD" ]; then
         if [ "${{ contains(inputs.cts_testing_mode, 'build-only')  }}" = "true" ]; then
-          cts_exclude_filter=$PWD/devops/cts_exclude_filter_compfails
+          cts_exclude_filter=$PWD/sycl/cts_exclude_filter/compfails
         elif [ "${{ contains(inputs.target_devices, 'opencl:cpu')  }}" = "true" ]; then
-          cts_exclude_filter=$PWD/devops/cts_exclude_filter_OCL_CPU
+          cts_exclude_filter=$PWD/sycl/cts_exclude_filter/OCL_CPU
         elif [ "${{ contains(inputs.target_devices, 'level_zero:gpu')  }}" = "true" ]; then
-          cts_exclude_filter=$PWD/devops/cts_exclude_filter_L0_GPU
+          cts_exclude_filter=$PWD/sycl/cts_exclude_filter/L0_GPU
         fi
 
         # List excluded SYCL CTS categories:
@@ -105,8 +119,8 @@ runs:
 
   # If the suite was built on another machine then the build contains the full
   # set of tests. We have special files to filter out some test categories,
-  # see "devops/cts_exclude_filter_*". Each configuration has its own file, e.g.
-  # there is "cts_exclude_filter_OCL_CPU" for opencl:cpu device. Therefore,
+  # see "sycl/cts_exclude_filter/*". Each configuration has its own file, e.g.
+  # there is "cts_exclude_filter/OCL_CPU" for opencl:cpu device. Therefore,
   # these files may differ from each other, so when there is a pre-built set of
   # tests, we need to filter it according to the filter-file.
   - name: Filter SYCL CTS test categories
@@ -115,9 +129,9 @@ runs:
     run: |
       cts_exclude_filter=""
       if [ "${{ contains(inputs.target_devices, 'opencl:cpu')  }}" = "true" ]; then
-        cts_exclude_filter=$PWD/devops/cts_exclude_filter_OCL_CPU
+        cts_exclude_filter=$PWD/sycl/cts_exclude_filter/OCL_CPU
       elif [ "${{ contains(inputs.target_devices, 'level_zero:gpu')  }}" = "true" ]; then
-        cts_exclude_filter=$PWD/devops/cts_exclude_filter_L0_GPU
+        cts_exclude_filter=$PWD/sycl/cts_exclude_filter/L0_GPU
       fi
 
       while IFS= read -r line; do

libclc/libspirv/lib/generic/math/clc_exp10.cl

@@ -123,22 +123,27 @@ _CLC_DEF _CLC_OVERLOAD double __clc_exp10(double x) {
   int m = n >> 6;
 
   double r =
-      R_LN10 * fma(-R_LOG10_2_BY_64_TL, dn, fma(-R_LOG10_2_BY_64_LD, dn, x));
+      R_LN10 * __spirv_ocl_fma(-R_LOG10_2_BY_64_TL, dn,
+                               __spirv_ocl_fma(-R_LOG10_2_BY_64_LD, dn, x));
 
   // 6 term tail of Taylor expansion of e^r
   double z2 =
-      r *
-      fma(r,
-          fma(r,
-              fma(r,
-                  fma(r, fma(r, 0x1.6c16c16c16c17p-10, 0x1.1111111111111p-7),
-                      0x1.5555555555555p-5),
-                  0x1.5555555555555p-3),
-              0x1.0000000000000p-1),
-          1.0);
+      r * __spirv_ocl_fma(
+              r,
+              __spirv_ocl_fma(
+                  r,
+                  __spirv_ocl_fma(
+                      r,
+                      __spirv_ocl_fma(r,
+                                      __spirv_ocl_fma(r, 0x1.6c16c16c16c17p-10,
+                                                      0x1.1111111111111p-7),
+                                      0x1.5555555555555p-5),
+                      0x1.5555555555555p-3),
+                  0x1.0000000000000p-1),
+              1.0);
 
   double2 tv = USE_TABLE(two_to_jby64_ep_tbl, j);
-  z2 = fma(tv.s0 + tv.s1, z2, tv.s1) + tv.s0;
+  z2 = __spirv_ocl_fma(tv.s0 + tv.s1, z2, tv.s1) + tv.s0;
 
   int small_value = (m < -1022) || ((m == -1022) && (z2 < 1.0));
 
@@ -147,7 +152,7 @@ _CLC_DEF _CLC_OVERLOAD double __clc_exp10(double x) {
   double z3 = z2 * as_double(((long)n1 + 1023) << 52);
   z3 *= as_double(((long)n2 + 1023) << 52);
 
-  z2 = ldexp(z2, m);
+  z2 = __spirv_ocl_ldexp(z2, m);
   z2 = small_value ? z3 : z2;
 
   z2 = __clc_isnan(x) ? x : z2;

libdevice/fallback-cmath.cpp

@@ -168,8 +168,13 @@ float __devicelib_fmaf(float x, float y, float z) {
   return __spirv_ocl_fma(x, y, z);
 }
 
+#if defined(__SPIR__) || defined(__SPIRV__)
+DEVICE_EXTERN_C_INLINE
+float __devicelib_sinf(float x) { return (x == 0.0f) ? x : __spirv_ocl_sin(x); }
+#else
 DEVICE_EXTERN_C_INLINE
 float __devicelib_sinf(float x) { return __spirv_ocl_sin(x); }
+#endif
 
 DEVICE_EXTERN_C_INLINE
 float __devicelib_cosf(float x) { return __spirv_ocl_cos(x); }