diff --git a/clang/lib/CodeGen/CodeGenTypes.cpp b/clang/lib/CodeGen/CodeGenTypes.cpp index b35275e4dd24c..2e7ea8d5dbbd8 100644 --- a/clang/lib/CodeGen/CodeGenTypes.cpp +++ b/clang/lib/CodeGen/CodeGenTypes.cpp @@ -52,20 +52,6 @@ void CodeGenTypes::addRecordTypeName(const RecordDecl *RD, llvm::raw_svector_ostream OS(TypeName); OS << RD->getKindName() << '.'; - // NOTE: The following block of code is copied from CLANG-3.6 with - // support of OpenCLCPlusPlus. It is rather the temporary solution - // that is going to be used until the general solution is ported/developed - // in the latest llvm trunk. - // - // For SYCL, the mangled type name is attached, so it can be - // reflown to proper name later. - if (getContext().getLangOpts().SYCLIsDevice) { - std::unique_ptr MC(getContext().createMangleContext()); - auto RDT = getContext().getRecordType(RD); - MC->mangleCXXRTTIName(RDT, OS); - OS << "."; - } - // FIXME: We probably want to make more tweaks to the printing policy. For // example, we should probably enable PrintCanonicalTypes and // FullyQualifiedNames. diff --git a/clang/test/CodeGenSYCL/address-space-cond-op.cpp b/clang/test/CodeGenSYCL/address-space-cond-op.cpp index 986d7d426361d..caa2126be30b2 100644 --- a/clang/test/CodeGenSYCL/address-space-cond-op.cpp +++ b/clang/test/CodeGenSYCL/address-space-cond-op.cpp @@ -14,7 +14,7 @@ struct S { // CHECK-NEXT: [[FROMBOOL:%.*]] = zext i1 [[COND:%.*]] to i8 // CHECK-NEXT: store i8 [[FROMBOOL]], i8 addrspace(4)* [[COND_ADDR_ASCAST]], align 1, [[TBAA12:!tbaa !.*]] // CHECK-NEXT: store [[STRUCT__ZTS1S_S]] addrspace(4)* [[LHS:%.*]], [[STRUCT__ZTS1S_S]] addrspace(4)* addrspace(4)* [[LHS_ADDR_ASCAST]], align 8, [[TBAA5:!tbaa !.*]] -// CHECK-NEXT: [[RHS_ASCAST:%.*]] = addrspacecast %struct._ZTS1S.S* [[RHS:%.*]] to [[STRUCT__ZTS1S_S]] addrspace(4)* +// CHECK-NEXT: [[RHS_ASCAST:%.*]] = addrspacecast %struct.S* [[RHS:%.*]] to [[STRUCT__ZTS1S_S]] addrspace(4)* // CHECK-NEXT: [[TMP0:%.*]] = load i8, i8 addrspace(4)* [[COND_ADDR_ASCAST]], align 1, [[TBAA12]], [[RNG14:!range !.*]] // CHECK-NEXT: [[TOBOOL:%.*]] = trunc i8 [[TMP0]] to i1 // CHECK-NEXT: br i1 [[TOBOOL]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]] diff --git a/clang/test/CodeGenSYCL/address-space-new.cpp b/clang/test/CodeGenSYCL/address-space-new.cpp index 8e05cd321f9ff..6610dc76da44b 100644 --- a/clang/test/CodeGenSYCL/address-space-new.cpp +++ b/clang/test/CodeGenSYCL/address-space-new.cpp @@ -100,10 +100,10 @@ void test() { Y yy; baz(yy); // CHECK: define {{.*}}spir_func void @{{.*}}baz{{.*}} - // CHECK: %[[FIRST:[a-zA-Z0-9]+]] = bitcast %struct.{{.*}}.Y addrspace(4)* %{{.*}} to i8 addrspace(4)* + // CHECK: %[[FIRST:[a-zA-Z0-9]+]] = bitcast %struct.Y addrspace(4)* %{{.*}} to i8 addrspace(4)* // CHECK: %[[OFFSET:[a-zA-Z0-9]+]].ptr = getelementptr inbounds i8, i8 addrspace(4)* %[[FIRST]], i64 8 - // CHECK: %[[SECOND:[a-zA-Z0-9]+]] = bitcast i8 addrspace(4)* %[[OFFSET]].ptr to %struct.{{.*}}.HasX addrspace(4)* - // CHECK: call spir_func void @{{.*}}bar{{.*}}(%struct.{{.*}}.HasX addrspace(4)* align 4 dereferenceable(4) %[[SECOND]]) + // CHECK: %[[SECOND:[a-zA-Z0-9]+]] = bitcast i8 addrspace(4)* %[[OFFSET]].ptr to %struct.HasX addrspace(4)* + // CHECK: call spir_func void @{{.*}}bar{{.*}}(%struct.HasX addrspace(4)* align 4 dereferenceable(4) %[[SECOND]]) } template diff --git a/clang/test/CodeGenSYCL/basic-kernel-wrapper.cpp b/clang/test/CodeGenSYCL/basic-kernel-wrapper.cpp index 7133752dcf217..b098c01a8ba4f 100644 --- a/clang/test/CodeGenSYCL/basic-kernel-wrapper.cpp +++ b/clang/test/CodeGenSYCL/basic-kernel-wrapper.cpp @@ -21,21 +21,21 @@ int main() { // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_function // CHECK-SAME: i32 addrspace(1)* [[MEM_ARG:%[a-zA-Z0-9_]+]], -// CHECK-SAME: %"struct.{{.*}}.cl::sycl::range"* byval{{.*}}align 4 [[ACC_RANGE:%[a-zA-Z0-9_]+_1]], -// CHECK-SAME: %"struct.{{.*}}.cl::sycl::range"* byval{{.*}}align 4 [[MEM_RANGE:%[a-zA-Z0-9_]+_2]], -// CHECK-SAME: %"struct.{{.*}}.cl::sycl::id"* byval{{.*}}align 4 [[OFFSET:%[a-zA-Z0-9_]+]]) +// CHECK-SAME: %"struct.cl::sycl::range"* byval{{.*}}align 4 [[ACC_RANGE:%[a-zA-Z0-9_]+_1]], +// CHECK-SAME: %"struct.cl::sycl::range"* byval{{.*}}align 4 [[MEM_RANGE:%[a-zA-Z0-9_]+_2]], +// CHECK-SAME: %"struct.cl::sycl::id"* byval{{.*}}align 4 [[OFFSET:%[a-zA-Z0-9_]+]]) // Check alloca for pointer argument // CHECK: [[MEM_ARG]].addr = alloca i32 addrspace(1)* // Check lambda object alloca -// CHECK: [[ANONALLOCA:%[0-9]+]] = alloca %class.{{.*}}.anon -// CHECK: [[ANON:%[0-9]+]] = addrspacecast %class.{{.*}}.anon* [[ANONALLOCA]] to %class.{{.*}}.anon addrspace(4)* +// CHECK: [[ANONALLOCA:%[0-9]+]] = alloca %class.anon +// CHECK: [[ANON:%[0-9]+]] = addrspacecast %class.anon* [[ANONALLOCA]] to %class.anon addrspace(4)* // Check allocas for ranges -// CHECK: [[ARANGEA:%agg.tmp.*]] = alloca %"struct.{{.*}}.cl::sycl::range" -// CHECK: [[ARANGET:%agg.tmp.*]] = addrspacecast %"struct.{{.*}}.cl::sycl::range"* [[ARANGEA]] to %"struct.{{.*}}.cl::sycl::range" addrspace(4)* -// CHECK: [[MRANGEA:%agg.tmp.*]] = alloca %"struct.{{.*}}.cl::sycl::range" -// CHECK: [[MRANGET:%agg.tmp.*]] = addrspacecast %"struct.{{.*}}.cl::sycl::range"* [[MRANGEA]] to %"struct.{{.*}}.cl::sycl::range" addrspace(4)* -// CHECK: [[OIDA:%agg.tmp.*]] = alloca %"struct.{{.*}}.cl::sycl::id" -// CHECK: [[OIDT:%agg.tmp.*]] = addrspacecast %"struct.{{.*}}.cl::sycl::id"* [[OIDA]] to %"struct.{{.*}}.cl::sycl::id" addrspace(4)* +// CHECK: [[ARANGEA:%agg.tmp.*]] = alloca %"struct.cl::sycl::range" +// CHECK: [[ARANGET:%agg.tmp.*]] = addrspacecast %"struct.cl::sycl::range"* [[ARANGEA]] to %"struct.cl::sycl::range" addrspace(4)* +// CHECK: [[MRANGEA:%agg.tmp.*]] = alloca %"struct.cl::sycl::range" +// CHECK: [[MRANGET:%agg.tmp.*]] = addrspacecast %"struct.cl::sycl::range"* [[MRANGEA]] to %"struct.cl::sycl::range" addrspace(4)* +// CHECK: [[OIDA:%agg.tmp.*]] = alloca %"struct.cl::sycl::id" +// CHECK: [[OIDT:%agg.tmp.*]] = addrspacecast %"struct.cl::sycl::id"* [[OIDA]] to %"struct.cl::sycl::id" addrspace(4)* // // Check store of kernel pointer argument to alloca // CHECK: store i32 addrspace(1)* [[MEM_ARG]], i32 addrspace(1)* addrspace(4)* [[MEM_ARG]].addr.ascast, align 8 @@ -44,16 +44,16 @@ int main() { // CHECK: call spir_func {{.*}}accessor // Check accessor GEP -// CHECK: [[ACCESSOR:%[a-zA-Z0-9_]+]] = getelementptr inbounds %class.{{.*}}.anon, %class.{{.*}}.anon addrspace(4)* [[ANON]], i32 0, i32 0 +// CHECK: [[ACCESSOR:%[a-zA-Z0-9_]+]] = getelementptr inbounds %class.anon, %class.anon addrspace(4)* [[ANON]], i32 0, i32 0 // Check load from kernel pointer argument alloca // CHECK: [[MEM_LOAD:%[a-zA-Z0-9_]+]] = load i32 addrspace(1)*, i32 addrspace(1)* addrspace(4)* [[MEM_ARG]].addr.ascast // Check accessor __init method call -// CHECK: [[ARANGE:%agg.tmp.*]] = addrspacecast %"struct.{{.*}}.cl::sycl::range" addrspace(4)* [[ARANGET]] to %"struct.{{.*}}.cl::sycl::range"* -// CHECK: [[MRANGE:%agg.tmp.*]] = addrspacecast %"struct.{{.*}}.cl::sycl::range" addrspace(4)* [[MRANGET]] to %"struct.{{.*}}.cl::sycl::range"* -// CHECK: [[OID:%agg.tmp.*]] = addrspacecast %"struct.{{.*}}.cl::sycl::id" addrspace(4)* [[OIDT]] to %"struct.{{.*}}.cl::sycl::id"* -// CHECK: call spir_func void @{{.*}}__init{{.*}}(%"class.{{.*}}.cl::sycl::accessor" addrspace(4)* {{[^,]*}} [[ACCESSOR]], i32 addrspace(1)* [[MEM_LOAD]], %"struct.{{.*}}.cl::sycl::range"* byval({{.*}}) align 4 [[ARANGE]], %"struct.{{.*}}.cl::sycl::range"* byval({{.*}}) align 4 [[MRANGE]], %"struct.{{.*}}.cl::sycl::id"* byval({{.*}}) align 4 [[OID]]) +// CHECK: [[ARANGE:%agg.tmp.*]] = addrspacecast %"struct.cl::sycl::range" addrspace(4)* [[ARANGET]] to %"struct.cl::sycl::range"* +// CHECK: [[MRANGE:%agg.tmp.*]] = addrspacecast %"struct.cl::sycl::range" addrspace(4)* [[MRANGET]] to %"struct.cl::sycl::range"* +// CHECK: [[OID:%agg.tmp.*]] = addrspacecast %"struct.cl::sycl::id" addrspace(4)* [[OIDT]] to %"struct.cl::sycl::id"* +// CHECK: call spir_func void @{{.*}}__init{{.*}}(%"class.cl::sycl::accessor" addrspace(4)* {{[^,]*}} [[ACCESSOR]], i32 addrspace(1)* [[MEM_LOAD]], %"struct.cl::sycl::range"* byval({{.*}}) align 4 [[ARANGE]], %"struct.cl::sycl::range"* byval({{.*}}) align 4 [[MRANGE]], %"struct.cl::sycl::id"* byval({{.*}}) align 4 [[OID]]) // Check lambda "()" operator call -// CHECK: call spir_func void @{{.*}}(%class.{{.*}}.anon addrspace(4)* {{[^,]*}}) +// CHECK: call spir_func void @{{.*}}(%class.anon addrspace(4)* {{[^,]*}}) diff --git a/clang/test/CodeGenSYCL/check-direct-attribute-propagation.cpp b/clang/test/CodeGenSYCL/check-direct-attribute-propagation.cpp index 6176c5fb25dfb..7108ddadcb29c 100644 --- a/clang/test/CodeGenSYCL/check-direct-attribute-propagation.cpp +++ b/clang/test/CodeGenSYCL/check-direct-attribute-propagation.cpp @@ -304,7 +304,7 @@ int main() { // Test attribute is not propagated. // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name32() #0 !kernel_arg_buffer_location ![[NUM]] - // CHECK: define {{.*}}spir_func void @{{.*}}Functor10{{.*}}(%class._ZTS9Functor10.Functor10 addrspace(4)* align 1 dereferenceable_or_null(1) %this) #3 comdat align 2 + // CHECK: define {{.*}}spir_func void @{{.*}}Functor10{{.*}}(%class.Functor10 addrspace(4)* align 1 dereferenceable_or_null(1) %this) #3 comdat align 2 // CHECK-NOT: noalias // CHECK-SAME: { // CHECK: define dso_local spir_func void @_Z4foo8v() @@ -312,12 +312,12 @@ int main() { h.single_task(f10); // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name33() #0 !kernel_arg_buffer_location ![[NUM]] - // CHECK: define {{.*}}spir_func void @{{.*}}Foo8{{.*}}(%class.{{.*}}Foo8.Foo8 addrspace(4)* noalias align 1 dereferenceable_or_null(1) %this) #3 comdat align 2 + // CHECK: define {{.*}}spir_func void @{{.*}}Foo8{{.*}}(%class.Foo8 addrspace(4)* noalias align 1 dereferenceable_or_null(1) %this) #3 comdat align 2 Foo8 boo8; h.single_task(boo8); // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name34() #0 !kernel_arg_buffer_location ![[NUM]] - // CHECK: define {{.*}}spir_func void @{{.*}}(%class.{{.*}}.anon addrspace(4)* noalias align 1 dereferenceable_or_null(1) %this) #4 align 2 + // CHECK: define {{.*}}spir_func void @{{.*}}(%class.anon{{.*}} addrspace(4)* noalias align 1 dereferenceable_or_null(1) %this) #4 align 2 h.single_task( []() [[intel::kernel_args_restrict]]{}); }); diff --git a/clang/test/CodeGenSYCL/debug-info-srcpos-kernel.cpp b/clang/test/CodeGenSYCL/debug-info-srcpos-kernel.cpp index ece9d8ed3cec5..7b61d4bc32250 100644 --- a/clang/test/CodeGenSYCL/debug-info-srcpos-kernel.cpp +++ b/clang/test/CodeGenSYCL/debug-info-srcpos-kernel.cpp @@ -25,7 +25,7 @@ int main() { } // CHECK: define{{.*}} spir_kernel {{.*}}19use_kernel_for_test({{.*}}){{.*}} !dbg [[KERNEL:![0-9]+]] {{.*}}{ -// CHECK: getelementptr inbounds %class.{{.*}}.anon{{.*}} !dbg [[LINE_A0:![0-9]+]] +// CHECK: getelementptr inbounds %class.anon, %class.anon addrspace(4)* %{{[0-9]+}}, i32 0, i32 0, !dbg [[LINE_A0:![0-9]+]] // CHECK: call spir_func void {{.*}}6__init{{.*}} !dbg [[LINE_A0]] // CHECK: call spir_func void @_ZZ4mainENKUlvE_clEv{{.*}} !dbg [[LINE_B0:![0-9]+]] // CHECK: ret void, !dbg [[LINE_C0:![0-9]+]] diff --git a/clang/test/CodeGenSYCL/device-functions.cpp b/clang/test/CodeGenSYCL/device-functions.cpp index d19cbbb10f168..2366d51f6e2d8 100644 --- a/clang/test/CodeGenSYCL/device-functions.cpp +++ b/clang/test/CodeGenSYCL/device-functions.cpp @@ -22,6 +22,6 @@ int main() { return 0; } // CHECK: define {{.*}}spir_kernel void @_ZTSZ4mainE11fake_kernel() -// CHECK: define internal spir_func void @_ZZ4mainENKUlvE_clEv(%class.{{.*}}.anon addrspace(4)* {{[^,]*}} %this) +// CHECK: define internal spir_func void @_ZZ4mainENKUlvE_clEv(%class.anon addrspace(4)* {{[^,]*}} %this) // CHECK: define {{.*}}spir_func void @_Z3foov() // CHECK: define linkonce_odr spir_func i32 @_Z3barIiET_S0_(i32 %arg) diff --git a/clang/test/CodeGenSYCL/device-variables.cpp b/clang/test/CodeGenSYCL/device-variables.cpp index 25d5e18eff912..27811233a728b 100644 --- a/clang/test/CodeGenSYCL/device-variables.cpp +++ b/clang/test/CodeGenSYCL/device-variables.cpp @@ -32,10 +32,10 @@ int main() { // CHECK: store i32 1, i32 addrspace(4)* %b foo(local_value); // Local variables and constexprs captured by lambda - // CHECK: [[GEP:%[0-9]+]] = getelementptr inbounds %class.{{.*}}.anon, %class.{{.*}}.anon addrspace(4)* %{{.*}}, i32 0, i32 0 + // CHECK: [[GEP:%[0-9]+]] = getelementptr inbounds %class.anon, %class.anon addrspace(4)* %{{.*}}, i32 0, i32 0 // CHECK: call spir_func void @{{.*}}foo{{.*}}(i32 addrspace(4)* align 4 dereferenceable(4) [[GEP]]) int some_device_local_var = some_local_var; - // CHECK: [[GEP1:%[0-9]+]] = getelementptr inbounds %class.{{.*}}.anon, %class.{{.*}}.anon addrspace(4)* %{{.*}}, i32 0, i32 1 + // CHECK: [[GEP1:%[0-9]+]] = getelementptr inbounds %class.anon, %class.anon addrspace(4)* %{{.*}}, i32 0, i32 1 // CHECK: [[LOAD1:%[0-9]+]] = load i32, i32 addrspace(4)* [[GEP1]] // CHECK: store i32 [[LOAD1]], i32 addrspace(4)* %some_device_local_var }); diff --git a/clang/test/CodeGenSYCL/inheritance.cpp b/clang/test/CodeGenSYCL/inheritance.cpp index 00e1e504a5e4b..4173800fc1159 100644 --- a/clang/test/CodeGenSYCL/inheritance.cpp +++ b/clang/test/CodeGenSYCL/inheritance.cpp @@ -40,34 +40,34 @@ int main() { } // Check kernel paramters -// CHECK: define {{.*}}spir_kernel void @{{.*}}derived(%struct.{{.*}}.base* byval(%struct.{{.*}}.base) align 4 %_arg__base, %struct.{{.*}}.__wrapper_class* byval(%struct.{{.*}}.__wrapper_class) align 8 %_arg_e, i32 %_arg_a) +// CHECK: define {{.*}}spir_kernel void @{{.*}}derived(%struct.base* byval(%struct.base) align 4 %_arg__base, %struct.__wrapper_class* byval(%struct.__wrapper_class) align 8 %_arg_e, i32 %_arg_a) // Check alloca for kernel paramters // CHECK: %[[ARG_AA:[a-zA-Z0-9_.]+]] = alloca i32, align 4 // CHECK: %[[ARG_A:[a-zA-Z0-9_.]+]] = addrspacecast i32* %[[ARG_AA]] to i32 addrspace(4)* // Check alloca for local functor object -// CHECK: %[[LOCAL_OBJECT:[a-zA-Z0-9_.]+]] = alloca %struct.{{.*}}.derived, align 8 +// CHECK: %[[LOCAL_OBJECT:[a-zA-Z0-9_.]+]] = alloca %struct.derived, align 8 // CHECK: store i32 %_arg_a, i32 addrspace(4)* %[[ARG_A]], align 4 // Initialize 'base' subobject -// CHECK: %[[DERIVED_TO_BASE:.*]] = bitcast %struct.{{.*}}.derived addrspace(4)* %[[LOCAL_OBJECT]].ascast to %struct.{{.*}}.base addrspace(4)* -// CHECK: %[[BASE_TO_PTR:.*]] = bitcast %struct.{{.*}}.base addrspace(4)* %[[DERIVED_TO_BASE]] to i8 addrspace(4)* -// CHECK: %[[PARAM_TO_PTR:.*]] = bitcast %struct.{{.*}}.base addrspace(4)* %_arg__base.ascast to i8 addrspace(4)* +// CHECK: %[[DERIVED_TO_BASE:.*]] = bitcast %struct.derived addrspace(4)* %[[LOCAL_OBJECT]].ascast to %struct.base addrspace(4)* +// CHECK: %[[BASE_TO_PTR:.*]] = bitcast %struct.base addrspace(4)* %[[DERIVED_TO_BASE]] to i8 addrspace(4)* +// CHECK: %[[PARAM_TO_PTR:.*]] = bitcast %struct.base addrspace(4)* %_arg__base.ascast to i8 addrspace(4)* // CHECK: call void @llvm.memcpy.p4i8.p4i8.i64(i8 addrspace(4)* align 8 %[[BASE_TO_PTR]], i8 addrspace(4)* align 4 %[[PARAM_TO_PTR]], i64 12, i1 false) // Initialize 'second_base' subobject // First, derived-to-base cast with offset: -// CHECK: %[[DERIVED_PTR:.*]] = bitcast %struct.{{.*}}.derived addrspace(4)* %[[LOCAL_OBJECT]].ascast to i8 addrspace(4)* +// CHECK: %[[DERIVED_PTR:.*]] = bitcast %struct.derived addrspace(4)* %[[LOCAL_OBJECT]].ascast to i8 addrspace(4)* // CHECK: %[[OFFSET_CALC:.*]] = getelementptr inbounds i8, i8 addrspace(4)* %[[DERIVED_PTR]], i64 16 -// CHECK: %[[TO_SECOND_BASE:.*]] = bitcast i8 addrspace(4)* %[[OFFSET_CALC]] to %class.{{.*}}.second_base addrspace(4)* +// CHECK: %[[TO_SECOND_BASE:.*]] = bitcast i8 addrspace(4)* %[[OFFSET_CALC]] to %class.second_base addrspace(4)* // Initialize 'second_base::e' -// CHECK: %[[SECOND_BASE_PTR:.*]] = getelementptr inbounds %class.{{.*}}.second_base, %class.{{.*}}.second_base addrspace(4)* %[[TO_SECOND_BASE]], i32 0, i32 0 -// CHECK: %[[PTR_TO_WRAPPER:.*]] = getelementptr inbounds %struct.{{.*}}.__wrapper_class, %struct.{{.*}}.__wrapper_class addrspace(4)* %_arg_e.ascast, i32 0, i32 0 +// CHECK: %[[SECOND_BASE_PTR:.*]] = getelementptr inbounds %class.second_base, %class.second_base addrspace(4)* %[[TO_SECOND_BASE]], i32 0, i32 0 +// CHECK: %[[PTR_TO_WRAPPER:.*]] = getelementptr inbounds %struct.__wrapper_class, %struct.__wrapper_class addrspace(4)* %_arg_e.ascast, i32 0, i32 0 // CHECK: %[[LOAD_PTR:.*]] = load i32 addrspace(1)*, i32 addrspace(1)* addrspace(4)* %[[PTR_TO_WRAPPER]] // CHECK: %[[AS_CAST:.*]] = addrspacecast i32 addrspace(1)* %[[LOAD_PTR]] to i32 addrspace(4)* // CHECK: store i32 addrspace(4)* %[[AS_CAST]], i32 addrspace(4)* addrspace(4)* %[[SECOND_BASE_PTR]] // Initialize field 'a' -// CHECK: %[[GEP_A:[a-zA-Z0-9]+]] = getelementptr inbounds %struct.{{.*}}.derived, %struct.{{.*}}.derived addrspace(4)* %[[LOCAL_OBJECT]].ascast, i32 0, i32 3 +// CHECK: %[[GEP_A:[a-zA-Z0-9]+]] = getelementptr inbounds %struct.derived, %struct.derived addrspace(4)* %[[LOCAL_OBJECT]].ascast, i32 0, i32 3 // CHECK: %[[LOAD_A:[0-9]+]] = load i32, i32 addrspace(4)* %[[ARG_A]], align 4 // CHECK: store i32 %[[LOAD_A]], i32 addrspace(4)* %[[GEP_A]] diff --git a/clang/test/CodeGenSYCL/intel-fpga-ivdep-array.cpp b/clang/test/CodeGenSYCL/intel-fpga-ivdep-array.cpp index 193044a62e2e3..7925e6bb885f5 100644 --- a/clang/test/CodeGenSYCL/intel-fpga-ivdep-array.cpp +++ b/clang/test/CodeGenSYCL/intel-fpga-ivdep-array.cpp @@ -155,16 +155,16 @@ void ivdep_struct() { int *ptr; int arr[10]; } s; - // CHECK: %[[STRUCT:[0-9a-z]+]] = alloca %struct.{{.+}}.S + // CHECK: %[[STRUCT:[0-9a-z]+]] = alloca %struct.S [[intel::ivdep(s.arr, 5)]] for (int i = 0; i != 10; ++i) s.arr[i] = 0; - // CHECK: %[[STRUCT_ARR:[0-9a-z]+]] = getelementptr inbounds %struct.{{.+}}.S, %struct.{{.+}}.S addrspace(4)* %[[STRUCT]].ascast, i32 0, i32 1 + // CHECK: %[[STRUCT_ARR:[0-9a-z]+]] = getelementptr inbounds %struct.S, %struct.S addrspace(4)* %[[STRUCT]].ascast, i32 0, i32 1 // CHECK: %{{[0-9a-z]+}} = getelementptr inbounds [10 x i32], [10 x i32] addrspace(4)* %[[STRUCT_ARR]], i64 0, i64 %{{[0-9a-z]+}}, !llvm.index.group ![[IDX_GROUP_STRUCT_ARR:[0-9]+]] // CHECK: br label %for.cond, !llvm.loop ![[MD_LOOP_STRUCT_ARR:[0-9]+]] [[intel::ivdep(s.ptr, 5)]] for (int i = 0; i != 10; ++i) s.ptr[i] = 0; - // CHECK: %[[STRUCT_PTR:[0-9a-z]+]] = getelementptr inbounds %struct.{{.+}}.S, %struct.{{.+}}.S addrspace(4)* %[[STRUCT]].ascast, i32 0, i32 0 + // CHECK: %[[STRUCT_PTR:[0-9a-z]+]] = getelementptr inbounds %struct.S, %struct.S addrspace(4)* %[[STRUCT]].ascast, i32 0, i32 0 // CHECK: %[[LOAD_STRUCT_PTR:[0-9a-z]+]] = load i32 addrspace(4)*, i32 addrspace(4)* addrspace(4)* %[[STRUCT_PTR]] // CHECK: %{{[0-9a-z]+}} = getelementptr inbounds i32, i32 addrspace(4)* %[[LOAD_STRUCT_PTR]], i64 %{{[0-9a-z]+}}, !llvm.index.group ![[IDX_GROUP_STRUCT_PTR:[0-9]+]] // CHECK: br label %for.cond{{[0-9]*}}, !llvm.loop ![[MD_LOOP_STRUCT_PTR:[0-9]+]] diff --git a/clang/test/CodeGenSYCL/intel-fpga-local.cpp b/clang/test/CodeGenSYCL/intel-fpga-local.cpp index f48dee1319cc6..95fcac216ab8b 100644 --- a/clang/test/CodeGenSYCL/intel-fpga-local.cpp +++ b/clang/test/CodeGenSYCL/intel-fpga-local.cpp @@ -135,49 +135,49 @@ void attrs_on_struct() { int force_p2d [[intel::force_pow2_depth(1)]]; } s; - // CHECK-DEVICE: %[[FIELD_NUMBANKS:.*]] = getelementptr inbounds %struct.{{.*}}.attrs_on_struct{{.*}} + // CHECK-DEVICE: %[[FIELD_NUMBANKS:.*]] = getelementptr inbounds %struct.attrs_on_struct{{.*}} // CHECK-DEVICE: call i32 addrspace(4)* @llvm.ptr.annotation.p4i32{{.*}}%[[FIELD_NUMBANKS]]{{.*}}[[ANN_numbanks_4]] s.numbanks = 0; - // CHECK-DEVICE: %[[FIELD_REGISTER:.*]] = getelementptr inbounds %struct.{{.*}}.attrs_on_struct{{.*}} + // CHECK-DEVICE: %[[FIELD_REGISTER:.*]] = getelementptr inbounds %struct.attrs_on_struct{{.*}} // CHECK-DEVICE: call i32 addrspace(4)* @llvm.ptr.annotation.p4i32{{.*}}%[[FIELD_REGISTER]]{{.*}}[[ANN_register]] s.reg = 0; - // CHECK-DEVICE: %[[FIELD_MEM_DEFAULT:.*]] = getelementptr inbounds %struct.{{.*}}.attrs_on_struct{{.*}} + // CHECK-DEVICE: %[[FIELD_MEM_DEFAULT:.*]] = getelementptr inbounds %struct.attrs_on_struct{{.*}} // CHECK-DEVICE: call i32 addrspace(4)* @llvm.ptr.annotation.p4i32{{.*}}%[[FIELD_MEM_DEFAULT]]{{.*}}[[ANN_memory_default]] s.memory = 0; - // CHECK-DEVICE: %[[FIELD_MEM_BLOCKRAM:.*]] = getelementptr inbounds %struct.{{.*}}.attrs_on_struct{{.*}} + // CHECK-DEVICE: %[[FIELD_MEM_BLOCKRAM:.*]] = getelementptr inbounds %struct.attrs_on_struct{{.*}} // CHECK-DEVICE: call i32 addrspace(4)* @llvm.ptr.annotation.p4i32{{.*}}%[[FIELD_MEM_BLOCKRAM]]{{.*}}[[ANN_memory_blockram]] s.memory_blockram = 0; - // CHECK-DEVICE: %[[FIELD_MEM_MLAB:.*]] = getelementptr inbounds %struct.{{.*}}.attrs_on_struct{{.*}} + // CHECK-DEVICE: %[[FIELD_MEM_MLAB:.*]] = getelementptr inbounds %struct.attrs_on_struct{{.*}} // CHECK-DEVICE: call i32 addrspace(4)* @llvm.ptr.annotation.p4i32{{.*}}%[[FIELD_MEM_MLAB]]{{.*}}[[ANN_memory_mlab]] s.memory_mlab = 0; - // CHECK-DEVICE: %[[FIELD_BANKWIDTH:.*]] = getelementptr inbounds %struct.{{.*}}.attrs_on_struct{{.*}} + // CHECK-DEVICE: %[[FIELD_BANKWIDTH:.*]] = getelementptr inbounds %struct.attrs_on_struct{{.*}} // CHECK-DEVICE: call i32 addrspace(4)* @llvm.ptr.annotation.p4i32{{.*}}%[[FIELD_BANKWIDTH]]{{.*}}[[ANN_bankwidth_4]] s.bankwidth = 0; - // CHECK-DEVICE: %[[FIELD_PRIV_COPIES:.*]] = getelementptr inbounds %struct.{{.*}}.attrs_on_struct{{.*}} + // CHECK-DEVICE: %[[FIELD_PRIV_COPIES:.*]] = getelementptr inbounds %struct.attrs_on_struct{{.*}} // CHECK-DEVICE: call i32 addrspace(4)* @llvm.ptr.annotation.p4i32{{.*}}%[[FIELD_PRIV_COPIES]]{{.*}}[[ANN_private_copies_8]] s.privatecopies = 0; - // CHECK-DEVICE: %[[FIELD_SINGLEPUMP:.*]] = getelementptr inbounds %struct.{{.*}}.attrs_on_struct{{.*}} + // CHECK-DEVICE: %[[FIELD_SINGLEPUMP:.*]] = getelementptr inbounds %struct.attrs_on_struct{{.*}} // CHECK-DEVICE: call i32 addrspace(4)* @llvm.ptr.annotation.p4i32{{.*}}%[[FIELD_SINGLEPUMP]]{{.*}}[[ANN_singlepump]] s.singlepump = 0; - // CHECK-DEVICE: %[[FIELD_DOUBLEPUMP:.*]] = getelementptr inbounds %struct.{{.*}}.attrs_on_struct{{.*}} + // CHECK-DEVICE: %[[FIELD_DOUBLEPUMP:.*]] = getelementptr inbounds %struct.attrs_on_struct{{.*}} // CHECK-DEVICE: call i32 addrspace(4)* @llvm.ptr.annotation.p4i32{{.*}}%[[FIELD_DOUBLEPUMP]]{{.*}}[[ANN_doublepump]] s.doublepump = 0; - // CHECK-DEVICE: %[[FIELD_MERGE_DEPTH:.*]] = getelementptr inbounds %struct.{{.*}}.attrs_on_struct{{.*}} + // CHECK-DEVICE: %[[FIELD_MERGE_DEPTH:.*]] = getelementptr inbounds %struct.attrs_on_struct{{.*}} // CHECK-DEVICE: call i32 addrspace(4)* @llvm.ptr.annotation.p4i32{{.*}}%[[FIELD_MERGE_DEPTH]]{{.*}}[[ANN_merge_depth]] s.merge_depth = 0; - // CHECK-DEVICE: %[[FIELD_MERGE_WIDTH:.*]] = getelementptr inbounds %struct.{{.*}}.attrs_on_struct{{.*}} + // CHECK-DEVICE: %[[FIELD_MERGE_WIDTH:.*]] = getelementptr inbounds %struct.attrs_on_struct{{.*}} // CHECK-DEVICE: call i32 addrspace(4)* @llvm.ptr.annotation.p4i32{{.*}}%[[FIELD_MERGE_WIDTH]]{{.*}}[[ANN_merge_width]] s.merge_width = 0; - // CHECK-DEVICE: %[[FIELD_MAX_REPLICATES:.*]] = getelementptr inbounds %struct.{{.*}}.attrs_on_struct{{.*}} + // CHECK-DEVICE: %[[FIELD_MAX_REPLICATES:.*]] = getelementptr inbounds %struct.attrs_on_struct{{.*}} // CHECK-DEVICE: call i32 addrspace(4)* @llvm.ptr.annotation.p4i32{{.*}}%[[FIELD_MAX_REPLICATES]]{{.*}}[[ANN_max_replicates_2]] s.maxreplicates = 0; - // CHECK-DEVICE: %[[FIELD_DUALPORT:.*]] = getelementptr inbounds %struct.{{.*}}.attrs_on_struct{{.*}} + // CHECK-DEVICE: %[[FIELD_DUALPORT:.*]] = getelementptr inbounds %struct.attrs_on_struct{{.*}} // CHECK-DEVICE: call i32 addrspace(4)* @llvm.ptr.annotation.p4i32{{.*}}%[[FIELD_DUALPORT]]{{.*}}[[ANN_simple_dual_port]] s.dualport = 0; - // CHECK-DEVICE: %[[FIELD_BANKBITS:.*]] = getelementptr inbounds %struct.{{.*}}.attrs_on_struct{{.*}} + // CHECK-DEVICE: %[[FIELD_BANKBITS:.*]] = getelementptr inbounds %struct.attrs_on_struct{{.*}} // CHECK-DEVICE: call i32 addrspace(4)* @llvm.ptr.annotation.p4i32{{.*}}%[[FIELD_BANKBITS]]{{.*}}[[ANN_bankbits_4_5]] s.bankbits = 0; - // CHECK-DEVICE: %[[FIELD_FP2D:.*]] = getelementptr inbounds %struct.{{.*}}.attrs_on_struct{{.*}} + // CHECK-DEVICE: %[[FIELD_FP2D:.*]] = getelementptr inbounds %struct.attrs_on_struct{{.*}} // CHECK-DEVICE: call i32 addrspace(4)* @llvm.ptr.annotation.p4i32{{.*}}%[[FIELD_FP2D]]{{.*}}[[ANN_force_pow2_depth_1]] s.force_p2d = 0; } @@ -221,22 +221,22 @@ void attrs_with_template_param() { int force_p2d [[intel::force_pow2_depth(C)]]; } s; - // CHECK-DEVICE: %[[FIELD_NUMBANKS:.*]] = getelementptr inbounds %struct.{{.*}}.templ_on_struct_fields{{.*}} + // CHECK-DEVICE: %[[FIELD_NUMBANKS:.*]] = getelementptr inbounds %struct.templ_on_struct_fields{{.*}} // CHECK-DEVICE: call i32 addrspace(4)* @llvm.ptr.annotation.p4i32{{.*}}%[[FIELD_NUMBANKS]]{{.*}}[[ANN_numbanks_4]] s.numbanks = 0; - // CHECK-DEVICE: %[[FIELD_BANKWIDTH:.*]] = getelementptr inbounds %struct.{{.*}}.templ_on_struct_fields{{.*}} + // CHECK-DEVICE: %[[FIELD_BANKWIDTH:.*]] = getelementptr inbounds %struct.templ_on_struct_fields{{.*}} // CHECK-DEVICE: call i32 addrspace(4)* @llvm.ptr.annotation.p4i32{{.*}}%[[FIELD_BANKWIDTH]]{{.*}}[[ANN_bankwidth_4]] s.bankwidth = 0; - // CHECK-DEVICE: %[[FIELD_PRIV_COPIES:.*]] = getelementptr inbounds %struct.{{.*}}.templ_on_struct_fields{{.*}} + // CHECK-DEVICE: %[[FIELD_PRIV_COPIES:.*]] = getelementptr inbounds %struct.templ_on_struct_fields{{.*}} // CHECK-DEVICE: call i32 addrspace(4)* @llvm.ptr.annotation.p4i32{{.*}}%[[FIELD_PRIV_COPIES]]{{.*}}[[ANN_private_copies_4]] s.privatecopies = 0; - // CHECK-DEVICE: %[[FIELD_MAX_REPLICATES:.*]] = getelementptr inbounds %struct.{{.*}}.templ_on_struct_fields{{.*}} + // CHECK-DEVICE: %[[FIELD_MAX_REPLICATES:.*]] = getelementptr inbounds %struct.templ_on_struct_fields{{.*}} // CHECK-DEVICE: call i32 addrspace(4)* @llvm.ptr.annotation.p4i32{{.*}}%[[FIELD_MAX_REPLICATES]]{{.*}}[[ANN_max_replicates_4]] s.maxreplicates = 0; - // CHECK-DEVICE: %[[FIELD_BANKBITS:.*]] = getelementptr inbounds %struct.{{.*}}.templ_on_struct_fields{{.*}} + // CHECK-DEVICE: %[[FIELD_BANKBITS:.*]] = getelementptr inbounds %struct.templ_on_struct_fields{{.*}} // CHECK-DEVICE: call i32 addrspace(4)* @llvm.ptr.annotation.p4i32{{.*}}%[[FIELD_BANKBITS]]{{.*}}[[ANN_bankbits_4_5]] s.bankbits = 0; - // CHECK-DEVICE: %[[FIELD_FP2D:.*]] = getelementptr inbounds %struct.{{.*}}.templ_on_struct_fields{{.*}} + // CHECK-DEVICE: %[[FIELD_FP2D:.*]] = getelementptr inbounds %struct.templ_on_struct_fields{{.*}} // CHECK-DEVICE: call i32 addrspace(4)* @llvm.ptr.annotation.p4i32{{.*}}%[[FIELD_FP2D]]{{.*}}[[ANN_force_pow2_depth_1]] s.force_p2d = 0; } @@ -254,7 +254,7 @@ void field_addrspace_cast() { } } state_var; // CHECK-DEVICE: define internal {{.*}} @_ZZ20field_addrspace_castvEN5stateC2Ev - // CHECK-DEVICE: %[[MEM:[a-zA-Z0-9]+]] = getelementptr inbounds %{{.*}}, %struct._ZTSZ20field_addrspace_castvE5state.state addrspace(4)* %{{.*}}, i32 0, i32 0 + // CHECK-DEVICE: %[[MEM:[a-zA-Z0-9]+]] = getelementptr inbounds %{{.*}}, %struct.state addrspace(4)* %{{.*}}, i32 0, i32 0 // CHECK-DEVICE: %[[BITCAST:[0-9]+]] = bitcast [8 x i32] addrspace(4)* %[[MEM]] to i8 addrspace(4)* // CHECK-DEVICE: %[[ANN:[0-9]+]] = call i8 addrspace(4)* @llvm.ptr.annotation.p4i8(i8 addrspace(4)* %[[BITCAST]], {{.*}}, {{.*}}) // CHECK-DEVICE: %{{[0-9]+}} = bitcast i8 addrspace(4)* %[[ANN]] to [8 x i32] addrspace(4) diff --git a/clang/test/CodeGenSYCL/kernel-handler.cpp b/clang/test/CodeGenSYCL/kernel-handler.cpp index 7046916224e6b..508f49a2983b5 100644 --- a/clang/test/CodeGenSYCL/kernel-handler.cpp +++ b/clang/test/CodeGenSYCL/kernel-handler.cpp @@ -24,16 +24,16 @@ void test(int val) { // NONATIVESUPPORT: define dso_local void @{{.*}}test_kernel_handler{{[^(]*}} // NONATIVESUPPORT-SAME: (i32 %_arg_, i8 addrspace(1)* %_arg__specialization_constants_buffer) -// NONATIVESUPPORT: %kh = alloca %"class.[[MANGLEDCLASS:[a-zA-Z0-9_]+]].cl::sycl::kernel_handler", align 1 +// NONATIVESUPPORT: %kh = alloca %"class.cl::sycl::kernel_handler", align 1 // NONATIVESUPPORT: %[[KH:[0-9]+]] = load i8 addrspace(1)*, i8 addrspace(1)** %_arg__specialization_constants_buffer.addr, align 8 // NONATIVESUPPORT: %[[ADDRSPACECAST:[0-9]+]] = addrspacecast i8 addrspace(1)* %[[KH]] to i8* -// NONATIVESUPPORT: call void @{{.*}}__init_specialization_constants_buffer{{.*}}(%"class.[[MANGLEDCLASS]].cl::sycl::kernel_handler"* nonnull align 1 dereferenceable(1) %kh, i8* %[[ADDRSPACECAST]]) +// NONATIVESUPPORT: call void @{{.*}}__init_specialization_constants_buffer{{.*}}(%"class.cl::sycl::kernel_handler"* nonnull align 1 dereferenceable(1) %kh, i8* %[[ADDRSPACECAST]]) // NONATIVESUPPORT: void @[[MANGLEDKERNELCALL:[a-zA-Z0-9_$]+]] -// NONATIVESUPPORT-SAME: byval(%"class.[[MANGLEDCLASS]].cl::sycl::kernel_handler") +// NONATIVESUPPORT-SAME: byval(%"class.cl::sycl::kernel_handler") // NATIVESUPPORT: define dso_local spir_kernel void @{{.*}}test_kernel_handler{{[^(]*}} // NATIVESUPPORT-SAME: (i32 %_arg_) -// NATIVESUPPORT: %kh = alloca %"class.[[MANGLEDCLASS:[a-zA-Z0-9_]+]].cl::sycl::kernel_handler" +// NATIVESUPPORT: %kh = alloca %"class.cl::sycl::kernel_handler" // NATIVESUPPORT-NOT: __init_specialization_constants_buffer // NATIVE-SUPPORT: call spir_func void @"[[MANGLEDKERNELCALL:[a-zA-Z0-9_$]+]]" -// NATIVE-SUPPORT-SAME: byval(%"class.[[MANGLEDCLASS]].cl::sycl::kernel_handler") +// NATIVE-SUPPORT-SAME: byval(%"class.cl::sycl::kernel_handler") diff --git a/clang/test/CodeGenSYCL/kernel-param-acc-array.cpp b/clang/test/CodeGenSYCL/kernel-param-acc-array.cpp index 4633ce6cede59..64003b9f9921d 100644 --- a/clang/test/CodeGenSYCL/kernel-param-acc-array.cpp +++ b/clang/test/CodeGenSYCL/kernel-param-acc-array.cpp @@ -26,38 +26,38 @@ int main() { // Check kernel_A parameters // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_A // CHECK-SAME: i32 addrspace(1)* [[MEM_ARG1:%[a-zA-Z0-9_]+]], -// CHECK-SAME: %"struct.{{.*}}.cl::sycl::range"* byval{{.*}}align 4 [[ACC_RANGE1:%[a-zA-Z0-9_]+_1]], -// CHECK-SAME: %"struct.{{.*}}.cl::sycl::range"* byval{{.*}}align 4 [[MEM_RANGE1:%[a-zA-Z0-9_]+_2]], -// CHECK-SAME: %"struct.{{.*}}.cl::sycl::id"* byval{{.*}}align 4 [[OFFSET1:%[a-zA-Z0-9_]+_3]], +// CHECK-SAME: %"struct.cl::sycl::range"* byval{{.*}}align 4 [[ACC_RANGE1:%[a-zA-Z0-9_]+_1]], +// CHECK-SAME: %"struct.cl::sycl::range"* byval{{.*}}align 4 [[MEM_RANGE1:%[a-zA-Z0-9_]+_2]], +// CHECK-SAME: %"struct.cl::sycl::id"* byval{{.*}}align 4 [[OFFSET1:%[a-zA-Z0-9_]+_3]], // CHECK-SAME: i32 addrspace(1)* [[MEM_ARG2:%[a-zA-Z0-9_]+_4]], -// CHECK-SAME: %"struct.{{.*}}.cl::sycl::range"* byval{{.*}}align 4 [[ACC_RANGE2:%[a-zA-Z0-9_]+_6]], -// CHECK-SAME: %"struct.{{.*}}.cl::sycl::range"* byval{{.*}}align 4 [[MEM_RANGE2:%[a-zA-Z0-9_]+_7]], -// CHECK-SAME: %"struct.{{.*}}.cl::sycl::id"* byval{{.*}}align 4 [[OFFSET2:%[a-zA-Z0-9_]+_8]]) +// CHECK-SAME: %"struct.cl::sycl::range"* byval{{.*}}align 4 [[ACC_RANGE2:%[a-zA-Z0-9_]+_6]], +// CHECK-SAME: %"struct.cl::sycl::range"* byval{{.*}}align 4 [[MEM_RANGE2:%[a-zA-Z0-9_]+_7]], +// CHECK-SAME: %"struct.cl::sycl::id"* byval{{.*}}align 4 [[OFFSET2:%[a-zA-Z0-9_]+_8]]) // CHECK alloca for pointer arguments // CHECK: [[MEM_ARG1:%[a-zA-Z0-9_.]+]] = alloca i32 addrspace(1)*, align 8 // CHECK: [[MEM_ARG2:%[a-zA-Z0-9_.]+]] = alloca i32 addrspace(1)*, align 8 // CHECK lambda object alloca -// CHECK: [[LOCAL_OBJECTA:%0]] = alloca %class.{{.*}}.anon, align 4 -// CHECK: [[LOCAL_OBJECT:%.*]] = addrspacecast %class.{{.*}}.anon* [[LOCAL_OBJECTA]] to %class.{{.*}}.anon addrspace(4)* +// CHECK: [[LOCAL_OBJECTA:%0]] = alloca %class.anon, align 4 +// CHECK: [[LOCAL_OBJECT:%.*]] = addrspacecast %class.anon* [[LOCAL_OBJECTA]] to %class.anon addrspace(4)* // CHECK allocas for ranges -// CHECK: [[ACC_RANGE1A:%[a-zA-Z0-9_.]+]] = alloca %"struct.{{.*}}.cl::sycl::range" -// CHECK: [[ACC_RANGE1AS:%.*]] = addrspacecast %"struct.{{.*}}.cl::sycl::range"* [[ACC_RANGE1A]] to %"struct.{{.*}}.cl::sycl::range" addrspace(4)* -// CHECK: [[MEM_RANGE1A:%[a-zA-Z0-9_.]+]] = alloca %"struct.{{.*}}.cl::sycl::range" -// CHECK: [[MEM_RANGE1AS:%.*]] = addrspacecast %"struct.{{.*}}.cl::sycl::range"* [[MEM_RANGE1A]] to %"struct.{{.*}}.cl::sycl::range" addrspace(4)* -// CHECK: [[OFFSET1A:%[a-zA-Z0-9_.]+]] = alloca %"struct.{{.*}}.cl::sycl::id" -// CHECK: [[OFFSET1AS:%.*]] = addrspacecast %"struct.{{.*}}.cl::sycl::id"* [[OFFSET1A]] to %"struct.{{.*}}.cl::sycl::id" addrspace(4)* -// CHECK: [[ACC_RANGE2A:%[a-zA-Z0-9_.]+]] = alloca %"struct.{{.*}}.cl::sycl::range" -// CHECK: [[ACC_RANGE2AS:%.*]] = addrspacecast %"struct.{{.*}}.cl::sycl::range"* [[ACC_RANGE2A]] to %"struct.{{.*}}.cl::sycl::range" addrspace(4)* -// CHECK: [[MEM_RANGE2A:%[a-zA-Z0-9_.]+]] = alloca %"struct.{{.*}}.cl::sycl::range" -// CHECK: [[MEM_RANGE2AS:%.*]] = addrspacecast %"struct.{{.*}}.cl::sycl::range"* [[MEM_RANGE2A]] to %"struct.{{.*}}.cl::sycl::range" addrspace(4)* -// CHECK: [[OFFSET2A:%[a-zA-Z0-9_.]+]] = alloca %"struct.{{.*}}.cl::sycl::id" -// CHECK: [[OFFSET2AS:%.*]] = addrspacecast %"struct.{{.*}}.cl::sycl::id"* [[OFFSET2A]] to %"struct.{{.*}}.cl::sycl::id" addrspace(4)* +// CHECK: [[ACC_RANGE1A:%[a-zA-Z0-9_.]+]] = alloca %"struct.cl::sycl::range" +// CHECK: [[ACC_RANGE1AS:%.*]] = addrspacecast %"struct.cl::sycl::range"* [[ACC_RANGE1A]] to %"struct.cl::sycl::range" addrspace(4)* +// CHECK: [[MEM_RANGE1A:%[a-zA-Z0-9_.]+]] = alloca %"struct.cl::sycl::range" +// CHECK: [[MEM_RANGE1AS:%.*]] = addrspacecast %"struct.cl::sycl::range"* [[MEM_RANGE1A]] to %"struct.cl::sycl::range" addrspace(4)* +// CHECK: [[OFFSET1A:%[a-zA-Z0-9_.]+]] = alloca %"struct.cl::sycl::id" +// CHECK: [[OFFSET1AS:%.*]] = addrspacecast %"struct.cl::sycl::id"* [[OFFSET1A]] to %"struct.cl::sycl::id" addrspace(4)* +// CHECK: [[ACC_RANGE2A:%[a-zA-Z0-9_.]+]] = alloca %"struct.cl::sycl::range" +// CHECK: [[ACC_RANGE2AS:%.*]] = addrspacecast %"struct.cl::sycl::range"* [[ACC_RANGE2A]] to %"struct.cl::sycl::range" addrspace(4)* +// CHECK: [[MEM_RANGE2A:%[a-zA-Z0-9_.]+]] = alloca %"struct.cl::sycl::range" +// CHECK: [[MEM_RANGE2AS:%.*]] = addrspacecast %"struct.cl::sycl::range"* [[MEM_RANGE2A]] to %"struct.cl::sycl::range" addrspace(4)* +// CHECK: [[OFFSET2A:%[a-zA-Z0-9_.]+]] = alloca %"struct.cl::sycl::id" +// CHECK: [[OFFSET2AS:%.*]] = addrspacecast %"struct.cl::sycl::id"* [[OFFSET2A]] to %"struct.cl::sycl::id" addrspace(4)* // CHECK accessor array default inits -// CHECK: [[ACCESSOR_ARRAY1:%[a-zA-Z0-9_]+]] = getelementptr inbounds %class.{{.*}}.anon, %class.{{.*}}.anon addrspace(4)* [[LOCAL_OBJECT]], i32 0, i32 0 +// CHECK: [[ACCESSOR_ARRAY1:%[a-zA-Z0-9_]+]] = getelementptr inbounds %class.anon, %class.anon addrspace(4)* [[LOCAL_OBJECT]], i32 0, i32 0 // CHECK: [[BEGIN:%[a-zA-Z0-9._]*]] = getelementptr inbounds [2 x [[ACCESSOR:.*]]], [2 x [[ACCESSOR]]] addrspace(4)* [[ACCESSOR_ARRAY1]], i64 0, i64 0 // Clang takes advantage of element 1 having the same address as the array, so it doesn't do a GEP. // CTOR Call #1 @@ -67,21 +67,21 @@ int main() { // CHECK: call spir_func void @{{.+}}([[ACCESSOR]] addrspace(4)* {{[^,]*}} [[ELEM2_GEP]]) // CHECK acc[0] __init method call -// CHECK: [[ACCESSOR_ARRAY1:%[a-zA-Z0-9_]+]] = getelementptr inbounds %class.{{.*}}.anon, %class.{{.*}}.anon addrspace(4)* [[LOCAL_OBJECT]], i32 0, i32 0 +// CHECK: [[ACCESSOR_ARRAY1:%[a-zA-Z0-9_]+]] = getelementptr inbounds %class.anon, %class.anon addrspace(4)* [[LOCAL_OBJECT]], i32 0, i32 0 // CHECK: [[INDEX1:%[a-zA-Z0-9._]*]] = getelementptr inbounds [2 x [[ACCESSOR]]], [2 x [[ACCESSOR]]] addrspace(4)* [[ACCESSOR_ARRAY1]], i64 0, i64 0 // CHECK load from kernel pointer argument alloca // CHECK: [[MEM_LOAD1:%[a-zA-Z0-9_]+]] = load i32 addrspace(1)*, i32 addrspace(1)* addrspace(4)* [[MEM_ARG1]] -// CHECK: [[ACC_RANGE1:%.*]] = addrspacecast %"struct.{{.*}}.cl::sycl::range" addrspace(4)* [[ACC_RANGE1AS]] to %"struct.{{.*}}.cl::sycl::range"* -// CHECK: [[MEM_RANGE1:%.*]] = addrspacecast %"struct.{{.*}}.cl::sycl::range" addrspace(4)* [[MEM_RANGE1AS]] to %"struct.{{.*}}.cl::sycl::range"* -// CHECK: [[OFFSET1:%.*]] = addrspacecast %"struct.{{.*}}.cl::sycl::id" addrspace(4)* [[OFFSET1AS]] to %"struct.{{.*}}.cl::sycl::id"* -// CHECK: call spir_func void @{{.*}}__init{{.*}}(%"class.{{.*}}.cl::sycl::accessor" addrspace(4)* {{[^,]*}} [[INDEX1]], i32 addrspace(1)* [[MEM_LOAD1]], %"struct.{{.*}}.cl::sycl::range"* byval({{.*}}) align 4 [[ACC_RANGE1]], %"struct.{{.*}}.cl::sycl::range"* byval({{.*}}) align 4 [[MEM_RANGE1]], %"struct.{{.*}}.cl::sycl::id"* byval({{.*}}) align 4 [[OFFSET1]]) +// CHECK: [[ACC_RANGE1:%.*]] = addrspacecast %"struct.cl::sycl::range" addrspace(4)* [[ACC_RANGE1AS]] to %"struct.cl::sycl::range"* +// CHECK: [[MEM_RANGE1:%.*]] = addrspacecast %"struct.cl::sycl::range" addrspace(4)* [[MEM_RANGE1AS]] to %"struct.cl::sycl::range"* +// CHECK: [[OFFSET1:%.*]] = addrspacecast %"struct.cl::sycl::id" addrspace(4)* [[OFFSET1AS]] to %"struct.cl::sycl::id"* +// CHECK: call spir_func void @{{.*}}__init{{.*}}(%"class.cl::sycl::accessor" addrspace(4)* {{[^,]*}} [[INDEX1]], i32 addrspace(1)* [[MEM_LOAD1]], %"struct.cl::sycl::range"* byval({{.*}}) align 4 [[ACC_RANGE1]], %"struct.cl::sycl::range"* byval({{.*}}) align 4 [[MEM_RANGE1]], %"struct.cl::sycl::id"* byval({{.*}}) align 4 [[OFFSET1]]) // CHECK acc[1] __init method call -// CHECK: [[ACCESSOR_ARRAY2:%[a-zA-Z0-9_]+]] = getelementptr inbounds %class.{{.*}}.anon, %class.{{.*}}.anon addrspace(4)* [[LOCAL_OBJECT]], i32 0, i32 0 +// CHECK: [[ACCESSOR_ARRAY2:%[a-zA-Z0-9_]+]] = getelementptr inbounds %class.anon, %class.anon addrspace(4)* [[LOCAL_OBJECT]], i32 0, i32 0 // CHECK: [[INDEX2:%[a-zA-Z0-9._]*]] = getelementptr inbounds [2 x [[ACCESSOR]]], [2 x [[ACCESSOR]]] addrspace(4)* [[ACCESSOR_ARRAY2]], i64 0, i64 1 // CHECK load from kernel pointer argument alloca // CHECK: [[MEM_LOAD2:%[a-zA-Z0-9_]+]] = load i32 addrspace(1)*, i32 addrspace(1)* addrspace(4)* [[MEM_ARG2]] -// CHECK: [[ACC_RANGE2:%.*]] = addrspacecast %"struct.{{.*}}.cl::sycl::range" addrspace(4)* [[ACC_RANGE2AS]] to %"struct.{{.*}}.cl::sycl::range"* -// CHECK: [[MEM_RANGE2:%.*]] = addrspacecast %"struct.{{.*}}.cl::sycl::range" addrspace(4)* [[MEM_RANGE2AS]] to %"struct.{{.*}}.cl::sycl::range"* -// CHECK: [[OFFSET2:%.*]] = addrspacecast %"struct.{{.*}}.cl::sycl::id" addrspace(4)* [[OFFSET2AS]] to %"struct.{{.*}}.cl::sycl::id"* -// CHECK: call spir_func void @{{.*}}__init{{.*}}(%"class.{{.*}}.cl::sycl::accessor" addrspace(4)* {{[^,]*}} [[INDEX2]], i32 addrspace(1)* [[MEM_LOAD2]], %"struct.{{.*}}.cl::sycl::range"* byval({{.*}}) align 4 [[ACC_RANGE2]], %"struct.{{.*}}.cl::sycl::range"* byval({{.*}}) align 4 [[MEM_RANGE2]], %"struct.{{.*}}.cl::sycl::id"* byval({{.*}}) align 4 [[OFFSET2]]) +// CHECK: [[ACC_RANGE2:%.*]] = addrspacecast %"struct.cl::sycl::range" addrspace(4)* [[ACC_RANGE2AS]] to %"struct.cl::sycl::range"* +// CHECK: [[MEM_RANGE2:%.*]] = addrspacecast %"struct.cl::sycl::range" addrspace(4)* [[MEM_RANGE2AS]] to %"struct.cl::sycl::range"* +// CHECK: [[OFFSET2:%.*]] = addrspacecast %"struct.cl::sycl::id" addrspace(4)* [[OFFSET2AS]] to %"struct.cl::sycl::id"* +// CHECK: call spir_func void @{{.*}}__init{{.*}}(%"class.cl::sycl::accessor" addrspace(4)* {{[^,]*}} [[INDEX2]], i32 addrspace(1)* [[MEM_LOAD2]], %"struct.cl::sycl::range"* byval({{.*}}) align 4 [[ACC_RANGE2]], %"struct.cl::sycl::range"* byval({{.*}}) align 4 [[MEM_RANGE2]], %"struct.cl::sycl::id"* byval({{.*}}) align 4 [[OFFSET2]]) diff --git a/clang/test/CodeGenSYCL/kernel-param-member-acc-array.cpp b/clang/test/CodeGenSYCL/kernel-param-member-acc-array.cpp index 5d1ce4dcbed8d..c135e89a3845b 100644 --- a/clang/test/CodeGenSYCL/kernel-param-member-acc-array.cpp +++ b/clang/test/CodeGenSYCL/kernel-param-member-acc-array.cpp @@ -46,18 +46,18 @@ int main() { // CHECK: [[LOCAL_OBJECT:%.*]] = addrspacecast %class{{.*}}.anon* %0 to %class{{.*}}.anon addrspace(4)* // Check allocas for ranges -// CHECK: [[ACC_RANGE1A:%[a-zA-Z0-9_.]+]] = alloca %"struct.{{.*}}.cl::sycl::range" -// CHECK: [[ACC_RANGE1AS:%.*]] = addrspacecast %"struct.{{.*}}.cl::sycl::range"* [[ACC_RANGE1A]] to %"struct.{{.*}}.cl::sycl::range" addrspace(4)* -// CHECK: [[MEM_RANGE1A:%[a-zA-Z0-9_.]+]] = alloca %"struct.{{.*}}.cl::sycl::range" -// CHECK: [[MEM_RANGE1AS:%.*]] = addrspacecast %"struct.{{.*}}.cl::sycl::range"* [[MEM_RANGE1A]] to %"struct.{{.*}}.cl::sycl::range" addrspace(4)* -// CHECK: [[OFFSET1A:%[a-zA-Z0-9_.]+]] = alloca %"struct.{{.*}}.cl::sycl::id" -// CHECK: [[OFFSET1AS:%.*]] = addrspacecast %"struct.{{.*}}.cl::sycl::id"* [[OFFSET1A]] to %"struct.{{.*}}.cl::sycl::id" addrspace(4)* -// CHECK: [[ACC_RANGE2A:%[a-zA-Z0-9_.]+]] = alloca %"struct.{{.*}}.cl::sycl::range" -// CHECK: [[ACC_RANGE2AS:%.*]] = addrspacecast %"struct.{{.*}}.cl::sycl::range"* [[ACC_RANGE2A]] to %"struct.{{.*}}.cl::sycl::range" addrspace(4)* -// CHECK: [[MEM_RANGE2A:%[a-zA-Z0-9_.]+]] = alloca %"struct.{{.*}}.cl::sycl::range" -// CHECK: [[MEM_RANGE2AS:%.*]] = addrspacecast %"struct.{{.*}}.cl::sycl::range"* [[MEM_RANGE2A]] to %"struct.{{.*}}.cl::sycl::range" addrspace(4)* -// CHECK: [[OFFSET2A:%[a-zA-Z0-9_.]+]] = alloca %"struct.{{.*}}.cl::sycl::id" -// CHECK: [[OFFSET2AS:%.*]] = addrspacecast %"struct.{{.*}}.cl::sycl::id"* [[OFFSET2A]] to %"struct.{{.*}}.cl::sycl::id" addrspace(4)* +// CHECK: [[ACC_RANGE1A:%[a-zA-Z0-9_.]+]] = alloca %"struct.cl::sycl::range" +// CHECK: [[ACC_RANGE1AS:%.*]] = addrspacecast %"struct.cl::sycl::range"* [[ACC_RANGE1A]] to %"struct.cl::sycl::range" addrspace(4)* +// CHECK: [[MEM_RANGE1A:%[a-zA-Z0-9_.]+]] = alloca %"struct.cl::sycl::range" +// CHECK: [[MEM_RANGE1AS:%.*]] = addrspacecast %"struct.cl::sycl::range"* [[MEM_RANGE1A]] to %"struct.cl::sycl::range" addrspace(4)* +// CHECK: [[OFFSET1A:%[a-zA-Z0-9_.]+]] = alloca %"struct.cl::sycl::id" +// CHECK: [[OFFSET1AS:%.*]] = addrspacecast %"struct.cl::sycl::id"* [[OFFSET1A]] to %"struct.cl::sycl::id" addrspace(4)* +// CHECK: [[ACC_RANGE2A:%[a-zA-Z0-9_.]+]] = alloca %"struct.cl::sycl::range" +// CHECK: [[ACC_RANGE2AS:%.*]] = addrspacecast %"struct.cl::sycl::range"* [[ACC_RANGE2A]] to %"struct.cl::sycl::range" addrspace(4)* +// CHECK: [[MEM_RANGE2A:%[a-zA-Z0-9_.]+]] = alloca %"struct.cl::sycl::range" +// CHECK: [[MEM_RANGE2AS:%.*]] = addrspacecast %"struct.cl::sycl::range"* [[MEM_RANGE2A]] to %"struct.cl::sycl::range" addrspace(4)* +// CHECK: [[OFFSET2A:%[a-zA-Z0-9_.]+]] = alloca %"struct.cl::sycl::id" +// CHECK: [[OFFSET2AS:%.*]] = addrspacecast %"struct.cl::sycl::id"* [[OFFSET2A]] to %"struct.cl::sycl::id" addrspace(4)* // CHECK accessor array default inits // CHECK: [[ACCESSOR_WRAPPER:%[a-zA-Z0-9_]+]] = getelementptr inbounds %class{{.*}}.anon, %class{{.*}}.anon addrspace(4)* [[LOCAL_OBJECT]], i32 0, i32 0 @@ -74,9 +74,9 @@ int main() { // CHECK: [[GEP_MEMBER_ACC1:%[a-zA-Z0-9_]+]] = getelementptr inbounds %struct{{.*}}.struct_acc_t, %struct{{.*}}.struct_acc_t addrspace(4)* [[GEP_LAMBDA1]], i32 0, i32 0 // CHECK: [[ARRAY_IDX1:%[a-zA-Z0-9._]*]] = getelementptr inbounds [2 x [[ACCESSOR]]], [2 x [[ACCESSOR]]] addrspace(4)* [[GEP_MEMBER_ACC1]], i64 0, i64 0 // CHECK: [[MEM_LOAD1:%[a-zA-Z0-9_]+]] = load i32 addrspace(1)*, i32 addrspace(1)* addrspace(4)* [[MEM_ARG1]].addr -// CHECK: [[ACC_RANGE1:%.*]] = addrspacecast %"struct.{{.*}}.cl::sycl::range" addrspace(4)* [[ACC_RANGE1AS]] to %"struct.{{.*}}.cl::sycl::range"* -// CHECK: [[MEM_RANGE1:%.*]] = addrspacecast %"struct.{{.*}}.cl::sycl::range" addrspace(4)* [[MEM_RANGE1AS]] to %"struct.{{.*}}.cl::sycl::range"* -// CHECK: [[OFFSET1:%.*]] = addrspacecast %"struct.{{.*}}.cl::sycl::id" addrspace(4)* [[OFFSET1AS]] to %"struct.{{.*}}.cl::sycl::id"* +// CHECK: [[ACC_RANGE1:%.*]] = addrspacecast %"struct.cl::sycl::range" addrspace(4)* [[ACC_RANGE1AS]] to %"struct.cl::sycl::range"* +// CHECK: [[MEM_RANGE1:%.*]] = addrspacecast %"struct.cl::sycl::range" addrspace(4)* [[MEM_RANGE1AS]] to %"struct.cl::sycl::range"* +// CHECK: [[OFFSET1:%.*]] = addrspacecast %"struct.cl::sycl::id" addrspace(4)* [[OFFSET1AS]] to %"struct.cl::sycl::id"* // CHECK: call spir_func void @{{.*}}__init{{.*}}([[ACCESSOR]] addrspace(4)* {{[^,]*}} [[ARRAY_IDX1]], i32 addrspace(1)* [[MEM_LOAD1]], %"struct{{.*}}.cl::sycl::range"* byval({{.*}}) align 4 [[ACC_RANGE1]], %"struct{{.*}}.cl::sycl::range"* byval({{.*}}) align 4 [[MEM_RANGE1]], %"struct{{.*}}.cl::sycl::id"* byval({{.*}}) align 4 [[OFFSET1]]) // Check acc[1] __init method call @@ -84,7 +84,7 @@ int main() { // CHECK: [[GEP_MEMBER_ACC2:%[a-zA-Z0-9_]+]] = getelementptr inbounds %struct{{.*}}.struct_acc_t, %struct{{.*}}.struct_acc_t addrspace(4)* [[GEP_LAMBDA2]], i32 0, i32 0 // CHECK: [[ARRAY_IDX2:%[a-zA-Z0-9_]*]] = getelementptr inbounds [2 x [[ACCESSOR]]], [2 x [[ACCESSOR]]] addrspace(4)* [[GEP_MEMBER_ACC2]], i64 0, i64 1 // CHECK: [[MEM_LOAD2:%[a-zA-Z0-9_]+]] = load i32 addrspace(1)*, i32 addrspace(1)* addrspace(4)* [[MEM_ARG1]].addr -// CHECK: [[ACC_RANGE2:%.*]] = addrspacecast %"struct.{{.*}}.cl::sycl::range" addrspace(4)* [[ACC_RANGE2AS]] to %"struct.{{.*}}.cl::sycl::range"* -// CHECK: [[MEM_RANGE2:%.*]] = addrspacecast %"struct.{{.*}}.cl::sycl::range" addrspace(4)* [[MEM_RANGE2AS]] to %"struct.{{.*}}.cl::sycl::range"* -// CHECK: [[OFFSET2:%.*]] = addrspacecast %"struct.{{.*}}.cl::sycl::id" addrspace(4)* [[OFFSET2AS]] to %"struct.{{.*}}.cl::sycl::id"* +// CHECK: [[ACC_RANGE2:%.*]] = addrspacecast %"struct.cl::sycl::range" addrspace(4)* [[ACC_RANGE2AS]] to %"struct.cl::sycl::range"* +// CHECK: [[MEM_RANGE2:%.*]] = addrspacecast %"struct.cl::sycl::range" addrspace(4)* [[MEM_RANGE2AS]] to %"struct.cl::sycl::range"* +// CHECK: [[OFFSET2:%.*]] = addrspacecast %"struct.cl::sycl::id" addrspace(4)* [[OFFSET2AS]] to %"struct.cl::sycl::id"* // CHECK: call spir_func void @{{.*}}__init{{.*}}([[ACCESSOR]] addrspace(4)* {{[^,]*}} [[ARRAY_IDX2]], i32 addrspace(1)* [[MEM_LOAD2]], %"struct{{.*}}.cl::sycl::range"* byval({{.*}}) align 4 [[ACC_RANGE2]], %"struct{{.*}}.cl::sycl::range"* byval({{.*}}) align 4 [[MEM_RANGE2]], %"struct{{.*}}.cl::sycl::id"* byval({{.*}}) align 4 [[OFFSET2]]) diff --git a/clang/test/CodeGenSYCL/sampler.cpp b/clang/test/CodeGenSYCL/sampler.cpp index 72cbc0b4f3ed9..7330e17c5f54e 100644 --- a/clang/test/CodeGenSYCL/sampler.cpp +++ b/clang/test/CodeGenSYCL/sampler.cpp @@ -2,14 +2,14 @@ // CHECK: define {{.*}}spir_kernel void @{{[a-zA-Z0-9_]+}}(%opencl.sampler_t addrspace(2)* [[SAMPLER_ARG:%[a-zA-Z0-9_]+]]) // CHECK-NEXT: entry: // CHECK-NEXT: [[SAMPLER_ARG]].addr = alloca %opencl.sampler_t addrspace(2)*, align 8 -// CHECK: [[ANON:%[0-9]+]] = alloca %class.{{.*}}.anon, align 8 -// CHECK: [[ANONCAST:%[0-9]+]] = addrspacecast %class.{{.*}}.anon* [[ANON]] to %class.{{.*}}.anon addrspace(4)* +// CHECK: [[ANON:%[0-9]+]] = alloca %class.anon, align 8 +// CHECK: [[ANONCAST:%[0-9]+]] = addrspacecast %class.anon* [[ANON]] to %class.anon addrspace(4)* // CHECK: store %opencl.sampler_t addrspace(2)* [[SAMPLER_ARG]], %opencl.sampler_t addrspace(2)* addrspace(4)* [[SAMPLER_ARG]].addr.ascast, align 8 -// CHECK-NEXT: [[BITCAST:%[0-9]+]] = bitcast %class.{{.*}}.anon* [[ANON]] to i8* +// CHECK-NEXT: [[BITCAST:%[0-9]+]] = bitcast %class.anon* [[ANON]] to i8* // CHECK-NEXT: call void @llvm.lifetime.start.p0i8(i64 8, i8* [[BITCAST]]) #4 -// CHECK-NEXT: [[GEP:%[0-9]+]] = getelementptr inbounds %class.{{.*}}.anon, %class.{{.*}}.anon addrspace(4)* [[ANONCAST]], i32 0, i32 0 +// CHECK-NEXT: [[GEP:%[0-9]+]] = getelementptr inbounds %class.anon, %class.anon addrspace(4)* [[ANONCAST]], i32 0, i32 0 // CHECK-NEXT: [[LOAD_SAMPLER_ARG:%[0-9]+]] = load %opencl.sampler_t addrspace(2)*, %opencl.sampler_t addrspace(2)* addrspace(4)* [[SAMPLER_ARG]].addr.ascast, align 8 -// CHECK-NEXT: call spir_func void @{{[a-zA-Z0-9_]+}}(%"class.{{.*}}.cl::sycl::sampler" addrspace(4)* {{[^,]*}} [[GEP]], %opencl.sampler_t addrspace(2)* [[LOAD_SAMPLER_ARG]]) +// CHECK-NEXT: call spir_func void @{{[a-zA-Z0-9_]+}}(%"class.cl::sycl::sampler" addrspace(4)* {{[^,]*}} [[GEP]], %opencl.sampler_t addrspace(2)* [[LOAD_SAMPLER_ARG]]) // // CHECK: define {{.*}}spir_kernel void @{{[a-zA-Z0-9_]+}}(%opencl.sampler_t addrspace(2)* [[SAMPLER_ARG_WRAPPED:%[a-zA-Z0-9_]+]], i32 [[ARG_A:%[a-zA-Z0-9_]+]]) @@ -17,24 +17,24 @@ // Check alloca // CHECK: [[SAMPLER_ARG_WRAPPED]].addr = alloca %opencl.sampler_t addrspace(2)*, align 8 // CHECK: [[ARG_A]].addr = alloca i32, align 4 -// CHECK: [[LAMBDAA:%[0-9]+]] = alloca %class.{{.*}}.anon, align 8 -// CHECK: [[LAMBDA:%[0-9]+]] = addrspacecast %class.{{.*}}.anon* [[LAMBDAA]] to %class.{{.*}}.anon addrspace(4)* +// CHECK: [[LAMBDAA:%[0-9]+]] = alloca %class.anon.0, align 8 +// CHECK: [[LAMBDA:%[0-9]+]] = addrspacecast %class.anon.0* [[LAMBDAA]] to %class.anon.0 addrspace(4)* // Check argument store // CHECK: store %opencl.sampler_t addrspace(2)* [[SAMPLER_ARG_WRAPPED]], %opencl.sampler_t addrspace(2)* addrspace(4)* [[SAMPLER_ARG_WRAPPED]].addr.ascast, align 8 // CHECK: store i32 [[ARG_A]], i32 addrspace(4)* [[ARG_A]].addr.ascast, align 4 // Initialize 'a' -// CHECK: [[GEP_LAMBDA:%[0-9]+]] = getelementptr inbounds %class.{{.*}}.anon, %class.{{.*}}.anon addrspace(4)* [[LAMBDA]], i32 0, i32 0 -// CHECK: [[GEP_A:%[a-zA-Z0-9]+]] = getelementptr inbounds %struct.{{.*}}.sampler_wrapper, %struct.{{.*}}.sampler_wrapper addrspace(4)* [[GEP_LAMBDA]], i32 0, i32 1 +// CHECK: [[GEP_LAMBDA:%[0-9]+]] = getelementptr inbounds %class.anon.0, %class.anon.0 addrspace(4)* [[LAMBDA]], i32 0, i32 0 +// CHECK: [[GEP_A:%[a-zA-Z0-9]+]] = getelementptr inbounds %struct.sampler_wrapper, %struct.sampler_wrapper addrspace(4)* [[GEP_LAMBDA]], i32 0, i32 1 // CHECK: [[LOAD_A:%[0-9]+]] = load i32, i32 addrspace(4)* [[ARG_A]].addr.ascast, align 4 // CHECK: store i32 [[LOAD_A]], i32 addrspace(4)* [[GEP_A]], align 8 // Initialize wrapped sampler 'smpl' -// CHECK: [[GEP_LAMBDA_0:%[0-9]+]] = getelementptr inbounds %class.{{.*}}.anon, %class.{{.*}}.anon addrspace(4)* [[LAMBDA]], i32 0, i32 0 -// CHECK: [[GEP_SMPL:%[a-zA-Z0-9]+]] = getelementptr inbounds %struct.{{.*}}.sampler_wrapper, %struct.{{.*}}.sampler_wrapper addrspace(4)* [[GEP_LAMBDA_0]], i32 0, i32 0 +// CHECK: [[GEP_LAMBDA_0:%[0-9]+]] = getelementptr inbounds %class.anon.0, %class.anon.0 addrspace(4)* [[LAMBDA]], i32 0, i32 0 +// CHECK: [[GEP_SMPL:%[a-zA-Z0-9]+]] = getelementptr inbounds %struct.sampler_wrapper, %struct.sampler_wrapper addrspace(4)* [[GEP_LAMBDA_0]], i32 0, i32 0 // CHECK: [[LOAD_SMPL:%[0-9]+]] = load %opencl.sampler_t addrspace(2)*, %opencl.sampler_t addrspace(2)* addrspace(4)* [[SAMPLER_ARG_WRAPPED]].addr.ascast, align 8 -// CHECK: call spir_func void @{{[a-zA-Z0-9_]+}}(%"class.{{.*}}.cl::sycl::sampler" addrspace(4)* {{.*}}, %opencl.sampler_t addrspace(2)* [[LOAD_SMPL]]) +// CHECK: call spir_func void @{{[a-zA-Z0-9_]+}}(%"class.cl::sycl::sampler" addrspace(4)* {{.*}}, %opencl.sampler_t addrspace(2)* [[LOAD_SMPL]]) // #include "Inputs/sycl.hpp" diff --git a/clang/test/CodeGenSYCL/spir-calling-conv.cpp b/clang/test/CodeGenSYCL/spir-calling-conv.cpp index ebfb025707b0b..e9bb67528fd4f 100644 --- a/clang/test/CodeGenSYCL/spir-calling-conv.cpp +++ b/clang/test/CodeGenSYCL/spir-calling-conv.cpp @@ -9,7 +9,7 @@ int main() { // CHECK: define {{.*}}spir_kernel void @_ZTSZ4mainE15kernel_function() - // CHECK: call spir_func void @_ZZ4mainENKUlvE_clEv(%class.{{.*}}.anon addrspace(4)* {{[^,]*}} %{{.+}}) + // CHECK: call spir_func void @_ZZ4mainENKUlvE_clEv(%class.anon addrspace(4)* {{[^,]*}} %{{.+}}) // CHECK: define internal spir_func void @_ZZ4mainENKUlvE_clEv(%class.{{.*}}anon addrspace(4)* {{[^,]*}} %this) diff --git a/clang/test/CodeGenSYCL/spir-enum.cpp b/clang/test/CodeGenSYCL/spir-enum.cpp index 501e760234103..c7c392eefaee8 100644 --- a/clang/test/CodeGenSYCL/spir-enum.cpp +++ b/clang/test/CodeGenSYCL/spir-enum.cpp @@ -22,8 +22,8 @@ int main() { // CHECK: define {{.*}}spir_kernel void @_ZTSZ4test9enum_typeE15kernel_function(i32 %_arg_) - // CHECK: getelementptr inbounds %class.{{.*}}.anon, %class.{{.*}}.anon addrspace(4)* - // CHECK: call spir_func void @_ZZ4test9enum_typeENKUlvE_clEv(%class.{{.*}}.anon addrspace(4)* {{[^,]*}} %{{.+}}) + // CHECK: getelementptr inbounds %class.anon, %class.anon addrspace(4)* + // CHECK: call spir_func void @_ZZ4test9enum_typeENKUlvE_clEv(%class.anon addrspace(4)* {{[^,]*}} %{{.+}}) test( enum_type::B ); return 0; diff --git a/clang/test/CodeGenSYCL/stall_enable_device.cpp b/clang/test/CodeGenSYCL/stall_enable_device.cpp index 5a0368bebaa22..1b91bc160b2aa 100644 --- a/clang/test/CodeGenSYCL/stall_enable_device.cpp +++ b/clang/test/CodeGenSYCL/stall_enable_device.cpp @@ -47,7 +47,7 @@ int main() { // CHECK: define {{.*}}spir_kernel void @{{.*}}test_kernel4() // CHECK-NOT: !stall_enable // CHECK-SAME: { - // CHECK: define {{.*}}spir_func void @{{.*}}func1{{.*}}(%class.{{.*}}func1{{.*}}.anon addrspace(4)* align 1 dereferenceable_or_null(1) %this) #4 align 2 !stall_enable ![[NUM4]] + // CHECK: define {{.*}}spir_func void @{{.*}}func1{{.*}}(%class.anon{{.*}} addrspace(4)* align 1 dereferenceable_or_null(1) %this) #4 align 2 !stall_enable ![[NUM4]] h.single_task( []() { func1(); }); diff --git a/clang/test/CodeGenSYCL/union-kernel-param.cpp b/clang/test/CodeGenSYCL/union-kernel-param.cpp index fadcf7ce074ea..197e9e29f3ec5 100644 --- a/clang/test/CodeGenSYCL/union-kernel-param.cpp +++ b/clang/test/CodeGenSYCL/union-kernel-param.cpp @@ -28,15 +28,15 @@ int main() { } // CHECK kernel_A parameters -// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_A(%union.{{.*}}.MyUnion* byval(%union.{{.*}}.MyUnion) align 4 [[MEM_ARG:%[a-zA-Z0-9_]+]]) +// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_A(%union.MyUnion* byval(%union.MyUnion) align 4 [[MEM_ARG:%[a-zA-Z0-9_]+]]) // Check lambda object alloca -// CHECK: [[LOCAL_OBJECT:%0]] = alloca %class.{{.*}}.anon, align 4 +// CHECK: [[LOCAL_OBJECT:%0]] = alloca %class.anon, align 4 -// CHECK: [[LOCAL_OBJECTAS:%.*]] = addrspacecast %class.{{.*}}.anon* [[LOCAL_OBJECT]] to %class.{{.*}}.anon addrspace(4)* -// CHECK: [[MEM_ARGAS:%.*]] = addrspacecast %union.{{.*}}.MyUnion* [[MEM_ARG]] to %union.{{.*}}.MyUnion addrspace(4)* -// CHECK: [[L_STRUCT_ADDR:%[a-zA-Z0-9_]+]] = getelementptr inbounds %class.{{.*}}.anon, %class.{{.*}}.anon addrspace(4)* [[LOCAL_OBJECTAS]], i32 0, i32 0 +// CHECK: [[LOCAL_OBJECTAS:%.*]] = addrspacecast %class.anon* [[LOCAL_OBJECT]] to %class.anon addrspace(4)* +// CHECK: [[MEM_ARGAS:%.*]] = addrspacecast %union.MyUnion* [[MEM_ARG]] to %union.MyUnion addrspace(4)* +// CHECK: [[L_STRUCT_ADDR:%[a-zA-Z0-9_]+]] = getelementptr inbounds %class.anon, %class.anon addrspace(4)* [[LOCAL_OBJECTAS]], i32 0, i32 0 // CHECK: [[MEMCPY_DST:%[0-9a-zA-Z_]+]] = bitcast %union.{{.*}}MyUnion addrspace(4)* [[L_STRUCT_ADDR]] to i8 addrspace(4)* // CHECK: [[MEMCPY_SRC:%[0-9a-zA-Z_]+]] = bitcast %union.{{.*}}MyUnion addrspace(4)* [[MEM_ARGAS]] to i8 addrspace(4)* // CHECK: call void @llvm.memcpy.p4i8.p4i8.i64(i8 addrspace(4)* align 4 [[MEMCPY_DST]], i8 addrspace(4)* align 4 [[MEMCPY_SRC]], i64 12, i1 false) -// CHECK: call spir_func void @{{.*}}(%class.{{.*}}.anon addrspace(4)* {{[^,]*}} [[LOCAL_OBJECTAS]]) +// CHECK: call spir_func void @{{.*}}(%class.anon addrspace(4)* {{[^,]*}} [[LOCAL_OBJECTAS]]) diff --git a/clang/test/CodeGenSYCL/virtual-types.cpp b/clang/test/CodeGenSYCL/virtual-types.cpp index a723f4579965c..fb69a06b63c45 100644 --- a/clang/test/CodeGenSYCL/virtual-types.cpp +++ b/clang/test/CodeGenSYCL/virtual-types.cpp @@ -18,9 +18,9 @@ int main() { // Struct layout big enough for vtable. -// CHECK: %struct._ZTS6Struct.Struct = type { i32 (...)** } +// CHECK: %struct.Struct = type { i32 (...)** } // VTable: -// CHECK: @_ZTV6Struct = linkonce_odr unnamed_addr constant { [3 x i8*] } { [3 x i8*] [i8* null, i8* bitcast ({ i8*, i8* }* @_ZTI6Struct to i8*), i8* bitcast (void (%struct._ZTS6Struct.Struct addrspace(4)*)* @_ZN6Struct3fooEv to i8*)] }, comdat, align 8 +// CHECK: @_ZTV6Struct = linkonce_odr unnamed_addr constant { [3 x i8*] } { [3 x i8*] [i8* null, i8* bitcast ({ i8*, i8* }* @_ZTI6Struct to i8*), i8* bitcast (void (%struct.Struct addrspace(4)*)* @_ZN6Struct3fooEv to i8*)] }, comdat, align 8 // CHECK: @[[TYPEINFO:.+]] = external global i8* // TypeInfo Name: // CHECK: @_ZTS6Struct = linkonce_odr constant [8 x i8] c"6Struct\00", comdat, align 1 diff --git a/llvm/lib/SYCLLowerIR/LowerESIMD.cpp b/llvm/lib/SYCLLowerIR/LowerESIMD.cpp index 029263c0dc35c..e82d153eae43b 100644 --- a/llvm/lib/SYCLLowerIR/LowerESIMD.cpp +++ b/llvm/lib/SYCLLowerIR/LowerESIMD.cpp @@ -1351,15 +1351,17 @@ SmallPtrSet collectGenXVolatileTypes(Module &M) { continue; auto GTy = dyn_cast(PTy->getPointerElementType()); // TODO FIXME relying on type name in LLVM IR is fragile, needs rework - if (!GTy || - !GTy->getName().endswith("sycl::ext::intel::experimental::esimd::simd")) + if (!GTy || !GTy->getName() + .rtrim(".0123456789") + .endswith("sycl::ext::intel::experimental::esimd::simd")) continue; assert(GTy->getNumContainedTypes() == 1); auto VTy = GTy->getContainedType(0); if ((GTy = dyn_cast(VTy))) { - assert( - GTy->getName().endswith( - "sycl::ext::intel::experimental::esimd::detail::simd_obj_impl")); + assert(GTy->getName() + .rtrim(".0123456789") + .endswith("sycl::ext::intel::experimental::esimd::detail::" + "simd_obj_impl")); VTy = GTy->getContainedType(0); } assert(VTy->isVectorTy()); diff --git a/llvm/lib/SYCLLowerIR/LowerESIMDVLoadVStore.cpp b/llvm/lib/SYCLLowerIR/LowerESIMDVLoadVStore.cpp index a3d795ec7a66f..2b4df4cc2b70a 100644 --- a/llvm/lib/SYCLLowerIR/LowerESIMDVLoadVStore.cpp +++ b/llvm/lib/SYCLLowerIR/LowerESIMDVLoadVStore.cpp @@ -68,7 +68,7 @@ char ESIMDLowerLoadStore::ID = 0; INITIALIZE_PASS(ESIMDLowerLoadStore, "ESIMDLowerLoadStore", "Lower ESIMD reference loads and stores", false, false) -// Lower non-volatilE vload/vstore intrinsic calls into normal load/store +// Lower non-volatile vload/vstore intrinsic calls into normal load/store // instructions. PreservedAnalyses ESIMDLowerLoadStorePass::run(Function &F, FunctionAnalysisManager &FAM) { diff --git a/llvm/lib/SYCLLowerIR/LowerESIMDVecArg.cpp b/llvm/lib/SYCLLowerIR/LowerESIMDVecArg.cpp index 848b289352a03..53fd3795ddba9 100644 --- a/llvm/lib/SYCLLowerIR/LowerESIMDVecArg.cpp +++ b/llvm/lib/SYCLLowerIR/LowerESIMDVecArg.cpp @@ -15,7 +15,7 @@ // Parameter %0 is of type simd* // define dso_local spir_func void @_Z3fooPiN2cm3gen4simdIiLi16EEE(i32 // addrspace(4)* %C, -// "class._ZTSN2cm3gen4simdIiLi16EEE.cm::gen::simd" * %0) +// "class.cm::gen::simd" * %0) // local_unnamed_addr #2 { // // New IR: @@ -29,7 +29,7 @@ // <16 x i32>* %0) local_unnamed_addr #2 { // entry: // % 1 = bitcast<16 x i32> * % 0 to % -// "class._ZTSN2cm3gen4simdIiLi16EEE.cm::gen::simd" * +// "class.cm::gen::simd" * // // It is OK not to rewrite a function (for example, when its address is taken) // since it does not affect correctness. But that may lead to vector backend @@ -39,15 +39,15 @@ // // Old IR: // ====== -// @vc = global %"class._ZTSN2cm3gen4simdIiLi16EEE.cm::gen::simd" +// @vc = global %"class.cm::gen::simd" // zeroinitializer, align 64 #0 // // % call.cm.i.i = tail call<16 x i32> @llvm.genx.vload.v16i32.p4v16i32( // <16 x i32> addrspace(4) * getelementptr( -// % "class._ZTSN2cm3gen4simdIiLi16EEE.cm::gen::simd", -// % "class._ZTSN2cm3gen4simdIiLi16EEE.cm::gen::simd" addrspace(4) * -// addrspacecast(% "class._ZTSN2cm3gen4simdIiLi16EEE.cm::gen::simd" * @vc to -// % "class._ZTSN2cm3gen4simdIiLi16EEE.cm::gen::simd" addrspace(4) *), i64 0, +// % "class.cm::gen::simd", +// % "class.cm::gen::simd" addrspace(4) * +// addrspacecast(% "class.cm::gen::simd" * @vc to +// % "class.cm::gen::simd" addrspace(4) *), i64 0, // i32 0)) // // New IR: @@ -58,12 +58,12 @@ // // % call.cm.i.i = tail call<16 x i32> @llvm.genx.vload.v16i32.p4v16i32( // <16 x i32> addrspace(4) * getelementptr( -// % "class._ZTSN2cm3gen4simdIiLi16EEE.cm::gen::simd", -// % "class._ZTSN2cm3gen4simdIiLi16EEE.cm::gen::simd" addrspace(4) * -// addrspacecast(% "class._ZTSN2cm3gen4simdIiLi16EEE.cm::gen::simd" * +// % "class.cm::gen::simd", +// % "class.cm::gen::simd" addrspace(4) * +// addrspacecast(% "class.cm::gen::simd" * // bitcast(<16 x i32> * @0 to -// %"class._ZTSN2cm3gen4simdIiLi16EEE.cm::gen::simd" *) to % -// "class._ZTSN2cm3gen4simdIiLi16EEE.cm::gen::simd" addrspace(4) *), +// %"class.cm::gen::simd" *) to % +// "class.cm::gen::simd" addrspace(4) *), // i64 0, i32 0)) //===----------------------------------------------------------------------===// diff --git a/sycl/test/check_device_code/kernel_arguments_as.cpp b/sycl/test/check_device_code/kernel_arguments_as.cpp index de2f10eedd1dc..752ff98f2ad7e 100644 --- a/sycl/test/check_device_code/kernel_arguments_as.cpp +++ b/sycl/test/check_device_code/kernel_arguments_as.cpp @@ -5,12 +5,12 @@ // // Check the address space of the pointer in accessor class. // -// CHECK: %struct{{.*}}AccWrapper = type { %"class{{.*}}cl::sycl::accessor" } -// CHECK: %"class{{.*}}cl::sycl::accessor" = type { %"class{{.*}}AccessorImplDevice", %[[UNION:.*]] } +// CHECK: %struct.AccWrapper = type { %"class.cl::sycl::accessor.1" } +// CHECK: %"class.cl::sycl::accessor.1" = type { %"class{{.*}}AccessorImplDevice", %[[UNION:.*]] } // CHECK-DISABLE: %[[UNION]] = type { i32 addrspace(1)* } // CHECK-ENABLE: %[[UNION]] = type { i32 addrspace(5)* } -// CHECK: %struct{{.*}}AccWrapper = type { %"class{{.*}}cl::sycl::accessor" } -// CHECK-NEXT: %"class{{.*}}cl::sycl::accessor" = type { %"class{{.*}}LocalAccessorBaseDevice", i32 addrspace(3)* } +// CHECK: %struct.AccWrapper.{{[0-9]+}} = type { %"class.cl::sycl::accessor.[[NUM:[0-9]+]]" } +// CHECK-NEXT: %"class.cl::sycl::accessor.[[NUM]]" = type { %"class{{.*}}LocalAccessorBaseDevice", i32 addrspace(3)* } // // Check that kernel arguments doesn't have generic address space. // @@ -37,7 +37,7 @@ int main() { auto acc_wrapped = AccWrapper{acc}; auto local_acc_wrapped = AccWrapper{local_acc}; cgh.parallel_for( - cl::sycl::range<1>(buf.get_count()), [=](cl::sycl::item<1> it) { + cl::sycl::range<1>(buf.size()), [=](cl::sycl::item<1> it) { auto idx = it.get_linear_id(); acc_wrapped.accessor[idx] = local_acc_wrapped.accessor[idx]; }); diff --git a/sycl/test/esimd/intrins_trans.cpp b/sycl/test/esimd/intrins_trans.cpp index 563bc65b8beac..d848a9d7bd15e 100644 --- a/sycl/test/esimd/intrins_trans.cpp +++ b/sycl/test/esimd/intrins_trans.cpp @@ -45,14 +45,14 @@ SYCL_ESIMD_FUNCTION SYCL_EXTERNAL simd foo() { simd_mask pred; v_addr += offsets; - __esimd_flat_atomic0( + __esimd_flat_atomic0( v_addr.data(), pred.data()); // CHECK: %{{[0-9a-zA-Z_.]+}} = call <32 x i32> @llvm.genx.svm.atomic.inc.v32i32.v32i1.v32i64(<32 x i1> %{{[0-9a-zA-Z_.]+}}, <32 x i64> %{{[0-9a-zA-Z_.]+}}, <32 x i32> undef) - __esimd_flat_atomic1( + __esimd_flat_atomic1( v_addr.data(), v1.data(), pred.data()); // CHECK: %{{[0-9a-zA-Z_.]+}} = call <32 x i32> @llvm.genx.svm.atomic.add.v32i32.v32i1.v32i64(<32 x i1> %{{[0-9a-zA-Z_.]+}}, <32 x i64> %{{[0-9a-zA-Z_.]+}}, <32 x i32> %{{[0-9a-zA-Z_.]+}}, <32 x i32> undef) - __esimd_flat_atomic2( + __esimd_flat_atomic2( v_addr.data(), v1.data(), v1.data(), pred.data()); // CHECK: %{{[0-9a-zA-Z_.]+}} = call <32 x i32> @llvm.genx.svm.atomic.cmpxchg.v32i32.v32i1.v32i64(<32 x i1> %{{[0-9a-zA-Z_.]+}}, <32 x i64> %{{[0-9a-zA-Z_.]+}}, <32 x i32> %{{[0-9a-zA-Z_.]+}}, <32 x i32> %{{[0-9a-zA-Z_.]+}}, <32 x i32> undef)