|
58 | 58 | // OPT-NEXT: ret void |
59 | 59 | // |
60 | 60 | // OPT-SPIRV-LABEL: define spir_kernel void @_Z7kernel1Pi( |
61 | | -// OPT-SPIRV-SAME: ptr addrspace(1) noundef [[X_COERCE:%.*]]) local_unnamed_addr addrspace(4) #[[ATTR0:[0-9]+]] !max_work_group_size [[META5:![0-9]+]] { |
| 61 | +// OPT-SPIRV-SAME: ptr addrspace(1) noundef captures(none) [[X_COERCE:%.*]]) local_unnamed_addr addrspace(4) #[[ATTR0:[0-9]+]] !max_work_group_size [[META5:![0-9]+]] { |
62 | 62 | // OPT-SPIRV-NEXT: [[ENTRY:.*:]] |
63 | 63 | // OPT-SPIRV-NEXT: [[TMP0:%.*]] = load i32, ptr addrspace(1) [[X_COERCE]], align 4 |
64 | 64 | // OPT-SPIRV-NEXT: [[INC:%.*]] = add nsw i32 [[TMP0]], 1 |
@@ -124,7 +124,7 @@ __global__ void kernel1(int *x) { |
124 | 124 | // OPT-NEXT: ret void |
125 | 125 | // |
126 | 126 | // OPT-SPIRV-LABEL: define spir_kernel void @_Z7kernel2Ri( |
127 | | -// OPT-SPIRV-SAME: ptr addrspace(1) noundef align 4 dereferenceable(4) [[X_COERCE:%.*]]) local_unnamed_addr addrspace(4) #[[ATTR0]] !max_work_group_size [[META5]] { |
| 127 | +// OPT-SPIRV-SAME: ptr addrspace(1) noundef align 4 captures(none) dereferenceable(4) [[X_COERCE:%.*]]) local_unnamed_addr addrspace(4) #[[ATTR0]] !max_work_group_size [[META5]] { |
128 | 128 | // OPT-SPIRV-NEXT: [[ENTRY:.*:]] |
129 | 129 | // OPT-SPIRV-NEXT: [[TMP0:%.*]] = load i32, ptr addrspace(1) [[X_COERCE]], align 4 |
130 | 130 | // OPT-SPIRV-NEXT: [[INC:%.*]] = add nsw i32 [[TMP0]], 1 |
@@ -191,7 +191,7 @@ __global__ void kernel2(int &x) { |
191 | 191 | // OPT-NEXT: ret void |
192 | 192 | // |
193 | 193 | // OPT-SPIRV-LABEL: define spir_kernel void @_Z7kernel3PU3AS2iPU3AS1i( |
194 | | -// OPT-SPIRV-SAME: ptr addrspace(2) noundef readonly captures(none) [[X:%.*]], ptr addrspace(1) noundef writeonly captures(none) initializes((0, 4)) [[Y:%.*]]) local_unnamed_addr addrspace(4) #[[ATTR1:[0-9]+]] !max_work_group_size [[META5]] { |
| 194 | +// OPT-SPIRV-SAME: ptr addrspace(2) noundef readonly captures(none) [[X:%.*]], ptr addrspace(1) noundef writeonly captures(none) initializes((0, 4)) [[Y:%.*]]) local_unnamed_addr addrspace(4) #[[ATTR0]] !max_work_group_size [[META5]] { |
195 | 195 | // OPT-SPIRV-NEXT: [[ENTRY:.*:]] |
196 | 196 | // OPT-SPIRV-NEXT: [[TMP0:%.*]] = load i32, ptr addrspace(2) [[X]], align 4 |
197 | 197 | // OPT-SPIRV-NEXT: store i32 [[TMP0]], ptr addrspace(1) [[Y]], align 4 |
@@ -257,7 +257,7 @@ __global__ void kernel3(__attribute__((address_space(2))) int *x, |
257 | 257 | // OPT-NEXT: ret void |
258 | 258 | // |
259 | 259 | // OPT-SPIRV-LABEL: define spir_func void @_Z4funcPi( |
260 | | -// OPT-SPIRV-SAME: ptr addrspace(4) noundef captures(none) [[X:%.*]]) local_unnamed_addr addrspace(4) #[[ATTR2:[0-9]+]] { |
| 260 | +// OPT-SPIRV-SAME: ptr addrspace(4) noundef captures(none) [[X:%.*]]) local_unnamed_addr addrspace(4) #[[ATTR1:[0-9]+]] { |
261 | 261 | // OPT-SPIRV-NEXT: [[ENTRY:.*:]] |
262 | 262 | // OPT-SPIRV-NEXT: [[TMP0:%.*]] = load i32, ptr addrspace(4) [[X]], align 4 |
263 | 263 | // OPT-SPIRV-NEXT: [[INC:%.*]] = add nsw i32 [[TMP0]], 1 |
@@ -339,7 +339,7 @@ struct S { |
339 | 339 | // OPT-NEXT: ret void |
340 | 340 | // |
341 | 341 | // OPT-SPIRV-LABEL: define spir_kernel void @_Z7kernel41S( |
342 | | -// OPT-SPIRV-SAME: [[STRUCT_S:%.*]] [[S_COERCE:%.*]]) local_unnamed_addr addrspace(4) #[[ATTR0]] !max_work_group_size [[META5]] { |
| 342 | +// OPT-SPIRV-SAME: [[STRUCT_S:%.*]] [[S_COERCE:%.*]]) local_unnamed_addr addrspace(4) #[[ATTR2:[0-9]+]] !max_work_group_size [[META5]] { |
343 | 343 | // OPT-SPIRV-NEXT: [[ENTRY:.*:]] |
344 | 344 | // OPT-SPIRV-NEXT: [[TMP0:%.*]] = extractvalue [[STRUCT_S]] [[S_COERCE]], 0 |
345 | 345 | // OPT-SPIRV-NEXT: [[TMP1:%.*]] = extractvalue [[STRUCT_S]] [[S_COERCE]], 1 |
@@ -442,13 +442,13 @@ __global__ void kernel4(struct S s) { |
442 | 442 | // OPT-NEXT: ret void |
443 | 443 | // |
444 | 444 | // OPT-SPIRV-LABEL: define spir_kernel void @_Z7kernel5P1S( |
445 | | -// OPT-SPIRV-SAME: ptr addrspace(1) noundef [[S_COERCE:%.*]]) local_unnamed_addr addrspace(4) #[[ATTR0]] !max_work_group_size [[META5]] { |
| 445 | +// OPT-SPIRV-SAME: ptr addrspace(1) noundef readonly captures(none) [[S_COERCE:%.*]]) local_unnamed_addr addrspace(4) #[[ATTR2]] !max_work_group_size [[META5]] { |
446 | 446 | // OPT-SPIRV-NEXT: [[ENTRY:.*:]] |
447 | 447 | // OPT-SPIRV-NEXT: [[TMP0:%.*]] = load ptr addrspace(4), ptr addrspace(1) [[S_COERCE]], align 8 |
448 | 448 | // OPT-SPIRV-NEXT: [[TMP1:%.*]] = load i32, ptr addrspace(4) [[TMP0]], align 4 |
449 | 449 | // OPT-SPIRV-NEXT: [[INC:%.*]] = add nsw i32 [[TMP1]], 1 |
450 | 450 | // OPT-SPIRV-NEXT: store i32 [[INC]], ptr addrspace(4) [[TMP0]], align 4 |
451 | | -// OPT-SPIRV-NEXT: [[Y:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[S_COERCE]], i64 8 |
| 451 | +// OPT-SPIRV-NEXT: [[Y:%.*]] = getelementptr inbounds nuw i8, ptr addrspace(1) [[S_COERCE]], i64 8 |
452 | 452 | // OPT-SPIRV-NEXT: [[TMP2:%.*]] = load ptr addrspace(4), ptr addrspace(1) [[Y]], align 8 |
453 | 453 | // OPT-SPIRV-NEXT: [[TMP3:%.*]] = load float, ptr addrspace(4) [[TMP2]], align 4 |
454 | 454 | // OPT-SPIRV-NEXT: [[ADD:%.*]] = fadd contract float [[TMP3]], 1.000000e+00 |
@@ -545,7 +545,7 @@ struct T { |
545 | 545 | // OPT-NEXT: ret void |
546 | 546 | // |
547 | 547 | // OPT-SPIRV-LABEL: define spir_kernel void @_Z7kernel61T( |
548 | | -// OPT-SPIRV-SAME: [[STRUCT_T:%.*]] [[T_COERCE:%.*]]) local_unnamed_addr addrspace(4) #[[ATTR0]] !max_work_group_size [[META5]] { |
| 548 | +// OPT-SPIRV-SAME: [[STRUCT_T:%.*]] [[T_COERCE:%.*]]) local_unnamed_addr addrspace(4) #[[ATTR2]] !max_work_group_size [[META5]] { |
549 | 549 | // OPT-SPIRV-NEXT: [[ENTRY:.*:]] |
550 | 550 | // OPT-SPIRV-NEXT: [[TMP0:%.*]] = extractvalue [[STRUCT_T]] [[T_COERCE]], 0 |
551 | 551 | // OPT-SPIRV-NEXT: [[DOTFCA_0_EXTRACT:%.*]] = extractvalue [2 x ptr addrspace(4)] [[TMP0]], 0 |
@@ -625,7 +625,7 @@ __global__ void kernel6(struct T t) { |
625 | 625 | // OPT-NEXT: ret void |
626 | 626 | // |
627 | 627 | // OPT-SPIRV-LABEL: define spir_kernel void @_Z7kernel7Pi( |
628 | | -// OPT-SPIRV-SAME: ptr addrspace(1) noalias noundef [[X_COERCE:%.*]]) local_unnamed_addr addrspace(4) #[[ATTR0]] !max_work_group_size [[META5]] { |
| 628 | +// OPT-SPIRV-SAME: ptr addrspace(1) noalias noundef captures(none) [[X_COERCE:%.*]]) local_unnamed_addr addrspace(4) #[[ATTR0]] !max_work_group_size [[META5]] { |
629 | 629 | // OPT-SPIRV-NEXT: [[ENTRY:.*:]] |
630 | 630 | // OPT-SPIRV-NEXT: [[TMP0:%.*]] = load i32, ptr addrspace(1) [[X_COERCE]], align 4 |
631 | 631 | // OPT-SPIRV-NEXT: [[INC:%.*]] = add nsw i32 [[TMP0]], 1 |
@@ -692,7 +692,7 @@ struct SS { |
692 | 692 | // OPT-NEXT: ret void |
693 | 693 | // |
694 | 694 | // OPT-SPIRV-LABEL: define spir_kernel void @_Z7kernel82SS( |
695 | | -// OPT-SPIRV-SAME: [[STRUCT_SS:%.*]] [[A_COERCE:%.*]]) local_unnamed_addr addrspace(4) #[[ATTR0]] !max_work_group_size [[META5]] { |
| 695 | +// OPT-SPIRV-SAME: [[STRUCT_SS:%.*]] [[A_COERCE:%.*]]) local_unnamed_addr addrspace(4) #[[ATTR2]] !max_work_group_size [[META5]] { |
696 | 696 | // OPT-SPIRV-NEXT: [[ENTRY:.*:]] |
697 | 697 | // OPT-SPIRV-NEXT: [[TMP0:%.*]] = extractvalue [[STRUCT_SS]] [[A_COERCE]], 0 |
698 | 698 | // OPT-SPIRV-NEXT: [[TMP1:%.*]] = load float, ptr addrspace(4) [[TMP0]], align 4 |
|
0 commit comments