Skip to content

Commit 430a40d

Browse files
AlexMacleanAlexMaclean
authored
[NVPTX] extend type support for nvvm.{min,max,mulhi,sad} (llvm#78385)
Ensure intrinsics and auto-upgrades support i16, i32, and i64 for for `nvvm.{min,max,mulhi,sad}` - `nvvm.min` and `nvvm.max`: These are auto-upgraded to `select` instructions but it is still nice to support the 16 bit variants just in case any generators of IR are still trying to use these intrinsics. - `nvvm.sad` added both the 16 and 64 bit variants, also marked this instruction as speculateble. These directly correspond to the PTX `sad.{u16,s16,u64,s64}` instructions. - `nvvm.mulhi` added the 16 bit variants. These directly correspond to the PTX `mul.hi.{s,u}16` instructions.
1 parent f2b5a31 commit 430a40d

File tree

6 files changed

+312
-42
lines changed

6 files changed

+312
-42
lines changed

‎llvm/include/llvm/IR/IntrinsicsNVVM.td

Lines changed: 24 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -615,6 +615,13 @@ let TargetPrefix = "nvvm" in {
615615
// Multiplication
616616
//
617617

618+
def int_nvvm_mulhi_s : ClangBuiltin<"__nvvm_mulhi_s">,
619+
DefaultAttrsIntrinsic<[llvm_i16_ty], [llvm_i16_ty, llvm_i16_ty],
620+
[IntrNoMem, IntrSpeculatable, Commutative]>;
621+
def int_nvvm_mulhi_us : ClangBuiltin<"__nvvm_mulhi_us">,
622+
DefaultAttrsIntrinsic<[llvm_i16_ty], [llvm_i16_ty, llvm_i16_ty],
623+
[IntrNoMem, IntrSpeculatable, Commutative]>;
624+
618625
def int_nvvm_mulhi_i : ClangBuiltin<"__nvvm_mulhi_i">,
619626
DefaultAttrsIntrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty],
620627
[IntrNoMem, IntrSpeculatable, Commutative]>;
@@ -730,12 +737,27 @@ let TargetPrefix = "nvvm" in {
730737
// Sad
731738
//
732739

740+
def int_nvvm_sad_s : ClangBuiltin<"__nvvm_sad_s">,
741+
DefaultAttrsIntrinsic<[llvm_i16_ty], [llvm_i16_ty, llvm_i16_ty, llvm_i16_ty],
742+
[IntrNoMem, Commutative, IntrSpeculatable]>;
743+
def int_nvvm_sad_us : ClangBuiltin<"__nvvm_sad_us">,
744+
DefaultAttrsIntrinsic<[llvm_i16_ty], [llvm_i16_ty, llvm_i16_ty, llvm_i16_ty],
745+
[IntrNoMem, Commutative, IntrSpeculatable]>;
746+
733747
def int_nvvm_sad_i : ClangBuiltin<"__nvvm_sad_i">,
734748
DefaultAttrsIntrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
735-
[IntrNoMem, Commutative]>;
749+
[IntrNoMem, Commutative, IntrSpeculatable]>;
736750
def int_nvvm_sad_ui : ClangBuiltin<"__nvvm_sad_ui">,
737751
DefaultAttrsIntrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
738-
[IntrNoMem, Commutative]>;
752+
[IntrNoMem, Commutative, IntrSpeculatable]>;
753+
754+
def int_nvvm_sad_ll : ClangBuiltin<"__nvvm_sad_ll">,
755+
DefaultAttrsIntrinsic<[llvm_i64_ty], [llvm_i64_ty, llvm_i64_ty, llvm_i64_ty],
756+
[IntrNoMem, Commutative, IntrSpeculatable]>;
757+
def int_nvvm_sad_ull : ClangBuiltin<"__nvvm_sad_ull">,
758+
DefaultAttrsIntrinsic<[llvm_i64_ty], [llvm_i64_ty, llvm_i64_ty, llvm_i64_ty],
759+
[IntrNoMem, Commutative, IntrSpeculatable]>;
760+
739761

740762
//
741763
// Floor Ceil

‎llvm/lib/IR/AutoUpgrade.cpp

Lines changed: 10 additions & 7 deletions
Original file line numberDiff line numberDiff line change
@@ -1209,7 +1209,8 @@ static bool UpgradeIntrinsicFunction1(Function *F, Function *&NewFn) {
12091209
Expand = true;
12101210
else if (Name.consume_front("max.") || Name.consume_front("min."))
12111211
// nvvm.{min,max}.{i,ii,ui,ull}
1212-
Expand = Name == "i" || Name == "ll" || Name == "ui" || Name == "ull";
1212+
Expand = Name == "s" || Name == "i" || Name == "ll" || Name == "us" ||
1213+
Name == "ui" || Name == "ull";
12131214
else if (Name.consume_front("atomic.load.add."))
12141215
// nvvm.atomic.load.add.{f32.p,f64.p}
12151216
Expand = Name.starts_with("f32.p") || Name.starts_with("f64.p");
@@ -4132,19 +4133,21 @@ void llvm::UpgradeIntrinsicCall(CallBase *CI, Function *NewFn) {
41324133
Value *Val = CI->getArgOperand(1);
41334134
Rep = Builder.CreateAtomicRMW(AtomicRMWInst::FAdd, Ptr, Val, MaybeAlign(),
41344135
AtomicOrdering::SequentiallyConsistent);
4135-
} else if (IsNVVM && (Name == "max.i" || Name == "max.ll" ||
4136-
Name == "max.ui" || Name == "max.ull")) {
4136+
} else if (IsNVVM && Name.consume_front("max.") &&
4137+
(Name == "s" || Name == "i" || Name == "ll" || Name == "us" ||
4138+
Name == "ui" || Name == "ull")) {
41374139
Value *Arg0 = CI->getArgOperand(0);
41384140
Value *Arg1 = CI->getArgOperand(1);
4139-
Value *Cmp = Name.ends_with(".ui") || Name.ends_with(".ull")
4141+
Value *Cmp = Name.starts_with("u")
41404142
? Builder.CreateICmpUGE(Arg0, Arg1, "max.cond")
41414143
: Builder.CreateICmpSGE(Arg0, Arg1, "max.cond");
41424144
Rep = Builder.CreateSelect(Cmp, Arg0, Arg1, "max");
4143-
} else if (IsNVVM && (Name == "min.i" || Name == "min.ll" ||
4144-
Name == "min.ui" || Name == "min.ull")) {
4145+
} else if (IsNVVM && Name.consume_front("min.") &&
4146+
(Name == "s" || Name == "i" || Name == "ll" || Name == "us" ||
4147+
Name == "ui" || Name == "ull")) {
41454148
Value *Arg0 = CI->getArgOperand(0);
41464149
Value *Arg1 = CI->getArgOperand(1);
4147-
Value *Cmp = Name.ends_with(".ui") || Name.ends_with(".ull")
4150+
Value *Cmp = Name.starts_with("u")
41484151
? Builder.CreateICmpULE(Arg0, Arg1, "min.cond")
41494152
: Builder.CreateICmpSLE(Arg0, Arg1, "min.cond");
41504153
Rep = Builder.CreateSelect(Cmp, Arg0, Arg1, "min");

‎llvm/lib/Target/NVPTX/NVPTXIntrinsics.td

Lines changed: 12 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -770,11 +770,14 @@ defm INT_NVVM_FMAN : MIN_MAX<"max">;
770770
// Multiplication
771771
//
772772

773+
def INT_NVVM_MULHI_S : F_MATH_2<"mul.hi.s16 \t$dst, $src0, $src1;", Int16Regs,
774+
Int16Regs, Int16Regs, int_nvvm_mulhi_s>;
775+
def INT_NVVM_MULHI_US : F_MATH_2<"mul.hi.u16 \t$dst, $src0, $src1;", Int16Regs,
776+
Int16Regs, Int16Regs, int_nvvm_mulhi_us>;
773777
def INT_NVVM_MULHI_I : F_MATH_2<"mul.hi.s32 \t$dst, $src0, $src1;", Int32Regs,
774778
Int32Regs, Int32Regs, int_nvvm_mulhi_i>;
775779
def INT_NVVM_MULHI_UI : F_MATH_2<"mul.hi.u32 \t$dst, $src0, $src1;", Int32Regs,
776780
Int32Regs, Int32Regs, int_nvvm_mulhi_ui>;
777-
778781
def INT_NVVM_MULHI_LL : F_MATH_2<"mul.hi.s64 \t$dst, $src0, $src1;", Int64Regs,
779782
Int64Regs, Int64Regs, int_nvvm_mulhi_ll>;
780783
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;",
851854
// Sad
852855
//
853856

857+
def INT_NVVM_SAD_S : F_MATH_3<"sad.s16 \t$dst, $src0, $src1, $src2;",
858+
Int16Regs, Int16Regs, Int16Regs, Int16Regs, int_nvvm_sad_s>;
859+
def INT_NVVM_SAD_US : F_MATH_3<"sad.u16 \t$dst, $src0, $src1, $src2;",
860+
Int16Regs, Int16Regs, Int16Regs, Int16Regs, int_nvvm_sad_us>;
854861
def INT_NVVM_SAD_I : F_MATH_3<"sad.s32 \t$dst, $src0, $src1, $src2;",
855862
Int32Regs, Int32Regs, Int32Regs, Int32Regs, int_nvvm_sad_i>;
856863
def INT_NVVM_SAD_UI : F_MATH_3<"sad.u32 \t$dst, $src0, $src1, $src2;",
857864
Int32Regs, Int32Regs, Int32Regs, Int32Regs, int_nvvm_sad_ui>;
865+
def INT_NVVM_SAD_LL : F_MATH_3<"sad.s64 \t$dst, $src0, $src1, $src2;",
866+
Int64Regs, Int64Regs, Int64Regs, Int64Regs, int_nvvm_sad_ll>;
867+
def INT_NVVM_SAD_ULL : F_MATH_3<"sad.u64 \t$dst, $src0, $src1, $src2;",
868+
Int64Regs, Int64Regs, Int64Regs, Int64Regs, int_nvvm_sad_ull>;
858869

859870
//
860871
// Floor Ceil

‎llvm/test/Assembler/auto_upgrade_nvvm_intrinsics.ll

Lines changed: 52 additions & 32 deletions
Original file line numberDiff line numberDiff line change
@@ -13,12 +13,16 @@ declare float @llvm.nvvm.h2f(i16)
1313
declare i32 @llvm.nvvm.abs.i(i32)
1414
declare i64 @llvm.nvvm.abs.ll(i64)
1515

16+
declare i16 @llvm.nvvm.max.s(i16, i16)
1617
declare i32 @llvm.nvvm.max.i(i32, i32)
1718
declare i64 @llvm.nvvm.max.ll(i64, i64)
19+
declare i16 @llvm.nvvm.max.us(i16, i16)
1820
declare i32 @llvm.nvvm.max.ui(i32, i32)
1921
declare i64 @llvm.nvvm.max.ull(i64, i64)
22+
declare i16 @llvm.nvvm.min.s(i16, i16)
2023
declare i32 @llvm.nvvm.min.i(i32, i32)
2124
declare i64 @llvm.nvvm.min.ll(i64, i64)
25+
declare i16 @llvm.nvvm.min.us(i16, i16)
2226
declare i32 @llvm.nvvm.min.ui(i32, i32)
2327
declare i64 @llvm.nvvm.min.ull(i64, i64)
2428

@@ -65,38 +69,54 @@ define void @abs(i32 %a, i64 %b) {
6569
}
6670

6771
; CHECK-LABEL: @min_max
68-
define void @min_max(i32 %a1, i32 %a2, i64 %b1, i64 %b2) {
69-
; CHECK: [[maxi:%[a-zA-Z0-9.]+]] = icmp sge i32 %a1, %a2
70-
; CHECK: select i1 [[maxi]], i32 %a1, i32 %a2
71-
%r1 = call i32 @llvm.nvvm.max.i(i32 %a1, i32 %a2)
72-
73-
; CHECK: [[maxll:%[a-zA-Z0-9.]+]] = icmp sge i64 %b1, %b2
74-
; CHECK: select i1 [[maxll]], i64 %b1, i64 %b2
75-
%r2 = call i64 @llvm.nvvm.max.ll(i64 %b1, i64 %b2)
76-
77-
; CHECK: [[maxui:%[a-zA-Z0-9.]+]] = icmp uge i32 %a1, %a2
78-
; CHECK: select i1 [[maxui]], i32 %a1, i32 %a2
79-
%r3 = call i32 @llvm.nvvm.max.ui(i32 %a1, i32 %a2)
80-
81-
; CHECK: [[maxull:%[a-zA-Z0-9.]+]] = icmp uge i64 %b1, %b2
82-
; CHECK: select i1 [[maxull]], i64 %b1, i64 %b2
83-
%r4 = call i64 @llvm.nvvm.max.ull(i64 %b1, i64 %b2)
84-
85-
; CHECK: [[mini:%[a-zA-Z0-9.]+]] = icmp sle i32 %a1, %a2
86-
; CHECK: select i1 [[mini]], i32 %a1, i32 %a2
87-
%r5 = call i32 @llvm.nvvm.min.i(i32 %a1, i32 %a2)
88-
89-
; CHECK: [[minll:%[a-zA-Z0-9.]+]] = icmp sle i64 %b1, %b2
90-
; CHECK: select i1 [[minll]], i64 %b1, i64 %b2
91-
%r6 = call i64 @llvm.nvvm.min.ll(i64 %b1, i64 %b2)
92-
93-
; CHECK: [[minui:%[a-zA-Z0-9.]+]] = icmp ule i32 %a1, %a2
94-
; CHECK: select i1 [[minui]], i32 %a1, i32 %a2
95-
%r7 = call i32 @llvm.nvvm.min.ui(i32 %a1, i32 %a2)
96-
97-
; CHECK: [[minull:%[a-zA-Z0-9.]+]] = icmp ule i64 %b1, %b2
98-
; CHECK: select i1 [[minull]], i64 %b1, i64 %b2
99-
%r8 = call i64 @llvm.nvvm.min.ull(i64 %b1, i64 %b2)
72+
define void @min_max(i16 %a1, i16 %a2, i32 %b1, i32 %b2, i64 %c1, i64 %c2) {
73+
; CHECK: [[maxs:%[a-zA-Z0-9.]+]] = icmp sge i16 %a1, %a2
74+
; CHECK: select i1 [[maxs]], i16 %a1, i16 %a2
75+
%r1 = call i16 @llvm.nvvm.max.s(i16 %a1, i16 %a2)
76+
77+
; CHECK: [[maxi:%[a-zA-Z0-9.]+]] = icmp sge i32 %b1, %b2
78+
; CHECK: select i1 [[maxi]], i32 %b1, i32 %b2
79+
%r2 = call i32 @llvm.nvvm.max.i(i32 %b1, i32 %b2)
80+
81+
; CHECK: [[maxll:%[a-zA-Z0-9.]+]] = icmp sge i64 %c1, %c2
82+
; CHECK: select i1 [[maxll]], i64 %c1, i64 %c2
83+
%r3 = call i64 @llvm.nvvm.max.ll(i64 %c1, i64 %c2)
84+
85+
; CHECK: [[maxus:%[a-zA-Z0-9.]+]] = icmp uge i16 %a1, %a2
86+
; CHECK: select i1 [[maxus]], i16 %a1, i16 %a2
87+
%r4 = call i16 @llvm.nvvm.max.us(i16 %a1, i16 %a2)
88+
89+
; CHECK: [[maxui:%[a-zA-Z0-9.]+]] = icmp uge i32 %b1, %b2
90+
; CHECK: select i1 [[maxui]], i32 %b1, i32 %b2
91+
%r5 = call i32 @llvm.nvvm.max.ui(i32 %b1, i32 %b2)
92+
93+
; CHECK: [[maxull:%[a-zA-Z0-9.]+]] = icmp uge i64 %c1, %c2
94+
; CHECK: select i1 [[maxull]], i64 %c1, i64 %c2
95+
%r6 = call i64 @llvm.nvvm.max.ull(i64 %c1, i64 %c2)
96+
97+
; CHECK: [[mins:%[a-zA-Z0-9.]+]] = icmp sle i16 %a1, %a2
98+
; CHECK: select i1 [[mins]], i16 %a1, i16 %a2
99+
%r7 = call i16 @llvm.nvvm.min.s(i16 %a1, i16 %a2)
100+
101+
; CHECK: [[mini:%[a-zA-Z0-9.]+]] = icmp sle i32 %b1, %b2
102+
; CHECK: select i1 [[mini]], i32 %b1, i32 %b2
103+
%r8 = call i32 @llvm.nvvm.min.i(i32 %b1, i32 %b2)
104+
105+
; CHECK: [[minll:%[a-zA-Z0-9.]+]] = icmp sle i64 %c1, %c2
106+
; CHECK: select i1 [[minll]], i64 %c1, i64 %c2
107+
%r9 = call i64 @llvm.nvvm.min.ll(i64 %c1, i64 %c2)
108+
109+
; CHECK: [[minus:%[a-zA-Z0-9.]+]] = icmp ule i16 %a1, %a2
110+
; CHECK: select i1 [[minus]], i16 %a1, i16 %a2
111+
%r10 = call i16 @llvm.nvvm.min.us(i16 %a1, i16 %a2)
112+
113+
; CHECK: [[minui:%[a-zA-Z0-9.]+]] = icmp ule i32 %b1, %b2
114+
; CHECK: select i1 [[minui]], i32 %b1, i32 %b2
115+
%r11 = call i32 @llvm.nvvm.min.ui(i32 %b1, i32 %b2)
116+
117+
; CHECK: [[minull:%[a-zA-Z0-9.]+]] = icmp ule i64 %c1, %c2
118+
; CHECK: select i1 [[minull]], i64 %c1, i64 %c2
119+
%r12 = call i64 @llvm.nvvm.min.ull(i64 %c1, i64 %c2)
100120

101121
ret void
102122
}

‎llvm/test/CodeGen/NVPTX/mulhi-intrins.ll

Lines changed: 104 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,104 @@
1+
; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 4
2+
; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_50 | FileCheck %s
3+
; RUN: %if ptxas %{ llc < %s -mtriple=nvptx64 -mcpu=sm_50 | %ptxas-verify %}
4+
5+
define i16 @test_mulhi_i16(i16 %x, i16 %y) {
6+
; CHECK-LABEL: test_mulhi_i16(
7+
; CHECK: {
8+
; CHECK-NEXT: .reg .b16 %rs<4>;
9+
; CHECK-NEXT: .reg .b32 %r<2>;
10+
; CHECK-EMPTY:
11+
; CHECK-NEXT: // %bb.0:
12+
; CHECK-NEXT: ld.param.u16 %rs1, [test_mulhi_i16_param_0];
13+
; CHECK-NEXT: ld.param.u16 %rs2, [test_mulhi_i16_param_1];
14+
; CHECK-NEXT: mul.hi.s16 %rs3, %rs1, %rs2;
15+
; CHECK-NEXT: cvt.u32.u16 %r1, %rs3;
16+
; CHECK-NEXT: st.param.b32 [func_retval0+0], %r1;
17+
; CHECK-NEXT: ret;
18+
%1 = call i16 @llvm.nvvm.mulhi.s(i16 %x, i16 %y)
19+
ret i16 %1
20+
}
21+
22+
define i16 @test_mulhi_u16(i16 %x, i16 %y) {
23+
; CHECK-LABEL: test_mulhi_u16(
24+
; CHECK: {
25+
; CHECK-NEXT: .reg .b16 %rs<4>;
26+
; CHECK-NEXT: .reg .b32 %r<2>;
27+
; CHECK-EMPTY:
28+
; CHECK-NEXT: // %bb.0:
29+
; CHECK-NEXT: ld.param.u16 %rs1, [test_mulhi_u16_param_0];
30+
; CHECK-NEXT: ld.param.u16 %rs2, [test_mulhi_u16_param_1];
31+
; CHECK-NEXT: mul.hi.u16 %rs3, %rs1, %rs2;
32+
; CHECK-NEXT: cvt.u32.u16 %r1, %rs3;
33+
; CHECK-NEXT: st.param.b32 [func_retval0+0], %r1;
34+
; CHECK-NEXT: ret;
35+
%1 = call i16 @llvm.nvvm.mulhi.us(i16 %x, i16 %y)
36+
ret i16 %1
37+
}
38+
39+
define i32 @test_mulhi_i32(i32 %x, i32 %y) {
40+
; CHECK-LABEL: test_mulhi_i32(
41+
; CHECK: {
42+
; CHECK-NEXT: .reg .b32 %r<4>;
43+
; CHECK-EMPTY:
44+
; CHECK-NEXT: // %bb.0:
45+
; CHECK-NEXT: ld.param.u32 %r1, [test_mulhi_i32_param_0];
46+
; CHECK-NEXT: ld.param.u32 %r2, [test_mulhi_i32_param_1];
47+
; CHECK-NEXT: mul.hi.s32 %r3, %r1, %r2;
48+
; CHECK-NEXT: st.param.b32 [func_retval0+0], %r3;
49+
; CHECK-NEXT: ret;
50+
%1 = call i32 @llvm.nvvm.mulhi.i(i32 %x, i32 %y)
51+
ret i32 %1
52+
}
53+
54+
define i32 @test_mulhi_u32(i32 %x, i32 %y) {
55+
; CHECK-LABEL: test_mulhi_u32(
56+
; CHECK: {
57+
; CHECK-NEXT: .reg .b32 %r<4>;
58+
; CHECK-EMPTY:
59+
; CHECK-NEXT: // %bb.0:
60+
; CHECK-NEXT: ld.param.u32 %r1, [test_mulhi_u32_param_0];
61+
; CHECK-NEXT: ld.param.u32 %r2, [test_mulhi_u32_param_1];
62+
; CHECK-NEXT: mul.hi.u32 %r3, %r1, %r2;
63+
; CHECK-NEXT: st.param.b32 [func_retval0+0], %r3;
64+
; CHECK-NEXT: ret;
65+
%1 = call i32 @llvm.nvvm.mulhi.ui(i32 %x, i32 %y)
66+
ret i32 %1
67+
}
68+
69+
define i64 @test_mulhi_i64(i64 %x, i64 %y) {
70+
; CHECK-LABEL: test_mulhi_i64(
71+
; CHECK: {
72+
; CHECK-NEXT: .reg .b64 %rd<4>;
73+
; CHECK-EMPTY:
74+
; CHECK-NEXT: // %bb.0:
75+
; CHECK-NEXT: ld.param.u64 %rd1, [test_mulhi_i64_param_0];
76+
; CHECK-NEXT: ld.param.u64 %rd2, [test_mulhi_i64_param_1];
77+
; CHECK-NEXT: mul.hi.s64 %rd3, %rd1, %rd2;
78+
; CHECK-NEXT: st.param.b64 [func_retval0+0], %rd3;
79+
; CHECK-NEXT: ret;
80+
%1 = call i64 @llvm.nvvm.mulhi.ll(i64 %x, i64 %y)
81+
ret i64 %1
82+
}
83+
84+
define i64 @test_mulhi_u64(i64 %x, i64 %y) {
85+
; CHECK-LABEL: test_mulhi_u64(
86+
; CHECK: {
87+
; CHECK-NEXT: .reg .b64 %rd<4>;
88+
; CHECK-EMPTY:
89+
; CHECK-NEXT: // %bb.0:
90+
; CHECK-NEXT: ld.param.u64 %rd1, [test_mulhi_u64_param_0];
91+
; CHECK-NEXT: ld.param.u64 %rd2, [test_mulhi_u64_param_1];
92+
; CHECK-NEXT: mul.hi.u64 %rd3, %rd1, %rd2;
93+
; CHECK-NEXT: st.param.b64 [func_retval0+0], %rd3;
94+
; CHECK-NEXT: ret;
95+
%1 = call i64 @llvm.nvvm.mulhi.ull(i64 %x, i64 %y)
96+
ret i64 %1
97+
}
98+
99+
declare i16 @llvm.nvvm.mulhi.s(i16, i16)
100+
declare i16 @llvm.nvvm.mulhi.us(i16, i16)
101+
declare i32 @llvm.nvvm.mulhi.i(i32, i32)
102+
declare i32 @llvm.nvvm.mulhi.ui(i32, i32)
103+
declare i64 @llvm.nvvm.mulhi.ll(i64, i64)
104+
declare i64 @llvm.nvvm.mulhi.ull(i64, i64)

0 commit comments

Comments
 (0)