diff --git a/clang/lib/AST/ASTContext.cpp b/clang/lib/AST/ASTContext.cpp index e9b8b7404718b..5d7e50688f354 100644 --- a/clang/lib/AST/ASTContext.cpp +++ b/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()) { + (D->hasAttr() && + D->getAttr()->isImplicit())) { if (L == GVA_DiscardableODR) return GVA_StrongODR; } diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index 8db2ba2e9903a..18759a0f3ee37 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/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); diff --git a/clang/test/CodeGenSYCL/device_has.cpp b/clang/test/CodeGenSYCL/device_has.cpp index a7c1b01fd2bca..6848c2284c904 100644 --- a/clang/test/CodeGenSYCL/device_has.cpp +++ b/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(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([]() [[sycl::device_has(sycl::aspect::gpu)]] {}); }); } diff --git a/clang/test/CodeGenSYCL/dynamic_local_accessor.cpp b/clang/test/CodeGenSYCL/dynamic_local_accessor.cpp index ec1c110585fb0..2d11ff72156b6 100644 --- a/clang/test/CodeGenSYCL/dynamic_local_accessor.cpp +++ b/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 diff --git a/clang/test/CodeGenSYCL/dynamic_work_group_memory.cpp b/clang/test/CodeGenSYCL/dynamic_work_group_memory.cpp index 4e2602b5cbaa7..bd989342fb621 100644 --- a/clang/test/CodeGenSYCL/dynamic_work_group_memory.cpp +++ b/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 diff --git a/clang/test/CodeGenSYCL/generated-types-initialization.cpp b/clang/test/CodeGenSYCL/generated-types-initialization.cpp index 91c13fa271222..b5e7f305a93ae 100644 --- a/clang/test/CodeGenSYCL/generated-types-initialization.cpp +++ b/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 diff --git a/clang/test/CodeGenSYCL/kernel-handler.cpp b/clang/test/CodeGenSYCL/kernel-handler.cpp index c0fe1c12e934d..6d20b328bbd39 100644 --- a/clang/test/CodeGenSYCL/kernel-handler.cpp +++ b/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 diff --git a/clang/test/CodeGenSYCL/kernel-op-calls.cpp b/clang/test/CodeGenSYCL/kernel-op-calls.cpp index 1dee50ba4b3ad..60a6077bd877c 100644 --- a/clang/test/CodeGenSYCL/kernel-op-calls.cpp +++ b/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 - [[sycl::work_group_size_hint(1, 2, 3)]] void operator()(sycl::id item) const {} - + template + [[sycl::work_group_size_hint(1, 2, 3)]] + void operator()(sycl::id 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); }); diff --git a/clang/test/CodeGenSYCL/kernel-param-acc-array.cpp b/clang/test/CodeGenSYCL/kernel-param-acc-array.cpp index b6f82a3b8c224..ca93a24b6cdc2 100644 --- a/clang/test/CodeGenSYCL/kernel-param-acc-array.cpp +++ b/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 diff --git a/clang/test/CodeGenSYCL/kernel-param-member-acc-array.cpp b/clang/test/CodeGenSYCL/kernel-param-member-acc-array.cpp index ffdc51506ff86..ef30002d20bcd 100644 --- a/clang/test/CodeGenSYCL/kernel-param-member-acc-array.cpp +++ b/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 diff --git a/clang/test/CodeGenSYCL/kernel-simple-instrumentation.cpp b/clang/test/CodeGenSYCL/kernel-simple-instrumentation.cpp index 7728e3c8c490d..39449070f11be 100644 --- a/clang/test/CodeGenSYCL/kernel-simple-instrumentation.cpp +++ b/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() diff --git a/clang/test/CodeGenSYCL/launch_bounds_nvptx.cpp b/clang/test/CodeGenSYCL/launch_bounds_nvptx.cpp index 886e9a94ceec1..e3902d340d385 100644 --- a/clang/test/CodeGenSYCL/launch_bounds_nvptx.cpp +++ b/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} diff --git a/clang/test/CodeGenSYCL/max-concurrency.cpp b/clang/test/CodeGenSYCL/max-concurrency.cpp index 5c92bcfbc7fba..7e17e89ba546e 100644 --- a/clang/test/CodeGenSYCL/max-concurrency.cpp +++ b/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) diff --git a/clang/test/CodeGenSYCL/odr-kernel.cpp b/clang/test/CodeGenSYCL/odr-kernel.cpp new file mode 100644 index 0000000000000..a598033d0f12f --- /dev/null +++ b/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 + __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(f); + } +} diff --git a/clang/test/CodeGenSYCL/pipeline_kernel.cpp b/clang/test/CodeGenSYCL/pipeline_kernel.cpp index f3aace06fe804..24a2f2e38b7ec 100644 --- a/clang/test/CodeGenSYCL/pipeline_kernel.cpp +++ b/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} diff --git a/clang/test/CodeGenSYCL/reqd-sub-group-size-spirv-intrin.cpp b/clang/test/CodeGenSYCL/reqd-sub-group-size-spirv-intrin.cpp index fe1c242fb7bcd..e6220cae6c4fd 100644 --- a/clang/test/CodeGenSYCL/reqd-sub-group-size-spirv-intrin.cpp +++ b/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{{.*}}({{.*}}) diff --git a/clang/test/CodeGenSYCL/stream.cpp b/clang/test/CodeGenSYCL/stream.cpp index 520f3c398b762..98e54f6f482f8 100644 --- a/clang/test/CodeGenSYCL/stream.cpp +++ b/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_]+]], diff --git a/clang/test/CodeGenSYCL/sub-group-size.cpp b/clang/test/CodeGenSYCL/sub-group-size.cpp index 3f4cf91818f82..ed6e23a4adfe0 100644 --- a/clang/test/CodeGenSYCL/sub-group-size.cpp +++ b/clang/test/CodeGenSYCL/sub-group-size.cpp @@ -23,7 +23,7 @@ void default_behavior() { kernel_single_task([]() { }); } -// 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]] diff --git a/clang/test/CodeGenSYCL/sycl-intelfpga-field.cpp b/clang/test/CodeGenSYCL/sycl-intelfpga-field.cpp index ca2d9d572c522..209f32cda79c4 100644 --- a/clang/test/CodeGenSYCL/sycl-intelfpga-field.cpp +++ b/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 diff --git a/clang/test/CodeGenSYCL/work_group_memory.cpp b/clang/test/CodeGenSYCL/work_group_memory.cpp index a4ebb72ad862d..babaf960aa9bb 100644 --- a/clang/test/CodeGenSYCL/work_group_memory.cpp +++ b/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 diff --git a/sycl/test-e2e/Matrix/joint_matrix_bf16_fill_k_cache_unroll.cpp b/sycl/test-e2e/Matrix/joint_matrix_bf16_fill_k_cache_unroll.cpp index 12f3a62aed608..68092f76a0a31 100644 --- a/sycl/test-e2e/Matrix/joint_matrix_bf16_fill_k_cache_unroll.cpp +++ b/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 diff --git a/sycl/test-e2e/Matrix/joint_matrix_bf16_fill_k_cache_unroll_init.cpp b/sycl/test-e2e/Matrix/joint_matrix_bf16_fill_k_cache_unroll_init.cpp index 6586aed6b4617..0372f7e27d5c6 100644 --- a/sycl/test-e2e/Matrix/joint_matrix_bf16_fill_k_cache_unroll_init.cpp +++ b/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 diff --git a/sycl/test/check_device_code/esimd/NBarrierAttr.cpp b/sycl/test/check_device_code/esimd/NBarrierAttr.cpp index 174b870c9b965..0ec0959cf7aaa 100644 --- a/sycl/test/check_device_code/esimd/NBarrierAttr.cpp +++ b/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([=]() 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) } diff --git a/sycl/test/check_device_code/esimd/dae.cpp b/sycl/test/check_device_code/esimd/dae.cpp index 222a8628647a8..dcbf40734979b 100644 --- a/sycl/test/check_device_code/esimd/dae.cpp +++ b/sycl/test/check_device_code/esimd/dae.cpp @@ -15,7 +15,7 @@ __attribute__((sycl_kernel)) void my_kernel(Func kernelFunc) { SYCL_EXTERNAL SYCL_ESIMD_FUNCTION ESIMD_NOINLINE void callee(int x) {} -// CHECK: define dso_local spir_kernel {{.*}} !sycl_kernel_omit_args ![[#MD:]] +// CHECK: define {{.*}}spir_kernel {{.*}} !sycl_kernel_omit_args ![[#MD:]] SYCL_EXTERNAL void __attribute__((noinline)) caller(int x) { my_kernel([=]() SYCL_ESIMD_KERNEL { callee(x); }); } diff --git a/sycl/test/check_device_code/esimd/genx_func_attr.cpp b/sycl/test/check_device_code/esimd/genx_func_attr.cpp index f92c5f40c36a3..5a0aebb88d184 100644 --- a/sycl/test/check_device_code/esimd/genx_func_attr.cpp +++ b/sycl/test/check_device_code/esimd/genx_func_attr.cpp @@ -24,7 +24,7 @@ SYCL_ESIMD_FUNCTION SYCL_EXTERNAL ESIMD_NOINLINE void callee(int x) { // inherits SLMSize and NBarrierCount from callee void caller_abc(int x) { kernel([=]() SYCL_ESIMD_KERNEL { callee(x); }); - // CHECK: define dso_local spir_kernel void @_ZTSZ10caller_abciE10kernel_abc() local_unnamed_addr #[[ATTR:[0-9]+]] + // CHECK: define {{.*}}spir_kernel void @_ZTSZ10caller_abciE10kernel_abc() local_unnamed_addr #[[ATTR:[0-9]+]] } // inherits only NBarrierCount from callee @@ -33,7 +33,7 @@ void caller_xyz(int x) { slm_init(1235); // also works in non-O0 callee(x); }); - // CHECK: define dso_local spir_kernel void @_ZTSZ10caller_xyziE10kernel_xyz() local_unnamed_addr #[[ATTR]] + // CHECK: define {{.*}}spir_kernel void @_ZTSZ10caller_xyziE10kernel_xyz() local_unnamed_addr #[[ATTR]] } // CHECK: attributes #[[ATTR]] = { {{.*}} "VCNamedBarrierCount"="13" "VCSLMSize"="2469" diff --git a/sycl/test/extensions/properties/properties_kernel_device_has_warning.cpp b/sycl/test/extensions/properties/properties_kernel_device_has_warning.cpp index 150a77f8f014f..a3c095ae03975 100644 --- a/sycl/test/extensions/properties/properties_kernel_device_has_warning.cpp +++ b/sycl/test/extensions/properties/properties_kernel_device_has_warning.cpp @@ -1,7 +1,4 @@ // RUN: %clangxx -fsycl-device-only -Xclang -verify -Xclang -verify-ignore-unexpected=note %s -// -// XFAIL: * -// XFAIL-TRACKER: https://github.com/intel/llvm/issues/19409 // Tests for warnings when propagated aspects do not match the aspects available // in a function, as specified through the 'sycl::device_has' property.