Merge from 'sycl' to 'sycl-web'

Author: iclsrc

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

@@ -86,15 +86,13 @@ jobs:
             image_options: -u 1001 --device=/dev/dri --privileged --cap-add SYS_ADMIN
             target_devices: level_zero:gpu;opencl:gpu;opencl:cpu
             reset_intel_gpu: true
-            install_drivers: ${{ contains(needs.detect_changes.outputs.filters, 'drivers') }}
             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/ubuntu2204_intel_drivers:latest
             image_options: -u 1001 --device=/dev/dri --privileged --cap-add SYS_ADMIN
             target_devices: level_zero:gpu;opencl:gpu
             reset_intel_gpu: true
-            install_drivers: ${{ contains(needs.detect_changes.outputs.filters, 'drivers') }}
             extra_lit_opts: --param matrix-xmx8=True --param gpu-intel-dg2=True
             env: '{"LIT_FILTER":${{ needs.determine_arc_tests.outputs.arc_tests }} }'
           - name: E2E tests with dev igc on Intel Arc A-Series Graphics
@@ -103,14 +101,9 @@ jobs:
             image_options: -u 1001 --device=/dev/dri --privileged --cap-add SYS_ADMIN
             target_devices: level_zero:gpu;opencl:gpu
             reset_intel_gpu: true
-            install_drivers: >-
-              ${{ contains(needs.detect_changes.outputs.filters, 'drivers') ||
-              contains(needs.detect_changes.outputs.filters, 'devigccfg') }}
-            use_dev_igc: ${{ contains(needs.detect_changes.outputs.filters, 'devigccfg') }}
             extra_lit_opts: --param matrix-xmx8=True --param gpu-intel-dg2=True
             env: '{"LIT_FILTER":${{ needs.determine_arc_tests.outputs.arc_tests }} }'
-            # Run only if the PR does not have the 'ci-no-devigc' label.
-            skip_run: ${{ contains(github.event.pull_request.labels.*.name, 'ci-no-devigc') }}
+            use_igc_dev: true
 
     uses: ./.github/workflows/sycl-linux-run-tests.yml
     with:
@@ -120,11 +113,13 @@ jobs:
       image_options: ${{ matrix.image_options }}
       target_devices: ${{ matrix.target_devices }}
       reset_intel_gpu: ${{ matrix.reset_intel_gpu }}
-      install_drivers: ${{ matrix.install_drivers }}
-      use_dev_igc: ${{ matrix.use_dev_igc }}
       extra_lit_opts: ${{ matrix.extra_lit_opts }}
       env: ${{ matrix.env || '{}' }}
-      skip_run: ${{ matrix.skip_run || 'false' }}
+
+      install_igc_driver: ${{ contains(needs.detect_changes.outputs.filters, 'drivers') }}
+      install_dev_igc_driver: ${{ matrix.use_igc_dev && contains(needs.detect_changes.outputs.filters, 'devigccfg') || 'false' }}
+      # Run only if the PR does not have the 'ci-no-devigc' label.
+      skip_run: ${{matrix.use_igc_dev && contains(github.event.pull_request.labels.*.name, 'ci-no-devigc') || 'false'}}
 
       ref: ${{ github.sha }}
       merge_ref: ''

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

@@ -65,10 +65,10 @@ on:
       reset_intel_gpu:
         type: string
         required: False
-      install_drivers:
+      install_igc_driver:
         type: string
         required: False
-      use_dev_igc:
+      install_dev_igc_driver:
         type: string
         required: False
       env:
@@ -134,13 +134,13 @@ on:
           Extra options to be added to LIT_OPTS.
         default: ''
 
-      install_drivers:
+      install_igc_driver:
         type: choice
         options:
           - false
           - true
 
-      use_dev_igc:
+      install_dev_igc_driver:
         type: choice
         options:
           - false
@@ -193,15 +193,15 @@ jobs:
       run: |
         git -C khronos_sycl_cts submodule update --init
     - name: Install drivers
-      if: inputs.install_drivers == 'true'
+      if: inputs.install_igc_driver == 'true' || inputs.install_dev_igc_driver == 'true'
       env:
         GITHUB_TOKEN: ${{ github.token }}
       run: |
-        if [ "${{ inputs.use_dev_igc }}" = "true" ]; then
+        if [ "${{ inputs.install_dev_igc_driver }}" = "true" ]; then
             # If libllvm14 is already installed (dev igc docker), still return true.
             sudo apt-get install -yqq libllvm14 || true;
         fi
-        sudo -E bash devops/scripts/install_drivers.sh llvm/devops/dependencies.json ${{ inputs.use_dev_igc == 'true' && 'llvm/devops/dependencies-igc-dev.json --use-dev-igc' || '' }} --all
+        sudo -E bash devops/scripts/install_drivers.sh llvm/devops/dependencies.json ${{ inputs.install_dev_igc_driver == 'true' && 'llvm/devops/dependencies-igc-dev.json --use-dev-igc' || '' }} --all
     - name: Source OneAPI TBB vars.sh
       shell: bash
       run: |

clang/include/clang/Basic/DiagnosticSemaKinds.td

@@ -12374,6 +12374,8 @@ def err_bit_cast_type_size_mismatch : Error<
   "__builtin_bit_cast source size does not equal destination size (%0 vs %1)">;
 
 // SYCL-specific diagnostics
+def warn_sycl_incorrect_use_attribute_non_kernel_function : Warning<
+  "%0 attribute can only be applied to a SYCL kernel function">, InGroup<SyclStrict>;
 def warn_sycl_kernel_num_of_template_params : Warning<
   "'sycl_kernel' attribute only applies to a function template with at least"
   " two template parameters">, InGroup<IgnoredAttributes>;

clang/include/clang/Sema/SemaSYCL.h

@@ -259,6 +259,8 @@ class SemaSYCL : public SemaBase {
   // useful notes that shows where the kernel was called.
   bool DiagnosingSYCLKernel = false;
 
+  llvm::DenseSet<const FunctionDecl *> SYCLKernelFunctions;
+
 public:
   SemaSYCL(Sema &S);
 
@@ -300,6 +302,10 @@ class SemaSYCL : public SemaBase {
   void addSyclDeviceDecl(Decl *d) { SyclDeviceDecls.insert(d); }
   llvm::SetVector<Decl *> &syclDeviceDecls() { return SyclDeviceDecls; }
 
+  void addSYCLKernelFunction(const FunctionDecl *FD) {
+    SYCLKernelFunctions.insert(FD);
+  }
+
   /// Lazily creates and returns SYCL integration header instance.
   SYCLIntegrationHeader &getSyclIntegrationHeader() {
     if (SyclIntHeader == nullptr)
@@ -375,6 +381,8 @@ class SemaSYCL : public SemaBase {
                                    SourceLocation Loc,
                                    DeviceDiagnosticReason Reason);
 
+  void performSYCLDelayedAttributesAnalaysis(const FunctionDecl *FD);
+
   /// Tells whether given variable is a SYCL explicit SIMD extension's "private
   /// global" variable - global variable in the private address space.
   bool isSYCLEsimdPrivateGlobal(VarDecl *VDecl) {

clang/lib/Sema/Sema.cpp

@@ -1909,6 +1909,10 @@ class DeferredDiagnosticsEmitter
   void checkFunc(SourceLocation Loc, FunctionDecl *FD) {
     auto &Done = DoneMap[InOMPDeviceContext > 0 ? 1 : 0];
     FunctionDecl *Caller = UsePath.empty() ? nullptr : UsePath.back();
+
+    if (!Caller && S.LangOpts.SYCLIsDevice)
+      S.SYCL().performSYCLDelayedAttributesAnalaysis(FD);
+
     if ((!ShouldEmitRootNode && !S.getLangOpts().OpenMP && !Caller) ||
         S.shouldIgnoreInHostDeviceCheck(FD) || InUsePath.count(FD))
       return;

clang/lib/Sema/SemaSYCL.cpp

@@ -882,6 +882,8 @@ class SingleDeviceFunctionTracker {
         // having a kernel lambda with a lambda call inside of it.
         KernelBody = CurrentDecl;
       }
+      if (KernelBody)
+        Parent.SemaSYCLRef.addSYCLKernelFunction(KernelBody);
     }
 
     // Recurse.
@@ -6852,3 +6854,19 @@ ExprResult SemaSYCL::ActOnUniqueStableNameExpr(SourceLocation OpLoc,
 
   return BuildUniqueStableNameExpr(OpLoc, LParen, RParen, TSI);
 }
+
+void SemaSYCL::performSYCLDelayedAttributesAnalaysis(const FunctionDecl *FD) {
+  if (SYCLKernelFunctions.contains(FD))
+    return;
+
+  for (const auto *KernelAttr : std::vector<AttributeCommonInfo *>{
+           FD->getAttr<SYCLReqdWorkGroupSizeAttr>(),
+           FD->getAttr<IntelReqdSubGroupSizeAttr>(),
+           FD->getAttr<SYCLWorkGroupSizeHintAttr>(),
+           FD->getAttr<VecTypeHintAttr>()}) {
+    if (KernelAttr)
+      Diag(KernelAttr->getLoc(),
+           diag::warn_sycl_incorrect_use_attribute_non_kernel_function)
+          << KernelAttr;
+  }
+}

clang/test/Driver/sycl.c

@@ -1,3 +1,6 @@
+// Failing on Windows - temporarily disable
+// REQUIRES: system-linux
+
 // RUN: %clang -### -fsycl -c %s 2>&1 | FileCheck %s --check-prefix=ENABLED
 // RUN: %clang -### -fsycl %s 2>&1 | FileCheck %s --check-prefix=ENABLED
 // RUN: %clang -### -fno-sycl -fsycl %s 2>&1 | FileCheck %s --check-prefix=ENABLED

clang/test/SemaSYCL/check-work-group-size-hint-device.cpp

@@ -17,7 +17,8 @@
 
 // Produce a conflicting attribute warning when the args are different.
 [[sycl::work_group_size_hint(4, 1, 1)]] void f3();    // expected-note {{previous attribute is here}}
-[[sycl::work_group_size_hint(1, 1, 32)]] void f3() {} // expected-warning {{attribute 'work_group_size_hint' is already applied with different arguments}}
+[[sycl::work_group_size_hint(1, 1, 32)]] void f3() {} // expected-warning {{attribute 'work_group_size_hint' is already applied with different arguments}} \
+// expected-warning {{'work_group_size_hint' attribute can only be applied to a SYCL kernel function}}
 
 // 1 and 2 dim versions
 [[sycl::work_group_size_hint(2)]] void f4();    // ok
@@ -70,10 +71,13 @@ void instantiate() {
   f8<0>(); // expected-note {{in instantiation}}
 #endif
 
+  // expected-warning@#f9prev {{'work_group_size_hint' attribute can only be applied to a SYCL kernel function}}
   f9<1, 1, 1>(); // OK, args are the same on the redecl.
 
   // expected-warning@#f9 {{attribute 'work_group_size_hint' is already applied with different arguments}}
   // expected-note@#f9prev {{previous attribute is here}}
+  // expected-warning@#f9prev {{'work_group_size_hint' attribute can only be applied to a SYCL kernel function}}
+
   f9<1, 2, 3>(); // expected-note {{in instantiation}}
 }
 
@@ -97,14 +101,14 @@ class Functor16x2x1 {
 
 class Functor4x4x4 {
 public:
-  [[sycl::work_group_size_hint(4, 4, 4)]] void operator()() const {};
+  [[sycl::work_group_size_hint(4, 4, 4)]] void operator()() const {}; // expected-warning {{'work_group_size_hint' attribute can only be applied to a SYCL kernel function}}
 };
 
 // Checking whether propagation of the attribute happens or not, according to the SYCL version.
 #if defined(EXPECT_PROP) // if attribute is propagated, then we expect errors here
 void f8x8x8(){};
 #else // otherwise no error
-[[sycl::work_group_size_hint(8, 8, 8)]] void f8x8x8(){};
+[[sycl::work_group_size_hint(8, 8, 8)]] void f8x8x8(){}; // expected-warning {{'work_group_size_hint' attribute can only be applied to a SYCL kernel function}}
 #endif
 class FunctorNoProp {
 public:

clang/test/SemaSYCL/device_has.cpp

@@ -17,7 +17,6 @@ enum class aspect {
 
 [[sycl::device_has("123")]] void func1() {}                          // expected-error{{'device_has' attribute argument is invalid; argument must be device aspect of type sycl::aspect}}
 [[sycl::device_has(fake_cl::sycl::aspect::aspect1)]] void func2() {} // expected-error{{'device_has' attribute argument is invalid; argument must be device aspect of type sycl::aspect}}
-
 [[sycl::device_has(sycl::aspect::cpu)]] void func3();   // expected-note{{previous attribute is here}}
 [[sycl::device_has(sycl::aspect::gpu)]] void func3() {} // expected-warning{{attribute 'device_has' is already applied}}
 

clang/test/SemaSYCL/intel-max-work-group-size.cpp

@@ -70,15 +70,15 @@ void instantiate() {
 // a declaration along with [[sycl::reqd_work_group_size(X1, Y1, Z1)]]
 // attribute, check to see if values of reqd_work_group_size arguments are
 // equal or less than values coming from max_work_group_size attribute.
-[[sycl::reqd_work_group_size(64, 64, 64)]] // expected-note {{conflicting attribute is here}}
+[[sycl::reqd_work_group_size(64, 64, 64)]] // expected-note {{conflicting attribute is here}} // expected-warning {{'reqd_work_group_size' attribute can only be applied to a SYCL kernel function}}
 [[intel::max_work_group_size(64, 16, 64)]] // expected-error {{'max_work_group_size' attribute conflicts with 'reqd_work_group_size' attribute}}
 void
 f9() {}
 
 [[intel::max_work_group_size(4, 4, 4)]] void f10();
 [[sycl::reqd_work_group_size(2, 2, 2)]] void f10(); // OK
 
-[[sycl::reqd_work_group_size(2, 2, 2)]] [[intel::max_work_group_size(4, 4, 4)]] void f11() {} // OK
+[[sycl::reqd_work_group_size(2, 2, 2)]] [[intel::max_work_group_size(4, 4, 4)]] void f11() {} // expected-warning {{'reqd_work_group_size' attribute can only be applied to a SYCL kernel function}}
 
 [[sycl::reqd_work_group_size(64, 64, 64)]] void f12(); // expected-note {{conflicting attribute is here}}
 [[intel::max_work_group_size(16, 16, 16)]] void f12(); // expected-error {{'max_work_group_size' attribute conflicts with 'reqd_work_group_size' attribute}}
@@ -91,14 +91,14 @@ f13() {}
 [[sycl::reqd_work_group_size(64, 64, 64)]] void f14(); // expected-error{{'reqd_work_group_size' attribute conflicts with 'max_work_group_size' attribute}}
 
 [[cl::reqd_work_group_size(1, 2, 3)]] // expected-warning {{attribute 'cl::reqd_work_group_size' is deprecated}} \
-		                      // expected-note {{did you mean to use 'sycl::reqd_work_group_size' instead?}}
+		                      // expected-note {{did you mean to use 'sycl::reqd_work_group_size' instead?}} // expected-warning {{'reqd_work_group_size' attribute can only be applied to a SYCL kernel function}}
 [[intel::max_work_group_size(1, 2, 3)]] void
 f15() {} // OK
 
 [[intel::max_work_group_size(2, 3, 7)]] void f16(); // expected-note {{conflicting attribute is here}}
 [[sycl::reqd_work_group_size(7, 3, 2)]] void f16(); // expected-error{{'reqd_work_group_size' attribute conflicts with 'max_work_group_size' attribute}}
 
-[[intel::max_work_group_size(1, 2, 3)]] [[sycl::reqd_work_group_size(1, 2, 3)]] void f17(){}; // OK
+[[intel::max_work_group_size(1, 2, 3)]] [[sycl::reqd_work_group_size(1, 2, 3)]] void f17(){}; // expected-warning {{'reqd_work_group_size' attribute can only be applied to a SYCL kernel function}}
 
 [[sycl::reqd_work_group_size(16)]]            // expected-note {{conflicting attribute is here}}
 [[intel::max_work_group_size(16, 1, 1)]] void // expected-error {{'max_work_group_size' attribute conflicts with 'reqd_work_group_size' attribute}}

clang/test/SemaSYCL/reqd-sub-group-size.cpp

@@ -33,17 +33,19 @@ int main() {
   return 0;
 }
 [[intel::reqd_sub_group_size(16)]] SYCL_EXTERNAL void B();
-[[intel::reqd_sub_group_size(16)]] void A() {
+[[intel::reqd_sub_group_size(16)]] void A() // expected-warning {{'reqd_sub_group_size' attribute can only be applied to a SYCL kernel function}} 
+{ 
 }
 
-[[intel::reqd_sub_group_size(16)]] SYCL_EXTERNAL void B() {
+[[intel::reqd_sub_group_size(16)]] SYCL_EXTERNAL void B() { // expected-warning {{'reqd_sub_group_size' attribute can only be applied to a SYCL kernel function}}
   A();
 }
 // expected-note@+1 {{conflicting attribute is here}}
-[[intel::reqd_sub_group_size(2)]] void sg_size2() {}
+[[intel::reqd_sub_group_size(2)]] void sg_size2() {} // expected-warning {{'reqd_sub_group_size' attribute can only be applied to a SYCL kernel function}}
 
-// expected-note@+2 {{conflicting attribute is here}}
-// expected-error@+1 {{conflicting attributes applied to a SYCL kernel}}
+// expected-note@+3 {{conflicting attribute is here}}
+// expected-error@+2 {{conflicting attributes applied to a SYCL kernel}}
+// expected-warning@+1 {{'reqd_sub_group_size' attribute can only be applied to a SYCL kernel function}}
 [[intel::reqd_sub_group_size(4)]] __attribute__((sycl_device)) void sg_size4() {
   sg_size2();
 }
@@ -67,7 +69,7 @@ int main() {
 
 // No diagnostic is emitted because the arguments match.
 [[intel::reqd_sub_group_size(12)]] void same();
-[[intel::reqd_sub_group_size(12)]] void same() {} // OK
+[[intel::reqd_sub_group_size(12)]] void same() {} // expected-warning {{'reqd_sub_group_size' attribute can only be applied to a SYCL kernel function}}
 
 // No diagnostic because the attributes are synonyms with identical behavior.
 [[sycl::reqd_sub_group_size(12)]] void same(); // OK
@@ -117,10 +119,12 @@ int check() {
 
 // Test that checks template parameter support on function.
 template <int N>
-// expected-error@+1{{'reqd_sub_group_size' attribute requires a positive integral compile time constant expression}}
+// expected-error@+2{{'reqd_sub_group_size' attribute requires a positive integral compile time constant expression}}
+// expected-warning@+1 {{'reqd_sub_group_size' attribute can only be applied to a SYCL kernel function}}
 [[intel::reqd_sub_group_size(N)]] void func3() {}
 
 template <int N>
+// expected-warning@+1 {{'reqd_sub_group_size' attribute can only be applied to a SYCL kernel function}}
 [[intel::reqd_sub_group_size(4)]] void func4(); // expected-note {{previous attribute is here}}
 
 template <int N>