Skip to content

Commit 71f39f9

Browse files
enferexenferex
authored
Revert "[SelectionDAG] Preserve volatile undef stores. (#99918)"
This reverts commit 404071b.
1 parent 2bb18e2 commit 71f39f9

File tree

7 files changed

+31
-99
lines changed

7 files changed

+31
-99
lines changed

‎llvm/lib/CodeGen/SelectionDAG/DAGCombiner.cpp

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -21413,7 +21413,7 @@ SDValue DAGCombiner::visitSTORE(SDNode *N) {
2141321413
}
2141421414

2141521415
// Turn 'store undef, Ptr' -> nothing.
21416-
if (Value.isUndef() && ST->isUnindexed() && !ST->isVolatile())
21416+
if (Value.isUndef() && ST->isUnindexed())
2141721417
return Chain;
2141821418

2141921419
// Try to infer better alignment information than the store already has.

‎llvm/test/CodeGen/AMDGPU/abi-attribute-hints-undefined-behavior.ll

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -255,6 +255,7 @@ define amdgpu_kernel void @marked_kernel_use_other_sgpr(ptr addrspace(1) %ptr) #
255255
%queue.load = load volatile i8, ptr addrspace(4) %queue.ptr
256256
%implicitarg.load = load volatile i8, ptr addrspace(4) %implicitarg.ptr
257257
%dispatch.load = load volatile i8, ptr addrspace(4) %dispatch.ptr
258+
store volatile i64 %dispatch.id, ptr addrspace(1) %ptr
258259
ret void
259260
}
260261

‎llvm/test/CodeGen/AMDGPU/bitcast-vector-extract.ll

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -73,7 +73,7 @@ define amdgpu_kernel void @store_bitcast_constant_v8i32_to_v16i16(ptr addrspace(
7373
define amdgpu_kernel void @store_value_lowered_to_undef_bitcast_source(ptr addrspace(1) %out, i64 %a, i64 %b) #0 {
7474
%undef = call i64 @llvm.amdgcn.icmp.i64(i64 %a, i64 %b, i32 999) #1
7575
%bc = bitcast i64 %undef to <2 x i32>
76-
store <2 x i32> %bc, ptr addrspace(1) %out
76+
store volatile <2 x i32> %bc, ptr addrspace(1) %out
7777
ret void
7878
}
7979

@@ -83,7 +83,7 @@ define amdgpu_kernel void @store_value_lowered_to_undef_bitcast_source_extractel
8383
%undef = call i64 @llvm.amdgcn.icmp.i64(i64 %a, i64 %b, i32 9999) #1
8484
%bc = bitcast i64 %undef to <2 x i32>
8585
%elt1 = extractelement <2 x i32> %bc, i32 1
86-
store i32 %elt1, ptr addrspace(1) %out
86+
store volatile i32 %elt1, ptr addrspace(1) %out
8787
ret void
8888
}
8989

‎llvm/test/CodeGen/AMDGPU/lds-global-non-entry-func.ll

Lines changed: 3 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -98,8 +98,8 @@ define void @func_use_lds_global() {
9898
ret void
9999
}
100100

101-
; ERR: warning: <unknown>:0:0: in function func_use_lds_global_constexpr_cast void (ptr addrspace(1)): local memory global used by non-kernel function
102-
define void @func_use_lds_global_constexpr_cast(ptr addrspace(1) %out) {
101+
; ERR: warning: <unknown>:0:0: in function func_use_lds_global_constexpr_cast void (): local memory global used by non-kernel function
102+
define void @func_use_lds_global_constexpr_cast() {
103103
; GFX8-SDAG-LABEL: func_use_lds_global_constexpr_cast:
104104
; GFX8-SDAG: ; %bb.0:
105105
; GFX8-SDAG-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
@@ -153,7 +153,7 @@ define void @func_use_lds_global_constexpr_cast(ptr addrspace(1) %out) {
153153
; GISEL-NEXT: s_setpc_b64 s[30:31]
154154
; GISEL-NEXT: .LBB1_2:
155155
; GISEL-NEXT: s_endpgm
156-
store i32 ptrtoint (ptr addrspace(3) @lds to i32), ptr addrspace(1) %out, align 4
156+
store volatile i32 ptrtoint (ptr addrspace(3) @lds to i32), ptr addrspace(1) poison, align 4
157157
ret void
158158
}
159159

‎llvm/test/CodeGen/AMDGPU/lower-work-group-id-intrinsics-pal.ll

Lines changed: 6 additions & 17 deletions
Original file line numberDiff line numberDiff line change
@@ -1,6 +1,6 @@
11
; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py
2-
; RUN: llc -mtriple=amdgcn-amd-amdpal -mcpu=gfx900 -mattr=-architected-sgprs -global-isel=0 -verify-machineinstrs < %s | FileCheck -check-prefixes=GFX9-SDAG %s
3-
; RUN: llc -mtriple=amdgcn-amd-amdpal -mcpu=gfx900 -mattr=-architected-sgprs -global-isel=1 -verify-machineinstrs < %s | FileCheck -check-prefixes=GFX9-GISEL %s
2+
; RUN: llc -mtriple=amdgcn-amd-amdpal -mcpu=gfx900 -mattr=-architected-sgprs -global-isel=0 -verify-machineinstrs < %s | FileCheck -check-prefixes=GFX9 %s
3+
; RUN: llc -mtriple=amdgcn-amd-amdpal -mcpu=gfx900 -mattr=-architected-sgprs -global-isel=1 -verify-machineinstrs < %s | FileCheck -check-prefixes=GFX9 %s
44
; RUN: llc -mtriple=amdgcn-amd-amdpal -mcpu=gfx900 -mattr=+architected-sgprs -global-isel=0 -verify-machineinstrs < %s | FileCheck -check-prefixes=GFX9ARCH-SDAG %s
55
; RUN: llc -mtriple=amdgcn-amd-amdpal -mcpu=gfx900 -mattr=+architected-sgprs -global-isel=1 -verify-machineinstrs < %s | FileCheck -check-prefixes=GFX9ARCH-GISEL %s
66
; RUN: llc -mtriple=amdgcn-amd-amdpal -mcpu=gfx1200 -global-isel=0 -verify-machineinstrs < %s | FileCheck -check-prefixes=GFX12,GFX12-SDAG %s
@@ -126,21 +126,10 @@ define amdgpu_cs void @caller() {
126126
declare amdgpu_gfx void @callee(i32)
127127

128128
define amdgpu_gfx void @workgroup_ids_gfx(ptr addrspace(1) %outx, ptr addrspace(1) %outy, ptr addrspace(1) %outz) {
129-
; GFX9-SDAG-LABEL: workgroup_ids_gfx:
130-
; GFX9-SDAG: ; %bb.0:
131-
; GFX9-SDAG-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
132-
; GFX9-SDAG-NEXT: global_store_dword v[0:1], v0, off
133-
; GFX9-SDAG-NEXT: s_waitcnt vmcnt(0)
134-
; GFX9-SDAG-NEXT: global_store_dword v[2:3], v0, off
135-
; GFX9-SDAG-NEXT: s_waitcnt vmcnt(0)
136-
; GFX9-SDAG-NEXT: global_store_dword v[4:5], v0, off
137-
; GFX9-SDAG-NEXT: s_waitcnt vmcnt(0)
138-
; GFX9-SDAG-NEXT: s_setpc_b64 s[30:31]
139-
;
140-
; GFX9-GISEL-LABEL: workgroup_ids_gfx:
141-
; GFX9-GISEL: ; %bb.0:
142-
; GFX9-GISEL-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
143-
; GFX9-GISEL-NEXT: s_setpc_b64 s[30:31]
129+
; GFX9-LABEL: workgroup_ids_gfx:
130+
; GFX9: ; %bb.0:
131+
; GFX9-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
132+
; GFX9-NEXT: s_setpc_b64 s[30:31]
144133
;
145134
; GFX9ARCH-SDAG-LABEL: workgroup_ids_gfx:
146135
; GFX9ARCH-SDAG: ; %bb.0:

‎llvm/test/CodeGen/AMDGPU/mem-builtins.ll

Lines changed: 18 additions & 18 deletions
Original file line numberDiff line numberDiff line change
@@ -9,64 +9,64 @@ declare hidden i32 @strnlen(ptr nocapture, i32) #1
99
declare hidden i32 @strcmp(ptr nocapture, ptr nocapture) #1
1010

1111

12-
; ERROR: error: <unknown>:0:0: in function test_memcmp void (ptr addrspace(1), ptr addrspace(1), ptr, ptr addrspace(1)): unsupported call to function memcmp
12+
; ERROR: error: <unknown>:0:0: in function test_memcmp void (ptr addrspace(1), ptr addrspace(1), ptr): unsupported call to function memcmp
1313

1414
; GCN: s_add_u32 s{{[0-9]+}}, s{{[0-9]+}}, memcmp@rel32@lo+4
1515
; GCN: s_addc_u32 s{{[0-9]+}}, s{{[0-9]+}}, memcmp@rel32@hi+12
16-
define amdgpu_kernel void @test_memcmp(ptr addrspace(1) %x, ptr addrspace(1) %y, ptr nocapture %p, ptr addrspace(1) %out) #0 {
16+
define amdgpu_kernel void @test_memcmp(ptr addrspace(1) %x, ptr addrspace(1) %y, ptr nocapture %p) #0 {
1717
entry:
1818
%cmp = tail call i32 @memcmp(ptr addrspace(1) %x, ptr addrspace(1) %y, i64 2)
19-
store i32 %cmp, ptr addrspace(1) %out
19+
store volatile i32 %cmp, ptr addrspace(1) undef
2020
ret void
2121
}
2222

23-
; ERROR: error: <unknown>:0:0: in function test_memchr void (ptr addrspace(1), i32, i64, ptr addrspace(1)): unsupported call to function memchr
23+
; ERROR: error: <unknown>:0:0: in function test_memchr void (ptr addrspace(1), i32, i64): unsupported call to function memchr
2424

2525
; GCN: s_add_u32 s{{[0-9]+}}, s{{[0-9]+}}, memchr@rel32@lo+4
2626
; GCN: s_addc_u32 s{{[0-9]+}}, s{{[0-9]+}}, memchr@rel32@hi+12
27-
define amdgpu_kernel void @test_memchr(ptr addrspace(1) %src, i32 %char, i64 %len, ptr addrspace(1) %out) #0 {
27+
define amdgpu_kernel void @test_memchr(ptr addrspace(1) %src, i32 %char, i64 %len) #0 {
2828
%res = call ptr addrspace(1) @memchr(ptr addrspace(1) %src, i32 %char, i64 %len)
29-
store ptr addrspace(1) %res, ptr addrspace(1) %out
29+
store volatile ptr addrspace(1) %res, ptr addrspace(1) undef
3030
ret void
3131
}
3232

33-
; ERROR: error: <unknown>:0:0: in function test_strcpy void (ptr, ptr, ptr addrspace(1)): unsupported call to function strcpy
33+
; ERROR: error: <unknown>:0:0: in function test_strcpy void (ptr, ptr): unsupported call to function strcpy
3434

3535
; GCN: s_add_u32 s{{[0-9]+}}, s{{[0-9]+}}, strcpy@rel32@lo+4
3636
; GCN: s_addc_u32 s{{[0-9]+}}, s{{[0-9]+}}, strcpy@rel32@hi+12
37-
define amdgpu_kernel void @test_strcpy(ptr %dst, ptr %src, ptr addrspace(1) %out) #0 {
37+
define amdgpu_kernel void @test_strcpy(ptr %dst, ptr %src) #0 {
3838
%res = call ptr @strcpy(ptr %dst, ptr %src)
39-
store ptr %res, ptr addrspace(1) %out
39+
store volatile ptr %res, ptr addrspace(1) undef
4040
ret void
4141
}
4242

43-
; ERROR: error: <unknown>:0:0: in function test_strcmp void (ptr, ptr, ptr addrspace(1)): unsupported call to function strcmp
43+
; ERROR: error: <unknown>:0:0: in function test_strcmp void (ptr, ptr): unsupported call to function strcmp
4444

4545
; GCN: s_add_u32 s{{[0-9]+}}, s{{[0-9]+}}, strcmp@rel32@lo+4
4646
; GCN: s_addc_u32 s{{[0-9]+}}, s{{[0-9]+}}, strcmp@rel32@hi+12
47-
define amdgpu_kernel void @test_strcmp(ptr %src0, ptr %src1, ptr addrspace(1) %out) #0 {
47+
define amdgpu_kernel void @test_strcmp(ptr %src0, ptr %src1) #0 {
4848
%res = call i32 @strcmp(ptr %src0, ptr %src1)
49-
store i32 %res, ptr addrspace(1) %out
49+
store volatile i32 %res, ptr addrspace(1) undef
5050
ret void
5151
}
5252

53-
; ERROR: error: <unknown>:0:0: in function test_strlen void (ptr, ptr addrspace(1)): unsupported call to function strlen
53+
; ERROR: error: <unknown>:0:0: in function test_strlen void (ptr): unsupported call to function strlen
5454

5555
; GCN: s_add_u32 s{{[0-9]+}}, s{{[0-9]+}}, strlen@rel32@lo+4
5656
; GCN: s_addc_u32 s{{[0-9]+}}, s{{[0-9]+}}, strlen@rel32@hi+12
57-
define amdgpu_kernel void @test_strlen(ptr %src, ptr addrspace(1) %out) #0 {
57+
define amdgpu_kernel void @test_strlen(ptr %src) #0 {
5858
%res = call i32 @strlen(ptr %src)
59-
store i32 %res, ptr addrspace(1) %out
59+
store volatile i32 %res, ptr addrspace(1) undef
6060
ret void
6161
}
6262

63-
; ERROR: error: <unknown>:0:0: in function test_strnlen void (ptr, i32, ptr addrspace(1)): unsupported call to function strnlen
63+
; ERROR: error: <unknown>:0:0: in function test_strnlen void (ptr, i32): unsupported call to function strnlen
6464

6565
; GCN: s_add_u32 s{{[0-9]+}}, s{{[0-9]+}}, strnlen@rel32@lo+4
6666
; GCN: s_addc_u32 s{{[0-9]+}}, s{{[0-9]+}}, strnlen@rel32@hi+12
67-
define amdgpu_kernel void @test_strnlen(ptr %src, i32 %size, ptr addrspace(1) %out) #0 {
67+
define amdgpu_kernel void @test_strnlen(ptr %src, i32 %size) #0 {
6868
%res = call i32 @strnlen(ptr %src, i32 %size)
69-
store i32 %res, ptr addrspace(1) %out
69+
store volatile i32 %res, ptr addrspace(1) undef
7070
ret void
7171
}
7272

‎llvm/test/CodeGen/NVPTX/store-undef.ll

Lines changed: 0 additions & 58 deletions
Original file line numberDiff line numberDiff line change
@@ -90,61 +90,3 @@ define void @test_store_def(i64 %param0, i32 %param1, ptr %out) {
9090
store %struct.T %S2, ptr %out
9191
ret void
9292
}
93-
94-
define void @test_store_volatile_undef(ptr %out, <8 x i32> %vec) {
95-
; CHECK-LABEL: test_store_volatile_undef(
96-
; CHECK: {
97-
; CHECK-NEXT: .reg .b32 %r<23>;
98-
; CHECK-NEXT: .reg .b64 %rd<5>;
99-
; CHECK-EMPTY:
100-
; CHECK-NEXT: // %bb.0:
101-
; CHECK-NEXT: ld.param.u64 %rd1, [test_store_volatile_undef_param_0];
102-
; CHECK-NEXT: st.volatile.v4.u32 [%rd1+16], {%r1, %r2, %r3, %r4};
103-
; CHECK-NEXT: st.volatile.v2.u32 [%rd1+8], {%r5, %r6};
104-
; CHECK-NEXT: st.volatile.u64 [%rd1], %rd2;
105-
; CHECK-NEXT: ld.param.v4.u32 {%r7, %r8, %r9, %r10}, [test_store_volatile_undef_param_1];
106-
; CHECK-NEXT: ld.param.v4.u32 {%r11, %r12, %r13, %r14}, [test_store_volatile_undef_param_1+16];
107-
; CHECK-NEXT: st.volatile.v4.u32 [%rd3], {%r11, %r12, %r13, %r14};
108-
; CHECK-NEXT: st.volatile.v4.u32 [%rd4], {%r7, %r8, %r9, %r10};
109-
; CHECK-NEXT: st.volatile.v4.u32 [%rd1+16], {%r15, %r16, %r17, %r18};
110-
; CHECK-NEXT: st.volatile.v4.u32 [%rd1], {%r19, %r20, %r21, %r22};
111-
; CHECK-NEXT: ret;
112-
store volatile %struct.T undef, ptr %out
113-
store volatile <8 x i32> %vec, ptr undef
114-
store volatile <8 x i32> undef, ptr %out
115-
ret void
116-
}
117-
118-
define void @test_store_volatile_of_poison(ptr %out) {
119-
; CHECK-LABEL: test_store_volatile_of_poison(
120-
; CHECK: {
121-
; CHECK-NEXT: .reg .b32 %r<7>;
122-
; CHECK-NEXT: .reg .b64 %rd<3>;
123-
; CHECK-EMPTY:
124-
; CHECK-NEXT: // %bb.0:
125-
; CHECK-NEXT: ld.param.u64 %rd1, [test_store_volatile_of_poison_param_0];
126-
; CHECK-NEXT: st.volatile.v4.u32 [%rd1+16], {%r1, %r2, %r3, %r4};
127-
; CHECK-NEXT: st.volatile.v2.u32 [%rd1+8], {%r5, %r6};
128-
; CHECK-NEXT: st.volatile.u64 [%rd1], %rd2;
129-
; CHECK-NEXT: ret;
130-
store volatile %struct.T poison, ptr %out
131-
ret void
132-
}
133-
134-
define void @test_store_volatile_to_poison(%struct.T %param) {
135-
; CHECK-LABEL: test_store_volatile_to_poison(
136-
; CHECK: {
137-
; CHECK-NEXT: .reg .b32 %r<7>;
138-
; CHECK-NEXT: .reg .b64 %rd<5>;
139-
; CHECK-EMPTY:
140-
; CHECK-NEXT: // %bb.0:
141-
; CHECK-NEXT: ld.param.u64 %rd1, [test_store_volatile_to_poison_param_0];
142-
; CHECK-NEXT: ld.param.v2.u32 {%r1, %r2}, [test_store_volatile_to_poison_param_0+8];
143-
; CHECK-NEXT: ld.param.v4.u32 {%r3, %r4, %r5, %r6}, [test_store_volatile_to_poison_param_0+16];
144-
; CHECK-NEXT: st.volatile.v4.u32 [%rd2], {%r3, %r4, %r5, %r6};
145-
; CHECK-NEXT: st.volatile.v2.u32 [%rd3], {%r1, %r2};
146-
; CHECK-NEXT: st.volatile.u64 [%rd4], %rd1;
147-
; CHECK-NEXT: ret;
148-
store volatile %struct.T %param, ptr poison
149-
ret void
150-
}

0 commit comments

Comments
 (0)