Skip to content

[NFC] Remove LLVM type name altering #3915

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
merged 10 commits into from
Oct 9, 2021
Merged

[NFC] Remove LLVM type name altering #3915

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
10 commits
Select commit Hold shift + click to select a range
884d951
[NFC] Remove LLVM type name altering
bader Jun 10, 2021
6ec44c9
Fix code formatting.
bader Jun 10, 2021
117ddea
Improve clang/test/CodeGenSYCL/debug-info-srcpos-kernel.cpp test check.
bader Jun 22, 2021
25c4e4a
Improve clang/test/CodeGenSYCL/intel-fpga-local.cpp test check.
bader Jun 22, 2021
4275bf9
Merge remote-tracking branch 'intel/sycl' into llvm-type-altering
bader Jun 22, 2021
73f2434
Fix clang/test/CodeGenSYCL/stall_enable_device.cpp test check.
bader Jun 22, 2021
7d74a47
Merge remote-tracking branch 'intel/sycl' into llvm-type-altering
bader Oct 8, 2021
29a97c9
Fix clang and sycl tests after merge.
bader Oct 8, 2021
6f99a1d
Apply clang-format.
bader Oct 8, 2021
183fab2
Fix kernel_arugments_as test on Windows.
bader Oct 8, 2021
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
14 changes: 0 additions & 14 deletions clang/lib/CodeGen/CodeGenTypes.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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<MangleContext> 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.
Expand Down
2 changes: 1 addition & 1 deletion clang/test/CodeGenSYCL/address-space-cond-op.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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:%.*]]
Expand Down
6 changes: 3 additions & 3 deletions clang/test/CodeGenSYCL/address-space-new.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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 <typename name, typename Func>
Expand Down
34 changes: 17 additions & 17 deletions clang/test/CodeGenSYCL/basic-kernel-wrapper.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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
Expand All @@ -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)* {{[^,]*}})
6 changes: 3 additions & 3 deletions clang/test/CodeGenSYCL/check-direct-attribute-propagation.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -304,20 +304,20 @@ 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()
Functor10 f10;
h.single_task<class kernel_name32>(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<class kernel_name33>(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<class kernel_name34>(
[]() [[intel::kernel_args_restrict]]{});
});
Expand Down
2 changes: 1 addition & 1 deletion clang/test/CodeGenSYCL/debug-info-srcpos-kernel.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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]+]]
Expand Down
2 changes: 1 addition & 1 deletion clang/test/CodeGenSYCL/device-functions.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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)
4 changes: 2 additions & 2 deletions clang/test/CodeGenSYCL/device-variables.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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
});
Expand Down
20 changes: 10 additions & 10 deletions clang/test/CodeGenSYCL/inheritance.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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]]
6 changes: 3 additions & 3 deletions clang/test/CodeGenSYCL/intel-fpga-ivdep-array.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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]+]]
Expand Down
Loading