[llvm] [NVPTX] Fix BFE selection for i64 values (PR #228685)
via llvm-commits
llvm-commits at lists.llvm.org
Sat Oct 3 03:26:57 PDT 2026
https://github.com/wenhu1024 updated https://github.com/llvm/llvm-project/pull/228685
>From 3d8794ae701216820798534ec7050db02a99b1c2 Mon Sep 17 00:00:00 2001
From: wenhu1024 <1340773297 at qq.com>
Date: Sat, 3 Oct 2026 17:41:41 +0800
Subject: [PATCH] [NVPTX] Fix BFE selection for i64 values
Use the source value width when checking the BFE start position.
The shift amount is i32 even for i64 values, so using its type width
incorrectly rejects valid bit extracts starting at bit 32 or above.
---
llvm/lib/Target/NVPTX/NVPTXISelDAGToDAG.cpp | 2 +-
llvm/test/CodeGen/NVPTX/bfe.ll | 31 ++++++++-
llvm/test/CodeGen/NVPTX/convert-fp-i128.ll | 64 +++++++++----------
.../CodeGen/NVPTX/float-to-arbitrary-fp.ll | 18 +++---
llvm/test/CodeGen/NVPTX/i128.ll | 8 +--
5 files changed, 75 insertions(+), 48 deletions(-)
diff --git a/llvm/lib/Target/NVPTX/NVPTXISelDAGToDAG.cpp b/llvm/lib/Target/NVPTX/NVPTXISelDAGToDAG.cpp
index 186bfa185e135a..b048362d985117 100644
--- a/llvm/lib/Target/NVPTX/NVPTXISelDAGToDAG.cpp
+++ b/llvm/lib/Target/NVPTX/NVPTXISelDAGToDAG.cpp
@@ -1816,7 +1816,7 @@ bool NVPTXDAGToDAGISel::tryBFE(SDNode *N) {
uint64_t StartVal = StartConst->getZExtValue();
// How many "good" bits do we have left? "good" is defined here as bits
// that exist in the original value, not shifted in.
- int64_t GoodBits = Start.getValueSizeInBits() - StartVal;
+ int64_t GoodBits = Val.getValueSizeInBits() - StartVal;
if (NumBits > GoodBits) {
// Do not handle the case where bits have been shifted in. In theory
// we could handle this, but the cost is likely higher than just
diff --git a/llvm/test/CodeGen/NVPTX/bfe.ll b/llvm/test/CodeGen/NVPTX/bfe.ll
index 644bf3606e8f64..515529cc0461b1 100644
--- a/llvm/test/CodeGen/NVPTX/bfe.ll
+++ b/llvm/test/CodeGen/NVPTX/bfe.ll
@@ -51,6 +51,36 @@ define i32 @bfe2(i32 %a) {
ret i32 %val1
}
+define i64 @bfe_i64_start_above_32(i64 %a, i64 %b) {
+; CHECK-O3-LABEL: bfe_i64_start_above_32(
+; CHECK-O3: {
+; CHECK-O3-NEXT: .reg .b64 %rd<5>;
+; CHECK-O3-EMPTY:
+; CHECK-O3-NEXT: // %bb.0:
+; CHECK-O3-NEXT: ld.param.b64 %rd1, [bfe_i64_start_above_32_param_0];
+; CHECK-O3-NEXT: ld.param.b64 %rd2, [bfe_i64_start_above_32_param_1];
+; CHECK-O3-NEXT: xor.b64 %rd3, %rd1, %rd2;
+; CHECK-O3-NEXT: bfe.u64 %rd4, %rd3, 40, 8;
+; CHECK-O3-NEXT: st.param.b64 [func_retval0], %rd4;
+; CHECK-O3-NEXT: ret;
+;
+; CHECK-O0-LABEL: bfe_i64_start_above_32(
+; CHECK-O0: {
+; CHECK-O0-NEXT: .reg .b64 %rd<5>;
+; CHECK-O0-EMPTY:
+; CHECK-O0-NEXT: // %bb.0:
+; CHECK-O0-NEXT: ld.param.b64 %rd2, [bfe_i64_start_above_32_param_1];
+; CHECK-O0-NEXT: ld.param.b64 %rd1, [bfe_i64_start_above_32_param_0];
+; CHECK-O0-NEXT: xor.b64 %rd3, %rd1, %rd2;
+; CHECK-O0-NEXT: bfe.u64 %rd4, %rd3, 40, 8;
+; CHECK-O0-NEXT: st.param.b64 [func_retval0], %rd4;
+; CHECK-O0-NEXT: ret;
+ %x = xor i64 %a, %b
+ %val0 = lshr i64 %x, 40
+ %val1 = and i64 %val0, 255
+ ret i64 %val1
+}
+
define i32 @no_bfe_on_32bit_overflow(i32 %a) {
; CHECK-LABEL: no_bfe_on_32bit_overflow(
; CHECK: {
@@ -218,4 +248,3 @@ define i64 @bfe4(i128 %a) {
%shr = ashr exact i64 %and, 17
ret i64 %shr
}
-
diff --git a/llvm/test/CodeGen/NVPTX/convert-fp-i128.ll b/llvm/test/CodeGen/NVPTX/convert-fp-i128.ll
index f9eff6a8d6cd16..54305b95539455 100644
--- a/llvm/test/CodeGen/NVPTX/convert-fp-i128.ll
+++ b/llvm/test/CodeGen/NVPTX/convert-fp-i128.ll
@@ -74,38 +74,37 @@ define i128 @cvt_u128_f64(double %x) {
; CHECK: {
; CHECK-NEXT: .reg .pred %p<4>;
; CHECK-NEXT: .reg .b32 %r<7>;
-; CHECK-NEXT: .reg .b64 %rd<10>;
+; CHECK-NEXT: .reg .b64 %rd<9>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0: // %fp-to-i-entry
; CHECK-NEXT: ld.param.b64 %rd3, [cvt_u128_f64_param_0];
-; CHECK-NEXT: shr.u64 %rd4, %rd3, 52;
-; CHECK-NEXT: and.b64 %rd1, %rd4, 2047;
+; CHECK-NEXT: bfe.u64 %rd1, %rd3, 52, 11;
; CHECK-NEXT: setp.lt.u64 %p1, %rd1, 1023;
-; CHECK-NEXT: mov.b64 %rd9, 0;
-; CHECK-NEXT: mov.b64 %rd8, %rd9;
+; CHECK-NEXT: mov.b64 %rd8, 0;
+; CHECK-NEXT: mov.b64 %rd7, %rd8;
; CHECK-NEXT: @%p1 bra $L__BB2_4;
; CHECK-NEXT: // %bb.1: // %fp-to-i-if-check.exp.size
-; CHECK-NEXT: and.b64 %rd5, %rd3, 4503599627370495;
-; CHECK-NEXT: or.b64 %rd2, %rd5, 4503599627370496;
+; CHECK-NEXT: and.b64 %rd4, %rd3, 4503599627370495;
+; CHECK-NEXT: or.b64 %rd2, %rd4, 4503599627370496;
; CHECK-NEXT: setp.gt.u64 %p2, %rd1, 1074;
; CHECK-NEXT: @%p2 bra $L__BB2_3;
; CHECK-NEXT: // %bb.2: // %fp-to-i-if-exp.small
; CHECK-NEXT: cvt.u32.u64 %r5, %rd1;
; CHECK-NEXT: sub.s32 %r6, 1075, %r5;
-; CHECK-NEXT: shr.u64 %rd8, %rd2, %r6;
+; CHECK-NEXT: shr.u64 %rd7, %rd2, %r6;
; CHECK-NEXT: bra.uni $L__BB2_4;
; CHECK-NEXT: $L__BB2_3: // %fp-to-i-if-exp.large
; CHECK-NEXT: cvt.u32.u64 %r1, %rd1;
; CHECK-NEXT: sub.s32 %r2, 1139, %r1;
-; CHECK-NEXT: shr.u64 %rd6, %rd2, %r2;
+; CHECK-NEXT: shr.u64 %rd5, %rd2, %r2;
; CHECK-NEXT: add.s32 %r3, %r1, -1139;
-; CHECK-NEXT: shl.b64 %rd7, %rd2, %r3;
+; CHECK-NEXT: shl.b64 %rd6, %rd2, %r3;
; CHECK-NEXT: add.s32 %r4, %r1, -1075;
; CHECK-NEXT: setp.gt.s32 %p3, %r4, 63;
-; CHECK-NEXT: selp.b64 %rd9, %rd7, %rd6, %p3;
-; CHECK-NEXT: shl.b64 %rd8, %rd2, %r4;
+; CHECK-NEXT: selp.b64 %rd8, %rd6, %rd5, %p3;
+; CHECK-NEXT: shl.b64 %rd7, %rd2, %r4;
; CHECK-NEXT: $L__BB2_4: // %fp-to-i-cleanup
-; CHECK-NEXT: st.param.v2.b64 [func_retval0], {%rd8, %rd9};
+; CHECK-NEXT: st.param.v2.b64 [func_retval0], {%rd7, %rd8};
; CHECK-NEXT: ret;
%a = fptoui double %x to i128
ret i128 %a
@@ -482,47 +481,46 @@ define i128 @cvt_s128_f64(double %x) {
; CHECK: {
; CHECK-NEXT: .reg .pred %p<4>;
; CHECK-NEXT: .reg .b32 %r<7>;
-; CHECK-NEXT: .reg .b64 %rd<18>;
+; CHECK-NEXT: .reg .b64 %rd<17>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0: // %fp-to-i-entry
; CHECK-NEXT: ld.param.b64 %rd5, [cvt_s128_f64_param_0];
-; CHECK-NEXT: shr.u64 %rd6, %rd5, 52;
-; CHECK-NEXT: and.b64 %rd3, %rd6, 2047;
+; CHECK-NEXT: bfe.u64 %rd3, %rd5, 52, 11;
; CHECK-NEXT: setp.lt.u64 %p1, %rd3, 1023;
-; CHECK-NEXT: mov.b64 %rd16, 0;
-; CHECK-NEXT: mov.b64 %rd17, %rd16;
+; CHECK-NEXT: mov.b64 %rd15, 0;
+; CHECK-NEXT: mov.b64 %rd16, %rd15;
; CHECK-NEXT: @%p1 bra $L__BB8_4;
; CHECK-NEXT: // %bb.1: // %fp-to-i-if-check.exp.size
; CHECK-NEXT: shr.s64 %rd2, %rd5, 63;
; CHECK-NEXT: or.b64 %rd1, %rd2, 1;
-; CHECK-NEXT: and.b64 %rd7, %rd5, 4503599627370495;
-; CHECK-NEXT: or.b64 %rd4, %rd7, 4503599627370496;
+; CHECK-NEXT: and.b64 %rd6, %rd5, 4503599627370495;
+; CHECK-NEXT: or.b64 %rd4, %rd6, 4503599627370496;
; CHECK-NEXT: setp.gt.u64 %p2, %rd3, 1074;
; CHECK-NEXT: @%p2 bra $L__BB8_3;
; CHECK-NEXT: // %bb.2: // %fp-to-i-if-exp.small
; CHECK-NEXT: cvt.u32.u64 %r5, %rd3;
; CHECK-NEXT: sub.s32 %r6, 1075, %r5;
-; CHECK-NEXT: shr.u64 %rd14, %rd4, %r6;
-; CHECK-NEXT: mul.hi.u64 %rd15, %rd14, %rd1;
-; CHECK-NEXT: mad.lo.s64 %rd17, %rd14, %rd2, %rd15;
-; CHECK-NEXT: mul.lo.s64 %rd16, %rd14, %rd1;
+; CHECK-NEXT: shr.u64 %rd13, %rd4, %r6;
+; CHECK-NEXT: mul.hi.u64 %rd14, %rd13, %rd1;
+; CHECK-NEXT: mad.lo.s64 %rd16, %rd13, %rd2, %rd14;
+; CHECK-NEXT: mul.lo.s64 %rd15, %rd13, %rd1;
; CHECK-NEXT: bra.uni $L__BB8_4;
; CHECK-NEXT: $L__BB8_3: // %fp-to-i-if-exp.large
; CHECK-NEXT: cvt.u32.u64 %r1, %rd3;
; CHECK-NEXT: sub.s32 %r2, 1139, %r1;
-; CHECK-NEXT: shr.u64 %rd8, %rd4, %r2;
+; CHECK-NEXT: shr.u64 %rd7, %rd4, %r2;
; CHECK-NEXT: add.s32 %r3, %r1, -1139;
-; CHECK-NEXT: shl.b64 %rd9, %rd4, %r3;
+; CHECK-NEXT: shl.b64 %rd8, %rd4, %r3;
; CHECK-NEXT: add.s32 %r4, %r1, -1075;
; CHECK-NEXT: setp.gt.s32 %p3, %r4, 63;
-; CHECK-NEXT: selp.b64 %rd10, %rd9, %rd8, %p3;
-; CHECK-NEXT: shl.b64 %rd11, %rd4, %r4;
-; CHECK-NEXT: mul.hi.u64 %rd12, %rd11, %rd1;
-; CHECK-NEXT: mad.lo.s64 %rd13, %rd11, %rd2, %rd12;
-; CHECK-NEXT: mad.lo.s64 %rd17, %rd10, %rd1, %rd13;
-; CHECK-NEXT: mul.lo.s64 %rd16, %rd11, %rd1;
+; CHECK-NEXT: selp.b64 %rd9, %rd8, %rd7, %p3;
+; CHECK-NEXT: shl.b64 %rd10, %rd4, %r4;
+; CHECK-NEXT: mul.hi.u64 %rd11, %rd10, %rd1;
+; CHECK-NEXT: mad.lo.s64 %rd12, %rd10, %rd2, %rd11;
+; CHECK-NEXT: mad.lo.s64 %rd16, %rd9, %rd1, %rd12;
+; CHECK-NEXT: mul.lo.s64 %rd15, %rd10, %rd1;
; CHECK-NEXT: $L__BB8_4: // %fp-to-i-cleanup
-; CHECK-NEXT: st.param.v2.b64 [func_retval0], {%rd16, %rd17};
+; CHECK-NEXT: st.param.v2.b64 [func_retval0], {%rd15, %rd16};
; CHECK-NEXT: ret;
%a = fptosi double %x to i128
ret i128 %a
diff --git a/llvm/test/CodeGen/NVPTX/float-to-arbitrary-fp.ll b/llvm/test/CodeGen/NVPTX/float-to-arbitrary-fp.ll
index 16ac8c8634db5f..eaf30c67bf84fc 100644
--- a/llvm/test/CodeGen/NVPTX/float-to-arbitrary-fp.ll
+++ b/llvm/test/CodeGen/NVPTX/float-to-arbitrary-fp.ll
@@ -2012,15 +2012,15 @@ define i8 @to_f8e5m2_from_f64(double %x) {
; CHECK-NEXT: and.b64 %rd34, %rd33, 128;
; CHECK-NEXT: or.b64 %rd35, %rd34, %rd32;
; CHECK-NEXT: or.b64 %rd36, %rd35, %rd31;
-; CHECK-NEXT: bfe.u64 %rd37, %rd8, 50, 2;
-; CHECK-NEXT: and.b64 %rd38, %rd37, 1;
-; CHECK-NEXT: and.b64 %rd39, %rd8, 562949953421311;
-; CHECK-NEXT: setp.ne.b64 %p6, %rd39, 0;
-; CHECK-NEXT: selp.b64 %rd40, 1, 0, %p6;
-; CHECK-NEXT: or.b64 %rd41, %rd40, %rd38;
-; CHECK-NEXT: shr.u64 %rd42, %rd8, 49;
-; CHECK-NEXT: and.b64 %rd43, %rd42, %rd41;
-; CHECK-NEXT: add.s64 %rd44, %rd37, %rd43;
+; CHECK-NEXT: bfe.u64 %rd37, %rd9, 50, 1;
+; CHECK-NEXT: and.b64 %rd38, %rd8, 562949953421311;
+; CHECK-NEXT: setp.ne.b64 %p6, %rd38, 0;
+; CHECK-NEXT: selp.b64 %rd39, 1, 0, %p6;
+; CHECK-NEXT: or.b64 %rd40, %rd39, %rd37;
+; CHECK-NEXT: shr.u64 %rd41, %rd8, 49;
+; CHECK-NEXT: and.b64 %rd42, %rd41, %rd40;
+; CHECK-NEXT: bfe.u64 %rd43, %rd8, 50, 2;
+; CHECK-NEXT: add.s64 %rd44, %rd43, %rd42;
; CHECK-NEXT: setp.gt.s64 %p7, %rd44, 3;
; CHECK-NEXT: selp.b64 %rd45, 0, %rd44, %p7;
; CHECK-NEXT: selp.b64 %rd46, 1, 0, %p7;
diff --git a/llvm/test/CodeGen/NVPTX/i128.ll b/llvm/test/CodeGen/NVPTX/i128.ll
index df1214c3f5e006..cff23fe3117810 100644
--- a/llvm/test/CodeGen/NVPTX/i128.ll
+++ b/llvm/test/CodeGen/NVPTX/i128.ll
@@ -108,7 +108,7 @@ define i128 @srem_i128(i128 %lhs, i128 %rhs) {
; CHECK-NEXT: sub.cc.s64 %rd50, %rd6, %rd45;
; CHECK-NEXT: subc.cc.s64 %rd51, %rd7, %rd42;
; CHECK-NEXT: shr.s64 %rd52, %rd51, 63;
-; CHECK-NEXT: and.b64 %rd69, %rd52, 1;
+; CHECK-NEXT: bfe.u64 %rd69, %rd51, 63, 1;
; CHECK-NEXT: and.b64 %rd53, %rd52, %rd4;
; CHECK-NEXT: and.b64 %rd54, %rd52, %rd5;
; CHECK-NEXT: sub.cc.s64 %rd73, %rd45, %rd53;
@@ -236,7 +236,7 @@ define i128 @urem_i128(i128 %lhs, i128 %rhs) {
; CHECK-NEXT: sub.cc.s64 %rd41, %rd3, %rd36;
; CHECK-NEXT: subc.cc.s64 %rd42, %rd4, %rd33;
; CHECK-NEXT: shr.s64 %rd43, %rd42, 63;
-; CHECK-NEXT: and.b64 %rd56, %rd43, 1;
+; CHECK-NEXT: bfe.u64 %rd56, %rd42, 63, 1;
; CHECK-NEXT: and.b64 %rd44, %rd43, %rd1;
; CHECK-NEXT: and.b64 %rd45, %rd43, %rd2;
; CHECK-NEXT: sub.cc.s64 %rd60, %rd36, %rd44;
@@ -406,7 +406,7 @@ define i128 @sdiv_i128(i128 %lhs, i128 %rhs) {
; CHECK-NEXT: sub.cc.s64 %rd51, %rd6, %rd46;
; CHECK-NEXT: subc.cc.s64 %rd52, %rd7, %rd43;
; CHECK-NEXT: shr.s64 %rd53, %rd52, 63;
-; CHECK-NEXT: and.b64 %rd64, %rd53, 1;
+; CHECK-NEXT: bfe.u64 %rd64, %rd52, 63, 1;
; CHECK-NEXT: and.b64 %rd54, %rd53, %rd3;
; CHECK-NEXT: and.b64 %rd55, %rd53, %rd4;
; CHECK-NEXT: sub.cc.s64 %rd68, %rd46, %rd54;
@@ -528,7 +528,7 @@ define i128 @udiv_i128(i128 %lhs, i128 %rhs) {
; CHECK-NEXT: sub.cc.s64 %rd41, %rd1, %rd36;
; CHECK-NEXT: subc.cc.s64 %rd42, %rd2, %rd33;
; CHECK-NEXT: shr.s64 %rd43, %rd42, 63;
-; CHECK-NEXT: and.b64 %rd50, %rd43, 1;
+; CHECK-NEXT: bfe.u64 %rd50, %rd42, 63, 1;
; CHECK-NEXT: and.b64 %rd44, %rd43, %rd5;
; CHECK-NEXT: and.b64 %rd45, %rd43, %rd6;
; CHECK-NEXT: sub.cc.s64 %rd54, %rd36, %rd44;
More information about the llvm-commits
mailing list