[clang] [llvm] [NVPTX] Auto-upgrade NVVM mulhi intrinsics to generic mulh (PR #226578)
Alex MacLean via llvm-commits
llvm-commits at lists.llvm.org
Fri Sep 25 12:55:53 PDT 2026
https://github.com/AlexMaclean created https://github.com/llvm/llvm-project/pull/226578
Remove the six llvm.nvvm.mulhi intrinsics and auto-upgrade them to llvm.smulh or llvm.umulh. Update the Clang builtins and NVPTX tests to use the generic intrinsics, which already lower through the existing multiply-high patterns.
Add auto-upgrade coverage for all six variants and Clang codegen coverage for the four existing builtins.
>From 2f77949a1cc92b35d3d960bb841bf15d3ecca0fa Mon Sep 17 00:00:00 2001
From: Alex Maclean <amaclean at nvidia.com>
Date: Fri, 25 Sep 2026 12:42:25 -0700
Subject: [PATCH] [NVPTX] Auto-upgrade NVVM mulhi intrinsics to generic mulh
Remove the six llvm.nvvm.mulhi intrinsics and auto-upgrade them to llvm.smulh or llvm.umulh. Update the NVPTX codegen tests and Clang builtins to use the generic intrinsics, which already lower through the existing multiply-high patterns.
Add auto-upgrade coverage for all six variants and Clang codegen coverage for the four existing builtins.
---
clang/lib/CodeGen/TargetBuiltins/NVPTX.cpp | 8 +++++
clang/test/CodeGen/builtins-nvptx.c | 24 +++++++++++++++
llvm/include/llvm/IR/IntrinsicsNVVM.td | 15 ++++------
llvm/lib/IR/AutoUpgrade.cpp | 2 ++
llvm/lib/Target/NVPTX/NVPTXIntrinsics.td | 7 -----
.../Assembler/auto_upgrade_nvvm_intrinsics.ll | 30 +++++++++++++++++++
llvm/test/CodeGen/NVPTX/mulhi-intrins.ll | 24 +++++++--------
7 files changed, 82 insertions(+), 28 deletions(-)
diff --git a/clang/lib/CodeGen/TargetBuiltins/NVPTX.cpp b/clang/lib/CodeGen/TargetBuiltins/NVPTX.cpp
index 76f6757326eca..ad5631c890c94 100644
--- a/clang/lib/CodeGen/TargetBuiltins/NVPTX.cpp
+++ b/clang/lib/CodeGen/TargetBuiltins/NVPTX.cpp
@@ -446,6 +446,14 @@ static Value *MakeFAdd(unsigned IntrinsicID, APFloat::roundingMode RM,
Value *CodeGenFunction::EmitNVPTXBuiltinExpr(unsigned BuiltinID,
const CallExpr *E) {
switch (BuiltinID) {
+ case NVPTX::BI__nvvm_mulhi_i:
+ case NVPTX::BI__nvvm_mulhi_ui:
+ case NVPTX::BI__nvvm_mulhi_ll:
+ case NVPTX::BI__nvvm_mulhi_ull:
+ return Builder.CreateBinaryIntrinsic(
+ E->getType()->hasSignedIntegerRepresentation() ? Intrinsic::smulh
+ : Intrinsic::umulh,
+ EmitScalarExpr(E->getArg(0)), EmitScalarExpr(E->getArg(1)));
case NVPTX::BI__nvvm_atom_add_gen_i:
case NVPTX::BI__nvvm_atom_add_gen_l:
case NVPTX::BI__nvvm_atom_add_gen_ll:
diff --git a/clang/test/CodeGen/builtins-nvptx.c b/clang/test/CodeGen/builtins-nvptx.c
index cd3eaff9d4c53..c50569308febf 100644
--- a/clang/test/CodeGen/builtins-nvptx.c
+++ b/clang/test/CodeGen/builtins-nvptx.c
@@ -72,6 +72,30 @@
#define __shared__ __attribute__((shared))
#define __constant__ __attribute__((constant))
+__device__ int mulhi_i(int a, int b) {
+ // CHECK-LABEL: define{{.*}} @{{.*}}mulhi_i
+ // CHECK: call i32 @llvm.smulh.i32(
+ return __nvvm_mulhi_i(a, b);
+}
+
+__device__ unsigned int mulhi_ui(unsigned int a, unsigned int b) {
+ // CHECK-LABEL: define{{.*}} @{{.*}}mulhi_ui
+ // CHECK: call i32 @llvm.umulh.i32(
+ return __nvvm_mulhi_ui(a, b);
+}
+
+__device__ long long mulhi_ll(long long a, long long b) {
+ // CHECK-LABEL: define{{.*}} @{{.*}}mulhi_ll
+ // CHECK: call i64 @llvm.smulh.i64(
+ return __nvvm_mulhi_ll(a, b);
+}
+
+__device__ unsigned long long mulhi_ull(unsigned long long a, unsigned long long b) {
+ // CHECK-LABEL: define{{.*}} @{{.*}}mulhi_ull
+ // CHECK: call i64 @llvm.umulh.i64(
+ return __nvvm_mulhi_ull(a, b);
+}
+
__device__ int read_tid() {
// CHECK: call i32 @llvm.nvvm.read.ptx.sreg.tid.x()
diff --git a/llvm/include/llvm/IR/IntrinsicsNVVM.td b/llvm/include/llvm/IR/IntrinsicsNVVM.td
index 002a8864d3959..308f2b1abc076 100644
--- a/llvm/include/llvm/IR/IntrinsicsNVVM.td
+++ b/llvm/include/llvm/IR/IntrinsicsNVVM.td
@@ -97,6 +97,12 @@
// * llvm.nvvm.clz.ll --> trunc i64 llvm.ctlz.i64(x) to i32
// * llvm.nvvm.popc.i --> llvm.ctpop.i32
// * llvm.nvvm.popc.ll --> trunc i64 llvm.ctpop.i64 to i32
+// * llvm.nvvm.mulhi.s --> llvm.smulh.i16
+// * llvm.nvvm.mulhi.i --> llvm.smulh.i32
+// * llvm.nvvm.mulhi.ll --> llvm.smulh.i64
+// * llvm.nvvm.mulhi.us --> llvm.umulh.i16
+// * llvm.nvvm.mulhi.ui --> llvm.umulh.i32
+// * llvm.nvvm.mulhi.ull --> llvm.umulh.i64
// * llvm.nvvm.abs.i --> llvm.abs.i32(x, true)
// * llvm.nvvm.abs.ll --> llvm.abs.i64(x, true)
// * llvm.nvvm.max.i --> select(x sge y, x, y)
@@ -1667,15 +1673,6 @@ let TargetPrefix = "nvvm" in {
let IntrProperties = [IntrNoMem, IntrSpeculatable, Commutative,
IntrNoCreateUndefOrPoison] in {
foreach sign = ["", "u"] in {
- def int_nvvm_mulhi_ # sign # s : NVVMBuiltin,
- DefaultAttrsIntrinsic<[llvm_i16_ty], [llvm_i16_ty, llvm_i16_ty]>;
-
- def int_nvvm_mulhi_ # sign # i : NVVMBuiltin,
- DefaultAttrsIntrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty]>;
-
- def int_nvvm_mulhi_ # sign # ll : NVVMBuiltin,
- DefaultAttrsIntrinsic<[llvm_i64_ty], [llvm_i64_ty, llvm_i64_ty]>;
-
def int_nvvm_mul24_ # sign # i : NVVMBuiltin,
DefaultAttrsIntrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty]>;
}
diff --git a/llvm/lib/IR/AutoUpgrade.cpp b/llvm/lib/IR/AutoUpgrade.cpp
index cb0e0690a261f..78eb59f9b2e47 100644
--- a/llvm/lib/IR/AutoUpgrade.cpp
+++ b/llvm/lib/IR/AutoUpgrade.cpp
@@ -2068,6 +2068,8 @@ static bool upgradeIntrinsicFunction1(Function *F, Function *&NewFn,
.Cases({"min.s", "min.i", "min.ll"}, Intrinsic::smin)
.Cases({"max.us", "max.ui", "max.ull"}, Intrinsic::umax)
.Cases({"min.us", "min.ui", "min.ull"}, Intrinsic::umin)
+ .Cases({"mulhi.s", "mulhi.i", "mulhi.ll"}, Intrinsic::smulh)
+ .Cases({"mulhi.us", "mulhi.ui", "mulhi.ull"}, Intrinsic::umulh)
.Default(Intrinsic::not_intrinsic);
if (IID != Intrinsic::not_intrinsic) {
NewFn = Intrinsic::getOrInsertDeclaration(F->getParent(), IID,
diff --git a/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td b/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
index d5108e1f32582..b622910978161 100644
--- a/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
+++ b/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
@@ -2174,13 +2174,6 @@ defm INT_NVVM_FMAN : MIN_MAX<"max">;
// Multiplication
//
-def : Pat<(int_nvvm_mulhi_s i16:$a, i16:$b), (MUL_HI_S16rr $a, $b)>;
-def : Pat<(int_nvvm_mulhi_us i16:$a, i16:$b), (MUL_HI_U16rr $a, $b)>;
-def : Pat<(int_nvvm_mulhi_i i32:$a, i32:$b), (MUL_HI_S32rr $a, $b)>;
-def : Pat<(int_nvvm_mulhi_ui i32:$a, i32:$b), (MUL_HI_U32rr $a, $b)>;
-def : Pat<(int_nvvm_mulhi_ll i64:$a, i64:$b), (MUL_HI_S64rr $a, $b)>;
-def : Pat<(int_nvvm_mulhi_ull i64:$a, i64:$b), (MUL_HI_U64rr $a, $b)>;
-
def INT_NVVM_MUL_RN_FTZ_F : F_MATH_2<"mul.rn.ftz.f32", B32, B32, B32, int_nvvm_mul_rn_ftz_f>;
def INT_NVVM_MUL_RN_F : F_MATH_2<"mul.rn.f32", B32, B32, B32, int_nvvm_mul_rn_f>;
def INT_NVVM_MUL_RZ_FTZ_F : F_MATH_2<"mul.rz.ftz.f32", B32, B32, B32, int_nvvm_mul_rz_ftz_f>;
diff --git a/llvm/test/Assembler/auto_upgrade_nvvm_intrinsics.ll b/llvm/test/Assembler/auto_upgrade_nvvm_intrinsics.ll
index 05afba0452bc8..4561c8b136e39 100644
--- a/llvm/test/Assembler/auto_upgrade_nvvm_intrinsics.ll
+++ b/llvm/test/Assembler/auto_upgrade_nvvm_intrinsics.ll
@@ -32,6 +32,13 @@ declare i16 @llvm.nvvm.min.us(i16, i16)
declare i32 @llvm.nvvm.min.ui(i32, i32)
declare i64 @llvm.nvvm.min.ull(i64, i64)
+declare i16 @llvm.nvvm.mulhi.s(i16, i16)
+declare i32 @llvm.nvvm.mulhi.i(i32, i32)
+declare i64 @llvm.nvvm.mulhi.ll(i64, i64)
+declare i16 @llvm.nvvm.mulhi.us(i16, i16)
+declare i32 @llvm.nvvm.mulhi.ui(i32, i32)
+declare i64 @llvm.nvvm.mulhi.ull(i64, i64)
+
declare i32 @llvm.nvvm.bitcast.f2i(float)
declare float @llvm.nvvm.bitcast.i2f(i32)
declare i64 @llvm.nvvm.bitcast.d2ll(double)
@@ -281,6 +288,29 @@ define void @min_max(i16 %a1, i16 %a2, i32 %b1, i32 %b2, i64 %c1, i64 %c2) {
ret void
}
+; CHECK-LABEL: @mulhi
+define void @mulhi(i16 %a1, i16 %a2, i32 %b1, i32 %b2, i64 %c1, i64 %c2) {
+; CHECK: %r1 = call i16 @llvm.smulh.i16(i16 %a1, i16 %a2)
+ %r1 = call i16 @llvm.nvvm.mulhi.s(i16 %a1, i16 %a2)
+
+; CHECK: %r2 = call i32 @llvm.smulh.i32(i32 %b1, i32 %b2)
+ %r2 = call i32 @llvm.nvvm.mulhi.i(i32 %b1, i32 %b2)
+
+; CHECK: %r3 = call i64 @llvm.smulh.i64(i64 %c1, i64 %c2)
+ %r3 = call i64 @llvm.nvvm.mulhi.ll(i64 %c1, i64 %c2)
+
+; CHECK: %r4 = call i16 @llvm.umulh.i16(i16 %a1, i16 %a2)
+ %r4 = call i16 @llvm.nvvm.mulhi.us(i16 %a1, i16 %a2)
+
+; CHECK: %r5 = call i32 @llvm.umulh.i32(i32 %b1, i32 %b2)
+ %r5 = call i32 @llvm.nvvm.mulhi.ui(i32 %b1, i32 %b2)
+
+; CHECK: %r6 = call i64 @llvm.umulh.i64(i64 %c1, i64 %c2)
+ %r6 = call i64 @llvm.nvvm.mulhi.ull(i64 %c1, i64 %c2)
+
+ ret void
+}
+
; CHECK-LABEL: @bitcast
define void @bitcast(i32 %a, i64 %b, float %c, double %d) {
; CHECK: bitcast float %c to i32
diff --git a/llvm/test/CodeGen/NVPTX/mulhi-intrins.ll b/llvm/test/CodeGen/NVPTX/mulhi-intrins.ll
index 8a88e1b26c7ff..65e1f99caa5f7 100644
--- a/llvm/test/CodeGen/NVPTX/mulhi-intrins.ll
+++ b/llvm/test/CodeGen/NVPTX/mulhi-intrins.ll
@@ -15,7 +15,7 @@ define i16 @test_mulhi_i16(i16 %x, i16 %y) {
; CHECK-NEXT: cvt.u32.u16 %r1, %rs3;
; CHECK-NEXT: st.param.b32 [func_retval0], %r1;
; CHECK-NEXT: ret;
- %1 = call i16 @llvm.nvvm.mulhi.s(i16 %x, i16 %y)
+ %1 = call i16 @llvm.smulh.i16(i16 %x, i16 %y)
ret i16 %1
}
@@ -32,7 +32,7 @@ define i16 @test_mulhi_u16(i16 %x, i16 %y) {
; CHECK-NEXT: cvt.u32.u16 %r1, %rs3;
; CHECK-NEXT: st.param.b32 [func_retval0], %r1;
; CHECK-NEXT: ret;
- %1 = call i16 @llvm.nvvm.mulhi.us(i16 %x, i16 %y)
+ %1 = call i16 @llvm.umulh.i16(i16 %x, i16 %y)
ret i16 %1
}
@@ -47,7 +47,7 @@ define i32 @test_mulhi_i32(i32 %x, i32 %y) {
; CHECK-NEXT: mul.hi.s32 %r3, %r1, %r2;
; CHECK-NEXT: st.param.b32 [func_retval0], %r3;
; CHECK-NEXT: ret;
- %1 = call i32 @llvm.nvvm.mulhi.i(i32 %x, i32 %y)
+ %1 = call i32 @llvm.smulh.i32(i32 %x, i32 %y)
ret i32 %1
}
@@ -62,7 +62,7 @@ define i32 @test_mulhi_u32(i32 %x, i32 %y) {
; CHECK-NEXT: mul.hi.u32 %r3, %r1, %r2;
; CHECK-NEXT: st.param.b32 [func_retval0], %r3;
; CHECK-NEXT: ret;
- %1 = call i32 @llvm.nvvm.mulhi.ui(i32 %x, i32 %y)
+ %1 = call i32 @llvm.umulh.i32(i32 %x, i32 %y)
ret i32 %1
}
@@ -77,7 +77,7 @@ define i64 @test_mulhi_i64(i64 %x, i64 %y) {
; CHECK-NEXT: mul.hi.s64 %rd3, %rd1, %rd2;
; CHECK-NEXT: st.param.b64 [func_retval0], %rd3;
; CHECK-NEXT: ret;
- %1 = call i64 @llvm.nvvm.mulhi.ll(i64 %x, i64 %y)
+ %1 = call i64 @llvm.smulh.i64(i64 %x, i64 %y)
ret i64 %1
}
@@ -92,13 +92,13 @@ define i64 @test_mulhi_u64(i64 %x, i64 %y) {
; CHECK-NEXT: mul.hi.u64 %rd3, %rd1, %rd2;
; CHECK-NEXT: st.param.b64 [func_retval0], %rd3;
; CHECK-NEXT: ret;
- %1 = call i64 @llvm.nvvm.mulhi.ull(i64 %x, i64 %y)
+ %1 = call i64 @llvm.umulh.i64(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)
+declare i16 @llvm.smulh.i16(i16, i16)
+declare i16 @llvm.umulh.i16(i16, i16)
+declare i32 @llvm.smulh.i32(i32, i32)
+declare i32 @llvm.umulh.i32(i32, i32)
+declare i64 @llvm.smulh.i64(i64, i64)
+declare i64 @llvm.umulh.i64(i64, i64)
More information about the llvm-commits
mailing list