[llvm] [DAG] canCreateUndefOrPoison/isGuaranteedNotToBeUndefOrPoison - SCALAR_TO_VECTOR upper elements are poison (PR #217185)
Cyrus Ding via llvm-commits
llvm-commits at lists.llvm.org
Thu Aug 20 01:12:34 PDT 2026
https://github.com/dingcyrus updated https://github.com/llvm/llvm-project/pull/217185
>From b8570e2706aeab56cefc6fd30ede336298c81cea 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.
X86/bfloat.ll (PR92471) now materializes the poison padding lane
explicitly instead of folding it to zero.
Fixes #217028
---
llvm/lib/CodeGen/SelectionDAG/SelectionDAG.cpp | 8 ++++----
llvm/lib/CodeGen/SelectionDAG/TargetLowering.cpp | 12 +++++-------
llvm/test/CodeGen/AArch64/neon-dotreduce.ll | 7 ++++---
.../CodeGen/AMDGPU/shufflevector.v3bf16.v2bf16.ll | 6 ++++++
.../test/CodeGen/AMDGPU/shufflevector.v3f16.v2f16.ll | 6 ++++++
.../test/CodeGen/AMDGPU/shufflevector.v3i16.v2i16.ll | 6 ++++++
.../CodeGen/NVPTX/insert-vector-elt-shuffle-i8.ll | 2 +-
llvm/test/CodeGen/X86/bfloat.ll | 2 ++
8 files changed, 34 insertions(+), 15 deletions(-)
diff --git a/llvm/lib/CodeGen/SelectionDAG/SelectionDAG.cpp b/llvm/lib/CodeGen/SelectionDAG/SelectionDAG.cpp
index 00cef99069347..98daaacc3fba9 100644
--- a/llvm/lib/CodeGen/SelectionDAG/SelectionDAG.cpp
+++ b/llvm/lib/CodeGen/SelectionDAG/SelectionDAG.cpp
@@ -5819,8 +5819,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] &&
@@ -6097,8 +6097,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 cca7cdad0e2c8..2dbcb1a35a23d 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: {
diff --git a/llvm/test/CodeGen/AArch64/neon-dotreduce.ll b/llvm/test/CodeGen/AArch64/neon-dotreduce.ll
index f5312828cfd78..06d90d1e548f9 100644
--- a/llvm/test/CodeGen/AArch64/neon-dotreduce.ll
+++ b/llvm/test/CodeGen/AArch64/neon-dotreduce.ll
@@ -4906,16 +4906,17 @@ entry:
define i32 @test_udot_v33i8_nomla(ptr nocapture readonly %a1) {
; CHECK-SD-LABEL: test_udot_v33i8_nomla:
; CHECK-SD: // %bb.0: // %entry
-; CHECK-SD-NEXT: ldr b1, [x0, #32]
+; CHECK-SD-NEXT: ldrb w8, [x0, #32]
; CHECK-SD-NEXT: ldp q3, q2, [x0]
; CHECK-SD-NEXT: movi v0.2d, #0000000000000000
-; CHECK-SD-NEXT: ushll v1.8h, v1.8b, #0
+; CHECK-SD-NEXT: fmov s1, w8
; CHECK-SD-NEXT: ushll v4.8h, v2.8b, #0
; CHECK-SD-NEXT: ushll v5.8h, v3.8b, #0
; CHECK-SD-NEXT: ushll2 v2.8h, v2.16b, #0
; CHECK-SD-NEXT: ushll2 v3.8h, v3.16b, #0
-; CHECK-SD-NEXT: ushll v1.4s, v1.4h, #0
+; CHECK-SD-NEXT: ushll v1.8h, v1.8b, #0
; CHECK-SD-NEXT: uaddl2 v6.4s, v5.8h, v4.8h
+; CHECK-SD-NEXT: ushll v1.4s, v1.4h, #0
; CHECK-SD-NEXT: mov v0.s[0], v1.s[0]
; CHECK-SD-NEXT: uaddl2 v1.4s, v3.8h, v2.8h
; CHECK-SD-NEXT: uaddl v2.4s, v3.4h, v2.4h
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 7bccd6ba088ac..a91742c2375b9 100644
--- a/llvm/test/CodeGen/X86/bfloat.ll
+++ b/llvm/test/CodeGen/X86/bfloat.ll
@@ -2194,6 +2194,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