|
1 | 1 | // NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --version 5
|
2 |
| -// RUN: %clangxx -O3 -fsycl -fsycl-device-only -fno-discard-value-names -S -emit-llvm -fno-sycl-instrument-device-code -o - %s | FileCheck %s |
| 2 | +// RUN: %clangxx -D__ENABLE_USM_ADDR_SPACE__ -O3 -fsycl -fsycl-device-only -fno-discard-value-names -S -emit-llvm -fno-sycl-instrument-device-code -o - %s | FileCheck %s |
3 | 3 |
|
4 | 4 | // Linux/Windows have minor differences in the generated IR (e.g. TBAA
|
5 | 5 | // metadata). Having linux-only checks eases the maintenance without sacrifising
|
@@ -59,45 +59,67 @@ SYCL_EXTERNAL auto to_generic_decorated(decorated_generic_ptr<int> p) {
|
59 | 59 | SYCL_EXTERNAL auto to_generic_not_decorated(int *p) {
|
60 | 60 | return static_address_cast<access::address_space::generic_space>(p);
|
61 | 61 | }
|
| 62 | + |
| 63 | +// CHECK-LABEL: define dso_local spir_func void @_ZN14static_as_cast16to_global_deviceEPi( |
| 64 | +// CHECK-SAME: ptr addrspace(4) dead_on_unwind noalias nocapture writable writeonly sret(%"class.sycl::_V1::multi_ptr.3") align 8 [[AGG_RESULT:%.*]], ptr addrspace(4) noundef [[P:%.*]]) local_unnamed_addr #[[ATTR4]] !srcloc [[META35:![0-9]+]] !sycl_fixed_targets [[META7]] { |
| 65 | +// CHECK-NEXT: [[ENTRY:.*:]] |
| 66 | +// CHECK-NEXT: [[TMP0:%.*]] = addrspacecast ptr addrspace(4) [[P]] to ptr addrspace(5) |
| 67 | +// CHECK-NEXT: store ptr addrspace(5) [[TMP0]], ptr addrspace(4) [[AGG_RESULT]], align 8, !tbaa [[TBAA36:![0-9]+]], !alias.scope [[META38:![0-9]+]] |
| 68 | +// CHECK-NEXT: ret void |
| 69 | +// |
| 70 | +SYCL_EXTERNAL auto to_global_device(int *p) { |
| 71 | + return static_address_cast<access::address_space::ext_intel_global_device_space>(p); |
| 72 | +} |
| 73 | + |
| 74 | +// CHECK-LABEL: define dso_local spir_func void @_ZN14static_as_cast14to_global_hostEPi( |
| 75 | +// CHECK-SAME: ptr addrspace(4) dead_on_unwind noalias nocapture writable writeonly sret(%"class.sycl::_V1::multi_ptr.4") align 8 [[AGG_RESULT:%.*]], ptr addrspace(4) noundef [[P:%.*]]) local_unnamed_addr #[[ATTR4]] !srcloc [[META41:![0-9]+]] !sycl_fixed_targets [[META7]] { |
| 76 | +// CHECK-NEXT: [[ENTRY:.*:]] |
| 77 | +// CHECK-NEXT: [[TMP0:%.*]] = addrspacecast ptr addrspace(4) [[P]] to ptr addrspace(6) |
| 78 | +// CHECK-NEXT: store ptr addrspace(6) [[TMP0]], ptr addrspace(4) [[AGG_RESULT]], align 8, !tbaa [[TBAA42:![0-9]+]], !alias.scope [[META44:![0-9]+]] |
| 79 | +// CHECK-NEXT: ret void |
| 80 | +// |
| 81 | +SYCL_EXTERNAL auto to_global_host(int *p) { |
| 82 | + return static_address_cast<access::address_space::ext_intel_global_host_space>(p); |
| 83 | +} |
62 | 84 | } // namespace static_as_cast
|
63 | 85 |
|
64 | 86 | namespace dynamic_as_cast {
|
65 | 87 | // CHECK-LABEL: define dso_local spir_func void @_ZN15dynamic_as_cast19to_global_decoratedEN4sycl3_V19multi_ptrIiLNS1_6access13address_spaceE6ELNS3_9decoratedE1EEE(
|
66 |
| -// CHECK-SAME: ptr addrspace(4) dead_on_unwind noalias nocapture writable writeonly sret(%"class.sycl::_V1::multi_ptr") align 8 [[AGG_RESULT:%.*]], ptr nocapture noundef readonly byval(%"class.sycl::_V1::multi_ptr.0") align 8 [[P:%.*]]) local_unnamed_addr #[[ATTR0]] !srcloc [[META35:![0-9]+]] !sycl_fixed_targets [[META7]] { |
| 88 | +// CHECK-SAME: ptr addrspace(4) dead_on_unwind noalias nocapture writable writeonly sret(%"class.sycl::_V1::multi_ptr") align 8 [[AGG_RESULT:%.*]], ptr nocapture noundef readonly byval(%"class.sycl::_V1::multi_ptr.0") align 8 [[P:%.*]]) local_unnamed_addr #[[ATTR0]] !srcloc [[META47:![0-9]+]] !sycl_fixed_targets [[META7]] { |
67 | 89 | // CHECK-NEXT: [[ENTRY:.*:]]
|
68 | 90 | // CHECK-NEXT: [[TMP0:%.*]] = load i64, ptr [[P]], align 8, !tbaa [[TBAA8]]
|
69 | 91 | // CHECK-NEXT: [[TMP1:%.*]] = inttoptr i64 [[TMP0]] to ptr addrspace(4)
|
70 | 92 | // CHECK-NEXT: [[CALL_I_I_I_I:%.*]] = tail call spir_func noundef ptr addrspace(1) @_Z41__spirv_GenericCastToPtrExplicit_ToGlobalPvi(ptr addrspace(4) noundef [[TMP1]], i32 noundef 5) #[[ATTR5]]
|
71 |
| -// CHECK-NEXT: store ptr addrspace(1) [[CALL_I_I_I_I]], ptr addrspace(4) [[AGG_RESULT]], align 8, !tbaa [[TBAA12]], !alias.scope [[META36:![0-9]+]] |
| 93 | +// CHECK-NEXT: store ptr addrspace(1) [[CALL_I_I_I_I]], ptr addrspace(4) [[AGG_RESULT]], align 8, !tbaa [[TBAA12]], !alias.scope [[META48:![0-9]+]] |
72 | 94 | // CHECK-NEXT: ret void
|
73 | 95 | //
|
74 | 96 | SYCL_EXTERNAL auto to_global_decorated(decorated_generic_ptr<int> p) {
|
75 | 97 | return dynamic_address_cast<access::address_space::global_space>(p);
|
76 | 98 | }
|
77 | 99 | // CHECK-LABEL: define dso_local spir_func void @_ZN15dynamic_as_cast23to_global_not_decoratedEPi(
|
78 |
| -// CHECK-SAME: ptr addrspace(4) dead_on_unwind noalias nocapture writable writeonly sret(%"class.sycl::_V1::multi_ptr.1") align 8 [[AGG_RESULT:%.*]], ptr addrspace(4) noundef [[P:%.*]]) local_unnamed_addr #[[ATTR2]] !srcloc [[META41:![0-9]+]] !sycl_fixed_targets [[META7]] { |
| 100 | +// CHECK-SAME: ptr addrspace(4) dead_on_unwind noalias nocapture writable writeonly sret(%"class.sycl::_V1::multi_ptr.1") align 8 [[AGG_RESULT:%.*]], ptr addrspace(4) noundef [[P:%.*]]) local_unnamed_addr #[[ATTR2]] !srcloc [[META53:![0-9]+]] !sycl_fixed_targets [[META7]] { |
79 | 101 | // CHECK-NEXT: [[ENTRY:.*:]]
|
80 | 102 | // CHECK-NEXT: [[CALL_I_I_I:%.*]] = tail call spir_func noundef ptr addrspace(1) @_Z41__spirv_GenericCastToPtrExplicit_ToGlobalPvi(ptr addrspace(4) noundef [[P]], i32 noundef 5) #[[ATTR5]]
|
81 |
| -// CHECK-NEXT: store ptr addrspace(1) [[CALL_I_I_I]], ptr addrspace(4) [[AGG_RESULT]], align 8, !tbaa [[TBAA20]], !alias.scope [[META42:![0-9]+]] |
| 103 | +// CHECK-NEXT: store ptr addrspace(1) [[CALL_I_I_I]], ptr addrspace(4) [[AGG_RESULT]], align 8, !tbaa [[TBAA20]], !alias.scope [[META54:![0-9]+]] |
82 | 104 | // CHECK-NEXT: ret void
|
83 | 105 | //
|
84 | 106 | SYCL_EXTERNAL auto to_global_not_decorated(int *p) {
|
85 | 107 | return dynamic_address_cast<access::address_space::global_space>(p);
|
86 | 108 | }
|
87 | 109 | // CHECK-LABEL: define dso_local spir_func void @_ZN15dynamic_as_cast20to_generic_decoratedEN4sycl3_V19multi_ptrIiLNS1_6access13address_spaceE6ELNS3_9decoratedE1EEE(
|
88 |
| -// CHECK-SAME: ptr addrspace(4) dead_on_unwind noalias nocapture writable writeonly sret(%"class.sycl::_V1::multi_ptr.0") align 8 [[AGG_RESULT:%.*]], ptr nocapture noundef readonly byval(%"class.sycl::_V1::multi_ptr.0") align 8 [[P:%.*]]) local_unnamed_addr #[[ATTR3]] !srcloc [[META45:![0-9]+]] !sycl_fixed_targets [[META7]] { |
| 110 | +// CHECK-SAME: ptr addrspace(4) dead_on_unwind noalias nocapture writable writeonly sret(%"class.sycl::_V1::multi_ptr.0") align 8 [[AGG_RESULT:%.*]], ptr nocapture noundef readonly byval(%"class.sycl::_V1::multi_ptr.0") align 8 [[P:%.*]]) local_unnamed_addr #[[ATTR3]] !srcloc [[META57:![0-9]+]] !sycl_fixed_targets [[META7]] { |
89 | 111 | // CHECK-NEXT: [[ENTRY:.*:]]
|
90 | 112 | // CHECK-NEXT: [[TMP0:%.*]] = load i64, ptr [[P]], align 8, !tbaa [[TBAA8]]
|
91 |
| -// CHECK-NEXT: store i64 [[TMP0]], ptr addrspace(4) [[AGG_RESULT]], align 8, !tbaa [[TBAA8]], !alias.scope [[META46:![0-9]+]] |
| 113 | +// CHECK-NEXT: store i64 [[TMP0]], ptr addrspace(4) [[AGG_RESULT]], align 8, !tbaa [[TBAA8]], !alias.scope [[META58:![0-9]+]] |
92 | 114 | // CHECK-NEXT: ret void
|
93 | 115 | //
|
94 | 116 | SYCL_EXTERNAL auto to_generic_decorated(decorated_generic_ptr<int> p) {
|
95 | 117 | return dynamic_address_cast<access::address_space::generic_space>(p);
|
96 | 118 | }
|
97 | 119 | // CHECK-LABEL: define dso_local spir_func void @_ZN15dynamic_as_cast24to_generic_not_decoratedEPi(
|
98 |
| -// CHECK-SAME: ptr addrspace(4) dead_on_unwind noalias nocapture writable writeonly sret(%"class.sycl::_V1::multi_ptr.2") align 8 [[AGG_RESULT:%.*]], ptr addrspace(4) noundef [[P:%.*]]) local_unnamed_addr #[[ATTR4]] !srcloc [[META49:![0-9]+]] !sycl_fixed_targets [[META7]] { |
| 120 | +// CHECK-SAME: ptr addrspace(4) dead_on_unwind noalias nocapture writable writeonly sret(%"class.sycl::_V1::multi_ptr.2") align 8 [[AGG_RESULT:%.*]], ptr addrspace(4) noundef [[P:%.*]]) local_unnamed_addr #[[ATTR4]] !srcloc [[META61:![0-9]+]] !sycl_fixed_targets [[META7]] { |
99 | 121 | // CHECK-NEXT: [[ENTRY:.*:]]
|
100 |
| -// CHECK-NEXT: store ptr addrspace(4) [[P]], ptr addrspace(4) [[AGG_RESULT]], align 8, !tbaa [[TBAA30]], !alias.scope [[META50:![0-9]+]] |
| 122 | +// CHECK-NEXT: store ptr addrspace(4) [[P]], ptr addrspace(4) [[AGG_RESULT]], align 8, !tbaa [[TBAA30]], !alias.scope [[META62:![0-9]+]] |
101 | 123 | // CHECK-NEXT: ret void
|
102 | 124 | //
|
103 | 125 | SYCL_EXTERNAL auto to_generic_not_decorated(int *p) {
|
|
0 commit comments