[llvm] [NVPTX] implement `TLI.isTruncateFree(EVT, EVT)` (PR #181484)
via llvm-commits
llvm-commits at lists.llvm.org
Sat Feb 14 07:20:00 PST 2026
llvmbot wrote:
<!--LLVM PR SUMMARY COMMENT-->
@llvm/pr-subscribers-backend-nvptx
Author: Aayush Sabharwal (AayushSabharwal)
<details>
<summary>Changes</summary>
Closes #<!-- -->114339
While the linked issue (if I understand correctly) suggests making `isZextFree` for `i32->i64`, I found no documentation to confirm this. https://github.com/llvm/llvm-project/pull/115139#discussion_r1831507650 also suggests that this isn't free in general. As such, this PR only implements the missing `isTruncateFree` method. The test case is taken from `llvm/test/CodeGen/RISCV/trunc-free.ll`. https://godbolt.org/z/W7Y1Ko95b is the codegen prior to this change.
---
Full diff: https://github.com/llvm/llvm-project/pull/181484.diff
6 Files Affected:
- (modified) llvm/lib/Target/NVPTX/NVPTXISelLowering.h (+6)
- (modified) llvm/test/CodeGen/NVPTX/cse-mov-sym.ll (+2-2)
- (modified) llvm/test/CodeGen/NVPTX/i128-array.ll (+6-6)
- (modified) llvm/test/CodeGen/NVPTX/masked-load-3xhalf.ll (+12-11)
- (modified) llvm/test/CodeGen/NVPTX/mulwide.ll (+12-12)
- (added) llvm/test/CodeGen/NVPTX/truncate-free.ll (+36)
``````````diff
diff --git a/llvm/lib/Target/NVPTX/NVPTXISelLowering.h b/llvm/lib/Target/NVPTX/NVPTXISelLowering.h
index 9f35fe1e866fa..0bbaea4a48b6a 100644
--- a/llvm/lib/Target/NVPTX/NVPTXISelLowering.h
+++ b/llvm/lib/Target/NVPTX/NVPTXISelLowering.h
@@ -68,6 +68,12 @@ class NVPTXTargetLowering : public TargetLowering {
unsigned AS,
Instruction *I = nullptr) const override;
+ bool isTruncateFree(EVT SrcTy, EVT DSTTy) const override {
+ if (!SrcTy.isScalarInteger() || !DSTTy.isScalarInteger())
+ return false;
+ return SrcTy.getSizeInBits() == 64 && DSTTy.getSizeInBits() == 32;
+ }
+
bool isTruncateFree(Type *SrcTy, Type *DstTy) const override {
// Truncating 64-bit to 32-bit is free in SASS.
if (!SrcTy->isIntegerTy() || !DstTy->isIntegerTy())
diff --git a/llvm/test/CodeGen/NVPTX/cse-mov-sym.ll b/llvm/test/CodeGen/NVPTX/cse-mov-sym.ll
index 2e68208786d24..2f3c82ceb278a 100644
--- a/llvm/test/CodeGen/NVPTX/cse-mov-sym.ll
+++ b/llvm/test/CodeGen/NVPTX/cse-mov-sym.ll
@@ -17,17 +17,17 @@ define i32 @test_mov_sym(i32 %offset1, i32 %offset2, i1 %cond) {
; CHECK-NEXT: .reg .b64 %rd<6>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0: // %entry
+; CHECK-NEXT: ld.param.s32 %rd1, [test_mov_sym_param_0];
; CHECK-NEXT: ld.param.b8 %rs1, [test_mov_sym_param_2];
; CHECK-NEXT: and.b16 %rs2, %rs1, 1;
; CHECK-NEXT: setp.ne.b16 %p1, %rs2, 0;
-; CHECK-NEXT: ld.param.b32 %r1, [test_mov_sym_param_0];
-; CHECK-NEXT: cvt.s64.s32 %rd1, %r1;
; CHECK-NEXT: mov.b64 %rd2, global_smem;
; CHECK-NEXT: add.s64 %rd3, %rd2, %rd1;
; CHECK-NEXT: ld.shared.b32 %r4, [%rd3];
; CHECK-NEXT: not.pred %p2, %p1;
; CHECK-NEXT: @%p2 bra $L__BB0_4;
; CHECK-NEXT: // %bb.1: // %if1.preheader
+; CHECK-NEXT: cvt.u32.u64 %r1, %rd1;
; CHECK-NEXT: ld.param.b32 %r2, [test_mov_sym_param_1];
; CHECK-NEXT: setp.ne.b32 %p3, %r1, %r2;
; CHECK-NEXT: $L__BB0_2: // %if1
diff --git a/llvm/test/CodeGen/NVPTX/i128-array.ll b/llvm/test/CodeGen/NVPTX/i128-array.ll
index 7bd8a0021f1b5..7cf69688b8716 100644
--- a/llvm/test/CodeGen/NVPTX/i128-array.ll
+++ b/llvm/test/CodeGen/NVPTX/i128-array.ll
@@ -8,13 +8,13 @@ define [2 x i128] @foo(i64 %a, i32 %b) {
; CHECK-NEXT: .reg .b64 %rd<5>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
-; CHECK-NEXT: ld.param.b32 %r1, [foo_param_1];
; CHECK-NEXT: ld.param.b64 %rd1, [foo_param_0];
-; CHECK-NEXT: shr.s64 %rd2, %rd1, 63;
-; CHECK-NEXT: cvt.s64.s32 %rd3, %r1;
-; CHECK-NEXT: shr.s64 %rd4, %rd3, 63;
-; CHECK-NEXT: st.param.v2.b64 [func_retval0], {%rd1, %rd2};
-; CHECK-NEXT: st.param.v2.b64 [func_retval0+16], {%rd3, %rd4};
+; CHECK-NEXT: ld.param.s32 %rd2, [foo_param_1];
+; CHECK-NEXT: cvt.u32.u64 %r1, %rd2;
+; CHECK-NEXT: shr.s64 %rd3, %rd1, 63;
+; CHECK-NEXT: shr.s64 %rd4, %rd2, 63;
+; CHECK-NEXT: st.param.v2.b64 [func_retval0], {%rd1, %rd3};
+; CHECK-NEXT: st.param.v2.b64 [func_retval0+16], {%rd2, %rd4};
; CHECK-NEXT: ret;
%1 = sext i64 %a to i128
%2 = sext i32 %b to i128
diff --git a/llvm/test/CodeGen/NVPTX/masked-load-3xhalf.ll b/llvm/test/CodeGen/NVPTX/masked-load-3xhalf.ll
index bba240c694040..0d1fed969cb89 100644
--- a/llvm/test/CodeGen/NVPTX/masked-load-3xhalf.ll
+++ b/llvm/test/CodeGen/NVPTX/masked-load-3xhalf.ll
@@ -47,29 +47,30 @@ define void @halfx3_extend_chain(ptr align 16 captures(none) %rd0) {
define void @halfx3_no_align(ptr align 4 captures(none) %rd0) {
; CHECK-LABEL: halfx3_no_align(
; CHECK: {
-; CHECK-NEXT: .reg .b16 %rs<7>;
+; CHECK-NEXT: .reg .b16 %rs<6>;
; CHECK-NEXT: .reg .b32 %r<10>;
; CHECK-NEXT: .reg .b64 %rd<2>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.b64 %rd1, [halfx3_no_align_param_0];
; CHECK-NEXT: ld.b16 %rs1, [%rd1+4];
-; CHECK-NEXT: mov.b32 %r1, {%rs1, %rs2};
-; CHECK-NEXT: ld.b32 %r2, [%rd1];
-; CHECK-NEXT: mov.b32 %r3, 0;
-; CHECK-NEXT: max.f16x2 %r4, %r1, %r3;
-; CHECK-NEXT: max.f16x2 %r5, %r2, %r3;
+; CHECK-NEXT: mov.b32 %r1, 0;
+; CHECK-NEXT: mov.b32 {_, %rs2}, %r1;
+; CHECK-NEXT: mov.b32 %r2, {%rs1, %rs2};
+; CHECK-NEXT: ld.b32 %r3, [%rd1];
+; CHECK-NEXT: max.f16x2 %r4, %r2, %r1;
+; CHECK-NEXT: max.f16x2 %r5, %r3, %r1;
; CHECK-NEXT: st.b32 [%rd1], %r5;
; CHECK-NEXT: mov.b32 {%rs3, _}, %r4;
; CHECK-NEXT: st.b16 [%rd1+4], %rs3;
; CHECK-NEXT: ld.b16 %rs4, [%rd1+10];
-; CHECK-NEXT: mov.b32 %r6, {%rs4, %rs5};
+; CHECK-NEXT: mov.b32 %r6, {%rs4, %rs2};
; CHECK-NEXT: ld.b32 %r7, [%rd1+6];
-; CHECK-NEXT: max.f16x2 %r8, %r6, %r3;
-; CHECK-NEXT: max.f16x2 %r9, %r7, %r3;
+; CHECK-NEXT: max.f16x2 %r8, %r6, %r1;
+; CHECK-NEXT: max.f16x2 %r9, %r7, %r1;
; CHECK-NEXT: st.b32 [%rd1+6], %r9;
-; CHECK-NEXT: mov.b32 {%rs6, _}, %r8;
-; CHECK-NEXT: st.b16 [%rd1+10], %rs6;
+; CHECK-NEXT: mov.b32 {%rs5, _}, %r8;
+; CHECK-NEXT: st.b16 [%rd1+10], %rs5;
; CHECK-NEXT: ret;
%load1 = load <3 x half>, ptr %rd0, align 4
%p1 = fcmp ogt <3 x half> %load1, zeroinitializer
diff --git a/llvm/test/CodeGen/NVPTX/mulwide.ll b/llvm/test/CodeGen/NVPTX/mulwide.ll
index bde57fb7b95b0..bc818ee63790a 100644
--- a/llvm/test/CodeGen/NVPTX/mulwide.ll
+++ b/llvm/test/CodeGen/NVPTX/mulwide.ll
@@ -153,10 +153,10 @@ define i64 @mulwide32(i32 %a, i32 %b) {
; NOOPT-NEXT: .reg .b64 %rd<4>;
; NOOPT-EMPTY:
; NOOPT-NEXT: // %bb.0:
-; NOOPT-NEXT: ld.param.b32 %r2, [mulwide32_param_1];
-; NOOPT-NEXT: ld.param.b32 %r1, [mulwide32_param_0];
-; NOOPT-NEXT: cvt.s64.s32 %rd1, %r1;
-; NOOPT-NEXT: cvt.s64.s32 %rd2, %r2;
+; NOOPT-NEXT: ld.param.s32 %rd1, [mulwide32_param_0];
+; NOOPT-NEXT: cvt.u32.u64 %r1, %rd1;
+; NOOPT-NEXT: ld.param.s32 %rd2, [mulwide32_param_1];
+; NOOPT-NEXT: cvt.u32.u64 %r2, %rd2;
; NOOPT-NEXT: mul.lo.s64 %rd3, %rd1, %rd2;
; NOOPT-NEXT: st.param.b64 [func_retval0], %rd3;
; NOOPT-NEXT: ret;
@@ -185,10 +185,10 @@ define i64 @mulwideu32(i32 %a, i32 %b) {
; NOOPT-NEXT: .reg .b64 %rd<4>;
; NOOPT-EMPTY:
; NOOPT-NEXT: // %bb.0:
-; NOOPT-NEXT: ld.param.b32 %r2, [mulwideu32_param_1];
-; NOOPT-NEXT: ld.param.b32 %r1, [mulwideu32_param_0];
-; NOOPT-NEXT: cvt.u64.u32 %rd1, %r1;
-; NOOPT-NEXT: cvt.u64.u32 %rd2, %r2;
+; NOOPT-NEXT: ld.param.b32 %rd1, [mulwideu32_param_0];
+; NOOPT-NEXT: cvt.u32.u64 %r1, %rd1;
+; NOOPT-NEXT: ld.param.b32 %rd2, [mulwideu32_param_1];
+; NOOPT-NEXT: cvt.u32.u64 %r2, %rd2;
; NOOPT-NEXT: mul.lo.s64 %rd3, %rd1, %rd2;
; NOOPT-NEXT: st.param.b64 [func_retval0], %rd3;
; NOOPT-NEXT: ret;
@@ -300,8 +300,8 @@ define i64 @shl30(i32 %a) {
; NOOPT-NEXT: .reg .b64 %rd<3>;
; NOOPT-EMPTY:
; NOOPT-NEXT: // %bb.0:
-; NOOPT-NEXT: ld.param.b32 %r1, [shl30_param_0];
-; NOOPT-NEXT: cvt.s64.s32 %rd1, %r1;
+; NOOPT-NEXT: ld.param.s32 %rd1, [shl30_param_0];
+; NOOPT-NEXT: cvt.u32.u64 %r1, %rd1;
; NOOPT-NEXT: shl.b64 %rd2, %rd1, 30;
; NOOPT-NEXT: st.param.b64 [func_retval0], %rd2;
; NOOPT-NEXT: ret;
@@ -327,8 +327,8 @@ define i64 @shl31(i32 %a) {
; NOOPT-NEXT: .reg .b64 %rd<3>;
; NOOPT-EMPTY:
; NOOPT-NEXT: // %bb.0:
-; NOOPT-NEXT: ld.param.b32 %r1, [shl31_param_0];
-; NOOPT-NEXT: cvt.s64.s32 %rd1, %r1;
+; NOOPT-NEXT: ld.param.s32 %rd1, [shl31_param_0];
+; NOOPT-NEXT: cvt.u32.u64 %r1, %rd1;
; NOOPT-NEXT: shl.b64 %rd2, %rd1, 31;
; NOOPT-NEXT: st.param.b64 [func_retval0], %rd2;
; NOOPT-NEXT: ret;
diff --git a/llvm/test/CodeGen/NVPTX/truncate-free.ll b/llvm/test/CodeGen/NVPTX/truncate-free.ll
new file mode 100644
index 0000000000000..804c33951446c
--- /dev/null
+++ b/llvm/test/CodeGen/NVPTX/truncate-free.ll
@@ -0,0 +1,36 @@
+; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 6
+; RUN: llc -march=nvptx64 < %s | FileCheck %s
+
+define void @test_select_truncate_free(ptr %p, ptr %q, ptr %r) {
+; CHECK-LABEL: test_select_truncate_free(
+; CHECK: {
+; CHECK-NEXT: .reg .pred %p<2>;
+; CHECK-NEXT: .reg .b32 %r<2>;
+; CHECK-NEXT: .reg .b64 %rd<5>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.b64 %rd2, [test_select_truncate_free_param_0];
+; CHECK-NEXT: ld.b32 %rd3, [%rd2];
+; CHECK-NEXT: setp.eq.b64 %p1, %rd3, 0;
+; CHECK-NEXT: ld.param.b64 %rd4, [test_select_truncate_free_param_1];
+; CHECK-NEXT: st.b64 [%rd4], %rd3;
+; CHECK-NEXT: @%p1 bra $L__BB0_2;
+; CHECK-NEXT: // %bb.1: // %if
+; CHECK-NEXT: ld.param.b64 %rd1, [test_select_truncate_free_param_2];
+; CHECK-NEXT: cvt.u32.u64 %r1, %rd3;
+; CHECK-NEXT: st.b32 [%rd1], %r1;
+; CHECK-NEXT: $L__BB0_2: // %end
+; CHECK-NEXT: ret;
+ %a = load i32, ptr %p
+ %b = zext i32 %a to i64
+ store i64 %b, ptr %q
+ %c = icmp ne i32 %a, 0
+ br i1 %c, label %if, label %end
+
+if:
+ store i32 %a, ptr %r
+ br label %end
+
+end:
+ ret void
+}
``````````
</details>
https://github.com/llvm/llvm-project/pull/181484
More information about the llvm-commits
mailing list