[llvm] [NVPTX] extend type support for nvvm.{min, max, mulhi, sad} (PR #78385)
Alex MacLean via llvm-commits
llvm-commits at lists.llvm.org
Wed Jan 17 11:57:20 PST 2024
https://github.com/AlexMaclean updated https://github.com/llvm/llvm-project/pull/78385
>From 301acfc9f02aa43176a3500f15bc7da54113b094 Mon Sep 17 00:00:00 2001
From: Alex MacLean <amaclean at nvidia.com>
Date: Fri, 12 Jan 2024 23:07:10 +0000
Subject: [PATCH 1/3] [NVPTX] extend type support for nvvm.{min,max,mulhi,sad}
---
llvm/include/llvm/IR/IntrinsicsNVVM.td | 26 +++++-
llvm/lib/IR/AutoUpgrade.cpp | 16 ++--
llvm/lib/Target/NVPTX/NVPTXIntrinsics.td | 13 ++-
.../Assembler/auto_upgrade_nvvm_intrinsics.ll | 84 ++++++++++++-------
llvm/test/CodeGen/NVPTX/mulhi-intrins.ll | 50 +++++++++++
llvm/test/CodeGen/NVPTX/sad-intrins.ll | 50 +++++++++++
6 files changed, 197 insertions(+), 42 deletions(-)
create mode 100644 llvm/test/CodeGen/NVPTX/mulhi-intrins.ll
create mode 100644 llvm/test/CodeGen/NVPTX/sad-intrins.ll
diff --git a/llvm/include/llvm/IR/IntrinsicsNVVM.td b/llvm/include/llvm/IR/IntrinsicsNVVM.td
index 4665a1169ef4ee..d4cee2ff7d92dd 100644
--- a/llvm/include/llvm/IR/IntrinsicsNVVM.td
+++ b/llvm/include/llvm/IR/IntrinsicsNVVM.td
@@ -615,6 +615,13 @@ let TargetPrefix = "nvvm" in {
// Multiplication
//
+ def int_nvvm_mulhi_s : ClangBuiltin<"__nvvm_mulhi_s">,
+ DefaultAttrsIntrinsic<[llvm_i16_ty], [llvm_i16_ty, llvm_i16_ty],
+ [IntrNoMem, IntrSpeculatable, Commutative]>;
+ def int_nvvm_mulhi_us : ClangBuiltin<"__nvvm_mulhi_us">,
+ DefaultAttrsIntrinsic<[llvm_i16_ty], [llvm_i16_ty, llvm_i16_ty],
+ [IntrNoMem, IntrSpeculatable, Commutative]>;
+
def int_nvvm_mulhi_i : ClangBuiltin<"__nvvm_mulhi_i">,
DefaultAttrsIntrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty],
[IntrNoMem, IntrSpeculatable, Commutative]>;
@@ -730,12 +737,27 @@ let TargetPrefix = "nvvm" in {
// Sad
//
+ def int_nvvm_sad_s : ClangBuiltin<"__nvvm_sad_s">,
+ DefaultAttrsIntrinsic<[llvm_i16_ty], [llvm_i16_ty, llvm_i16_ty, llvm_i16_ty],
+ [IntrNoMem, Commutative, IntrSpeculatable]>;
+ def int_nvvm_sad_us : ClangBuiltin<"__nvvm_sad_us">,
+ DefaultAttrsIntrinsic<[llvm_i16_ty], [llvm_i16_ty, llvm_i16_ty, llvm_i16_ty],
+ [IntrNoMem, Commutative, IntrSpeculatable]>;
+
def int_nvvm_sad_i : ClangBuiltin<"__nvvm_sad_i">,
DefaultAttrsIntrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
- [IntrNoMem, Commutative]>;
+ [IntrNoMem, Commutative, IntrSpeculatable]>;
def int_nvvm_sad_ui : ClangBuiltin<"__nvvm_sad_ui">,
DefaultAttrsIntrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
- [IntrNoMem, Commutative]>;
+ [IntrNoMem, Commutative, IntrSpeculatable]>;
+
+ def int_nvvm_sad_ll : ClangBuiltin<"__nvvm_sad_ll">,
+ DefaultAttrsIntrinsic<[llvm_i64_ty], [llvm_i64_ty, llvm_i64_ty, llvm_i64_ty],
+ [IntrNoMem, Commutative, IntrSpeculatable]>;
+ def int_nvvm_sad_ull : ClangBuiltin<"__nvvm_sad_ull">,
+ DefaultAttrsIntrinsic<[llvm_i64_ty], [llvm_i64_ty, llvm_i64_ty, llvm_i64_ty],
+ [IntrNoMem, Commutative, IntrSpeculatable]>;
+
//
// Floor Ceil
diff --git a/llvm/lib/IR/AutoUpgrade.cpp b/llvm/lib/IR/AutoUpgrade.cpp
index 1a9e474911b39c..69ea591a8a7550 100644
--- a/llvm/lib/IR/AutoUpgrade.cpp
+++ b/llvm/lib/IR/AutoUpgrade.cpp
@@ -1209,7 +1209,8 @@ static bool UpgradeIntrinsicFunction1(Function *F, Function *&NewFn) {
Expand = true;
else if (Name.consume_front("max.") || Name.consume_front("min."))
// nvvm.{min,max}.{i,ii,ui,ull}
- Expand = Name == "i" || Name == "ll" || Name == "ui" || Name == "ull";
+ Expand = Name == "s" || Name == "i" || Name == "ll" || Name == "us" ||
+ Name == "ui" || Name == "ull";
else if (Name.consume_front("atomic.load.add."))
// nvvm.atomic.load.add.{f32.p,f64.p}
Expand = Name.starts_with("f32.p") || Name.starts_with("f64.p");
@@ -4132,19 +4133,20 @@ void llvm::UpgradeIntrinsicCall(CallBase *CI, Function *NewFn) {
Value *Val = CI->getArgOperand(1);
Rep = Builder.CreateAtomicRMW(AtomicRMWInst::FAdd, Ptr, Val, MaybeAlign(),
AtomicOrdering::SequentiallyConsistent);
- } else if (IsNVVM && (Name == "max.i" || Name == "max.ll" ||
- Name == "max.ui" || Name == "max.ull")) {
+ } else if (IsNVVM && (Name == "max.s" || Name == "max.i" || Name == "max.ll" ||
+ Name == "max.us" || Name == "max.ui" || Name == "max.ull")) {
Value *Arg0 = CI->getArgOperand(0);
Value *Arg1 = CI->getArgOperand(1);
- Value *Cmp = Name.ends_with(".ui") || Name.ends_with(".ull")
+ Value *Cmp = Name.starts_with("max.u")
? Builder.CreateICmpUGE(Arg0, Arg1, "max.cond")
: Builder.CreateICmpSGE(Arg0, Arg1, "max.cond");
Rep = Builder.CreateSelect(Cmp, Arg0, Arg1, "max");
- } else if (IsNVVM && (Name == "min.i" || Name == "min.ll" ||
- Name == "min.ui" || Name == "min.ull")) {
+ } else if (IsNVVM &&
+ (Name == "min.s" || Name == "min.i" || Name == "min.ll" ||
+ Name == "min.us" || Name == "min.ui" || Name == "min.ull")) {
Value *Arg0 = CI->getArgOperand(0);
Value *Arg1 = CI->getArgOperand(1);
- Value *Cmp = Name.ends_with(".ui") || Name.ends_with(".ull")
+ Value *Cmp = Name.starts_with("min.u")
? Builder.CreateICmpULE(Arg0, Arg1, "min.cond")
: Builder.CreateICmpSLE(Arg0, Arg1, "min.cond");
Rep = Builder.CreateSelect(Cmp, Arg0, Arg1, "min");
diff --git a/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td b/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
index c5dbe350e44472..5bb6389cf6a42b 100644
--- a/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
+++ b/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
@@ -770,11 +770,14 @@ defm INT_NVVM_FMAN : MIN_MAX<"max">;
// Multiplication
//
+def INT_NVVM_MULHI_S : F_MATH_2<"mul.hi.s16 \t$dst, $src0, $src1;", Int16Regs,
+ Int16Regs, Int16Regs, int_nvvm_mulhi_s>;
+def INT_NVVM_MULHI_US : F_MATH_2<"mul.hi.u16 \t$dst, $src0, $src1;", Int16Regs,
+ Int16Regs, Int16Regs, int_nvvm_mulhi_us>;
def INT_NVVM_MULHI_I : F_MATH_2<"mul.hi.s32 \t$dst, $src0, $src1;", Int32Regs,
Int32Regs, Int32Regs, int_nvvm_mulhi_i>;
def INT_NVVM_MULHI_UI : F_MATH_2<"mul.hi.u32 \t$dst, $src0, $src1;", Int32Regs,
Int32Regs, Int32Regs, int_nvvm_mulhi_ui>;
-
def INT_NVVM_MULHI_LL : F_MATH_2<"mul.hi.s64 \t$dst, $src0, $src1;", Int64Regs,
Int64Regs, Int64Regs, int_nvvm_mulhi_ll>;
def INT_NVVM_MULHI_ULL : F_MATH_2<"mul.hi.u64 \t$dst, $src0, $src1;", Int64Regs,
@@ -851,10 +854,18 @@ def INT_NVVM_DIV_RP_D : F_MATH_2<"div.rp.f64 \t$dst, $src0, $src1;",
// Sad
//
+def INT_NVVM_SAD_S : F_MATH_3<"sad.s16 \t$dst, $src0, $src1, $src2;",
+ Int16Regs, Int16Regs, Int16Regs, Int16Regs, int_nvvm_sad_s>;
+def INT_NVVM_SAD_US : F_MATH_3<"sad.u16 \t$dst, $src0, $src1, $src2;",
+ Int16Regs, Int16Regs, Int16Regs, Int16Regs, int_nvvm_sad_us>;
def INT_NVVM_SAD_I : F_MATH_3<"sad.s32 \t$dst, $src0, $src1, $src2;",
Int32Regs, Int32Regs, Int32Regs, Int32Regs, int_nvvm_sad_i>;
def INT_NVVM_SAD_UI : F_MATH_3<"sad.u32 \t$dst, $src0, $src1, $src2;",
Int32Regs, Int32Regs, Int32Regs, Int32Regs, int_nvvm_sad_ui>;
+def INT_NVVM_SAD_LL : F_MATH_3<"sad.s64 \t$dst, $src0, $src1, $src2;",
+ Int64Regs, Int64Regs, Int64Regs, Int64Regs, int_nvvm_sad_ll>;
+def INT_NVVM_SAD_ULL : F_MATH_3<"sad.u64 \t$dst, $src0, $src1, $src2;",
+ Int64Regs, Int64Regs, Int64Regs, Int64Regs, int_nvvm_sad_ull>;
//
// Floor Ceil
diff --git a/llvm/test/Assembler/auto_upgrade_nvvm_intrinsics.ll b/llvm/test/Assembler/auto_upgrade_nvvm_intrinsics.ll
index a295e5ea7ff004..1c11e1221fef34 100644
--- a/llvm/test/Assembler/auto_upgrade_nvvm_intrinsics.ll
+++ b/llvm/test/Assembler/auto_upgrade_nvvm_intrinsics.ll
@@ -13,12 +13,16 @@ declare float @llvm.nvvm.h2f(i16)
declare i32 @llvm.nvvm.abs.i(i32)
declare i64 @llvm.nvvm.abs.ll(i64)
+declare i16 @llvm.nvvm.max.s(i16, i16)
declare i32 @llvm.nvvm.max.i(i32, i32)
declare i64 @llvm.nvvm.max.ll(i64, i64)
+declare i16 @llvm.nvvm.max.us(i16, i16)
declare i32 @llvm.nvvm.max.ui(i32, i32)
declare i64 @llvm.nvvm.max.ull(i64, i64)
+declare i16 @llvm.nvvm.min.s(i16, i16)
declare i32 @llvm.nvvm.min.i(i32, i32)
declare i64 @llvm.nvvm.min.ll(i64, i64)
+declare i16 @llvm.nvvm.min.us(i16, i16)
declare i32 @llvm.nvvm.min.ui(i32, i32)
declare i64 @llvm.nvvm.min.ull(i64, i64)
@@ -65,38 +69,54 @@ define void @abs(i32 %a, i64 %b) {
}
; CHECK-LABEL: @min_max
-define void @min_max(i32 %a1, i32 %a2, i64 %b1, i64 %b2) {
-; CHECK: [[maxi:%[a-zA-Z0-9.]+]] = icmp sge i32 %a1, %a2
-; CHECK: select i1 [[maxi]], i32 %a1, i32 %a2
- %r1 = call i32 @llvm.nvvm.max.i(i32 %a1, i32 %a2)
-
-; CHECK: [[maxll:%[a-zA-Z0-9.]+]] = icmp sge i64 %b1, %b2
-; CHECK: select i1 [[maxll]], i64 %b1, i64 %b2
- %r2 = call i64 @llvm.nvvm.max.ll(i64 %b1, i64 %b2)
-
-; CHECK: [[maxui:%[a-zA-Z0-9.]+]] = icmp uge i32 %a1, %a2
-; CHECK: select i1 [[maxui]], i32 %a1, i32 %a2
- %r3 = call i32 @llvm.nvvm.max.ui(i32 %a1, i32 %a2)
-
-; CHECK: [[maxull:%[a-zA-Z0-9.]+]] = icmp uge i64 %b1, %b2
-; CHECK: select i1 [[maxull]], i64 %b1, i64 %b2
- %r4 = call i64 @llvm.nvvm.max.ull(i64 %b1, i64 %b2)
-
-; CHECK: [[mini:%[a-zA-Z0-9.]+]] = icmp sle i32 %a1, %a2
-; CHECK: select i1 [[mini]], i32 %a1, i32 %a2
- %r5 = call i32 @llvm.nvvm.min.i(i32 %a1, i32 %a2)
-
-; CHECK: [[minll:%[a-zA-Z0-9.]+]] = icmp sle i64 %b1, %b2
-; CHECK: select i1 [[minll]], i64 %b1, i64 %b2
- %r6 = call i64 @llvm.nvvm.min.ll(i64 %b1, i64 %b2)
-
-; CHECK: [[minui:%[a-zA-Z0-9.]+]] = icmp ule i32 %a1, %a2
-; CHECK: select i1 [[minui]], i32 %a1, i32 %a2
- %r7 = call i32 @llvm.nvvm.min.ui(i32 %a1, i32 %a2)
-
-; CHECK: [[minull:%[a-zA-Z0-9.]+]] = icmp ule i64 %b1, %b2
-; CHECK: select i1 [[minull]], i64 %b1, i64 %b2
- %r8 = call i64 @llvm.nvvm.min.ull(i64 %b1, i64 %b2)
+define void @min_max(i16 %a1, i16 %a2, i32 %b1, i32 %b2, i64 %c1, i64 %c2) {
+; CHECK: [[maxs:%[a-zA-Z0-9.]+]] = icmp sge i16 %a1, %a2
+; CHECK: select i1 [[maxs]], i16 %a1, i16 %a2
+ %r1 = call i16 @llvm.nvvm.max.s(i16 %a1, i16 %a2)
+
+; CHECK: [[maxi:%[a-zA-Z0-9.]+]] = icmp sge i32 %b1, %b2
+; CHECK: select i1 [[maxi]], i32 %b1, i32 %b2
+ %r2 = call i32 @llvm.nvvm.max.i(i32 %b1, i32 %b2)
+
+; CHECK: [[maxll:%[a-zA-Z0-9.]+]] = icmp sge i64 %c1, %c2
+; CHECK: select i1 [[maxll]], i64 %c1, i64 %c2
+ %r3 = call i64 @llvm.nvvm.max.ll(i64 %c1, i64 %c2)
+
+; CHECK: [[maxus:%[a-zA-Z0-9.]+]] = icmp uge i16 %a1, %a2
+; CHECK: select i1 [[maxus]], i16 %a1, i16 %a2
+ %r4 = call i16 @llvm.nvvm.max.us(i16 %a1, i16 %a2)
+
+; CHECK: [[maxui:%[a-zA-Z0-9.]+]] = icmp uge i32 %b1, %b2
+; CHECK: select i1 [[maxui]], i32 %b1, i32 %b2
+ %r5 = call i32 @llvm.nvvm.max.ui(i32 %b1, i32 %b2)
+
+; CHECK: [[maxull:%[a-zA-Z0-9.]+]] = icmp uge i64 %c1, %c2
+; CHECK: select i1 [[maxull]], i64 %c1, i64 %c2
+ %r6 = call i64 @llvm.nvvm.max.ull(i64 %c1, i64 %c2)
+
+; CHECK: [[mins:%[a-zA-Z0-9.]+]] = icmp sle i16 %a1, %a2
+; CHECK: select i1 [[mins]], i16 %a1, i16 %a2
+ %r7 = call i16 @llvm.nvvm.min.s(i16 %a1, i16 %a2)
+
+; CHECK: [[mini:%[a-zA-Z0-9.]+]] = icmp sle i32 %b1, %b2
+; CHECK: select i1 [[mini]], i32 %b1, i32 %b2
+ %r8 = call i32 @llvm.nvvm.min.i(i32 %b1, i32 %b2)
+
+; CHECK: [[minll:%[a-zA-Z0-9.]+]] = icmp sle i64 %c1, %c2
+; CHECK: select i1 [[minll]], i64 %c1, i64 %c2
+ %r9 = call i64 @llvm.nvvm.min.ll(i64 %c1, i64 %c2)
+
+; CHECK: [[minus:%[a-zA-Z0-9.]+]] = icmp ule i16 %a1, %a2
+; CHECK: select i1 [[minus]], i16 %a1, i16 %a2
+ %r10 = call i16 @llvm.nvvm.min.us(i16 %a1, i16 %a2)
+
+; CHECK: [[minui:%[a-zA-Z0-9.]+]] = icmp ule i32 %b1, %b2
+; CHECK: select i1 [[minui]], i32 %b1, i32 %b2
+ %r11 = call i32 @llvm.nvvm.min.ui(i32 %b1, i32 %b2)
+
+; CHECK: [[minull:%[a-zA-Z0-9.]+]] = icmp ule i64 %c1, %c2
+; CHECK: select i1 [[minull]], i64 %c1, i64 %c2
+ %r12 = call i64 @llvm.nvvm.min.ull(i64 %c1, i64 %c2)
ret void
}
diff --git a/llvm/test/CodeGen/NVPTX/mulhi-intrins.ll b/llvm/test/CodeGen/NVPTX/mulhi-intrins.ll
new file mode 100644
index 00000000000000..e28f3e24388294
--- /dev/null
+++ b/llvm/test/CodeGen/NVPTX/mulhi-intrins.ll
@@ -0,0 +1,50 @@
+; RUN: llc < %s -march=nvptx -mcpu=sm_50 | FileCheck %s
+
+; CHECK-LABEL: test_mulhi_i16
+; CHECK: mul.hi.s16
+define i16 @test_mulhi_i16(i16 %x, i16 %y) {
+ %1 = call i16 @llvm.nvvm.mulhi.s(i16 %x, i16 %y)
+ ret i16 %1
+}
+
+; CHECK-LABEL: test_mulhi_u16
+; CHECK: mul.hi.u16
+define i16 @test_mulhi_u16(i16 %x, i16 %y) {
+ %1 = call i16 @llvm.nvvm.mulhi.us(i16 %x, i16 %y)
+ ret i16 %1
+}
+
+; CHECK-LABEL: test_mulhi_i32
+; CHECK: mul.hi.s32
+define i32 @test_mulhi_i32(i32 %x, i32 %y) {
+ %1 = call i32 @llvm.nvvm.mulhi.i(i32 %x, i32 %y)
+ ret i32 %1
+}
+
+; CHECK-LABEL: test_mulhi_u32
+; CHECK: mul.hi.u32
+define i32 @test_mulhi_u32(i32 %x, i32 %y) {
+ %1 = call i32 @llvm.nvvm.mulhi.ui(i32 %x, i32 %y)
+ ret i32 %1
+}
+
+; CHECK-LABEL: test_mulhi_i64
+; CHECK: mul.hi.s64
+define i64 @test_mulhi_i64(i64 %x, i64 %y) {
+ %1 = call i64 @llvm.nvvm.mulhi.ll(i64 %x, i64 %y)
+ ret i64 %1
+}
+
+; CHECK-LABEL: test_mulhi_u64
+; CHECK: mul.hi.u64
+define i64 @test_mulhi_u64(i64 %x, i64 %y) {
+ %1 = call i64 @llvm.nvvm.mulhi.ull(i64 %x, i64 %y)
+ ret i64 %1
+}
+
+declare i16 @llvm.nvvm.mulhi.s(i16, i16)
+declare i16 @llvm.nvvm.mulhi.us(i16, i16)
+declare i32 @llvm.nvvm.mulhi.i(i32, i32)
+declare i32 @llvm.nvvm.mulhi.ui(i32, i32)
+declare i64 @llvm.nvvm.mulhi.ll(i64, i64)
+declare i64 @llvm.nvvm.mulhi.ull(i64, i64)
diff --git a/llvm/test/CodeGen/NVPTX/sad-intrins.ll b/llvm/test/CodeGen/NVPTX/sad-intrins.ll
new file mode 100644
index 00000000000000..72596c50a30244
--- /dev/null
+++ b/llvm/test/CodeGen/NVPTX/sad-intrins.ll
@@ -0,0 +1,50 @@
+; RUN: llc < %s -march=nvptx -mcpu=sm_50 | FileCheck %s
+
+; CHECK-LABEL: test_sad_i16
+; CHECK: sad.s16
+define i16 @test_sad_i16(i16 %x, i16 %y, i16 %z) {
+ %1 = call i16 @llvm.nvvm.sad.s(i16 %x, i16 %y, i16 %z)
+ ret i16 %1
+}
+
+; CHECK-LABEL: test_sad_u16
+; CHECK: sad.u16
+define i16 @test_sad_u16(i16 %x, i16 %y, i16 %z) {
+ %1 = call i16 @llvm.nvvm.sad.us(i16 %x, i16 %y, i16 %z)
+ ret i16 %1
+}
+
+; CHECK-LABEL: test_sad_i32
+; CHECK: sad.s32
+define i32 @test_sad_i32(i32 %x, i32 %y, i32 %z) {
+ %1 = call i32 @llvm.nvvm.sad.i(i32 %x, i32 %y, i32 %z)
+ ret i32 %1
+}
+
+; CHECK-LABEL: test_sad_u32
+; CHECK: sad.u32
+define i32 @test_sad_u32(i32 %x, i32 %y, i32 %z) {
+ %1 = call i32 @llvm.nvvm.sad.ui(i32 %x, i32 %y, i32 %z)
+ ret i32 %1
+}
+
+; CHECK-LABEL: test_sad_i64
+; CHECK: sad.s64
+define i64 @test_sad_i64(i64 %x, i64 %y, i64 %z) {
+ %1 = call i64 @llvm.nvvm.sad.ll(i64 %x, i64 %y, i64 %z)
+ ret i64 %1
+}
+
+; CHECK-LABEL: test_sad_u64
+; CHECK: sad.u64
+define i64 @test_sad_u64(i64 %x, i64 %y, i64 %z) {
+ %1 = call i64 @llvm.nvvm.sad.ull(i64 %x, i64 %y, i64 %z)
+ ret i64 %1
+}
+
+declare i16 @llvm.nvvm.sad.s(i16, i16, i16)
+declare i16 @llvm.nvvm.sad.us(i16, i16, i16)
+declare i32 @llvm.nvvm.sad.i(i32, i32, i32)
+declare i32 @llvm.nvvm.sad.ui(i32, i32, i32)
+declare i64 @llvm.nvvm.sad.ll(i64, i64, i64)
+declare i64 @llvm.nvvm.sad.ull(i64, i64, i64)
>From 2abf8aca2593cdfb6150458b2ff6858e3517a2ca Mon Sep 17 00:00:00 2001
From: Alex MacLean <amaclean at nvidia.com>
Date: Wed, 17 Jan 2024 03:12:54 +0000
Subject: [PATCH 2/3] fixup clang-formatting
---
llvm/lib/IR/AutoUpgrade.cpp | 5 +++--
1 file changed, 3 insertions(+), 2 deletions(-)
diff --git a/llvm/lib/IR/AutoUpgrade.cpp b/llvm/lib/IR/AutoUpgrade.cpp
index 69ea591a8a7550..f199fef9eb63d1 100644
--- a/llvm/lib/IR/AutoUpgrade.cpp
+++ b/llvm/lib/IR/AutoUpgrade.cpp
@@ -4133,8 +4133,9 @@ void llvm::UpgradeIntrinsicCall(CallBase *CI, Function *NewFn) {
Value *Val = CI->getArgOperand(1);
Rep = Builder.CreateAtomicRMW(AtomicRMWInst::FAdd, Ptr, Val, MaybeAlign(),
AtomicOrdering::SequentiallyConsistent);
- } else if (IsNVVM && (Name == "max.s" || Name == "max.i" || Name == "max.ll" ||
- Name == "max.us" || Name == "max.ui" || Name == "max.ull")) {
+ } else if (IsNVVM &&
+ (Name == "max.s" || Name == "max.i" || Name == "max.ll" ||
+ Name == "max.us" || Name == "max.ui" || Name == "max.ull")) {
Value *Arg0 = CI->getArgOperand(0);
Value *Arg1 = CI->getArgOperand(1);
Value *Cmp = Name.starts_with("max.u")
>From 8eccccfd57e976ddfa2fe700e3d60d878127bce6 Mon Sep 17 00:00:00 2001
From: Alex MacLean <amaclean at nvidia.com>
Date: Wed, 17 Jan 2024 19:57:04 +0000
Subject: [PATCH 3/3] address comments
---
llvm/lib/IR/AutoUpgrade.cpp | 16 ++---
llvm/test/CodeGen/NVPTX/mulhi-intrins.ll | 79 ++++++++++++++++++----
llvm/test/CodeGen/NVPTX/sad-intrins.ll | 85 ++++++++++++++++++++----
3 files changed, 146 insertions(+), 34 deletions(-)
diff --git a/llvm/lib/IR/AutoUpgrade.cpp b/llvm/lib/IR/AutoUpgrade.cpp
index f199fef9eb63d1..d2338e0eaa8d4e 100644
--- a/llvm/lib/IR/AutoUpgrade.cpp
+++ b/llvm/lib/IR/AutoUpgrade.cpp
@@ -4133,21 +4133,21 @@ void llvm::UpgradeIntrinsicCall(CallBase *CI, Function *NewFn) {
Value *Val = CI->getArgOperand(1);
Rep = Builder.CreateAtomicRMW(AtomicRMWInst::FAdd, Ptr, Val, MaybeAlign(),
AtomicOrdering::SequentiallyConsistent);
- } else if (IsNVVM &&
- (Name == "max.s" || Name == "max.i" || Name == "max.ll" ||
- Name == "max.us" || Name == "max.ui" || Name == "max.ull")) {
+ } else if (IsNVVM && Name.consume_front("max.") &&
+ (Name == "s" || Name == "i" || Name == "ll" || Name == "us" ||
+ Name == "ui" || Name == "ull")) {
Value *Arg0 = CI->getArgOperand(0);
Value *Arg1 = CI->getArgOperand(1);
- Value *Cmp = Name.starts_with("max.u")
+ Value *Cmp = Name.starts_with("u")
? Builder.CreateICmpUGE(Arg0, Arg1, "max.cond")
: Builder.CreateICmpSGE(Arg0, Arg1, "max.cond");
Rep = Builder.CreateSelect(Cmp, Arg0, Arg1, "max");
- } else if (IsNVVM &&
- (Name == "min.s" || Name == "min.i" || Name == "min.ll" ||
- Name == "min.us" || Name == "min.ui" || Name == "min.ull")) {
+ } else if (IsNVVM && Name.consume_front("min.") &&
+ (Name == "s" || Name == "i" || Name == "ll" || Name == "us" ||
+ Name == "ui" || Name == "ull")) {
Value *Arg0 = CI->getArgOperand(0);
Value *Arg1 = CI->getArgOperand(1);
- Value *Cmp = Name.starts_with("min.u")
+ Value *Cmp = Name.starts_with("u")
? Builder.CreateICmpULE(Arg0, Arg1, "min.cond")
: Builder.CreateICmpSLE(Arg0, Arg1, "min.cond");
Rep = Builder.CreateSelect(Cmp, Arg0, Arg1, "min");
diff --git a/llvm/test/CodeGen/NVPTX/mulhi-intrins.ll b/llvm/test/CodeGen/NVPTX/mulhi-intrins.ll
index e28f3e24388294..fc38f642624274 100644
--- a/llvm/test/CodeGen/NVPTX/mulhi-intrins.ll
+++ b/llvm/test/CodeGen/NVPTX/mulhi-intrins.ll
@@ -1,43 +1,96 @@
-; RUN: llc < %s -march=nvptx -mcpu=sm_50 | FileCheck %s
+; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 4
+; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_50 | FileCheck %s
-; CHECK-LABEL: test_mulhi_i16
-; CHECK: mul.hi.s16
define i16 @test_mulhi_i16(i16 %x, i16 %y) {
+; CHECK-LABEL: test_mulhi_i16(
+; CHECK: {
+; CHECK-NEXT: .reg .b16 %rs<4>;
+; CHECK-NEXT: .reg .b32 %r<2>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.u16 %rs1, [test_mulhi_i16_param_0];
+; CHECK-NEXT: ld.param.u16 %rs2, [test_mulhi_i16_param_1];
+; CHECK-NEXT: mul.hi.s16 %rs3, %rs1, %rs2;
+; CHECK-NEXT: cvt.u32.u16 %r1, %rs3;
+; CHECK-NEXT: st.param.b32 [func_retval0+0], %r1;
+; CHECK-NEXT: ret;
%1 = call i16 @llvm.nvvm.mulhi.s(i16 %x, i16 %y)
ret i16 %1
}
-; CHECK-LABEL: test_mulhi_u16
-; CHECK: mul.hi.u16
define i16 @test_mulhi_u16(i16 %x, i16 %y) {
+; CHECK-LABEL: test_mulhi_u16(
+; CHECK: {
+; CHECK-NEXT: .reg .b16 %rs<4>;
+; CHECK-NEXT: .reg .b32 %r<2>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.u16 %rs1, [test_mulhi_u16_param_0];
+; CHECK-NEXT: ld.param.u16 %rs2, [test_mulhi_u16_param_1];
+; CHECK-NEXT: mul.hi.u16 %rs3, %rs1, %rs2;
+; CHECK-NEXT: cvt.u32.u16 %r1, %rs3;
+; CHECK-NEXT: st.param.b32 [func_retval0+0], %r1;
+; CHECK-NEXT: ret;
%1 = call i16 @llvm.nvvm.mulhi.us(i16 %x, i16 %y)
ret i16 %1
}
-; CHECK-LABEL: test_mulhi_i32
-; CHECK: mul.hi.s32
define i32 @test_mulhi_i32(i32 %x, i32 %y) {
+; CHECK-LABEL: test_mulhi_i32(
+; CHECK: {
+; CHECK-NEXT: .reg .b32 %r<4>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.u32 %r1, [test_mulhi_i32_param_0];
+; CHECK-NEXT: ld.param.u32 %r2, [test_mulhi_i32_param_1];
+; CHECK-NEXT: mul.hi.s32 %r3, %r1, %r2;
+; CHECK-NEXT: st.param.b32 [func_retval0+0], %r3;
+; CHECK-NEXT: ret;
%1 = call i32 @llvm.nvvm.mulhi.i(i32 %x, i32 %y)
ret i32 %1
}
-; CHECK-LABEL: test_mulhi_u32
-; CHECK: mul.hi.u32
define i32 @test_mulhi_u32(i32 %x, i32 %y) {
+; CHECK-LABEL: test_mulhi_u32(
+; CHECK: {
+; CHECK-NEXT: .reg .b32 %r<4>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.u32 %r1, [test_mulhi_u32_param_0];
+; CHECK-NEXT: ld.param.u32 %r2, [test_mulhi_u32_param_1];
+; CHECK-NEXT: mul.hi.u32 %r3, %r1, %r2;
+; CHECK-NEXT: st.param.b32 [func_retval0+0], %r3;
+; CHECK-NEXT: ret;
%1 = call i32 @llvm.nvvm.mulhi.ui(i32 %x, i32 %y)
ret i32 %1
}
-; CHECK-LABEL: test_mulhi_i64
-; CHECK: mul.hi.s64
define i64 @test_mulhi_i64(i64 %x, i64 %y) {
+; CHECK-LABEL: test_mulhi_i64(
+; CHECK: {
+; CHECK-NEXT: .reg .b64 %rd<4>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.u64 %rd1, [test_mulhi_i64_param_0];
+; CHECK-NEXT: ld.param.u64 %rd2, [test_mulhi_i64_param_1];
+; CHECK-NEXT: mul.hi.s64 %rd3, %rd1, %rd2;
+; CHECK-NEXT: st.param.b64 [func_retval0+0], %rd3;
+; CHECK-NEXT: ret;
%1 = call i64 @llvm.nvvm.mulhi.ll(i64 %x, i64 %y)
ret i64 %1
}
-; CHECK-LABEL: test_mulhi_u64
-; CHECK: mul.hi.u64
define i64 @test_mulhi_u64(i64 %x, i64 %y) {
+; CHECK-LABEL: test_mulhi_u64(
+; CHECK: {
+; CHECK-NEXT: .reg .b64 %rd<4>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.u64 %rd1, [test_mulhi_u64_param_0];
+; CHECK-NEXT: ld.param.u64 %rd2, [test_mulhi_u64_param_1];
+; CHECK-NEXT: mul.hi.u64 %rd3, %rd1, %rd2;
+; CHECK-NEXT: st.param.b64 [func_retval0+0], %rd3;
+; CHECK-NEXT: ret;
%1 = call i64 @llvm.nvvm.mulhi.ull(i64 %x, i64 %y)
ret i64 %1
}
diff --git a/llvm/test/CodeGen/NVPTX/sad-intrins.ll b/llvm/test/CodeGen/NVPTX/sad-intrins.ll
index 72596c50a30244..1b6b9c6711769b 100644
--- a/llvm/test/CodeGen/NVPTX/sad-intrins.ll
+++ b/llvm/test/CodeGen/NVPTX/sad-intrins.ll
@@ -1,43 +1,102 @@
-; RUN: llc < %s -march=nvptx -mcpu=sm_50 | FileCheck %s
+; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 4
+; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_50 | FileCheck %s
-; CHECK-LABEL: test_sad_i16
-; CHECK: sad.s16
define i16 @test_sad_i16(i16 %x, i16 %y, i16 %z) {
+; CHECK-LABEL: test_sad_i16(
+; CHECK: {
+; CHECK-NEXT: .reg .b16 %rs<5>;
+; CHECK-NEXT: .reg .b32 %r<2>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.u16 %rs1, [test_sad_i16_param_0];
+; CHECK-NEXT: ld.param.u16 %rs2, [test_sad_i16_param_1];
+; CHECK-NEXT: ld.param.u16 %rs3, [test_sad_i16_param_2];
+; CHECK-NEXT: sad.s16 %rs4, %rs1, %rs2, %rs3;
+; CHECK-NEXT: cvt.u32.u16 %r1, %rs4;
+; CHECK-NEXT: st.param.b32 [func_retval0+0], %r1;
+; CHECK-NEXT: ret;
%1 = call i16 @llvm.nvvm.sad.s(i16 %x, i16 %y, i16 %z)
ret i16 %1
}
-; CHECK-LABEL: test_sad_u16
-; CHECK: sad.u16
define i16 @test_sad_u16(i16 %x, i16 %y, i16 %z) {
+; CHECK-LABEL: test_sad_u16(
+; CHECK: {
+; CHECK-NEXT: .reg .b16 %rs<5>;
+; CHECK-NEXT: .reg .b32 %r<2>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.u16 %rs1, [test_sad_u16_param_0];
+; CHECK-NEXT: ld.param.u16 %rs2, [test_sad_u16_param_1];
+; CHECK-NEXT: ld.param.u16 %rs3, [test_sad_u16_param_2];
+; CHECK-NEXT: sad.u16 %rs4, %rs1, %rs2, %rs3;
+; CHECK-NEXT: cvt.u32.u16 %r1, %rs4;
+; CHECK-NEXT: st.param.b32 [func_retval0+0], %r1;
+; CHECK-NEXT: ret;
%1 = call i16 @llvm.nvvm.sad.us(i16 %x, i16 %y, i16 %z)
ret i16 %1
}
-; CHECK-LABEL: test_sad_i32
-; CHECK: sad.s32
define i32 @test_sad_i32(i32 %x, i32 %y, i32 %z) {
+; CHECK-LABEL: test_sad_i32(
+; CHECK: {
+; CHECK-NEXT: .reg .b32 %r<5>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.u32 %r1, [test_sad_i32_param_0];
+; CHECK-NEXT: ld.param.u32 %r2, [test_sad_i32_param_1];
+; CHECK-NEXT: ld.param.u32 %r3, [test_sad_i32_param_2];
+; CHECK-NEXT: sad.s32 %r4, %r1, %r2, %r3;
+; CHECK-NEXT: st.param.b32 [func_retval0+0], %r4;
+; CHECK-NEXT: ret;
%1 = call i32 @llvm.nvvm.sad.i(i32 %x, i32 %y, i32 %z)
ret i32 %1
}
-; CHECK-LABEL: test_sad_u32
-; CHECK: sad.u32
define i32 @test_sad_u32(i32 %x, i32 %y, i32 %z) {
+; CHECK-LABEL: test_sad_u32(
+; CHECK: {
+; CHECK-NEXT: .reg .b32 %r<5>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.u32 %r1, [test_sad_u32_param_0];
+; CHECK-NEXT: ld.param.u32 %r2, [test_sad_u32_param_1];
+; CHECK-NEXT: ld.param.u32 %r3, [test_sad_u32_param_2];
+; CHECK-NEXT: sad.u32 %r4, %r1, %r2, %r3;
+; CHECK-NEXT: st.param.b32 [func_retval0+0], %r4;
+; CHECK-NEXT: ret;
%1 = call i32 @llvm.nvvm.sad.ui(i32 %x, i32 %y, i32 %z)
ret i32 %1
}
-; CHECK-LABEL: test_sad_i64
-; CHECK: sad.s64
define i64 @test_sad_i64(i64 %x, i64 %y, i64 %z) {
+; CHECK-LABEL: test_sad_i64(
+; CHECK: {
+; CHECK-NEXT: .reg .b64 %rd<5>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.u64 %rd1, [test_sad_i64_param_0];
+; CHECK-NEXT: ld.param.u64 %rd2, [test_sad_i64_param_1];
+; CHECK-NEXT: ld.param.u64 %rd3, [test_sad_i64_param_2];
+; CHECK-NEXT: sad.s64 %rd4, %rd1, %rd2, %rd3;
+; CHECK-NEXT: st.param.b64 [func_retval0+0], %rd4;
+; CHECK-NEXT: ret;
%1 = call i64 @llvm.nvvm.sad.ll(i64 %x, i64 %y, i64 %z)
ret i64 %1
}
-; CHECK-LABEL: test_sad_u64
-; CHECK: sad.u64
define i64 @test_sad_u64(i64 %x, i64 %y, i64 %z) {
+; CHECK-LABEL: test_sad_u64(
+; CHECK: {
+; CHECK-NEXT: .reg .b64 %rd<5>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.u64 %rd1, [test_sad_u64_param_0];
+; CHECK-NEXT: ld.param.u64 %rd2, [test_sad_u64_param_1];
+; CHECK-NEXT: ld.param.u64 %rd3, [test_sad_u64_param_2];
+; CHECK-NEXT: sad.u64 %rd4, %rd1, %rd2, %rd3;
+; CHECK-NEXT: st.param.b64 [func_retval0+0], %rd4;
+; CHECK-NEXT: ret;
%1 = call i64 @llvm.nvvm.sad.ull(i64 %x, i64 %y, i64 %z)
ret i64 %1
}
More information about the llvm-commits
mailing list