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

  CONFLICT (content): Merge conflict in libclc/CMakeLists.txt
  CONFLICT (rename/delete): libclc/opencl/lib/generic/workitem/get_num_groups.cl renamed to libclc/clc/include/clc/common/clc_step.h in HEAD, but deleted in 19964c3.
  CONFLICT (modify/delete): libclc/clc/include/clc/common/clc_step.h deleted in 19964c3 and modified in HEAD.  Version HEAD of libclc/clc/include/clc/common/clc_step.h left in tree.
  CONFLICT (rename/delete): libclc/opencl/lib/generic/workitem/get_local_size.cl renamed to libclc/clc/include/clc/shared/clc_vload.h in HEAD, but deleted in 19964c3.
  CONFLICT (modify/delete): libclc/clc/include/clc/shared/clc_vload.h deleted in 19964c3 and modified in HEAD.  Version HEAD of libclc/clc/include/clc/shared/clc_vload.h left in tree.
  CONFLICT (rename/delete): libclc/opencl/lib/generic/workitem/get_local_id.cl renamed to libclc/clc/include/clc/shared/clc_vstore.h in HEAD, but deleted in 19964c3.
  CONFLICT (modify/delete): libclc/clc/include/clc/shared/clc_vstore.h deleted in 19964c3 and modified in HEAD.  Version HEAD of libclc/clc/include/clc/shared/clc_vstore.h left in tree.
  CONFLICT (rename/delete): libclc/opencl/lib/generic/workitem/get_group_id.cl renamed to libclc/clc/lib/generic/async/clc_prefetch.cl in HEAD, but deleted in 19964c3.
  CONFLICT (modify/delete): libclc/clc/lib/generic/async/clc_prefetch.cl deleted in 19964c3 and modified in HEAD.  Version HEAD of libclc/clc/lib/generic/async/clc_prefetch.cl left in tree.
  CONFLICT (rename/delete): libclc/opencl/lib/generic/workitem/get_work_dim.cl renamed to libclc/clc/lib/generic/common/clc_step.cl in HEAD, but deleted in 19964c3.
  CONFLICT (modify/delete): libclc/clc/lib/generic/common/clc_step.cl deleted in 19964c3 and modified in HEAD.  Version HEAD of libclc/clc/lib/generic/common/clc_step.cl left in tree.
  CONFLICT (rename/delete): libclc/opencl/lib/generic/workitem/get_global_offset.cl renamed to libclc/clc/lib/generic/common/clc_step.inc in HEAD, but deleted in 19964c3.
  CONFLICT (modify/delete): libclc/clc/lib/generic/common/clc_step.inc deleted in 19964c3 and modified in HEAD.  Version HEAD of libclc/clc/lib/generic/common/clc_step.inc left in tree.
  CONFLICT (content): Merge conflict in libclc/opencl/lib/generic/async/prefetch.inc
  CONFLICT (content): Merge conflict in libclc/opencl/lib/generic/shared/vload.cl
  CONFLICT (content): Merge conflict in libclc/opencl/lib/generic/shared/vstore.cl

Author: mikolaj-pirog

.github/workflows/sycl-rel-nightly.yml

@@ -26,7 +26,7 @@ jobs:
     steps:
     - uses: actions/checkout@v4
       with:
-        ref: ${{ inputs.testing_branch || 'sycl-rel-6_1_0'  }}
+        ref: ${{ inputs.testing_branch || 'sycl-rel-6_2'  }}
     - run: git show --quiet | tee -a $GITHUB_STEP_SUMMARY
 
     - id: is_new_commit
@@ -46,7 +46,7 @@ jobs:
       build_artifact_suffix: default
       build_configure_extra_args: '--disable-jit --no-assertions --add_security_flags=sanitize --hip --cuda'
       build_image: ghcr.io/intel/llvm/ubuntu2204_build:latest
-      build_ref: ${{ inputs.testing_branch || 'sycl-rel-6_1_0'  }}
+      build_ref: ${{ inputs.testing_branch || 'sycl-rel-6_2'  }}
       pack_release: 'true'
 
       # We upload the build for people to download/use, override its name and
@@ -94,7 +94,7 @@ jobs:
       target_devices: ${{ matrix.target_devices }}
       tests_selector: e2e
       extra_lit_opts: ${{ matrix.extra_lit_opts }}
-      repo_ref: ${{ inputs.testing_branch || 'sycl-rel-6_1_0'  }}
+      repo_ref: ${{ inputs.testing_branch || 'sycl-rel-6_2'  }}
       devops_ref: sycl
       sycl_toolchain_artifact: sycl_linux_default
       sycl_toolchain_archive: ${{ needs.ubuntu2204_build.outputs.artifact_archive_name }}
@@ -105,7 +105,7 @@ jobs:
     if: ${{ github.repository == 'intel/llvm' && needs.check_for_new_commits.outputs.is_new_commit != 'false' }}
     uses: ./.github/workflows/sycl-windows-build.yml
     with:
-      ref: ${{ inputs.testing_branch || 'sycl-rel-6_1_0'  }}
+      ref: ${{ inputs.testing_branch || 'sycl-rel-6_2'  }}
       build_configure_extra_args: '--disable-jit --no-assertions --add_security_flags=sanitize'
       pack_release: 'true'
 
@@ -137,7 +137,7 @@ jobs:
       target_devices: level_zero:gpu
       sycl_toolchain_archive: ${{ needs.build-win.outputs.artifact_archive_name }}
       extra_lit_opts: ${{ matrix.extra_lit_opts }}
-      repo_ref: ${{ inputs.testing_branch || 'sycl-rel-6_1_0'  }}
+      repo_ref: ${{ inputs.testing_branch || 'sycl-rel-6_2'  }}
       devops_ref: sycl
 
   cuda-aws-start:
@@ -147,7 +147,7 @@ jobs:
     secrets: inherit
     with:
       mode: start
-      ref: ${{ inputs.testing_branch || 'sycl-rel-6_1_0'  }}
+      ref: ${{ inputs.testing_branch || 'sycl-rel-6_2'  }}
 
   cuda-run-tests:
     needs: [ubuntu2204_build, cuda-aws-start]
@@ -159,7 +159,7 @@ jobs:
       image: ghcr.io/intel/llvm/ubuntu2204_build:latest
       image_options: -u 1001 --gpus all --cap-add SYS_ADMIN --env NVIDIA_DISABLE_REQUIRE=1
       target_devices: cuda:gpu
-      repo_ref: ${{ inputs.testing_branch || 'sycl-rel-6_1_0'  }}
+      repo_ref: ${{ inputs.testing_branch || 'sycl-rel-6_2'  }}
       devops_ref: sycl
 
       sycl_toolchain_artifact: sycl_linux_default
@@ -173,7 +173,7 @@ jobs:
     secrets: inherit
     with:
       mode: stop
-      ref: ${{ inputs.testing_branch || 'sycl-rel-6_1_0'  }}
+      ref: ${{ inputs.testing_branch || 'sycl-rel-6_2'  }}
 
   build-sycl-cts:
     needs: ubuntu2204_build
@@ -185,7 +185,7 @@ jobs:
       cts_testing_mode: 'build-only'
       image_options: -u 1001 --device=/dev/dri --privileged --cap-add SYS_ADMIN
       tests_selector: cts
-      repo_ref: ${{ inputs.testing_branch || 'sycl-rel-6_1_0'  }}
+      repo_ref: ${{ inputs.testing_branch || 'sycl-rel-6_2'  }}
       devops_ref: sycl
       tests_ref: 87f8677c4b67cfe56ad6f09246036f10c306c977
       sycl_toolchain_artifact: sycl_linux_default
@@ -217,7 +217,7 @@ jobs:
       image_options: ${{ matrix.image_options }}
       target_devices: ${{ matrix.target_devices }}
       tests_selector: cts
-      repo_ref: ${{ inputs.testing_branch || 'sycl-rel-6_1_0'  }}
+      repo_ref: ${{ inputs.testing_branch || 'sycl-rel-6_2'  }}
       devops_ref: sycl
       sycl_toolchain_artifact: sycl_linux_default
       sycl_toolchain_archive: ${{ needs.ubuntu2204_build.outputs.artifact_archive_name }}

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

@@ -164,7 +164,9 @@ jobs:
     - name: Test adapter specific
       env:
         ZE_ENABLE_LOADER_DEBUG_TRACE: 1
-        LIT_OPTS: "--timeout 120 --filter-out 'adapters/level_zero/memcheck.test'"
+        LIT_OPTS: "--timeout 120"
+        # These tests cause timeouts on CI
+        LIT_FILTER_OUT: "(adapters/level_zero/memcheck.test|adapters/level_zero/v2/deferred_kernel_memcheck.test)"
       run: cmake --build build -j $(nproc) -- check-unified-runtime-adapter
       # Don't run adapter specific tests when building multiple adapters
       if: ${{ matrix.adapter.other_name == '' }}

clang/lib/CodeGen/CodeGenModule.cpp

@@ -6902,7 +6902,8 @@ void CodeGenModule::EmitGlobalFunctionDefinition(GlobalDecl GD,
     if (GD.getKernelReferenceKind() == KernelReferenceKind::Stub &&
         !D->hasAttr<NoInlineAttr>() &&
         !Fn->hasFnAttribute(llvm::Attribute::NoInline) &&
-        !D->hasAttr<OptimizeNoneAttr>() &&
+        !D->hasAttr<OptimizeNoneAttr>() && !LangOpts.SYCLIsNativeCPU &&
+        !LangOpts.SYCLIsDevice &&
         !Fn->hasFnAttribute(llvm::Attribute::OptimizeNone) &&
         !ShouldAddOptNone) {
       Fn->addFnAttr(llvm::Attribute::AlwaysInline);

clang/lib/Sema/SemaDeclAttr.cpp

@@ -6961,10 +6961,8 @@ ProcessDeclAttribute(Sema &S, Scope *scope, Decl *D, const ParsedAttr &AL,
          Aux && AL.existsInTarget(*Aux)))) {
     if (AL.isRegularKeywordAttribute() || AL.isDeclspecAttribute()) {
       S.Diag(AL.getLoc(), AL.isRegularKeywordAttribute()
-               ? (unsigned)diag::err_keyword_not_supported_on_target
-           : AL.isDeclspecAttribute()
-               ? (unsigned)diag::warn_unhandled_ms_attribute_ignored
-               : (unsigned)diag::warn_unknown_attribute_ignored)
+                              ? diag::err_keyword_not_supported_on_target
+                              : diag::warn_unhandled_ms_attribute_ignored)
           << AL.getAttrName() << AL.getRange();
     } else {
       S.DiagnoseUnknownAttribute(AL);

clang/test/SemaSYCL/intel-fpga-local.cpp

@@ -9,7 +9,7 @@ sycl::queue deviceQueue;
 void diagnostics()
 {
   // **doublepump
-  //expected-warning@+1 {{unknown attribute 'doublepump' ignored}}
+  //expected-warning@+1 {{unknown attribute 'intelfpga::doublepump' ignored}}
   [[intelfpga::doublepump]] unsigned int doublepump_var[64];
 
   //expected-error@+2{{attributes are not compatible}}
@@ -29,7 +29,7 @@ void diagnostics()
   unsigned int dpump_reg[64];
 
   // **singlepump
-  //expected-warning@+1 {{unknown attribute 'singlepump' ignored}}
+  //expected-warning@+1 {{unknown attribute 'intelfpga::singlepump' ignored}}
   [[intelfpga::singlepump]] unsigned int singlepump_var[64];
 
   //expected-error@+1{{attributes are not compatible}}
@@ -48,7 +48,7 @@ void diagnostics()
   unsigned int spump_reg[64];
 
   // **fpga_register
-  //expected-warning@+1 {{unknown attribute 'register' ignored}}
+  //expected-warning@+1 {{unknown attribute 'intelfpga::register' ignored}}
   [[intelfpga::register]] unsigned int reg_var[64];
 
   //expected-warning@+2{{attribute 'fpga_register' is already applied}}
@@ -155,7 +155,7 @@ void diagnostics()
   [[intel::fpga_memory]]
   [[intel::fpga_memory]] unsigned int mem_mem[64]; // OK
 
-  //expected-warning@+1 {{unknown attribute 'memory' ignored}}
+  //expected-warning@+1 {{unknown attribute 'intelfpga::memory' ignored}}
   [[intelfpga::memory]] unsigned int memory_var[64];
 
   // Check to see if there's a duplicate attribute with different values
@@ -189,7 +189,7 @@ void diagnostics()
   [[intel::fpga_memory]] unsigned int mem_mlabs_block_ram[64];
 
   // **bankwidth
-  //expected-warning@+1 {{unknown attribute 'bankwidth' ignored}}
+  //expected-warning@+1 {{unknown attribute 'intelfpga::bankwidth' ignored}}
   [[intelfpga::bankwidth(4)]] unsigned int bankwidth_var[32];
 
   // Checking of incompatible attributes.
@@ -200,7 +200,7 @@ void diagnostics()
   unsigned int bankwidth_reg[64];
 
   // **max_replicates
-  //expected-warning@+1 {{unknown attribute 'max_replicates' ignored}}
+  //expected-warning@+1 {{unknown attribute 'intelfpga::max_replicates' ignored}}
   [[intelfpga::max_replicates(2)]] unsigned int max_replicates_var[64];
 
   // Checking of different argument values.
@@ -223,7 +223,7 @@ void diagnostics()
   //expected-error@+1{{'simple_dual_port' attribute takes no arguments}}
   [[intel::simple_dual_port(0)]] unsigned int sdp[64];
 
-  //expected-warning@+1 {{unknown attribute 'simple_dual_port' ignored}}
+  //expected-warning@+1 {{unknown attribute 'intelfpga::simple_dual_port' ignored}}
   [[intelfpga::simple_dual_port]] unsigned int dual_port_var[64];
 
   //expected-note@+1 {{conflicting attribute is here}}
@@ -259,7 +259,7 @@ void diagnostics()
   [[intel::bankwidth(0)]] unsigned int bw_zero[64];
 
   // **private_copies
-  //expected-warning@+1 {{unknown attribute 'private_copies' ignored}}
+  //expected-warning@+1 {{unknown attribute 'intelfpga::private_copies' ignored}}
   [[intelfpga::private_copies(8)]] unsigned int private_copies_var[64];
 
   // Checking of incompatible attributes.
@@ -316,11 +316,11 @@ void diagnostics()
   //expected-error@+1{{requires a positive integral compile time constant expression}}
   [[intel::numbanks(0)]] unsigned int nb_zero[64];
 
-  //expected-warning@+1 {{unknown attribute 'numbanks' ignored}}
+  //expected-warning@+1 {{unknown attribute 'intelfpga::numbanks' ignored}}
   [[intelfpga::numbanks(8)]] unsigned int numbanks_var[32];
 
   // **merge
-  //expected-warning@+1 {{unknown attribute 'merge' ignored}}
+  //expected-warning@+1 {{unknown attribute 'intelfpga::merge' ignored}}
   [[intelfpga::merge("mrg1", "depth")]] unsigned int merge_depth_var[64];
 
   //expected-error@+2{{attributes are not compatible}}
@@ -386,11 +386,11 @@ void diagnostics()
   //expected-error@+1{{'bank_bits' attribute requires a non-negative integral compile time constant expression}}
   [[intel::bank_bits(-1)]] unsigned int bb_negative_arg[4];
 
-  //expected-warning@+1 {{unknown attribute 'bank_bits' ignored}}
+  //expected-warning@+1 {{unknown attribute 'intelfpga::bank_bits' ignored}}
   [[intelfpga::bank_bits(2, 3, 4, 5)]] unsigned int bankbits_var[64];
 
   // **force_pow2_depth
-  //expected-warning@+1 {{unknown attribute 'force_pow2_depth' ignored}}
+  //expected-warning@+1 {{unknown attribute 'intelfpga::force_pow2_depth' ignored}}
   [[intelfpga::force_pow2_depth(0)]] unsigned int arr_force_p2d_0_var[64];
 
   //expected-error@+1{{'force_pow2_depth' attribute requires integer constant value 0 or 1}}
@@ -423,16 +423,16 @@ void check_gnu_style() {
   //expected-warning@+1{{unknown attribute 'register' ignored}}
   int __attribute__((register)) reg;
 
-  //expected-warning@+1{{unknown attribute '__singlepump__' ignored}}
+  //expected-warning@+1{{unknown attribute 'singlepump' ignored}}
   unsigned int __attribute__((__singlepump__)) singlepump;
 
-  //expected-warning@+1{{unknown attribute '__doublepump__' ignored}}
+  //expected-warning@+1{{unknown attribute 'doublepump' ignored}}
   unsigned int __attribute__((__doublepump__)) doublepump;
 
-  //expected-warning@+1{{unknown attribute '__private_copies__' ignored}}
+  //expected-warning@+1{{unknown attribute 'private_copies' ignored}}
   int __attribute__((__private_copies__(4))) private_copies;
 
-  //expected-warning@+1{{unknown attribute '__merge__' ignored}}
+  //expected-warning@+1{{unknown attribute 'merge' ignored}}
   int __attribute__((__merge__("mrg1","depth"))) merge;
 
   //expected-warning@+1{{unknown attribute 'max_replicates' ignored}}

clang/test/SemaSYCL/intel-fpga-no-global-work-offset.cpp

@@ -9,7 +9,7 @@
 
 [[intel::no_global_work_offset(0, 1)]] void test1() {} // expected-error{{'no_global_work_offset' attribute takes no more than 1 argument}}
 
-[[intelfpga::no_global_work_offset]] void RemovedSpell(); // expected-warning {{unknown attribute 'no_global_work_offset' ignored}}
+[[intelfpga::no_global_work_offset]] void RemovedSpell(); // expected-warning {{unknown attribute 'intelfpga::no_global_work_offset' ignored}}
 
 // Test that checks wrong function template instantiation and ensures that the type
 // is checked properly when instantiating from the template definition.

clang/test/SemaSYCL/intel-max-global-work-dim-device.cpp

@@ -8,7 +8,7 @@ using namespace sycl;
 queue q;
 
 struct Func {
-  // expected-warning@+1 {{unknown attribute 'max_global_work_dim' ignored}}
+  // expected-warning@+1 {{unknown attribute 'intelfpga::max_global_work_dim' ignored}}
   [[intelfpga::max_global_work_dim(2)]] void operator()() const {}
 };
 

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

@@ -141,6 +141,6 @@ struct DAFuncObj {
 };
 
 struct Func {
-  // expected-warning@+1 {{unknown attribute 'max_work_group_size' ignored}}
+  // expected-warning@+1 {{unknown attribute 'intelfpga::max_work_group_size' ignored}}
   [[intelfpga::max_work_group_size(1, 1, 1)]] void operator()() const {}
 };

clang/test/SemaSYCL/num_simd_work_items.cpp

@@ -7,7 +7,7 @@
 using namespace sycl;
 queue q;
 
-// expected-warning@+1 {{unknown attribute 'num_simd_work_items' ignored}}
+// expected-warning@+1 {{unknown attribute 'intelfpga::num_simd_work_items' ignored}}
 [[intelfpga::num_simd_work_items(22)]] void RemoveSpelling();
 
 // No diagnostic is emitted because the arguments match.

clang/test/SemaSYCL/scheduler_target_fmax_mhz.cpp

@@ -3,7 +3,7 @@
 // Test that checks scheduler_target_fmax_mhz attribute support on Function.
 
 // Test for removed deprecated spelling of scheduler_target_fmax_mhz attribute.
-// expected-warning@+1 {{unknown attribute 'scheduler_target_fmax_mhz' ignored}}
+// expected-warning@+1 {{unknown attribute 'intelfpga::scheduler_target_fmax_mhz' ignored}}
 [[intelfpga::scheduler_target_fmax_mhz(2)]] void deprecate() {}
 
 // Tests for incorrect argument values for Intel FPGA scheduler_target_fmax_mhz function attribute.