[llvm] [DAG] canCreateUndefOrPoison/isGuaranteedNotToBeUndefOrPoison - SCALAR_TO_VECTOR upper elements are poison (PR #217185)
Cyrus Ding via llvm-commits
llvm-commits at lists.llvm.org
Wed Sep 16 00:57:42 PDT 2026
https://github.com/dingcyrus updated https://github.com/llvm/llvm-project/pull/217185
>From c086b60e639efa46c34351f0c0e699bd4c4f9076 Mon Sep 17 00:00:00 2001
From: Cyrus Ding <785101675 at qq.com>
Date: Wed, 19 Aug 2026 09:28:51 +0800
Subject: [PATCH] [DAG] canCreateUndefOrPoison/isGuaranteedNotToBeUndefOrPoison
- SCALAR_TO_VECTOR upper elements are poison
ISD::SCALAR_TO_VECTOR documents its upper elements (1..N-1) as poison,
not undef. Update the code that inspects those elements to match:
- canCreateUndefOrPoison / isGuaranteedNotToBeUndefOrPoison now use
includesPoison(Kind) instead of includesUndef(Kind) when the upper
elements are demanded, so PoisonOnly queries correctly report that
they can create poison and are not guaranteed to be poison-free.
- SimplifyDemandedVectorElts no longer reports the upper elements as
KnownUndef, which let ZERO_EXTEND / ZERO_EXTEND_VECTOR_INREG fold
zext(undef) to zero on poison lanes. SimplifyDemandedBits only needed
a comment update.
With the corrected queries, DAGCombiner's build_vector-to-AND/OR-mask
fold now (correctly) freezes source vectors whose masked-off lanes
trace back to SCALAR_TO_VECTOR poison lanes. SimplifyDemandedVectorElts'
freeze(scalar_to_vector(x)) -> scalar_to_vector(freeze(x)) fold would
then sink that freeze below the SCALAR_TO_VECTOR. When the scalar is a
load of a promoted (wider than the element) type, freeze(load) can
never be folded away (a load is only poison-free with !noundef - the
loaded value may be poison in memory), and the resulting
scalar_to_vector(freeze(load)) no longer matches ISel patterns such as
the AArch64 scalar_to_vector(extload) -> ldr b/h forms (this regressed
AArch64/neon-dotreduce.ll from ldr b to ldrb + fmov). Skip the fold
for promoted loads: freeze(scalar_to_vector(load)) is equivalent for
the demanded element zero, and ISel selects the freeze as a plain
copy, keeping the load visible to the patterns. The sink is still
beneficial for well-typed scalar_to_vector (e.g. the scalar load
widening in X86/widen-load-of-small-alloca-with-zero-upper-half.ll),
so it is unchanged there.
X86/bfloat.ll (PR92471) now materializes the poison padding lane
explicitly instead of folding it to zero.
Fixes #217028
---
.../lib/CodeGen/SelectionDAG/SelectionDAG.cpp | 8 ++---
.../CodeGen/SelectionDAG/TargetLowering.cpp | 31 ++++++++++++-------
.../AMDGPU/shufflevector.v3bf16.v2bf16.ll | 6 ++++
.../AMDGPU/shufflevector.v3f16.v2f16.ll | 6 ++++
.../AMDGPU/shufflevector.v3i16.v2i16.ll | 6 ++++
.../NVPTX/insert-vector-elt-shuffle-i8.ll | 2 +-
llvm/test/CodeGen/X86/bfloat.ll | 2 ++
7 files changed, 45 insertions(+), 16 deletions(-)
diff --git a/llvm/lib/CodeGen/SelectionDAG/SelectionDAG.cpp b/llvm/lib/CodeGen/SelectionDAG/SelectionDAG.cpp
index 66a445ec0827e..2c0e5bed9a74d 100644
--- a/llvm/lib/CodeGen/SelectionDAG/SelectionDAG.cpp
+++ b/llvm/lib/CodeGen/SelectionDAG/SelectionDAG.cpp
@@ -5814,8 +5814,8 @@ bool SelectionDAG::isGuaranteedNotToBeUndefOrPoison(SDValue Op,
}
case ISD::SCALAR_TO_VECTOR:
- // Check upper (known undef) elements.
- if (DemandedElts.ugt(1) && includesUndef(Kind))
+ // Check upper (known poison) elements.
+ if (DemandedElts.ugt(1) && includesPoison(Kind))
return false;
// Check element zero.
if (DemandedElts[0] &&
@@ -6092,8 +6092,8 @@ bool SelectionDAG::canCreateUndefOrPoison(SDValue Op, const APInt &DemandedElts,
!isKnownNeverZero(Op.getOperand(0), Depth + 1);
case ISD::SCALAR_TO_VECTOR:
- // Check if we demand any upper (undef) elements.
- return includesUndef(Kind) && DemandedElts.ugt(1);
+ // Check if we demand any upper (poison) elements.
+ return includesPoison(Kind) && DemandedElts.ugt(1);
case ISD::INSERT_VECTOR_ELT:
case ISD::EXTRACT_VECTOR_ELT: {
diff --git a/llvm/lib/CodeGen/SelectionDAG/TargetLowering.cpp b/llvm/lib/CodeGen/SelectionDAG/TargetLowering.cpp
index b5f2cd0e61d06..3143f527e7602 100644
--- a/llvm/lib/CodeGen/SelectionDAG/TargetLowering.cpp
+++ b/llvm/lib/CodeGen/SelectionDAG/TargetLowering.cpp
@@ -1265,7 +1265,7 @@ bool TargetLowering::SimplifyDemandedBits(
if (VT.isScalableVector())
return false;
if (!DemandedElts[0])
- return TLO.CombineTo(Op, TLO.DAG.getUNDEF(VT));
+ return TLO.CombineTo(Op, TLO.DAG.getPOISON(VT));
KnownBits SrcKnown;
SDValue Src = Op.getOperand(0);
@@ -1274,7 +1274,7 @@ bool TargetLowering::SimplifyDemandedBits(
if (SimplifyDemandedBits(Src, SrcDemandedBits, SrcKnown, TLO, Depth + 1))
return true;
- // Upper elements are undef, so only get the knownbits if we just demand
+ // Upper elements are poison, so only get the knownbits if we just demand
// the bottom element.
if (DemandedElts == 1)
Known = SrcKnown.anyextOrTrunc(BitWidth);
@@ -3352,11 +3352,9 @@ bool TargetLowering::SimplifyDemandedVectorElts(
switch (Opcode) {
case ISD::SCALAR_TO_VECTOR: {
- if (!DemandedElts[0]) {
- KnownUndef.setAllBits();
- return TLO.CombineTo(Op, TLO.DAG.getUNDEF(VT));
- }
- KnownUndef.setHighBits(NumElts - 1);
+ if (!DemandedElts[0])
+ return TLO.CombineTo(Op, TLO.DAG.getPOISON(VT));
+ // Upper elements are poison, not undef - don't mark them as KnownUndef.
break;
}
case ISD::BITCAST: {
@@ -3476,10 +3474,21 @@ bool TargetLowering::SimplifyDemandedVectorElts(
// TODO: Replace this with the general fold from DAGCombiner::visitFREEZE
// freeze(op(x, ...)) -> op(freeze(x), ...).
- if (N0.getOpcode() == ISD::SCALAR_TO_VECTOR && DemandedElts == 1)
- return TLO.CombineTo(
- Op, TLO.DAG.getNode(ISD::SCALAR_TO_VECTOR, DL, VT,
- TLO.DAG.getFreeze(N0.getOperand(0))));
+ // Don't sink the freeze below SCALAR_TO_VECTOR when the scalar is a load
+ // of a promoted (wider than the element) type: freeze(load) can never be
+ // folded away (the loaded value may be poison in memory), and the extra
+ // freeze node then blocks ISel patterns matching scalar_to_vector of a
+ // load, e.g. the AArch64 scalar_to_vector(extload) -> ldr b/h forms.
+ // freeze(scalar_to_vector(load)) is equivalent for the demanded element
+ // zero, and ISel selects the freeze as a plain copy.
+ if (N0.getOpcode() == ISD::SCALAR_TO_VECTOR && DemandedElts == 1) {
+ SDValue Scalar = N0.getOperand(0);
+ bool IsPromotedLoad = Scalar.getOpcode() == ISD::LOAD &&
+ Scalar.getValueType() != VT.getVectorElementType();
+ if (!IsPromotedLoad)
+ return TLO.CombineTo(Op, TLO.DAG.getNode(ISD::SCALAR_TO_VECTOR, DL, VT,
+ TLO.DAG.getFreeze(Scalar)));
+ }
break;
}
case ISD::BUILD_VECTOR: {
diff --git a/llvm/test/CodeGen/AMDGPU/shufflevector.v3bf16.v2bf16.ll b/llvm/test/CodeGen/AMDGPU/shufflevector.v3bf16.v2bf16.ll
index 86dfc6f37680e..43a776703f4b9 100644
--- a/llvm/test/CodeGen/AMDGPU/shufflevector.v3bf16.v2bf16.ll
+++ b/llvm/test/CodeGen/AMDGPU/shufflevector.v3bf16.v2bf16.ll
@@ -23,6 +23,7 @@ define void @v_shuffle_v3bf16_v2bf16__0_u_u(ptr addrspace(1) inreg %ptr) {
; GFX900-NEXT: ;;#ASMSTART
; GFX900-NEXT: ; def v1
; GFX900-NEXT: ;;#ASMEND
+; GFX900-NEXT: global_store_short v0, v0, s[16:17] offset:4
; GFX900-NEXT: global_store_dword v0, v1, s[16:17]
; GFX900-NEXT: s_waitcnt vmcnt(0)
; GFX900-NEXT: s_setpc_b64 s[30:31]
@@ -34,6 +35,7 @@ define void @v_shuffle_v3bf16_v2bf16__0_u_u(ptr addrspace(1) inreg %ptr) {
; GFX90A-NEXT: ;;#ASMSTART
; GFX90A-NEXT: ; def v1
; GFX90A-NEXT: ;;#ASMEND
+; GFX90A-NEXT: global_store_short v0, v0, s[16:17] offset:4
; GFX90A-NEXT: global_store_dword v0, v1, s[16:17]
; GFX90A-NEXT: s_waitcnt vmcnt(0)
; GFX90A-NEXT: s_setpc_b64 s[30:31]
@@ -45,6 +47,7 @@ define void @v_shuffle_v3bf16_v2bf16__0_u_u(ptr addrspace(1) inreg %ptr) {
; GFX942-NEXT: ;;#ASMSTART
; GFX942-NEXT: ; def v1
; GFX942-NEXT: ;;#ASMEND
+; GFX942-NEXT: global_store_short v0, v0, s[0:1] offset:4
; GFX942-NEXT: global_store_dword v0, v1, s[0:1]
; GFX942-NEXT: s_waitcnt vmcnt(0)
; GFX942-NEXT: s_setpc_b64 s[30:31]
@@ -1413,6 +1416,7 @@ define void @v_shuffle_v3bf16_v2bf16__0_2_2(ptr addrspace(1) inreg %ptr) {
; GFX900-NEXT: ;;#ASMSTART
; GFX900-NEXT: ; def v1
; GFX900-NEXT: ;;#ASMEND
+; GFX900-NEXT: global_store_short v0, v0, s[16:17] offset:4
; GFX900-NEXT: global_store_dword v0, v1, s[16:17]
; GFX900-NEXT: s_waitcnt vmcnt(0)
; GFX900-NEXT: s_setpc_b64 s[30:31]
@@ -1424,6 +1428,7 @@ define void @v_shuffle_v3bf16_v2bf16__0_2_2(ptr addrspace(1) inreg %ptr) {
; GFX90A-NEXT: ;;#ASMSTART
; GFX90A-NEXT: ; def v1
; GFX90A-NEXT: ;;#ASMEND
+; GFX90A-NEXT: global_store_short v0, v0, s[16:17] offset:4
; GFX90A-NEXT: global_store_dword v0, v1, s[16:17]
; GFX90A-NEXT: s_waitcnt vmcnt(0)
; GFX90A-NEXT: s_setpc_b64 s[30:31]
@@ -1435,6 +1440,7 @@ define void @v_shuffle_v3bf16_v2bf16__0_2_2(ptr addrspace(1) inreg %ptr) {
; GFX942-NEXT: ;;#ASMSTART
; GFX942-NEXT: ; def v1
; GFX942-NEXT: ;;#ASMEND
+; GFX942-NEXT: global_store_short v0, v0, s[0:1] offset:4
; GFX942-NEXT: global_store_dword v0, v1, s[0:1]
; GFX942-NEXT: s_waitcnt vmcnt(0)
; GFX942-NEXT: s_setpc_b64 s[30:31]
diff --git a/llvm/test/CodeGen/AMDGPU/shufflevector.v3f16.v2f16.ll b/llvm/test/CodeGen/AMDGPU/shufflevector.v3f16.v2f16.ll
index 6aa9657fd06f3..5967eedc20ebd 100644
--- a/llvm/test/CodeGen/AMDGPU/shufflevector.v3f16.v2f16.ll
+++ b/llvm/test/CodeGen/AMDGPU/shufflevector.v3f16.v2f16.ll
@@ -23,6 +23,7 @@ define void @v_shuffle_v3f16_v2f16__0_u_u(ptr addrspace(1) inreg %ptr) {
; GFX900-NEXT: ;;#ASMSTART
; GFX900-NEXT: ; def v1
; GFX900-NEXT: ;;#ASMEND
+; GFX900-NEXT: global_store_short v0, v0, s[16:17] offset:4
; GFX900-NEXT: global_store_dword v0, v1, s[16:17]
; GFX900-NEXT: s_waitcnt vmcnt(0)
; GFX900-NEXT: s_setpc_b64 s[30:31]
@@ -34,6 +35,7 @@ define void @v_shuffle_v3f16_v2f16__0_u_u(ptr addrspace(1) inreg %ptr) {
; GFX90A-NEXT: ;;#ASMSTART
; GFX90A-NEXT: ; def v1
; GFX90A-NEXT: ;;#ASMEND
+; GFX90A-NEXT: global_store_short v0, v0, s[16:17] offset:4
; GFX90A-NEXT: global_store_dword v0, v1, s[16:17]
; GFX90A-NEXT: s_waitcnt vmcnt(0)
; GFX90A-NEXT: s_setpc_b64 s[30:31]
@@ -45,6 +47,7 @@ define void @v_shuffle_v3f16_v2f16__0_u_u(ptr addrspace(1) inreg %ptr) {
; GFX942-NEXT: ;;#ASMSTART
; GFX942-NEXT: ; def v1
; GFX942-NEXT: ;;#ASMEND
+; GFX942-NEXT: global_store_short v0, v0, s[0:1] offset:4
; GFX942-NEXT: global_store_dword v0, v1, s[0:1]
; GFX942-NEXT: s_waitcnt vmcnt(0)
; GFX942-NEXT: s_setpc_b64 s[30:31]
@@ -1413,6 +1416,7 @@ define void @v_shuffle_v3f16_v2f16__0_2_2(ptr addrspace(1) inreg %ptr) {
; GFX900-NEXT: ;;#ASMSTART
; GFX900-NEXT: ; def v1
; GFX900-NEXT: ;;#ASMEND
+; GFX900-NEXT: global_store_short v0, v0, s[16:17] offset:4
; GFX900-NEXT: global_store_dword v0, v1, s[16:17]
; GFX900-NEXT: s_waitcnt vmcnt(0)
; GFX900-NEXT: s_setpc_b64 s[30:31]
@@ -1424,6 +1428,7 @@ define void @v_shuffle_v3f16_v2f16__0_2_2(ptr addrspace(1) inreg %ptr) {
; GFX90A-NEXT: ;;#ASMSTART
; GFX90A-NEXT: ; def v1
; GFX90A-NEXT: ;;#ASMEND
+; GFX90A-NEXT: global_store_short v0, v0, s[16:17] offset:4
; GFX90A-NEXT: global_store_dword v0, v1, s[16:17]
; GFX90A-NEXT: s_waitcnt vmcnt(0)
; GFX90A-NEXT: s_setpc_b64 s[30:31]
@@ -1435,6 +1440,7 @@ define void @v_shuffle_v3f16_v2f16__0_2_2(ptr addrspace(1) inreg %ptr) {
; GFX942-NEXT: ;;#ASMSTART
; GFX942-NEXT: ; def v1
; GFX942-NEXT: ;;#ASMEND
+; GFX942-NEXT: global_store_short v0, v0, s[0:1] offset:4
; GFX942-NEXT: global_store_dword v0, v1, s[0:1]
; GFX942-NEXT: s_waitcnt vmcnt(0)
; GFX942-NEXT: s_setpc_b64 s[30:31]
diff --git a/llvm/test/CodeGen/AMDGPU/shufflevector.v3i16.v2i16.ll b/llvm/test/CodeGen/AMDGPU/shufflevector.v3i16.v2i16.ll
index bcfe492695da6..25ccbdd52c92c 100644
--- a/llvm/test/CodeGen/AMDGPU/shufflevector.v3i16.v2i16.ll
+++ b/llvm/test/CodeGen/AMDGPU/shufflevector.v3i16.v2i16.ll
@@ -23,6 +23,7 @@ define void @v_shuffle_v3i16_v2i16__0_u_u(ptr addrspace(1) inreg %ptr) {
; GFX900-NEXT: ;;#ASMSTART
; GFX900-NEXT: ; def v1
; GFX900-NEXT: ;;#ASMEND
+; GFX900-NEXT: global_store_short v0, v0, s[16:17] offset:4
; GFX900-NEXT: global_store_dword v0, v1, s[16:17]
; GFX900-NEXT: s_waitcnt vmcnt(0)
; GFX900-NEXT: s_setpc_b64 s[30:31]
@@ -34,6 +35,7 @@ define void @v_shuffle_v3i16_v2i16__0_u_u(ptr addrspace(1) inreg %ptr) {
; GFX90A-NEXT: ;;#ASMSTART
; GFX90A-NEXT: ; def v1
; GFX90A-NEXT: ;;#ASMEND
+; GFX90A-NEXT: global_store_short v0, v0, s[16:17] offset:4
; GFX90A-NEXT: global_store_dword v0, v1, s[16:17]
; GFX90A-NEXT: s_waitcnt vmcnt(0)
; GFX90A-NEXT: s_setpc_b64 s[30:31]
@@ -45,6 +47,7 @@ define void @v_shuffle_v3i16_v2i16__0_u_u(ptr addrspace(1) inreg %ptr) {
; GFX942-NEXT: ;;#ASMSTART
; GFX942-NEXT: ; def v1
; GFX942-NEXT: ;;#ASMEND
+; GFX942-NEXT: global_store_short v0, v0, s[0:1] offset:4
; GFX942-NEXT: global_store_dword v0, v1, s[0:1]
; GFX942-NEXT: s_waitcnt vmcnt(0)
; GFX942-NEXT: s_setpc_b64 s[30:31]
@@ -1406,6 +1409,7 @@ define void @v_shuffle_v3i16_v2i16__0_2_2(ptr addrspace(1) inreg %ptr) {
; GFX900-NEXT: ;;#ASMSTART
; GFX900-NEXT: ; def v1
; GFX900-NEXT: ;;#ASMEND
+; GFX900-NEXT: global_store_short v0, v0, s[16:17] offset:4
; GFX900-NEXT: global_store_dword v0, v1, s[16:17]
; GFX900-NEXT: s_waitcnt vmcnt(0)
; GFX900-NEXT: s_setpc_b64 s[30:31]
@@ -1417,6 +1421,7 @@ define void @v_shuffle_v3i16_v2i16__0_2_2(ptr addrspace(1) inreg %ptr) {
; GFX90A-NEXT: ;;#ASMSTART
; GFX90A-NEXT: ; def v1
; GFX90A-NEXT: ;;#ASMEND
+; GFX90A-NEXT: global_store_short v0, v0, s[16:17] offset:4
; GFX90A-NEXT: global_store_dword v0, v1, s[16:17]
; GFX90A-NEXT: s_waitcnt vmcnt(0)
; GFX90A-NEXT: s_setpc_b64 s[30:31]
@@ -1428,6 +1433,7 @@ define void @v_shuffle_v3i16_v2i16__0_2_2(ptr addrspace(1) inreg %ptr) {
; GFX942-NEXT: ;;#ASMSTART
; GFX942-NEXT: ; def v1
; GFX942-NEXT: ;;#ASMEND
+; GFX942-NEXT: global_store_short v0, v0, s[0:1] offset:4
; GFX942-NEXT: global_store_dword v0, v1, s[0:1]
; GFX942-NEXT: s_waitcnt vmcnt(0)
; GFX942-NEXT: s_setpc_b64 s[30:31]
diff --git a/llvm/test/CodeGen/NVPTX/insert-vector-elt-shuffle-i8.ll b/llvm/test/CodeGen/NVPTX/insert-vector-elt-shuffle-i8.ll
index 11ec3ff776726..a662af8920f26 100644
--- a/llvm/test/CodeGen/NVPTX/insert-vector-elt-shuffle-i8.ll
+++ b/llvm/test/CodeGen/NVPTX/insert-vector-elt-shuffle-i8.ll
@@ -19,7 +19,7 @@ define ptx_kernel void @pack_i8(<2 x i8> %a) {
; CHECK-EMPTY:
; CHECK-NEXT: // end inline asm
; CHECK-NEXT: cvt.u32.u16 %r2, %rs1;
-; CHECK-NEXT: prmt.b32 %r3, %r2, 0, 0x14U;
+; CHECK-NEXT: prmt.b32 %r3, %r2, 0, 0x3214U;
; CHECK-NEXT: mov.b64 %rd1, 0;
; CHECK-NEXT: st.shared.v4.b32 [%rd1], {%r3, 0, 0, 0};
; CHECK-NEXT: ret;
diff --git a/llvm/test/CodeGen/X86/bfloat.ll b/llvm/test/CodeGen/X86/bfloat.ll
index 4a234a5e80608..1d5e65afec89c 100644
--- a/llvm/test/CodeGen/X86/bfloat.ll
+++ b/llvm/test/CodeGen/X86/bfloat.ll
@@ -2006,6 +2006,8 @@ define void @PR92471(ptr %0, ptr %1) nounwind {
; SSE2-NEXT: movq {{.*#+}} xmm0 = mem[0],zero
; SSE2-NEXT: movd {{.*#+}} xmm1 = mem[0],zero,zero,zero
; SSE2-NEXT: pinsrw $2, 12(%rdi), %xmm1
+; SSE2-NEXT: pextrw $7, %xmm0, %eax
+; SSE2-NEXT: pinsrw $3, %eax, %xmm1
; SSE2-NEXT: pxor %xmm2, %xmm2
; SSE2-NEXT: pxor %xmm3, %xmm3
; SSE2-NEXT: punpcklwd {{.*#+}} xmm3 = xmm3[0],xmm1[0],xmm3[1],xmm1[1],xmm3[2],xmm1[2],xmm3[3],xmm1[3]
More information about the llvm-commits
mailing list