Skip to content

Commit 08d3b4c

Browse files
jsjijsji
committed
[SYCL][NFC] Update test after nocatpure change (#17080)
1 parent def5289 commit 08d3b4c

File tree

6 files changed

+77
-77
lines changed

6 files changed

+77
-77
lines changed

‎sycl-jit/test/internalization/promote-private-non-unit-cuda.ll‎

Lines changed: 11 additions & 11 deletions
Original file line numberDiff line numberDiff line change
@@ -18,18 +18,18 @@ declare noundef i32 @llvm.nvvm.read.ptx.sreg.ntid.x() #0
1818
declare noundef i32 @llvm.nvvm.read.ptx.sreg.tid.x() #0
1919
declare ptr @llvm.nvvm.implicit.offset() #1
2020

21-
define void @fused_0(ptr addrspace(1) nocapture noundef align 16 %KernelOne__arg_accTmp,
22-
ptr nocapture noundef readonly byval(%"class.sycl::_V1::range") align 8 %KernelOne__arg_accTmp3,
23-
ptr addrspace(1) nocapture noundef readonly align 4 %KernelOne__arg_accIn,
24-
ptr nocapture noundef readonly byval(%"class.sycl::_V1::range") align 8 %KernelOne__arg_accIn6,
25-
ptr addrspace(1) nocapture noundef align 1 %KernelOne__arg_accTmp27,
26-
ptr nocapture noundef readonly byval(%"class.sycl::_V1::range") align 8 %KernelOne__arg_accTmp210,
27-
ptr addrspace(1) nocapture noundef writeonly align 4 %KernelTwo__arg_accOut,
28-
ptr nocapture noundef readonly byval(%"class.sycl::_V1::range") align 8 %KernelTwo__arg_accOut3)
21+
define void @fused_0(ptr addrspace(1) noundef align 16 %KernelOne__arg_accTmp,
22+
ptr noundef readonly byval(%"class.sycl::_V1::range") align 8 captures(none)%KernelOne__arg_accTmp3,
23+
ptr addrspace(1) noundef readonly align 4 %KernelOne__arg_accIn,
24+
ptr noundef readonly byval(%"class.sycl::_V1::range") align 8 captures(none)%KernelOne__arg_accIn6,
25+
ptr addrspace(1) noundef align 1 %KernelOne__arg_accTmp27,
26+
ptr noundef readonly byval(%"class.sycl::_V1::range") align 8 captures(none)%KernelOne__arg_accTmp210,
27+
ptr addrspace(1) noundef writeonly align 4 %KernelTwo__arg_accOut,
28+
ptr noundef readonly byval(%"class.sycl::_V1::range") align 8 captures(none)%KernelTwo__arg_accOut3)
2929
local_unnamed_addr #3 !sycl.kernel.promote !13 !sycl.kernel.promote.localsize !14 !sycl.kernel.promote.elemsize !15 {
30-
; CHECK-LABEL: define void @fused_0(
31-
; CHECK-SAME: ptr nocapture noundef readonly byval(%"class.sycl::_V1::range") align 8 [[KERNELONE__ARG_ACCTMP3:%[^,]*accTmp3]],
32-
; CHECK-SAME: ptr nocapture noundef readonly byval(%"class.sycl::_V1::range") align 8 [[KERNELONE__ARG_ACCTMP210:%[^,]*accTmp210]]
30+
; CHECK-LABEL: define ptx_kernel void @fused_0(
31+
; CHECK-SAME: ptr noundef readonly byval(%"class.sycl::_V1::range") align 8 captures(none) [[KERNELONE__ARG_ACCTMP3:%[^,]*accTmp3]],
32+
; CHECK-SAME: ptr noundef readonly byval(%"class.sycl::_V1::range") align 8 captures(none) [[KERNELONE__ARG_ACCTMP210:%[^,]*accTmp210]]
3333
; CHECK: entry:
3434
; CHECK: [[TMP0:%.*]] = alloca i8, i64 3, align 1
3535
; CHECK: [[TMP1:%.*]] = alloca i8, i64 96, align 16

‎sycl-jit/test/internalization/promote-private-non-unit.ll‎

Lines changed: 7 additions & 7 deletions
Original file line numberDiff line numberDiff line change
@@ -15,17 +15,17 @@ target triple = "spir64-unknown-unknown"
1515
declare spir_func i64 @_Z33__spirv_BuiltInGlobalInvocationIdi(i32) local_unnamed_addr #1
1616

1717
define spir_kernel void @fused_0(ptr addrspace(1) nocapture align 16 %KernelOne__arg_accTmp,
18-
ptr nocapture readonly byval(%"class.sycl::_V1::range") align 8 %KernelOne__arg_accTmp3,
19-
ptr addrspace(1) nocapture readonly align 4 %KernelOne__arg_accIn,
20-
ptr nocapture readonly byval(%"class.sycl::_V1::range") align 8 %KernelOne__arg_accIn6,
18+
ptr readonly byval(%"class.sycl::_V1::range") align 8 captures(none) %KernelOne__arg_accTmp3,
19+
ptr addrspace(1) readonly align 4 %KernelOne__arg_accIn,
20+
ptr readonly byval(%"class.sycl::_V1::range") align 8 captures(none) %KernelOne__arg_accIn6,
2121
ptr addrspace(1) nocapture writeonly align 1 %KernelOne__arg_accTmp27,
22-
ptr nocapture readonly byval(%"class.sycl::_V1::range") align 8 %KernelOne__arg_accTmp210,
22+
ptr readonly byval(%"class.sycl::_V1::range") align 8 captures(none) %KernelOne__arg_accTmp210,
2323
ptr addrspace(1) nocapture writeonly align 4 %KernelTwo__arg_accOut,
24-
ptr nocapture readonly byval(%"class.sycl::_V1::range") align 8 %KernelTwo__arg_accOut3)
24+
ptr readonly byval(%"class.sycl::_V1::range") align 8 captures(none) %KernelTwo__arg_accOut3)
2525
local_unnamed_addr #2 !sycl.kernel.promote !11 !sycl.kernel.promote.localsize !12 !sycl.kernel.promote.elemsize !13 {
2626
; CHECK-LABEL: define spir_kernel void @fused_0(
27-
; CHECK-SAME: ptr nocapture readonly byval(%"class.sycl::_V1::range") align 8 [[KERNELONE__ARG_ACCTMP3:%[^,]*accTmp3]],
28-
; CHECK-SAME: ptr nocapture readonly byval(%"class.sycl::_V1::range") align 8 [[KERNELONE__ARG_ACCTMP210:%[^,]*accTmp210]]
27+
; CHECK-SAME: ptr readonly byval(%"class.sycl::_V1::range") align 8 captures(none) [[KERNELONE__ARG_ACCTMP3:%[^,]*accTmp3]],
28+
; CHECK-SAME: ptr readonly byval(%"class.sycl::_V1::range") align 8 captures(none) [[KERNELONE__ARG_ACCTMP210:%[^,]*accTmp210]]
2929
; CHECK: entry:
3030
; CHECK: [[TMP0:%.*]] = alloca i8, i64 3, align 1
3131
; CHECK: [[TMP1:%.*]] = alloca i8, i64 96, align 16

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

Lines changed: 23 additions & 23 deletions
Original file line numberDiff line numberDiff line change
@@ -1,4 +1,4 @@
1-
// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --version 5
1+
// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --check-globals none --version 5
22
// 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
33

44
// Linux/Windows have minor differences in the generated IR (e.g. TBAA
@@ -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 nocapture writable writeonly sret(%"class.sycl::_V1::multi_ptr") align 8 initializes((0, 8)) [[AGG_RESULT:%.*]], ptr nocapture noundef readonly byval(%"class.sycl::_V1::multi_ptr.0") align 8 [[P:%.*]])
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]+]] {
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 [[TBAA12:![0-9]+]], !alias.scope [[META14:![0-9]+]]
21+
// CHECK-NEXT: store ptr addrspace(1) [[TMP2]], ptr addrspace(4) [[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 nocapture writable writeonly sret(%"class.sycl::_V1::multi_ptr.1") align 8 initializes((0, 8)) [[AGG_RESULT:%.*]], ptr addrspace(4) noundef [[P:%.*]])
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]] {
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 [[TBAA20:![0-9]+]], !alias.scope [[META22:![0-9]+]]
31+
// CHECK-NEXT: store ptr addrspace(1) [[TMP0]], ptr addrspace(4) [[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 nocapture writable writeonly sret(%"class.sycl::_V1::multi_ptr.0") align 8 initializes((0, 8)) [[AGG_RESULT:%.*]], ptr nocapture noundef readonly byval(%"class.sycl::_V1::multi_ptr.0")
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]] {
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 [[META26:![0-9]+]]
41+
// CHECK-NEXT: store i64 [[TMP0]], ptr addrspace(4) [[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 nocapture writable writeonly sret(%"class.sycl::_V1::multi_ptr.2") align 8 initializes((0, 8)) [[AGG_RESULT:%.*]], ptr addrspace(4) noundef [[P:%.*]])
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]] {
4949
// CHECK-NEXT: [[ENTRY:.*:]]
50-
// CHECK-NEXT: store ptr addrspace(4) [[P]], ptr addrspace(4) [[AGG_RESULT]], align 8, !tbaa [[TBAA30:![0-9]+]], !alias.scope [[META32:![0-9]+]]
50+
// CHECK-NEXT: store ptr addrspace(4) [[P]], ptr addrspace(4) [[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 nocapture writable writeonly sret(%"class.sycl::_V1::multi_ptr.3") align 8 initializes((0, 8)) [[AGG_RESULT:%.*]], ptr addrspace(4) noundef [[P:%.*]])
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]] {
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 [[TBAA36:![0-9]+]], !alias.scope [[META38:![0-9]+]]
61+
// CHECK-NEXT: store ptr addrspace(5) [[TMP0]], ptr addrspace(4) [[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 nocapture writable writeonly sret(%"class.sycl::_V1::multi_ptr.4") align 8 initializes((0, 8)) [[AGG_RESULT:%.*]], ptr addrspace(4) noundef [[P:%.*]])
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]] {
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 [[TBAA42:![0-9]+]], !alias.scope [[META44:![0-9]+]]
72+
// CHECK-NEXT: store ptr addrspace(6) [[TMP0]], ptr addrspace(4) [[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 nocapture writable writeonly sret(%"class.sycl::_V1::multi_ptr") align 8 initializes((0, 8)) [[AGG_RESULT:%.*]], ptr nocapture noundef readonly byval(%"class.sycl::_V1::multi_ptr.0") align 8 [[P:%.*]])
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]] {
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)
86-
// 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)
87-
// 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]+]]
86+
// 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]+]]
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 nocapture writable writeonly sret(%"class.sycl::_V1::multi_ptr.1") align 8 initializes((0, 8)) [[AGG_RESULT:%.*]], ptr addrspace(4) noundef [[P:%.*]])
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]] {
9595
// CHECK-NEXT: [[ENTRY:.*:]]
96-
// 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)
97-
// CHECK-NEXT: store ptr addrspace(1) [[CALL_I_I_I]], ptr addrspace(4) [[AGG_RESULT]], align 8, !tbaa [[TBAA20]], !alias.scope [[META54:![0-9]+]]
96+
// 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]+]]
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 nocapture writable writeonly sret(%"class.sycl::_V1::multi_ptr.0") align 8 initializes((0, 8)) [[AGG_RESULT:%.*]], ptr nocapture noundef readonly byval(%"class.sycl::_V1::multi_ptr.0") align 8 [[P:%.*]])
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]] {
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 [[META58:![0-9]+]]
107+
// CHECK-NEXT: store i64 [[TMP0]], ptr addrspace(4) [[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 nocapture writable writeonly sret(%"class.sycl::_V1::multi_ptr.2") align 8 initializes((0, 8)) [[AGG_RESULT:%.*]], ptr addrspace(4) noundef [[P:%.*]])
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]] {
115115
// CHECK-NEXT: [[ENTRY:.*:]]
116-
// CHECK-NEXT: store ptr addrspace(4) [[P]], ptr addrspace(4) [[AGG_RESULT]], align 8, !tbaa [[TBAA30]], !alias.scope [[META62:![0-9]+]]
116+
// CHECK-NEXT: store ptr addrspace(4) [[P]], ptr addrspace(4) [[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) {

0 commit comments

Comments
 (0)