Skip to content

Commit 3de1c4b

Browse files
jsjijsji
committed
[NFC][SYCL] Regen test after sret addressspace change
1 parent d6d75e1 commit 3de1c4b

File tree

9 files changed

+171
-170
lines changed

9 files changed

+171
-170
lines changed

‎sycl/test/check_device_code/extensions/address_cast.cpp

Lines changed: 20 additions & 20 deletions
Original file line numberDiff line numberDiff line change
@@ -13,63 +13,63 @@ using namespace sycl::ext::oneapi::experimental;
1313

1414
namespace static_as_cast {
1515
// CHECK-LABEL: define dso_local spir_func void @_ZN14static_as_cast19to_global_decoratedEN4sycl3_V19multi_ptrIiLNS1_6access13address_spaceE6ELNS3_9decoratedE1EEE(
16-
// CHECK-SAME: ptr addrspace(4) dead_on_unwind noalias writable writeonly sret(%"class.sycl::_V1::multi_ptr") align 8 captures(none) initializes((0, 8)) [[AGG_RESULT:%.*]], ptr noundef readonly byval(%"class.sycl::_V1::multi_ptr.0") align 8 captures(none) [[P:%.*]]) local_unnamed_addr #[[ATTR0:[0-9]+]] !srcloc [[META6:![0-9]+]] !sycl_fixed_targets [[META7:![0-9]+]] {
16+
// CHECK-SAME: ptr dead_on_unwind noalias writable writeonly sret(%"class.sycl::_V1::multi_ptr") align 8 captures(none) [[AGG_RESULT:%.*]], ptr noundef readonly byval(%"class.sycl::_V1::multi_ptr.0") align 8 captures(none) [[P:%.*]]) local_unnamed_addr #[[ATTR0:[0-9]+]] {{.*}}{
1717
// CHECK-NEXT: [[ENTRY:.*:]]
1818
// CHECK-NEXT: [[TMP0:%.*]] = load i64, ptr [[P]], align 8, !tbaa [[TBAA8:![0-9]+]]
1919
// CHECK-NEXT: [[TMP1:%.*]] = inttoptr i64 [[TMP0]] to ptr addrspace(4)
2020
// CHECK-NEXT: [[TMP2:%.*]] = addrspacecast ptr addrspace(4) [[TMP1]] to ptr addrspace(1)
21-
// CHECK-NEXT: store ptr addrspace(1) [[TMP2]], ptr addrspace(4) [[AGG_RESULT]], align 8, !tbaa [[TBAA13:![0-9]+]], !alias.scope [[META15:![0-9]+]]
21+
// CHECK-NEXT: store ptr addrspace(1) [[TMP2]], ptr [[AGG_RESULT]], align 8, !tbaa [[TBAA13:![0-9]+]], !alias.scope [[META15:![0-9]+]]
2222
// CHECK-NEXT: ret void
2323
//
2424
SYCL_EXTERNAL auto to_global_decorated(decorated_generic_ptr<int> p) {
2525
return static_address_cast<access::address_space::global_space>(p);
2626
}
2727
// CHECK-LABEL: define dso_local spir_func void @_ZN14static_as_cast23to_global_not_decoratedEPi(
28-
// CHECK-SAME: ptr addrspace(4) dead_on_unwind noalias writable writeonly sret(%"class.sycl::_V1::multi_ptr.1") align 8 captures(none) initializes((0, 8)) [[AGG_RESULT:%.*]], ptr addrspace(4) noundef [[P:%.*]]) local_unnamed_addr #[[ATTR1:[0-9]+]] !srcloc [[META20:![0-9]+]] !sycl_fixed_targets [[META7]] {
28+
// CHECK-SAME: ptr dead_on_unwind noalias writable writeonly sret(%"class.sycl::_V1::multi_ptr.1") align 8 captures(none) [[AGG_RESULT:%.*]], ptr addrspace(4) noundef [[P:%.*]]) local_unnamed_addr #[[ATTR1:[0-9]+]] {{.*}}{
2929
// CHECK-NEXT: [[ENTRY:.*:]]
3030
// CHECK-NEXT: [[TMP0:%.*]] = addrspacecast ptr addrspace(4) [[P]] to ptr addrspace(1)
31-
// CHECK-NEXT: store ptr addrspace(1) [[TMP0]], ptr addrspace(4) [[AGG_RESULT]], align 8, !tbaa [[TBAA21:![0-9]+]], !alias.scope [[META23:![0-9]+]]
31+
// CHECK-NEXT: store ptr addrspace(1) [[TMP0]], ptr [[AGG_RESULT]], align 8, !tbaa [[TBAA21:![0-9]+]], !alias.scope [[META23:![0-9]+]]
3232
// CHECK-NEXT: ret void
3333
//
3434
SYCL_EXTERNAL auto to_global_not_decorated(int *p) {
3535
return static_address_cast<access::address_space::global_space>(p);
3636
}
3737
// CHECK-LABEL: define dso_local spir_func void @_ZN14static_as_cast20to_generic_decoratedEN4sycl3_V19multi_ptrIiLNS1_6access13address_spaceE6ELNS3_9decoratedE1EEE(
38-
// CHECK-SAME: ptr addrspace(4) dead_on_unwind noalias writable writeonly sret(%"class.sycl::_V1::multi_ptr.0") align 8 captures(none) initializes((0, 8)) [[AGG_RESULT:%.*]], ptr noundef readonly byval(%"class.sycl::_V1::multi_ptr.0") align 8 captures(none) [[P:%.*]]) local_unnamed_addr #[[ATTR0]] !srcloc [[META26:![0-9]+]] !sycl_fixed_targets [[META7]] {
38+
// CHECK-SAME: ptr dead_on_unwind noalias writable writeonly sret(%"class.sycl::_V1::multi_ptr.0") align 8 captures(none) initializes((0, 8)) [[AGG_RESULT:%.*]], ptr noundef readonly byval(%"class.sycl::_V1::multi_ptr.0") align 8 captures(none) [[P:%.*]]) local_unnamed_addr #[[ATTR0]] {{.*}}{
3939
// CHECK-NEXT: [[ENTRY:.*:]]
4040
// CHECK-NEXT: [[TMP0:%.*]] = load i64, ptr [[P]], align 8, !tbaa [[TBAA8]]
41-
// CHECK-NEXT: store i64 [[TMP0]], ptr addrspace(4) [[AGG_RESULT]], align 8, !tbaa [[TBAA8]], !alias.scope [[META27:![0-9]+]]
41+
// CHECK-NEXT: store i64 [[TMP0]], ptr [[AGG_RESULT]], align 8, !tbaa [[TBAA8]], !alias.scope [[META27:![0-9]+]]
4242
// CHECK-NEXT: ret void
4343
//
4444
SYCL_EXTERNAL auto to_generic_decorated(decorated_generic_ptr<int> p) {
4545
return static_address_cast<access::address_space::generic_space>(p);
4646
}
4747
// CHECK-LABEL: define dso_local spir_func void @_ZN14static_as_cast24to_generic_not_decoratedEPi(
48-
// CHECK-SAME: ptr addrspace(4) dead_on_unwind noalias writable writeonly sret(%"class.sycl::_V1::multi_ptr.2") align 8 captures(none) initializes((0, 8)) [[AGG_RESULT:%.*]], ptr addrspace(4) noundef [[P:%.*]]) local_unnamed_addr #[[ATTR1]] !srcloc [[META30:![0-9]+]] !sycl_fixed_targets [[META7]] {
48+
// CHECK-SAME: ptr dead_on_unwind noalias writable writeonly sret(%"class.sycl::_V1::multi_ptr.2") align 8 captures(none) [[AGG_RESULT:%.*]], ptr addrspace(4) noundef [[P:%.*]]) local_unnamed_addr #[[ATTR1]] {{.*}}{
4949
// CHECK-NEXT: [[ENTRY:.*:]]
50-
// CHECK-NEXT: store ptr addrspace(4) [[P]], ptr addrspace(4) [[AGG_RESULT]], align 8, !tbaa [[TBAA31:![0-9]+]], !alias.scope [[META33:![0-9]+]]
50+
// CHECK-NEXT: store ptr addrspace(4) [[P]], ptr [[AGG_RESULT]], align 8, !tbaa [[TBAA31:![0-9]+]], !alias.scope [[META33:![0-9]+]]
5151
// CHECK-NEXT: ret void
5252
//
5353
SYCL_EXTERNAL auto to_generic_not_decorated(int *p) {
5454
return static_address_cast<access::address_space::generic_space>(p);
5555
}
5656

5757
// CHECK-LABEL: define dso_local spir_func void @_ZN14static_as_cast16to_global_deviceEPi(
58-
// CHECK-SAME: ptr addrspace(4) dead_on_unwind noalias writable writeonly sret(%"class.sycl::_V1::multi_ptr.3") align 8 captures(none) initializes((0, 8)) [[AGG_RESULT:%.*]], ptr addrspace(4) noundef [[P:%.*]]) local_unnamed_addr #[[ATTR1]] !srcloc [[META36:![0-9]+]] !sycl_fixed_targets [[META7]] {
58+
// CHECK-SAME: ptr dead_on_unwind noalias writable writeonly sret(%"class.sycl::_V1::multi_ptr.3") align 8 captures(none) [[AGG_RESULT:%.*]], ptr addrspace(4) noundef [[P:%.*]]) local_unnamed_addr #[[ATTR1]] {{.*}}{
5959
// CHECK-NEXT: [[ENTRY:.*:]]
6060
// CHECK-NEXT: [[TMP0:%.*]] = addrspacecast ptr addrspace(4) [[P]] to ptr addrspace(5)
61-
// CHECK-NEXT: store ptr addrspace(5) [[TMP0]], ptr addrspace(4) [[AGG_RESULT]], align 8, !tbaa [[TBAA37:![0-9]+]], !alias.scope [[META39:![0-9]+]]
61+
// CHECK-NEXT: store ptr addrspace(5) [[TMP0]], ptr [[AGG_RESULT]], align 8, !tbaa [[TBAA37:![0-9]+]], !alias.scope [[META39:![0-9]+]]
6262
// CHECK-NEXT: ret void
6363
//
6464
SYCL_EXTERNAL auto to_global_device(int *p) {
6565
return static_address_cast<access::address_space::ext_intel_global_device_space>(p);
6666
}
6767

6868
// CHECK-LABEL: define dso_local spir_func void @_ZN14static_as_cast14to_global_hostEPi(
69-
// CHECK-SAME: ptr addrspace(4) dead_on_unwind noalias writable writeonly sret(%"class.sycl::_V1::multi_ptr.4") align 8 captures(none) initializes((0, 8)) [[AGG_RESULT:%.*]], ptr addrspace(4) noundef [[P:%.*]]) local_unnamed_addr #[[ATTR1]] !srcloc [[META42:![0-9]+]] !sycl_fixed_targets [[META7]] {
69+
// CHECK-SAME: ptr dead_on_unwind noalias writable writeonly sret(%"class.sycl::_V1::multi_ptr.4") align 8 captures(none) [[AGG_RESULT:%.*]], ptr addrspace(4) noundef [[P:%.*]]) local_unnamed_addr #[[ATTR1]] {{.*}}{
7070
// CHECK-NEXT: [[ENTRY:.*:]]
7171
// CHECK-NEXT: [[TMP0:%.*]] = addrspacecast ptr addrspace(4) [[P]] to ptr addrspace(6)
72-
// CHECK-NEXT: store ptr addrspace(6) [[TMP0]], ptr addrspace(4) [[AGG_RESULT]], align 8, !tbaa [[TBAA43:![0-9]+]], !alias.scope [[META45:![0-9]+]]
72+
// CHECK-NEXT: store ptr addrspace(6) [[TMP0]], ptr [[AGG_RESULT]], align 8, !tbaa [[TBAA43:![0-9]+]], !alias.scope [[META45:![0-9]+]]
7373
// CHECK-NEXT: ret void
7474
//
7575
SYCL_EXTERNAL auto to_global_host(int *p) {
@@ -79,41 +79,41 @@ SYCL_EXTERNAL auto to_global_host(int *p) {
7979

8080
namespace dynamic_as_cast {
8181
// CHECK-LABEL: define dso_local spir_func void @_ZN15dynamic_as_cast19to_global_decoratedEN4sycl3_V19multi_ptrIiLNS1_6access13address_spaceE6ELNS3_9decoratedE1EEE(
82-
// CHECK-SAME: ptr addrspace(4) dead_on_unwind noalias writable writeonly sret(%"class.sycl::_V1::multi_ptr") align 8 captures(none) initializes((0, 8)) [[AGG_RESULT:%.*]], ptr noundef readonly byval(%"class.sycl::_V1::multi_ptr.0") align 8 captures(none) [[P:%.*]]) local_unnamed_addr #[[ATTR2:[0-9]+]] !srcloc [[META48:![0-9]+]] !sycl_fixed_targets [[META7]] {
82+
// CHECK-SAME: ptr dead_on_unwind noalias writable writeonly sret(%"class.sycl::_V1::multi_ptr") align 8 captures(none) [[AGG_RESULT:%.*]], ptr noundef readonly byval(%"class.sycl::_V1::multi_ptr.0") align 8 captures(none) [[P:%.*]]) local_unnamed_addr #[[ATTR2:[0-9]+]] {{.*}}{
8383
// CHECK-NEXT: [[ENTRY:.*:]]
8484
// CHECK-NEXT: [[TMP0:%.*]] = load i64, ptr [[P]], align 8, !tbaa [[TBAA8]]
8585
// CHECK-NEXT: [[TMP1:%.*]] = inttoptr i64 [[TMP0]] to ptr addrspace(4)
8686
// CHECK-NEXT: [[CALL_I_I_I:%.*]] = tail call spir_func noundef ptr addrspace(1) @_Z41__spirv_GenericCastToPtrExplicit_ToGlobalPvi(ptr addrspace(4) noundef [[TMP1]], i32 noundef 5) #[[ATTR5:[0-9]+]]
87-
// CHECK-NEXT: store ptr addrspace(1) [[CALL_I_I_I]], ptr addrspace(4) [[AGG_RESULT]], align 8, !tbaa [[TBAA13]], !alias.scope [[META49:![0-9]+]]
87+
// CHECK-NEXT: store ptr addrspace(1) [[CALL_I_I_I]], ptr [[AGG_RESULT]], align 8, !tbaa [[TBAA13]], !alias.scope [[META49:![0-9]+]]
8888
// CHECK-NEXT: ret void
8989
//
9090
SYCL_EXTERNAL auto to_global_decorated(decorated_generic_ptr<int> p) {
9191
return dynamic_address_cast<access::address_space::global_space>(p);
9292
}
9393
// CHECK-LABEL: define dso_local spir_func void @_ZN15dynamic_as_cast23to_global_not_decoratedEPi(
94-
// CHECK-SAME: ptr addrspace(4) dead_on_unwind noalias writable writeonly sret(%"class.sycl::_V1::multi_ptr.1") align 8 captures(none) initializes((0, 8)) [[AGG_RESULT:%.*]], ptr addrspace(4) noundef [[P:%.*]]) local_unnamed_addr #[[ATTR4:[0-9]+]] !srcloc [[META54:![0-9]+]] !sycl_fixed_targets [[META7]] {
94+
// CHECK-SAME: ptr dead_on_unwind noalias writable writeonly sret(%"class.sycl::_V1::multi_ptr.1") align 8 captures(none) [[AGG_RESULT:%.*]], ptr addrspace(4) noundef [[P:%.*]]) local_unnamed_addr #[[ATTR4:[0-9]+]] {{.*}}{
9595
// CHECK-NEXT: [[ENTRY:.*:]]
9696
// CHECK-NEXT: [[CALL_I_I:%.*]] = tail call spir_func noundef ptr addrspace(1) @_Z41__spirv_GenericCastToPtrExplicit_ToGlobalPvi(ptr addrspace(4) noundef [[P]], i32 noundef 5) #[[ATTR5]]
97-
// CHECK-NEXT: store ptr addrspace(1) [[CALL_I_I]], ptr addrspace(4) [[AGG_RESULT]], align 8, !tbaa [[TBAA21]], !alias.scope [[META55:![0-9]+]]
97+
// CHECK-NEXT: store ptr addrspace(1) [[CALL_I_I]], ptr [[AGG_RESULT]], align 8, !tbaa [[TBAA21]], !alias.scope [[META55:![0-9]+]]
9898
// CHECK-NEXT: ret void
9999
//
100100
SYCL_EXTERNAL auto to_global_not_decorated(int *p) {
101101
return dynamic_address_cast<access::address_space::global_space>(p);
102102
}
103103
// CHECK-LABEL: define dso_local spir_func void @_ZN15dynamic_as_cast20to_generic_decoratedEN4sycl3_V19multi_ptrIiLNS1_6access13address_spaceE6ELNS3_9decoratedE1EEE(
104-
// CHECK-SAME: ptr addrspace(4) dead_on_unwind noalias writable writeonly sret(%"class.sycl::_V1::multi_ptr.0") align 8 captures(none) initializes((0, 8)) [[AGG_RESULT:%.*]], ptr noundef readonly byval(%"class.sycl::_V1::multi_ptr.0") align 8 captures(none) [[P:%.*]]) local_unnamed_addr #[[ATTR0]] !srcloc [[META58:![0-9]+]] !sycl_fixed_targets [[META7]] {
104+
// CHECK-SAME: ptr dead_on_unwind noalias writable writeonly sret(%"class.sycl::_V1::multi_ptr.0") align 8 captures(none) initializes((0, 8)) [[AGG_RESULT:%.*]], ptr noundef readonly byval(%"class.sycl::_V1::multi_ptr.0") align 8 captures(none) [[P:%.*]]) local_unnamed_addr #[[ATTR0]] {{.*}}{
105105
// CHECK-NEXT: [[ENTRY:.*:]]
106106
// CHECK-NEXT: [[TMP0:%.*]] = load i64, ptr [[P]], align 8, !tbaa [[TBAA8]]
107-
// CHECK-NEXT: store i64 [[TMP0]], ptr addrspace(4) [[AGG_RESULT]], align 8, !tbaa [[TBAA8]], !alias.scope [[META59:![0-9]+]]
107+
// CHECK-NEXT: store i64 [[TMP0]], ptr [[AGG_RESULT]], align 8, !tbaa [[TBAA8]], !alias.scope [[META59:![0-9]+]]
108108
// CHECK-NEXT: ret void
109109
//
110110
SYCL_EXTERNAL auto to_generic_decorated(decorated_generic_ptr<int> p) {
111111
return dynamic_address_cast<access::address_space::generic_space>(p);
112112
}
113113
// CHECK-LABEL: define dso_local spir_func void @_ZN15dynamic_as_cast24to_generic_not_decoratedEPi(
114-
// CHECK-SAME: ptr addrspace(4) dead_on_unwind noalias writable writeonly sret(%"class.sycl::_V1::multi_ptr.2") align 8 captures(none) initializes((0, 8)) [[AGG_RESULT:%.*]], ptr addrspace(4) noundef [[P:%.*]]) local_unnamed_addr #[[ATTR1]] !srcloc [[META62:![0-9]+]] !sycl_fixed_targets [[META7]] {
114+
// CHECK-SAME: ptr dead_on_unwind noalias writable writeonly sret(%"class.sycl::_V1::multi_ptr.2") align 8 captures(none) [[AGG_RESULT:%.*]], ptr addrspace(4) noundef [[P:%.*]]) local_unnamed_addr #[[ATTR1]] {{.*}}{
115115
// CHECK-NEXT: [[ENTRY:.*:]]
116-
// CHECK-NEXT: store ptr addrspace(4) [[P]], ptr addrspace(4) [[AGG_RESULT]], align 8, !tbaa [[TBAA31]], !alias.scope [[META63:![0-9]+]]
116+
// CHECK-NEXT: store ptr addrspace(4) [[P]], ptr [[AGG_RESULT]], align 8, !tbaa [[TBAA31]], !alias.scope [[META63:![0-9]+]]
117117
// CHECK-NEXT: ret void
118118
//
119119
SYCL_EXTERNAL auto to_generic_not_decorated(int *p) {

‎sycl/test/check_device_code/vector/as.cpp

Lines changed: 3 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -14,7 +14,8 @@
1414

1515
template SYCL_EXTERNAL sycl::vec<int, 4> sycl::vec<float, 4>::as<sycl::vec<int, 4>>() const;
1616
// CHECK-LABEL: define weak_odr dso_local spir_func void @_ZNK4sycl3_V13vecIfLi4EE2asINS1_IiLi4EEEEET_v(
17-
// CHECK-SAME: ptr addrspace(4) dead_on_unwind noalias writable sret(%"class.sycl::_V1::vec") align 16 [[AGG_RESULT:%.*]], ptr addrspace(4) noundef align 16 dereferenceable_or_null(16) [[THIS:%.*]]) local_unnamed_addr #[[ATTR0:[0-9]+]] comdat align 2 !srcloc [[META5:![0-9]+]] !sycl_fixed_targets [[META6:![0-9]+]] {
17+
// CHECK-SAME: ptr dead_on_unwind noalias writable sret(%"class.sycl::_V1::vec") align 16 [[AGG_RESULT:%.*]], ptr addrspace(4) noundef align 16 dereferenceable_or_null(16) [[THIS:%.*]]) local_unnamed_addr #[[ATTR0:[0-9]+]] comdat align 2 !srcloc [[META6:![0-9]+]] !sycl_fixed_targets [[META7:![0-9]+]] {
1818
// CHECK-NEXT: entry:
19-
// CHECK-NEXT: tail call void @llvm.memcpy.p4.p4.i64(ptr addrspace(4) noundef align 16 dereferenceable(16) [[AGG_RESULT]], ptr addrspace(4) noundef align 16 dereferenceable(16) [[THIS]], i64 16, i1 false)
19+
// CHECK-NEXT: tail call void @llvm.memcpy.p0.p4.i64(ptr noundef nonnull align 16 dereferenceable(16) [[AGG_RESULT]], ptr addrspace(4) noundef align 16 dereferenceable(16) [[THIS]], i64 16, i1 false)
2020
// CHECK-NEXT: ret void
21+
//

0 commit comments

Comments
 (0)