[SYCL] Fix linkage adjustment of kernels

clang/lib/AST/ASTContext.cpp

@@ -12824,7 +12824,8 @@ static GVALinkage adjustGVALinkageForAttributes(const ASTContext &Context,
     if (Context.shouldExternalize(D))
       return GVA_StrongExternal;
   } else if (Context.getLangOpts().SYCLIsDevice &&
-             D->hasAttr<DeviceKernelAttr>()) {
+             (D->hasAttr<DeviceKernelAttr>() &&
+              D->getAttr<DeviceKernelAttr>()->isImplicit())) {
     if (L == GVA_DiscardableODR)
       return GVA_StrongODR;
   }

clang/lib/Sema/SemaSYCL.cpp

@@ -5443,9 +5443,13 @@ void SemaSYCL::ConstructOpenCLKernel(FunctionDecl *KernelCallerFunc,
   ESIMDKernelDiagnostics esimdKernel(*this, KernelObj->getLocation(),
                                      IsSIMDKernel);
 
-  SyclKernelDeclCreator kernel_decl(*this, KernelObj->getLocation(),
-                                    KernelCallerFunc->isInlined(), IsSIMDKernel,
-                                    KernelCallerFunc);
+  // In case of syntax errors in input programs we are not able to access
+  // CallOperator. In this case the value of IsInlined doesn't matter, because
+  // compilation will fail with errors anyways.
+  const bool IsInlined =
+      CallOperator ? CallOperator->isInlined() : /* placeholder */ false;
+  SyclKernelDeclCreator kernel_decl(*this, KernelObj->getLocation(), IsInlined,
+                                    IsSIMDKernel, KernelCallerFunc);
   SyclKernelBodyCreator kernel_body(*this, kernel_decl, KernelObj,
                                     KernelCallerFunc, IsSIMDKernel,
                                     CallOperator);

clang/test/CodeGenSYCL/device_has.cpp

@@ -6,7 +6,7 @@
 using namespace sycl;
 queue q;
 
-// CHECK-DAG: define dso_local spir_kernel void @{{.*}}kernel_name_1{{.*}} !sycl_declared_aspects ![[ASPECTS1:[0-9]+]] !srcloc ![[SRCLOC1:[0-9]+]]
+// CHECK-DAG: define {{.*}}spir_kernel void @{{.*}}kernel_name_1{{.*}} !sycl_declared_aspects ![[ASPECTS1:[0-9]+]] !srcloc ![[SRCLOC1:[0-9]+]]
 
 // CHECK-DAG: define {{.*}}spir_func void @{{.*}}func1{{.*}} !sycl_declared_aspects ![[ASPECTS1]] !srcloc ![[SRCLOC2:[0-9]+]] {
 [[sycl::device_has(sycl::aspect::cpu)]] void func1() {}
@@ -67,7 +67,7 @@ void foo() {
   q.submit([&](handler &h) {
     KernelFunctor f1;
     h.single_task<class kernel_name_1>(f1);
-    // CHECK-DAG: define dso_local spir_kernel void @{{.*}}kernel_name_2{{.*}} !sycl_declared_aspects ![[ASPECTS4:[0-9]+]] !srcloc ![[SRCLOC8:[0-9]+]]
+    // CHECK-DAG: define {{.*}}spir_kernel void @{{.*}}kernel_name_2{{.*}} !sycl_declared_aspects ![[ASPECTS4:[0-9]+]] !srcloc ![[SRCLOC8:[0-9]+]]
     h.single_task<class kernel_name_2>([]() [[sycl::device_has(sycl::aspect::gpu)]] {});
   });
 }

clang/test/CodeGenSYCL/dynamic_local_accessor.cpp

@@ -7,7 +7,7 @@
 // The first two RUN commands verify that the init call is generated with the correct arguments in LLVM IR
 // and the second two RUN commands verify the contents of the integration header produced by the frontend.
 //
-// CHECK-IR: define dso_local spir_kernel void @
+// CHECK-IR: define {{.*}}spir_kernel void @
 // CHECK-IR-SAME: ptr addrspace(3) noundef align 4 [[PTR:%[a-zA-Z0-9_]+]]
 //
 // CHECK-IR: [[PTR]].addr = alloca ptr addrspace(3), align 8

clang/test/CodeGenSYCL/dynamic_work_group_memory.cpp

@@ -7,7 +7,7 @@
 // The first two RUN commands verify that the init call is generated with the correct arguments in LLVM IR
 // and the second two RUN commands verify the contents of the integration header produced by the frontend.
 //
-// CHECK-IR: define dso_local spir_kernel void @
+// CHECK-IR: define {{.*}}spir_kernel void @
 // CHECK-IR-SAME: ptr addrspace(3) noundef align 4 [[PTR:%[a-zA-Z0-9_]+]]
 //
 // CHECK-IR: [[PTR]].addr = alloca ptr addrspace(3), align 8

clang/test/CodeGenSYCL/generated-types-initialization.cpp

@@ -38,7 +38,7 @@ int main() {
   });
   return 0;
 }
-// CHECK: define dso_local spir_kernel void @{{.*}}basic(ptr noundef byval(%struct.__generated_B) align 8 %_arg_Obj)
+// CHECK: define {{.*}}spir_kernel void @{{.*}}basic(ptr noundef byval(%struct.__generated_B) align 8 %_arg_Obj)
 //
 // Kernel object clone.
 // CHECK: %[[K:[a-zA-Z0-9_.]+]] = alloca %class.anon
@@ -54,7 +54,7 @@ int main() {
 // Kernel body call.
 // CHECK: call spir_func void @_ZZZ4mainENKUlRN4sycl3_V17handlerEE_clES2_ENKUlvE_clEv(ptr addrspace(4) noundef align 8 dereferenceable_or_null(16) %[[K_as_cast]])
 
-// CHECK: define dso_local spir_kernel void @{{.*}}nns(ptr noundef byval(%struct.__generated_B.0) align 8 %_arg_NNSObj)
+// CHECK: define {{.*}}spir_kernel void @{{.*}}nns(ptr noundef byval(%struct.__generated_B.0) align 8 %_arg_NNSObj)
 //
 // Kernel object clone.
 // CHECK: %[[NNSK:[a-zA-Z0-9_.]+]] = alloca %class.anon.2

clang/test/CodeGenSYCL/kernel-handler.cpp

@@ -22,7 +22,7 @@ void test(int val) {
   });
 }
 
-// ALL: define dso_local{{ spir_kernel | ptx_kernel | }}void @{{.*}}test_kernel_handler{{[^(]*}}
+// ALL: define {{.*}}{{ spir_kernel | ptx_kernel | }}void @{{.*}}test_kernel_handler{{[^(]*}}
 // NONATIVESUPPORT-SAME: (ptr noundef byval(%class.anon) align 4 %_arg__sycl_functor, ptr addrspace(1) noundef align 1 %_arg__specialization_constants_buffer)
 // NATIVESUPPORT-SAME: (i32 noundef %_arg_a, ptr addrspace(1) noundef align 1 %_arg__specialization_constants_buffer)
 // ALL: %kh = alloca %"class.sycl::_V1::kernel_handler", align 1

clang/test/CodeGenSYCL/kernel-op-calls.cpp

@@ -14,46 +14,46 @@ class Functor1 {
  [[sycl::reqd_sub_group_size(4)]] void operator()(sycl::id<1> id) const {}
 
  [[sycl::work_group_size_hint(1, 2, 3)]] void operator()(sycl::id<2> id) const {}
-
 };
 
 class ESIMDFunctor {
 public:
- ESIMDFunctor(){}
+  ESIMDFunctor(){}
 
   [[intel::sycl_explicit_simd]] void operator()(sycl::id<2> id) const {}
 
-  [[sycl::work_group_size_hint(1, 2, 3)]][[intel::sycl_explicit_simd]] void operator()(sycl::id<1> id) const {}
-
+  [[sycl::work_group_size_hint(1, 2, 3)]] [[intel::sycl_explicit_simd]]
+  void operator()(sycl::id<1> id) const {}
 };
 
 // Check templated 'operator()()' call works.
 class kernels {
-public:  
+public:
   kernels(){}
 
-    template<int Dimensions = 1>
-    [[sycl::work_group_size_hint(1, 2, 3)]] void operator()(sycl::id<Dimensions> item) const {}
-    
+  template<int Dimensions = 1>
+  [[sycl::work_group_size_hint(1, 2, 3)]]
+  void operator()(sycl::id<Dimensions> item) const {}
 };
 
 int main() {
 
   Q.submit([&](sycl::handler& cgh) {
       Functor1 F;
-      // CHECK: define dso_local spir_kernel void @_ZTS8Functor1() {{.*}} !kernel_arg_buffer_location !{{[0-9]+}} !intel_reqd_sub_group_size !{{[0-9]+}} !sycl_fixed_targets !{{[0-9]+}} {
+      // CHECK: define weak_odr spir_kernel void @_ZTS8Functor1() {{.*}} !intel_reqd_sub_group_size
       cgh.parallel_for(sycl::range<1>(10), F);
     });
 
   Q.submit([&](sycl::handler& cgh) {
       kernels K;
-      // CHECK: define dso_local spir_kernel void @_ZTS7kernels() {{.*}} !kernel_arg_buffer_location !{{[0-9]+}} !work_group_size_hint !{{[0-9]+}} !sycl_fixed_targets !{{[0-9]+}} {
+      // CHECK: define weak_odr spir_kernel void @_ZTS7kernels() {{.*}} !work_group_size_hint !{{[0-9]+}}
       cgh.parallel_for(sycl::range<1>(10), K);
     });
 
   Q.submit([&](sycl::handler& cgh) {
       ESIMDFunctor EF;
-      // CHECK: define dso_local spir_kernel void @_ZTS12ESIMDFunctor() {{.*}} !sycl_explicit_simd !{{[0-9]+}} !intel_reqd_sub_group_size !{{[0-9]+}} !work_group_size_hint !{{[0-9]+}} !kernel_arg_accessor_ptr !{{[0-9]+}} !sycl_fixed_targets !{{[0-9]+}} {
+      // CHECK: define weak_odr spir_kernel void @_ZTS12ESIMDFunctor() {{.*}} !work_group_size_hint
+      // CHECK-SAME: !sycl_explicit_simd
       cgh.parallel_for(sycl::range<1>(10), EF);
     });
 

clang/test/CodeGenSYCL/kernel-param-acc-array.cpp

@@ -24,7 +24,7 @@ int main() {
         acc[1].use();
       });
 }
-// CHECK-LABEL: define dso_local spir_kernel void @_ZTSZ4mainE8kernel_A(
+// CHECK-LABEL: define {{.*}}spir_kernel void @_ZTSZ4mainE8kernel_A(
 // CHECK-SAME: ptr addrspace(1) noundef align 4 [[_ARG_ACC:%.*]], ptr noundef byval(%"struct.sycl::_V1::range") align 4 [[_ARG_ACC1:%.*]], ptr noundef byval(%"struct.sycl::_V1::range") align 4 [[_ARG_ACC2:%.*]], ptr noundef byval(%"struct.sycl::_V1::id") align 4 [[_ARG_ACC3:%.*]], ptr addrspace(1) noundef align 4 [[_ARG_ACC4:%.*]], ptr noundef byval(%"struct.sycl::_V1::range") align 4 [[_ARG_ACC6:%.*]], ptr noundef byval(%"struct.sycl::_V1::range") align 4 [[_ARG_ACC7:%.*]], ptr noundef byval(%"struct.sycl::_V1::id") align 4 [[_ARG_ACC8:%.*]]) #[[ATTR0:[0-9]+]]
 // CHECK-NEXT:  [[ENTRY:.*:]]
 // CHECK-NEXT:    [[_ARG_ACC_ADDR:%.*]] = alloca ptr addrspace(1), align 8

clang/test/CodeGenSYCL/kernel-param-member-acc-array.cpp

@@ -28,7 +28,7 @@ int main() {
       });
 }
 
-// CHECK-LABEL: define dso_local spir_kernel void @_ZTSZ4mainE8kernel_C(
+// CHECK-LABEL: define {{.*}}spir_kernel void @_ZTSZ4mainE8kernel_C(
 // CHECK-SAME: ptr addrspace(1) noundef align 4 [[_ARG_MEMBER_ACC:%.*]], ptr noundef byval(%"struct.sycl::_V1::range") align 4 [[_ARG_MEMBER_ACC1:%.*]], ptr noundef byval(%"struct.sycl::_V1::range") align 4 [[_ARG_MEMBER_ACC2:%.*]], ptr noundef byval(%"struct.sycl::_V1::id") align 4 [[_ARG_MEMBER_ACC3:%.*]], ptr addrspace(1) noundef align 4 [[_ARG_MEMBER_ACC4:%.*]], ptr noundef byval(%"struct.sycl::_V1::range") align 4 [[_ARG_MEMBER_ACC6:%.*]], ptr noundef byval(%"struct.sycl::_V1::range") align 4 [[_ARG_MEMBER_ACC7:%.*]], ptr noundef byval(%"struct.sycl::_V1::id") align 4 [[_ARG_MEMBER_ACC8:%.*]]) #[[ATTR0:[0-9]+]]
 // CHECK-NEXT:  [[ENTRY:.*:]]
 // CHECK-NEXT:    [[_ARG_MEMBER_ACC_ADDR:%.*]] = alloca ptr addrspace(1), align 8

clang/test/CodeGenSYCL/kernel-simple-instrumentation.cpp

@@ -3,7 +3,7 @@
 
 // RUN: %clang_cc1 -fsycl-is-device -fsycl-instrument-device-code -triple spir64-unknown-unknown -emit-llvm %s -o - | FileCheck %s
 
-// CHECK: kernel_function
+// CHECK: define {{.*}}kernel_function
 // CHECK-NEXT: entry:
 // CHECK-NEXT: call spir_func void @__itt_offload_wi_start_wrapper()
 // CHECK: call spir_func void @__itt_offload_wi_finish_wrapper()

clang/test/CodeGenSYCL/launch_bounds_nvptx.cpp

@@ -44,9 +44,9 @@ int main() {
   return 0;
 }
 
-// CHECK: define dso_local ptx_kernel void @{{.*}}kernel_name1() #0 {{.*}} !min_work_groups_per_cu ![[MWGPC:[0-9]+]] !max_work_groups_per_mp ![[MWGPM:[0-9]+]] !max_work_group_size ![[MWGS:[0-9]+]]
-// CHECK: define dso_local ptx_kernel void @{{.*}}kernel_name2() #0 {{.*}} !min_work_groups_per_cu ![[MWGPC:[0-9]+]] !max_work_groups_per_mp ![[MWGPM:[0-9]+]] !max_work_group_size ![[MWGS:[0-9]+]]
-// CHECK: define dso_local ptx_kernel void @{{.*}}kernel_name3() #0 {{.*}} !min_work_groups_per_cu ![[MWGPC_MWGPM:[0-9]+]] !max_work_groups_per_mp ![[MWGPC_MWGPM]] !max_work_group_size ![[MWGS_2:[0-9]+]]
+// CHECK: define {{.*}}ptx_kernel void @{{.*}}kernel_name1() #0 {{.*}} !min_work_groups_per_cu ![[MWGPC:[0-9]+]] !max_work_groups_per_mp ![[MWGPM:[0-9]+]] !max_work_group_size ![[MWGS:[0-9]+]]
+// CHECK: define {{.*}}ptx_kernel void @{{.*}}kernel_name2() #0 {{.*}} !min_work_groups_per_cu ![[MWGPC:[0-9]+]] !max_work_groups_per_mp ![[MWGPM:[0-9]+]] !max_work_group_size ![[MWGS:[0-9]+]]
+// CHECK: define {{.*}}ptx_kernel void @{{.*}}kernel_name3() #0 {{.*}} !min_work_groups_per_cu ![[MWGPC_MWGPM:[0-9]+]] !max_work_groups_per_mp ![[MWGPC_MWGPM]] !max_work_group_size ![[MWGS_2:[0-9]+]]
 
 // CHECK: ![[MWGPC]] = !{i32 2}
 // CHECK: ![[MWGPM]] = !{i32 4}

clang/test/CodeGenSYCL/max-concurrency.cpp

@@ -41,7 +41,7 @@
 // CHECK: %this1 = load ptr addrspace(4), ptr addrspace(4) [[ADDR1_CAST]], align 8
 // CHECK: ret void
 
-// CHECK: define dso_local spir_kernel void @_ZTSZZ4mainENKUlRN4sycl3_V17handlerEE_clES2_E12kernel_name5()
+// CHECK: define {{.*}}spir_kernel void @_ZTSZZ4mainENKUlRN4sycl3_V17handlerEE_clES2_E12kernel_name5()
 // CHECK: entry:
 // CHECK: [[H1:%.*]] = alloca [[H:%.*]], align 1
 // CHECK: [[H2:%.*]] = addrspacecast ptr [[H1]] to ptr addrspace(4)

clang/test/CodeGenSYCL/odr-kernel.cpp

@@ -0,0 +1,78 @@
+// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -triple spir64-unknown-unknown -sycl-std=2020 -emit-llvm -o - %s | FileCheck %s
+//
+// Kernel definition may be shared by multiple translation unit if a kernel is
+// defined as a functor in a header file. Therefore, we need to make sure that
+// the linkage for emitted kernel is correct, i.e. it allows to merge the same
+// symbols without triggering multiple definitions error.
+
+#include "sycl.hpp"
+
+// CHECK-DAG: define weak_odr spir_kernel void @_ZTS13FunctorInline
+// CHECK-DAG: define weak_odr spir_kernel void @_ZTS14FunctorInline2
+// CHECK-DAG: define dso_local spir_kernel void @_ZTS15FunctorNoInline
+// CHECK-DAG: define dso_local spir_kernel void @_ZTSZ4mainE10KernelName
+// CHECK-DAG: define dso_local spir_kernel void @_Z32__sycl_kernel_FreeFunctionKernelv
+// CHECK-DAG: define weak_odr spir_kernel void @_Z38__sycl_kernel_FreeFunctionKernelInlinev
+
+class FunctorInline {
+public:
+  void operator()(sycl::id<1>) const {}
+};
+
+class FunctorInline2 {
+public:
+  void operator()(sycl::id<1>) const;
+};
+inline void FunctorInline2::operator()(sycl::id<1>) const {}
+
+class FunctorNoInline {
+public:
+  void operator()(sycl::id<1>) const;
+};
+void FunctorNoInline::operator()(sycl::id<1>) const {}
+
+class FunctorNoInline2 {
+public:
+  void operator()() const;
+};
+void FunctorNoInline2::operator()() const {}
+
+
+[[__sycl_detail__::add_ir_attributes_function("sycl-nd-range-kernel", 2)]]
+void FreeFunctionKernel() {}
+
+[[__sycl_detail__::add_ir_attributes_function("sycl-nd-range-kernel", 2)]]
+inline void FreeFunctionKernelInline() {}
+
+
+struct KernelLaunchWrapper {
+  template <typename KernelName, typename KernelType>
+  __attribute__((sycl_kernel))
+  static void kernel_single_task(const KernelType &kernelFunc) {
+    kernelFunc();
+  }
+};
+
+int main() {
+  sycl::queue q;
+
+  q.submit([&](sycl::handler &cgh) {
+    FunctorInline f;
+    cgh.parallel_for(sycl::range<1>(1024), f);
+  });
+
+  q.submit([&](sycl::handler &cgh) {
+    FunctorInline2 f;
+    cgh.parallel_for(sycl::range<1>(1024), f);
+  });
+
+  q.submit([&](sycl::handler &cgh) {
+    FunctorNoInline f;
+    cgh.parallel_for(sycl::range<1>(1024), f);
+  });
+
+  {
+    FunctorNoInline2 f;
+    KernelLaunchWrapper::kernel_single_task<class KernelName>(f);
+  }
+}

clang/test/CodeGenSYCL/pipeline_kernel.cpp

@@ -29,8 +29,8 @@ int main() {
   return 0;
 }
 
-// CHECK: define dso_local spir_kernel void @{{.*}}test_kernel1() #0 {{.*}} !pipeline_kernel ![[NUM5:[0-9]+]]
-// CHECK: define dso_local spir_kernel void @{{.*}}test_kernel2() #0 {{.*}} ![[NUM4:[0-9]+]]
-// CHECK: define dso_local spir_kernel void @{{.*}}test_kernel3() #0 {{.*}} !pipeline_kernel ![[NUM5]]
+// CHECK: define weak_odr spir_kernel void @{{.*}}test_kernel1() #0 {{.*}} !pipeline_kernel ![[NUM5:[0-9]+]]
+// CHECK: define weak_odr spir_kernel void @{{.*}}test_kernel2() #0 {{.*}} ![[NUM4:[0-9]+]]
+// CHECK: define weak_odr spir_kernel void @{{.*}}test_kernel3() #0 {{.*}} !pipeline_kernel ![[NUM5]]
 // CHECK: ![[NUM4]] = !{}
 // CHECK: ![[NUM5]] = !{i32 0}

clang/test/CodeGenSYCL/reqd-sub-group-size-spirv-intrin.cpp

@@ -20,7 +20,7 @@ int main() {
   return 0;
 }
 
-// CHECK: define dso_local spir_kernel void @{{.*}}main{{.*}}kernel_class() {{.*}} !intel_reqd_sub_group_size ![[SUBGROUPSIZE:[0-9]+]]
+// CHECK: define {{.*}}spir_kernel void @{{.*}}main{{.*}}kernel_class() {{.*}} !intel_reqd_sub_group_size ![[SUBGROUPSIZE:[0-9]+]]
 // CHECK: call spir_func void @{{.*}}__spirv_ControlBarrier{{.*}}({{.*}})
 
 // CHECK: declare spir_func void @{{.*}}__spirv_ControlBarrier{{.*}}({{.*}})

clang/test/CodeGenSYCL/stream.cpp

@@ -4,7 +4,7 @@
 // CHECK: %[[RANGE_TYPE:"struct.*sycl::_V1::range"]]
 // CHECK: %[[ID_TYPE:"struct.*sycl::_V1::id"]]
 
-// CHECK: define dso_local spir_kernel void @{{.*}}StreamTester
+// CHECK: define {{.*}}spir_kernel void @{{.*}}StreamTester
 // CHECK-SAME: ptr addrspace(1) noundef align 1 [[ACC_DATA:%[a-zA-Z0-9_]+]],
 // CHECK-SAME: ptr noundef byval(%[[RANGE_TYPE]]) align 4 [[ACC_RANGE1:%[a-zA-Z0-9_]+]],
 // CHECK-SAME: ptr noundef byval(%[[RANGE_TYPE]]) align 4 [[ACC_RANGE2:%[a-zA-Z0-9_]+]],

clang/test/CodeGenSYCL/sub-group-size.cpp

@@ -23,7 +23,7 @@ void default_behavior() {
   kernel_single_task<class Kernel1>([]() {
   });
 }
-// NONE-DAG: define {{.*}}spir_kernel void @{{.*}}Kernel1() #{{[0-9]+}} !kernel_arg_buffer_location !{{[0-9]+}} !sycl_fixed_targets !{{[0-9]+}} {
+// NONE-DAG: define {{.*}}spir_kernel void @{{.*}}Kernel1() #{{[0-9]+}}{{.*}} !kernel_arg_buffer_location !{{[0-9]+}} !sycl_fixed_targets !{{[0-9]+}} {
 // PRIM_DEF-DAG: define {{.*}}spir_kernel void @{{.*}}Kernel1() #{{[0-9]+}}{{.*}} !intel_reqd_sub_group_size ![[PRIMARY]]
 // TEN_DEF-DAG: define {{.*}}spir_kernel void @{{.*}}Kernel1() #{{[0-9]+}}{{.*}} !intel_reqd_sub_group_size ![[TEN]]
 

clang/test/CodeGenSYCL/sycl-intelfpga-field.cpp

@@ -20,7 +20,7 @@
 struct fooA {
     int *p;
 
-// CHECK: define dso_local spir_kernel void @_ZTS4fooA(ptr addrspace(1) {{.*}}%[[ARG:.*]])
+// CHECK: define {{.*}}spir_kernel void @_ZTS4fooA(ptr addrspace(1) {{.*}}%[[ARG:.*]])
 // CHECK: %[[ARG_ADDR:.*]] = alloca ptr addrspace(1), align 8
 // CHECK: %[[ARG_ADDR_AS_CAST:.*]] = addrspacecast ptr %[[ARG_ADDR]] to ptr addrspace(4)
 // CHECK: store ptr addrspace(1) %[[ARG]], ptr addrspace(4) %[[ARG_ADDR_AS_CAST]], align 8
@@ -39,7 +39,7 @@ struct fooA {
 struct fooB {
     float f;
 
-// CHECK: define dso_local spir_kernel void @_ZTS4fooB({{.*}}%[[ARG:.*]])
+// CHECK: define {{.*}}spir_kernel void @_ZTS4fooB({{.*}}%[[ARG:.*]])
 // CHECK: %[[ARG_ADDR:.*]] = alloca float, align 4
 // CHECK: %[[ARG_ADDR_AS_CAST:.*]] = addrspacecast ptr %[[ARG_ADDR]] to ptr addrspace(4)
 // CHECK: store float %[[ARG]], ptr addrspace(4) %[[ARG_ADDR_AS_CAST]], align 4
@@ -58,7 +58,7 @@ struct bar {
 struct fooC {
     bar b;
 
-// CHECK: define dso_local spir_kernel void @_ZTS4fooC({{.*}}%[[ARG:.*]])
+// CHECK: define {{.*}}spir_kernel void @_ZTS4fooC({{.*}}%[[ARG:.*]])
 // CHECK: %[[ARG_AS_CAST:.*]] = addrspacecast ptr %[[ARG]] to ptr addrspace(4)
 // CHECK: %[[GEP:.*]] = getelementptr inbounds
 // CHECK: call void @llvm.memcpy.p4.p4.i64(ptr addrspace(4) align 1 %[[GEP]], ptr addrspace(4) align 1 %[[ARG_AS_CAST]], i64 1, i1 false)
@@ -71,7 +71,7 @@ struct fooD {
     [[clang::annotate("my_ann_1")]]
     int n;
 
-// CHECK: define dso_local spir_kernel void @_ZTS4fooD(i32 {{.*}}%[[ARG:.*]])
+// CHECK: define {{.*}}spir_kernel void @_ZTS4fooD(i32 {{.*}}%[[ARG:.*]])
 // CHECK: %[[ARG_ADDR:.*]] = alloca i32, align 4
 // CHECK: %[[ARG_ADDR_AS_CAST:.*]] = addrspacecast ptr %[[ARG_ADDR]] to ptr addrspace(4)
 // CHECK: store i32 %[[ARG]], ptr addrspace(4) %[[ARG_ADDR_AS_CAST]], align 4

clang/test/CodeGenSYCL/work_group_memory.cpp

@@ -7,7 +7,7 @@
 // The first two RUN commands verify that the init call is generated with the correct arguments in LLVM IR
 // and the second two RUN commands verify the contents of the integration header produced by the frontend.
 //
-// CHECK-IR: define dso_local spir_kernel void @
+// CHECK-IR: define {{.*}}spir_kernel void @
 // CHECK-IR-SAME: ptr addrspace(3) noundef align 4 [[PTR:%[a-zA-Z0-9_]+]]
 //
 // CHECK-IR: [[PTR]].addr = alloca ptr addrspace(3), align 8

sycl/test-e2e/Matrix/joint_matrix_bf16_fill_k_cache_unroll.cpp

@@ -6,8 +6,6 @@
 //
 //===----------------------------------------------------------------------===//
 // REQUIRES: target-spir
-// XFAIL: (arch-intel_gpu_acm_g10 || arch-intel_gpu_pvc || arch-intel_gpu_bmg_g21)
-// XFAIL-TRACKER: CMPLRLLVM-66371
 // REQUIRES: aspect-ext_intel_matrix
 
 // RUN: %{build} -mllvm -inline-threshold=2000 %fp-model-precise -o %t.out -DMANUAL_UNROLL -DVNNI

sycl/test-e2e/Matrix/joint_matrix_bf16_fill_k_cache_unroll_init.cpp

@@ -6,9 +6,6 @@
 //
 //===----------------------------------------------------------------------===//
 // REQUIRES: target-spir
-// XFAIL: (arch-intel_gpu_acm_g10 || arch-intel_gpu_pvc || arch-intel_gpu_bmg_g21)
-// XFAIL-TRACKER: CMPLRLLVM-66371
-
 // REQUIRES: aspect-ext_intel_matrix, gpu
 
 // RUN: %{build} -mllvm -inline-threshold=2000 %fp-model-precise -o %t_gpu.out -DINIT_LIST -DMANUAL_UNROLL -DVNNI

sycl/test/check_device_code/esimd/NBarrierAttr.cpp

@@ -23,7 +23,7 @@ SYCL_ESIMD_FUNCTION SYCL_EXTERNAL ESIMD_NOINLINE void callee(int x) {
 // inherits SLMSize and NBarrierCount from callee
 void caller_abc(int x) {
   kernel<class kernel_abc>([=]() SYCL_ESIMD_KERNEL { callee(x); });
-  // CHECK: define dso_local spir_kernel void @_ZTSZ10caller_abciE10kernel_abc() local_unnamed_addr #[[ATTR1:[0-9]+]]
+  // CHECK: define {{.*}}spir_kernel void @_ZTSZ10caller_abciE10kernel_abc() local_unnamed_addr #[[ATTR1:[0-9]+]]
 }
 
 // inherits only NBarrierCount from callee
@@ -33,7 +33,7 @@ void caller_xyz(int x) {
     auto y = __ESIMD_ENS::named_barrier_allocate<35>();
     __ESIMD_NS::named_barrier_wait(y);
   });
-  // CHECK: define dso_local spir_kernel void @_ZTSZ10caller_xyziE10kernel_xyz() local_unnamed_addr #[[ATTR2:[0-9]+]]
+  // CHECK: define {{.*}}spir_kernel void @_ZTSZ10caller_xyziE10kernel_xyz() local_unnamed_addr #[[ATTR2:[0-9]+]]
   // CHECK: call void @llvm.genx.nbarrier(i8 0, i8 13, i8 0)
 }