[llvm] [NVPTX] Added intrinsics for i16x2 ALU instructions (PR #194767)
via llvm-commits
llvm-commits at lists.llvm.org
Tue Apr 28 18:47:52 PDT 2026
llvmbot wrote:
<!--LLVM PR SUMMARY COMMENT-->
@llvm/pr-subscribers-llvm-ir
Author: yasmincs
<details>
<summary>Changes</summary>
Added intrinsics for packed 16-bit integer instructions with NVPTX support: llvm.nvvm.add.s16x2, llvm.nvvm.sub.s16x2, llvm.nvvm.min.s16x2, llvm.nvvm.max.s16x2, llvm.nvvm.min.u16x2, and llvm.nvvm.max.u16x2.
---
Full diff: https://github.com/llvm/llvm-project/pull/194767.diff
4 Files Affected:
- (modified) llvm/include/llvm/IR/IntrinsicsNVVM.td (+19)
- (modified) llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp (+18)
- (modified) llvm/lib/Target/NVPTX/NVPTXIntrinsics.td (+11)
- (added) llvm/test/CodeGen/NVPTX/i16x2-intrinsics.ll (+121)
``````````diff
diff --git a/llvm/include/llvm/IR/IntrinsicsNVVM.td b/llvm/include/llvm/IR/IntrinsicsNVVM.td
index b81d347eb3e32..0e3701f02e776 100644
--- a/llvm/include/llvm/IR/IntrinsicsNVVM.td
+++ b/llvm/include/llvm/IR/IntrinsicsNVVM.td
@@ -1439,6 +1439,25 @@ let TargetPrefix = "nvvm" in {
def int_nvvm_neg_bf16x2 : NVVMBuiltin,
PureIntrinsic<[llvm_v2bf16_ty], [llvm_v2bf16_ty]>;
+ //
+ // Packed i16x2 integer ALU
+ //
+ let IntrProperties = [IntrNoMem, IntrSpeculatable, Commutative,
+ IntrNoCreateUndefOrPoison] in {
+ def int_nvvm_add_s16x2 : NVVMBuiltin,
+ DefaultAttrsIntrinsic<[llvm_v2i16_ty], [llvm_v2i16_ty, llvm_v2i16_ty]>;
+ def int_nvvm_min_s16x2 : NVVMBuiltin,
+ DefaultAttrsIntrinsic<[llvm_v2i16_ty], [llvm_v2i16_ty, llvm_v2i16_ty]>;
+ def int_nvvm_max_s16x2 : NVVMBuiltin,
+ DefaultAttrsIntrinsic<[llvm_v2i16_ty], [llvm_v2i16_ty, llvm_v2i16_ty]>;
+ def int_nvvm_min_u16x2 : NVVMBuiltin,
+ DefaultAttrsIntrinsic<[llvm_v2i16_ty], [llvm_v2i16_ty, llvm_v2i16_ty]>;
+ def int_nvvm_max_u16x2 : NVVMBuiltin,
+ DefaultAttrsIntrinsic<[llvm_v2i16_ty], [llvm_v2i16_ty, llvm_v2i16_ty]>;
+ }
+ def int_nvvm_sub_s16x2 : NVVMBuiltin,
+ PureIntrinsic<[llvm_v2i16_ty], [llvm_v2i16_ty, llvm_v2i16_ty]>;
+
//
// Round
//
diff --git a/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp b/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
index a5fd0a8724762..b4ccf0d5b27d1 100644
--- a/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
+++ b/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
@@ -3020,6 +3020,22 @@ static SDValue lowerPrmtIntrinsic(SDValue Op, SelectionDAG &DAG) {
return getPRMT(A, B, Selector, DL, DAG, Mode);
}
+/// llvm.nvvm.sub.s16x2 → add.s16x2(a, -b) with per-lane neg.s16.
+static SDValue lowerNvvmSubS16x2Intrinsic(SDValue Op, SelectionDAG &DAG) {
+ SDLoc DL(Op);
+ SDValue A = Op.getOperand(1);
+ SDValue B = Op.getOperand(2);
+ SmallVector<SDValue, 2> NegLanes;
+ SDValue Zero = DAG.getConstant(0, DL, MVT::i16);
+ for (unsigned I = 0; I < 2; ++I) {
+ SDValue Lane = DAG.getNode(ISD::EXTRACT_VECTOR_ELT, DL, MVT::i16, B,
+ DAG.getConstant(I, DL, MVT::i32));
+ NegLanes.push_back(DAG.getNode(ISD::SUB, DL, MVT::i16, Zero, Lane));
+ }
+ SDValue NegB = DAG.getNode(ISD::BUILD_VECTOR, DL, MVT::v2i16, NegLanes);
+ return DAG.getNode(ISD::ADD, DL, MVT::v2i16, A, NegB);
+}
+
#define TCGEN05_LD_RED_INTR(SHAPE, NUM, TYPE) \
Intrinsic::nvvm_tcgen05_ld_red_##SHAPE##_x##NUM##_##TYPE
@@ -3195,6 +3211,8 @@ static SDValue lowerIntrinsicWOChain(SDValue Op, SelectionDAG &DAG) {
case Intrinsic::nvvm_f32x4_to_e2m1x4_rs_satfinite:
case Intrinsic::nvvm_f32x4_to_e2m1x4_rs_relu_satfinite:
return lowerCvtRSIntrinsics(Op, DAG);
+ case Intrinsic::nvvm_sub_s16x2:
+ return lowerNvvmSubS16x2Intrinsic(Op, DAG);
}
}
diff --git a/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td b/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
index 8979276bc5afb..9c0469425d7ff 100644
--- a/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
+++ b/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
@@ -1601,6 +1601,17 @@ def INT_NVVM_NEG_BF16 : F_MATH_1<"neg.bf16", BF16RT,
def INT_NVVM_NEG_BF16X2 : F_MATH_1<"neg.bf16x2", BF16X2RT,
BF16X2RT, int_nvvm_neg_bf16x2, [hasPTX<70>, hasSM<80>]>;
+//
+// Packed i16x2 integer ALU
+//
+let Predicates = [hasPTX<80>, hasSM<90>] in {
+ def : Pat<(int_nvvm_add_s16x2 v2i16:$a, v2i16:$b), (ADD16x2 $a, $b)>;
+ def : Pat<(int_nvvm_min_s16x2 v2i16:$a, v2i16:$b), (SMIN16x2 $a, $b)>;
+ def : Pat<(int_nvvm_max_s16x2 v2i16:$a, v2i16:$b), (SMAX16x2 $a, $b)>;
+ def : Pat<(int_nvvm_min_u16x2 v2i16:$a, v2i16:$b), (UMIN16x2 $a, $b)>;
+ def : Pat<(int_nvvm_max_u16x2 v2i16:$a, v2i16:$b), (UMAX16x2 $a, $b)>;
+}
+
//
// Round
//
diff --git a/llvm/test/CodeGen/NVPTX/i16x2-intrinsics.ll b/llvm/test/CodeGen/NVPTX/i16x2-intrinsics.ll
new file mode 100644
index 0000000000000..2e4a6e9e0364c
--- /dev/null
+++ b/llvm/test/CodeGen/NVPTX/i16x2-intrinsics.ll
@@ -0,0 +1,121 @@
+; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 6
+; RUN: llc < %s -mtriple=nvptx64-nvidia-cuda -mcpu=sm_90 -mattr=+ptx80 | FileCheck %s
+
+declare <2 x i16> @llvm.nvvm.add.s16x2(<2 x i16>, <2 x i16>) #0
+declare <2 x i16> @llvm.nvvm.sub.s16x2(<2 x i16>, <2 x i16>) #0
+declare <2 x i16> @llvm.nvvm.min.s16x2(<2 x i16>, <2 x i16>) #0
+declare <2 x i16> @llvm.nvvm.max.s16x2(<2 x i16>, <2 x i16>) #0
+declare <2 x i16> @llvm.nvvm.min.u16x2(<2 x i16>, <2 x i16>) #0
+declare <2 x i16> @llvm.nvvm.max.u16x2(<2 x i16>, <2 x i16>) #0
+
+define <2 x i16> @test_add(<2 x i16> %a, <2 x i16> %b) {
+; CHECK-LABEL: test_add(
+; CHECK: {
+; CHECK-NEXT: .reg .b32 %r<4>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.b32 %r1, [test_add_param_0];
+; CHECK-NEXT: ld.param.b32 %r2, [test_add_param_1];
+; CHECK-NEXT: add.s16x2 %r3, %r1, %r2;
+; CHECK-NEXT: st.param.b32 [func_retval0], %r3;
+; CHECK-NEXT: ret;
+ %r = call <2 x i16> @llvm.nvvm.add.s16x2(<2 x i16> %a, <2 x i16> %b)
+ ret <2 x i16> %r
+}
+
+define <2 x i16> @test_sub(<2 x i16> %a, <2 x i16> %b) {
+; CHECK-LABEL: test_sub(
+; CHECK: {
+; CHECK-NEXT: .reg .b16 %rs<5>;
+; CHECK-NEXT: .reg .b32 %r<4>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.b32 %r1, [test_sub_param_0];
+; CHECK-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [test_sub_param_1];
+; CHECK-NEXT: neg.s16 %rs3, %rs2;
+; CHECK-NEXT: neg.s16 %rs4, %rs1;
+; CHECK-NEXT: mov.b32 %r2, {%rs4, %rs3};
+; CHECK-NEXT: add.s16x2 %r3, %r1, %r2;
+; CHECK-NEXT: st.param.b32 [func_retval0], %r3;
+; CHECK-NEXT: ret;
+ %r = call <2 x i16> @llvm.nvvm.sub.s16x2(<2 x i16> %a, <2 x i16> %b)
+ ret <2 x i16> %r
+}
+
+define <2 x i16> @test_min_s(<2 x i16> %a, <2 x i16> %b) {
+; CHECK-LABEL: test_min_s(
+; CHECK: {
+; CHECK-NEXT: .reg .b32 %r<4>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.b32 %r1, [test_min_s_param_0];
+; CHECK-NEXT: ld.param.b32 %r2, [test_min_s_param_1];
+; CHECK-NEXT: min.s16x2 %r3, %r1, %r2;
+; CHECK-NEXT: st.param.b32 [func_retval0], %r3;
+; CHECK-NEXT: ret;
+ %r = call <2 x i16> @llvm.nvvm.min.s16x2(<2 x i16> %a, <2 x i16> %b)
+ ret <2 x i16> %r
+}
+
+define <2 x i16> @test_max_s(<2 x i16> %a, <2 x i16> %b) {
+; CHECK-LABEL: test_max_s(
+; CHECK: {
+; CHECK-NEXT: .reg .b32 %r<4>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.b32 %r1, [test_max_s_param_0];
+; CHECK-NEXT: ld.param.b32 %r2, [test_max_s_param_1];
+; CHECK-NEXT: max.s16x2 %r3, %r1, %r2;
+; CHECK-NEXT: st.param.b32 [func_retval0], %r3;
+; CHECK-NEXT: ret;
+ %r = call <2 x i16> @llvm.nvvm.max.s16x2(<2 x i16> %a, <2 x i16> %b)
+ ret <2 x i16> %r
+}
+
+define <2 x i16> @test_min_u(<2 x i16> %a, <2 x i16> %b) {
+; CHECK-LABEL: test_min_u(
+; CHECK: {
+; CHECK-NEXT: .reg .b32 %r<4>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.b32 %r1, [test_min_u_param_0];
+; CHECK-NEXT: ld.param.b32 %r2, [test_min_u_param_1];
+; CHECK-NEXT: min.u16x2 %r3, %r1, %r2;
+; CHECK-NEXT: st.param.b32 [func_retval0], %r3;
+; CHECK-NEXT: ret;
+ %r = call <2 x i16> @llvm.nvvm.min.u16x2(<2 x i16> %a, <2 x i16> %b)
+ ret <2 x i16> %r
+}
+
+define <2 x i16> @test_max_u(<2 x i16> %a, <2 x i16> %b) {
+; CHECK-LABEL: test_max_u(
+; CHECK: {
+; CHECK-NEXT: .reg .b32 %r<4>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.b32 %r1, [test_max_u_param_0];
+; CHECK-NEXT: ld.param.b32 %r2, [test_max_u_param_1];
+; CHECK-NEXT: max.u16x2 %r3, %r1, %r2;
+; CHECK-NEXT: st.param.b32 [func_retval0], %r3;
+; CHECK-NEXT: ret;
+ %r = call <2 x i16> @llvm.nvvm.max.u16x2(<2 x i16> %a, <2 x i16> %b)
+ ret <2 x i16> %r
+}
+
+define <2 x i16> @test_sub_dag(<2 x i16> %a, <2 x i16> %b) {
+; CHECK-LABEL: test_sub_dag(
+; CHECK: {
+; CHECK-NEXT: .reg .b16 %rs<7>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [test_sub_dag_param_0];
+; CHECK-NEXT: ld.param.v2.b16 {%rs3, %rs4}, [test_sub_dag_param_1];
+; CHECK-NEXT: sub.s16 %rs5, %rs2, %rs4;
+; CHECK-NEXT: sub.s16 %rs6, %rs1, %rs3;
+; CHECK-NEXT: st.param.v2.b16 [func_retval0], {%rs6, %rs5};
+; CHECK-NEXT: ret;
+ %r = sub nsw <2 x i16> %a, %b
+ ret <2 x i16> %r
+}
+
+attributes #0 = { nocallback nosync nounwind speculatable willreturn memory(none) }
``````````
</details>
https://github.com/llvm/llvm-project/pull/194767
More information about the llvm-commits
mailing list