[clang] [llvm] [mlir] [AMDGPU] [ROCDL] Added Intrinsics for smed, umed, to support ISA instructions from ROCDL (PR #157748)
via llvm-commits
llvm-commits at lists.llvm.org
Thu Sep 11 06:49:06 PDT 2025
llvmbot wrote:
<!--LLVM PR SUMMARY COMMENT-->
@llvm/pr-subscribers-backend-amdgpu
Author: Keshav Vinayak Jha (keshavvinayak01)
<details>
<summary>Changes</summary>
## Description
Added support for AMDGPU signed (med3) intrinsics. Implemented `smed3` and `umed3` ISA instructions from ROCDL with complete end-to-end support including LLVM intrinsics, Clang builtins, AMDGPU backend instruction selection, and MLIR ROCDL dialect operations.
## Testing
- 4 Lit test files in `llvm/test/CodeGen/AMDGPU/`
- ROCDL -> LLVMIR lit tests for new `rocdl.med3.<dtype>` ops in `/test/Target/LLVMIR/rocdl.mlir`
Addresses [#<!-- -->157052](https://github.com/llvm/llvm-project/issues/157052)
---
Patch is 27.89 KiB, truncated to 20.00 KiB below, full version: https://github.com/llvm/llvm-project/pull/157748.diff
14 Files Affected:
- (modified) clang/include/clang/Basic/BuiltinsAMDGPU.def (+4)
- (modified) clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp (+8)
- (modified) llvm/include/llvm/IR/IntrinsicsAMDGPU.td (+12)
- (modified) llvm/lib/Target/AMDGPU/AMDGPUGISel.td (+2-2)
- (modified) llvm/lib/Target/AMDGPU/AMDGPUInstCombineIntrinsic.cpp (+172)
- (modified) llvm/lib/Target/AMDGPU/AMDGPUInstrInfo.td (+12-7)
- (modified) llvm/lib/Target/AMDGPU/AMDGPULegalizerInfo.cpp (+22)
- (modified) llvm/lib/Target/AMDGPU/AMDGPURegisterBankInfo.cpp (+3)
- (added) llvm/test/CodeGen/AMDGPU/llvm.amdgcn.smed3.i16.ll (+27)
- (added) llvm/test/CodeGen/AMDGPU/llvm.amdgcn.smed3.ll (+42)
- (added) llvm/test/CodeGen/AMDGPU/llvm.amdgcn.umed3.i16.ll (+27)
- (added) llvm/test/CodeGen/AMDGPU/llvm.amdgcn.umed3.ll (+42)
- (modified) mlir/include/mlir/Dialect/LLVMIR/ROCDLOps.td (+88)
- (modified) mlir/test/Target/LLVMIR/rocdl.mlir (+42)
``````````diff
diff --git a/clang/include/clang/Basic/BuiltinsAMDGPU.def b/clang/include/clang/Basic/BuiltinsAMDGPU.def
index fda16e42d2c6b..4f6ab2a36cd85 100644
--- a/clang/include/clang/Basic/BuiltinsAMDGPU.def
+++ b/clang/include/clang/Basic/BuiltinsAMDGPU.def
@@ -139,6 +139,8 @@ BUILTIN(__builtin_amdgcn_ds_bpermute, "iii", "nc")
BUILTIN(__builtin_amdgcn_readfirstlane, "ii", "nc")
BUILTIN(__builtin_amdgcn_readlane, "iii", "nc")
BUILTIN(__builtin_amdgcn_fmed3f, "ffff", "nc")
+BUILTIN(__builtin_amdgcn_smed3, "iiii", "nc")
+BUILTIN(__builtin_amdgcn_umed3, "UiUiUiUi", "nc")
BUILTIN(__builtin_amdgcn_ds_faddf, "ff*3fIiIiIb", "n")
BUILTIN(__builtin_amdgcn_ds_fminf, "ff*3fIiIiIb", "n")
BUILTIN(__builtin_amdgcn_ds_fmaxf, "ff*3fIiIiIb", "n")
@@ -265,6 +267,8 @@ TARGET_BUILTIN(__builtin_amdgcn_perm, "UiUiUiUi", "nc", "gfx8-insts")
//===----------------------------------------------------------------------===//
TARGET_BUILTIN(__builtin_amdgcn_fmed3h, "hhhh", "nc", "gfx9-insts")
+TARGET_BUILTIN(__builtin_amdgcn_smed3h, "ssss", "nc", "gfx9-insts")
+TARGET_BUILTIN(__builtin_amdgcn_umed3h, "UsUsUsUs", "nc", "gfx9-insts")
TARGET_BUILTIN(__builtin_amdgcn_global_atomic_fadd_f64, "dd*1d", "t", "gfx90a-insts")
TARGET_BUILTIN(__builtin_amdgcn_global_atomic_fadd_f32, "ff*1f", "t", "atomic-fadd-rtn-insts")
diff --git a/clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp b/clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp
index 07cf08c54985a..5d4c980c7c63e 100644
--- a/clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp
+++ b/clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp
@@ -606,6 +606,14 @@ Value *CodeGenFunction::EmitAMDGPUBuiltinExpr(unsigned BuiltinID,
case AMDGPU::BI__builtin_amdgcn_fmed3h:
return emitBuiltinWithOneOverloadedType<3>(*this, E,
Intrinsic::amdgcn_fmed3);
+ case AMDGPU::BI__builtin_amdgcn_smed3:
+ case AMDGPU::BI__builtin_amdgcn_smed3h:
+ return emitBuiltinWithOneOverloadedType<3>(*this, E,
+ Intrinsic::amdgcn_smed3);
+ case AMDGPU::BI__builtin_amdgcn_umed3:
+ case AMDGPU::BI__builtin_amdgcn_umed3h:
+ return emitBuiltinWithOneOverloadedType<3>(*this, E,
+ Intrinsic::amdgcn_umed3);
case AMDGPU::BI__builtin_amdgcn_ds_append:
case AMDGPU::BI__builtin_amdgcn_ds_consume: {
Intrinsic::ID Intrin = BuiltinID == AMDGPU::BI__builtin_amdgcn_ds_append ?
diff --git a/llvm/include/llvm/IR/IntrinsicsAMDGPU.td b/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
index 030d01d7a5f3f..52ba06ed4be25 100644
--- a/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
+++ b/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
@@ -543,6 +543,18 @@ def int_amdgcn_fmed3 :
[IntrNoMem, IntrSpeculatable]
>;
+def int_amdgcn_smed3 :
+ DefaultAttrsIntrinsic<[llvm_anyint_ty],
+ [LLVMMatchType<0>, LLVMMatchType<0>, LLVMMatchType<0>],
+ [IntrNoMem, IntrSpeculatable]
+>;
+
+def int_amdgcn_umed3 :
+ DefaultAttrsIntrinsic<[llvm_anyint_ty],
+ [LLVMMatchType<0>, LLVMMatchType<0>, LLVMMatchType<0>],
+ [IntrNoMem, IntrSpeculatable]
+>;
+
def int_amdgcn_cubeid : ClangBuiltin<"__builtin_amdgcn_cubeid">,
DefaultAttrsIntrinsic<[llvm_float_ty],
[llvm_float_ty, llvm_float_ty, llvm_float_ty],
diff --git a/llvm/lib/Target/AMDGPU/AMDGPUGISel.td b/llvm/lib/Target/AMDGPU/AMDGPUGISel.td
index bb4bf742fb861..486ec90edcaef 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPUGISel.td
+++ b/llvm/lib/Target/AMDGPU/AMDGPUGISel.td
@@ -256,8 +256,8 @@ def : GINodeEquiv<G_AMDGPU_CVT_F32_UBYTE2, AMDGPUcvt_f32_ubyte2>;
def : GINodeEquiv<G_AMDGPU_CVT_F32_UBYTE3, AMDGPUcvt_f32_ubyte3>;
def : GINodeEquiv<G_AMDGPU_CVT_PK_I16_I32, AMDGPUpk_i16_i32_impl>;
-def : GINodeEquiv<G_AMDGPU_SMED3, AMDGPUsmed3>;
-def : GINodeEquiv<G_AMDGPU_UMED3, AMDGPUumed3>;
+def : GINodeEquiv<G_AMDGPU_SMED3, AMDGPUsmed3_impl>;
+def : GINodeEquiv<G_AMDGPU_UMED3, AMDGPUumed3_impl>;
def : GINodeEquiv<G_AMDGPU_FMED3, AMDGPUfmed3_impl>;
def : GINodeEquiv<G_AMDGPU_CLAMP, AMDGPUclamp>;
diff --git a/llvm/lib/Target/AMDGPU/AMDGPUInstCombineIntrinsic.cpp b/llvm/lib/Target/AMDGPU/AMDGPUInstCombineIntrinsic.cpp
index 4fe5d00679436..c6cb4736f95df 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPUInstCombineIntrinsic.cpp
+++ b/llvm/lib/Target/AMDGPU/AMDGPUInstCombineIntrinsic.cpp
@@ -60,6 +60,26 @@ static APFloat fmed3AMDGCN(const APFloat &Src0, const APFloat &Src1,
return maxnum(Src0, Src1);
}
+// Constant fold llvm.amdgcn.smed3 intrinsics for standard inputs.
+static APInt smed3AMDGCN(const APInt &Src0, const APInt &Src1, const APInt &Src2) {
+ APInt Max3 = Src0.sgt(Src1) ? (Src0.sgt(Src2) ? Src0 : Src2)
+ : (Src1.sgt(Src2) ? Src1 : Src2);
+
+ if (Max3 == Src0) return Src1.sgt(Src2) ? Src1 : Src2;
+ if (Max3 == Src1) return Src0.sgt(Src2) ? Src0 : Src2;
+ return Src0.sgt(Src1) ? Src0 : Src1;
+}
+
+// Constant fold llvm.amdgcn.umed3 intrinsics for standard inputs.
+static APInt umed3AMDGCN(const APInt &Src0, const APInt &Src1, const APInt &Src2) {
+ APInt Max3 = Src0.ugt(Src1) ? (Src0.ugt(Src2) ? Src0 : Src2)
+ : (Src1.ugt(Src2) ? Src1 : Src2);
+
+ if (Max3 == Src0) return Src1.ugt(Src2) ? Src1 : Src2;
+ if (Max3 == Src1) return Src0.ugt(Src2) ? Src0 : Src2;
+ return Src0.ugt(Src1) ? Src0 : Src1;
+}
+
// Check if a value can be converted to a 16-bit value without losing
// precision.
// The value is expected to be either a float (IsFloat = true) or an unsigned
@@ -427,6 +447,36 @@ static Value *matchFPExtFromF16(Value *Arg) {
return nullptr;
}
+/// Match an sext from i16 to i32, or a constant we can convert.
+static Value *matchSExtFromI16(Value *Arg) {
+ Value *Src = nullptr;
+ ConstantInt *CInt = nullptr;
+ if (match(Arg, m_OneUse(m_SExt(m_Value(Src))))) {
+ if (Src->getType()->isIntegerTy(16))
+ return Src;
+ } else if (match(Arg, m_ConstantInt(CInt))) {
+ // Check if the constant fits in i16
+ if (CInt->getValue().getActiveBits() <= 16)
+ return ConstantInt::get(Type::getInt16Ty(Arg->getContext()), CInt->getValue().trunc(16));
+ }
+ return nullptr;
+}
+
+/// Match a zext from i16 to i32, or a constant we can convert.
+static Value *matchZExtFromI16(Value *Arg) {
+ Value *Src = nullptr;
+ ConstantInt *CInt = nullptr;
+ if (match(Arg, m_OneUse(m_ZExt(m_Value(Src))))) {
+ if (Src->getType()->isIntegerTy(16))
+ return Src;
+ } else if (match(Arg, m_ConstantInt(CInt))) {
+ // Check if the constant fits in i16
+ if (CInt->getValue().getActiveBits() <= 16)
+ return ConstantInt::get(Type::getInt16Ty(Arg->getContext()), CInt->getValue().trunc(16));
+ }
+ return nullptr;
+}
+
// Trim all zero components from the end of the vector \p UseV and return
// an appropriate bitset with known elements.
static APInt trimTrailingZerosInVector(InstCombiner &IC, Value *UseV,
@@ -1174,6 +1224,128 @@ GCNTTIImpl::instCombineIntrinsic(InstCombiner &IC, IntrinsicInst &II) const {
break;
}
+ case Intrinsic::amdgcn_smed3: {
+ Value *Src0 = II.getArgOperand(0);
+ Value *Src1 = II.getArgOperand(1);
+ Value *Src2 = II.getArgOperand(2);
+
+ // Propagate poison values.
+ for (Value *Src : {Src0, Src1, Src2}) {
+ if (isa<PoisonValue>(Src))
+ return IC.replaceInstUsesWith(II, Src);
+ }
+
+ bool Swap = false;
+ // Canonicalize constants to RHS operands.
+ //
+ // smed3(c0, x, c1) -> smed3(x, c0, c1)
+ if (isa<Constant>(Src0) && !isa<Constant>(Src1)) {
+ std::swap(Src0, Src1);
+ Swap = true;
+ }
+
+ if (isa<Constant>(Src1) && !isa<Constant>(Src2)) {
+ std::swap(Src1, Src2);
+ Swap = true;
+ }
+
+ if (isa<Constant>(Src0) && !isa<Constant>(Src1)) {
+ std::swap(Src0, Src1);
+ Swap = true;
+ }
+
+ if (Swap) {
+ II.setArgOperand(0, Src0);
+ II.setArgOperand(1, Src1);
+ II.setArgOperand(2, Src2);
+ return &II;
+ }
+
+ // Constant fold smed3 with constant operands.
+ if (const ConstantInt *C0 = dyn_cast<ConstantInt>(Src0)) {
+ if (const ConstantInt *C1 = dyn_cast<ConstantInt>(Src1)) {
+ if (const ConstantInt *C2 = dyn_cast<ConstantInt>(Src2)) {
+ APInt Result = smed3AMDGCN(C0->getValue(), C1->getValue(), C2->getValue());
+ return IC.replaceInstUsesWith(II, ConstantInt::get(II.getType(), Result));
+ }
+ }
+ }
+
+ // Width reduction for integer extensions.
+ // smed3((sext X), (sext Y), (sext Z)) -> sext (smed3(X, Y, Z))
+ if (Value *X = matchSExtFromI16(Src0)) {
+ if (Value *Y = matchSExtFromI16(Src1)) {
+ if (Value *Z = matchSExtFromI16(Src2)) {
+ Value *NewCall = IC.Builder.CreateIntrinsic(
+ IID, {X->getType()}, {X, Y, Z}, &II, II.getName());
+ return new SExtInst(NewCall, II.getType());
+ }
+ }
+ }
+
+ break;
+ }
+ case Intrinsic::amdgcn_umed3: {
+ Value *Src0 = II.getArgOperand(0);
+ Value *Src1 = II.getArgOperand(1);
+ Value *Src2 = II.getArgOperand(2);
+
+ // Propagate poison values.
+ for (Value *Src : {Src0, Src1, Src2}) {
+ if (isa<PoisonValue>(Src))
+ return IC.replaceInstUsesWith(II, Src);
+ }
+
+ bool Swap = false;
+ // Canonicalize constants to RHS operands.
+ //
+ // umed3(c0, x, c1) -> umed3(x, c0, c1)
+ if (isa<Constant>(Src0) && !isa<Constant>(Src1)) {
+ std::swap(Src0, Src1);
+ Swap = true;
+ }
+
+ if (isa<Constant>(Src1) && !isa<Constant>(Src2)) {
+ std::swap(Src1, Src2);
+ Swap = true;
+ }
+
+ if (isa<Constant>(Src0) && !isa<Constant>(Src1)) {
+ std::swap(Src0, Src1);
+ Swap = true;
+ }
+
+ if (Swap) {
+ II.setArgOperand(0, Src0);
+ II.setArgOperand(1, Src1);
+ II.setArgOperand(2, Src2);
+ return &II;
+ }
+
+ // Constant fold umed3 with constant operands.
+ if (const ConstantInt *C0 = dyn_cast<ConstantInt>(Src0)) {
+ if (const ConstantInt *C1 = dyn_cast<ConstantInt>(Src1)) {
+ if (const ConstantInt *C2 = dyn_cast<ConstantInt>(Src2)) {
+ APInt Result = umed3AMDGCN(C0->getValue(), C1->getValue(), C2->getValue());
+ return IC.replaceInstUsesWith(II, ConstantInt::get(II.getType(), Result));
+ }
+ }
+ }
+
+ // Width reduction for integer extensions.
+ // umed3((zext X), (zext Y), (zext Z)) -> zext (umed3(X, Y, Z))
+ if (Value *X = matchZExtFromI16(Src0)) {
+ if (Value *Y = matchZExtFromI16(Src1)) {
+ if (Value *Z = matchZExtFromI16(Src2)) {
+ Value *NewCall = IC.Builder.CreateIntrinsic(
+ IID, {X->getType()}, {X, Y, Z}, &II, II.getName());
+ return new ZExtInst(NewCall, II.getType());
+ }
+ }
+ }
+
+ break;
+ }
case Intrinsic::amdgcn_icmp:
case Intrinsic::amdgcn_fcmp: {
const ConstantInt *CC = cast<ConstantInt>(II.getArgOperand(2));
diff --git a/llvm/lib/Target/AMDGPU/AMDGPUInstrInfo.td b/llvm/lib/Target/AMDGPU/AMDGPUInstrInfo.td
index b8fa6f3fc6867..e9680e062cffa 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPUInstrInfo.td
+++ b/llvm/lib/Target/AMDGPU/AMDGPUInstrInfo.td
@@ -334,16 +334,13 @@ def AMDGPUmad_i24 : SDNode<"AMDGPUISD::MAD_I24", AMDGPUDTIntTernaryOp,
[]
>;
-def AMDGPUsmed3 : SDNode<"AMDGPUISD::SMED3", AMDGPUDTIntTernaryOp,
- []
->;
-
-def AMDGPUumed3 : SDNode<"AMDGPUISD::UMED3", AMDGPUDTIntTernaryOp,
- []
->;
def AMDGPUfmed3_impl : SDNode<"AMDGPUISD::FMED3", SDTFPTernaryOp, []>;
+def AMDGPUsmed3_impl : SDNode<"AMDGPUISD::SMED3", AMDGPUDTIntTernaryOp, []>;
+
+def AMDGPUumed3_impl : SDNode<"AMDGPUISD::UMED3", AMDGPUDTIntTernaryOp, []>;
+
def AMDGPUfdot2_impl : SDNode<"AMDGPUISD::FDOT2",
SDTypeProfile<1, 4, [SDTCisSameAs<0, 3>, SDTCisSameAs<1, 2>,
SDTCisFP<0>, SDTCisVec<1>,
@@ -448,6 +445,14 @@ def AMDGPUfmed3 : PatFrags<(ops node:$src0, node:$src1, node:$src2),
[(int_amdgcn_fmed3 node:$src0, node:$src1, node:$src2),
(AMDGPUfmed3_impl node:$src0, node:$src1, node:$src2)]>;
+def AMDGPUsmed3 : PatFrags<(ops node:$src0, node:$src1, node:$src2),
+ [(int_amdgcn_smed3 node:$src0, node:$src1, node:$src2),
+ (AMDGPUsmed3_impl node:$src0, node:$src1, node:$src2)]>;
+
+def AMDGPUumed3 : PatFrags<(ops node:$src0, node:$src1, node:$src2),
+ [(int_amdgcn_umed3 node:$src0, node:$src1, node:$src2),
+ (AMDGPUumed3_impl node:$src0, node:$src1, node:$src2)]>;
+
def AMDGPUdiv_fixup : PatFrags<(ops node:$src0, node:$src1, node:$src2),
[(int_amdgcn_div_fixup node:$src0, node:$src1, node:$src2),
(AMDGPUdiv_fixup_impl node:$src0, node:$src1, node:$src2)]>;
diff --git a/llvm/lib/Target/AMDGPU/AMDGPULegalizerInfo.cpp b/llvm/lib/Target/AMDGPU/AMDGPULegalizerInfo.cpp
index f18536cd4ab93..5da1e04c58bae 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPULegalizerInfo.cpp
+++ b/llvm/lib/Target/AMDGPU/AMDGPULegalizerInfo.cpp
@@ -7798,6 +7798,28 @@ bool AMDGPULegalizerInfo::legalizeIntrinsic(LegalizerHelper &Helper,
Observer.changedInstr(MI);
return true;
}
+ case Intrinsic::amdgcn_smed3: {
+ GISelChangeObserver &Observer = Helper.Observer;
+
+ // FIXME: This is to workaround the inability of tablegen match combiners to
+ // match intrinsics in patterns.
+ Observer.changingInstr(MI);
+ MI.setDesc(B.getTII().get(AMDGPU::G_AMDGPU_SMED3));
+ MI.removeOperand(1);
+ Observer.changedInstr(MI);
+ return true;
+ }
+ case Intrinsic::amdgcn_umed3: {
+ GISelChangeObserver &Observer = Helper.Observer;
+
+ // FIXME: This is to workaround the inability of tablegen match combiners to
+ // match intrinsics in patterns.
+ Observer.changingInstr(MI);
+ MI.setDesc(B.getTII().get(AMDGPU::G_AMDGPU_UMED3));
+ MI.removeOperand(1);
+ Observer.changedInstr(MI);
+ return true;
+ }
case Intrinsic::amdgcn_readlane:
case Intrinsic::amdgcn_writelane:
case Intrinsic::amdgcn_readfirstlane:
diff --git a/llvm/lib/Target/AMDGPU/AMDGPURegisterBankInfo.cpp b/llvm/lib/Target/AMDGPU/AMDGPURegisterBankInfo.cpp
index 36b27bef350ed..63141d065bf65 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPURegisterBankInfo.cpp
+++ b/llvm/lib/Target/AMDGPU/AMDGPURegisterBankInfo.cpp
@@ -4136,6 +4136,7 @@ AMDGPURegisterBankInfo::getInstrMapping(const MachineInstr &MI) const {
case AMDGPU::G_AMDGPU_CVT_F32_UBYTE3:
case AMDGPU::G_AMDGPU_CVT_PK_I16_I32:
case AMDGPU::G_AMDGPU_SMED3:
+ case AMDGPU::G_AMDGPU_UMED3:
case AMDGPU::G_AMDGPU_FMED3:
return getDefaultMappingVOP(MI);
case AMDGPU::G_UMULH:
@@ -4660,6 +4661,8 @@ AMDGPURegisterBankInfo::getInstrMapping(const MachineInstr &MI) const {
case Intrinsic::amdgcn_cvt_scalef32_sr_pk16_bf6_bf16:
case Intrinsic::amdgcn_sat_pk4_i4_i8:
case Intrinsic::amdgcn_sat_pk4_u4_u8:
+ case Intrinsic::amdgcn_smed3:
+ case Intrinsic::amdgcn_umed3:
case Intrinsic::amdgcn_fmed3:
case Intrinsic::amdgcn_cubeid:
case Intrinsic::amdgcn_cubema:
diff --git a/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.smed3.i16.ll b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.smed3.i16.ll
new file mode 100644
index 0000000000000..0f6f00309401c
--- /dev/null
+++ b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.smed3.i16.ll
@@ -0,0 +1,27 @@
+; RUN: llc -mtriple=amdgcn -mcpu=gfx900 < %s | FileCheck -check-prefix=GCN %s
+
+; GCN-LABEL: {{^}}test_smed3_i16:
+; GCN: v_med3_i16 v{{[0-9]+}}, s{{[0-9]+}}, v{{[0-9]+}}, v{{[0-9]+}}
+define amdgpu_kernel void @test_smed3_i16(ptr addrspace(1) %out, i32 %src0.arg, i32 %src1.arg, i32 %src2.arg) #1 {
+ %src0.i16 = trunc i32 %src0.arg to i16
+ %src1.i16 = trunc i32 %src1.arg to i16
+ %src2.i16 = trunc i32 %src2.arg to i16
+ %med3 = call i16 @llvm.amdgcn.smed3.i16(i16 %src0.i16, i16 %src1.i16, i16 %src2.i16)
+ store i16 %med3, ptr addrspace(1) %out
+ ret void
+}
+
+; GCN-LABEL: {{^}}test_smed3_zero_i16:
+; GCN: v_med3_i16 v{{[0-9]+}}, s{{[0-9]+}}, v{{[0-9]+}}, 0
+define amdgpu_kernel void @test_smed3_zero_i16(ptr addrspace(1) %out, i32 %src0.arg, i32 %src1.arg) #1 {
+ %src0.i16 = trunc i32 %src0.arg to i16
+ %src1.i16 = trunc i32 %src1.arg to i16
+ %med3 = call i16 @llvm.amdgcn.smed3.i16(i16 %src0.i16, i16 %src1.i16, i16 0)
+ store i16 %med3, ptr addrspace(1) %out
+ ret void
+}
+
+declare i16 @llvm.amdgcn.smed3.i16(i16, i16, i16) #0
+
+attributes #0 = { nounwind readnone }
+attributes #1 = { nounwind }
diff --git a/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.smed3.ll b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.smed3.ll
new file mode 100644
index 0000000000000..250fdc0d2d78d
--- /dev/null
+++ b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.smed3.ll
@@ -0,0 +1,42 @@
+; RUN: llc -mtriple=amdgcn < %s | FileCheck -check-prefix=GCN %s
+; RUN: llc -mtriple=amdgcn -mcpu=tonga < %s | FileCheck -check-prefix=GCN %s
+
+; GCN-LABEL: {{^}}test_smed3:
+; GCN: v_med3_i32 v{{[0-9]+}}, s{{[0-9]+}}, v{{[0-9]+}}, v{{[0-9]+}}
+define amdgpu_kernel void @test_smed3(ptr addrspace(1) %out, i32 %src0, i32 %src1, i32 %src2) #1 {
+ %med3 = call i32 @llvm.amdgcn.smed3.i32(i32 %src0, i32 %src1, i32 %src2)
+ store i32 %med3, ptr addrspace(1) %out
+ ret void
+}
+
+; GCN-LABEL: {{^}}test_smed3_multi_use:
+; GCN: v_med3_i32 [[MED3:v[0-9]+]], s{{[0-9]+}}, v{{[0-9]+}}, v{{[0-9]+}}
+; GCN: v_mul_lo_i32 v{{[0-9]+}}, [[MED3]], s{{[0-9]+}}
+define amdgpu_kernel void @test_smed3_multi_use(ptr addrspace(1) %out, i32 %src0, i32 %src1, i32 %src2, i32 %mul.arg) #1 {
+ %med3 = call i32 @llvm.amdgcn.smed3.i32(i32 %src0, i32 %src1, i32 %src2)
+ %med3.user = mul i32 %med3, %mul.arg
+ store volatile i32 %med3.user, ptr addrspace(1) %out
+ store volatile i32 %med3, ptr addrspace(1) %out
+ ret void
+}
+
+; GCN-LABEL: {{^}}test_smed3_constants:
+; GCN: v_med3_i32 v{{[0-9]+}}, s{{[0-9]+}}, v{{[0-9]+}}, 42
+define amdgpu_kernel void @test_smed3_constants(ptr addrspace(1) %out, i32 %src0, i32 %src1) #1 {
+ %med3 = call i32 @llvm.amdgcn.smed3.i32(i32 %src0, i32 %src1, i32 42)
+ store i32 %med3, ptr addrspace(1) %out
+ ret void
+}
+
+; GCN-LABEL: {{^}}test_smed3_zero:
+; GCN: v_med3_i32 v{{[0-9]+}}, s{{[0-9]+}}, v{{[0-9]+}}, 0
+define amdgpu_kernel void @test_smed3_zero(ptr addrspace(1) %out, i32 %src0, i32 %src1) #1 {
+ %med3 = call i32 @llvm.amdgcn.smed3.i32(i32 %src0, i32 %src1, i32 0)
+ store i32 %med3, ptr addrspace(1) %out
+ ret void
+}
+
+declare i32 @llvm.amdgcn.smed3.i32(i32, i32, i32) #0
+
+attributes #0 = { nounwind readnone }
+attributes #1 = { nounwind }
diff --git a/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.umed3.i16.ll b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.umed3.i16.ll
new file mode 100644
index 0000000000000..d484e8a4b0804
--- /dev/null
+++ b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.umed3.i16.ll
@@ -0,0 +1,27 @@
+; RUN: llc -mtriple=amdgcn -mcpu=gfx900 < %s | FileCheck -check-prefix=GCN %s
+
+; GCN-LABEL: {{^}}test_umed3_i16:
+; GCN: v_med3_u16 v{{[0-9]+}}, s{{[0-9]+}}, v{{[0-9]+}}, v{{[0-9]+}}
+define amdgpu_kernel void @test_umed3_i16(ptr addrspace(1) %out, i32 %src0.arg, i32 %src1.arg, i32 %src2.arg) #1 {
+ %src0.i16 = trunc i32 %src0.arg to i16
+ %src1.i16 = trunc i32 %src1.arg to i16
+ %src2.i16 = trunc i32 %src2.arg to i16
+ %med3 = call i16 @llvm.amdgcn.umed3.i16(i16 %src0.i16, i16 %src1.i16, i16 %src2.i16)
+ store i16 %med3, ptr addrspace(1) %out
+ ret void
+}
+
+; GCN-LABEL: {{^}}test_umed3_zero_i16:
+; GCN: v_med3_u16 v{{[0-9]+}}, s{{[0-9]+}}, v{{[0-9]+}}, 0
+define amdgpu_kernel void @test_umed3_zero_i16(ptr addrspace(1) %out, i32 %src0.arg, i32 %src1.arg) #1 {
+ %src0.i16 = trunc i32 %src0.arg to i16
+ %src1.i16 = trunc i32 %src1.arg to i16
+ %med3 = call i16 @llvm.amdgcn.umed3.i16(i16 %src0.i16, i16 %src1.i16, i16 0)
+ store i16 %med3, ptr addrspace(1) %out
+ ret void
+}
+
+declare i16 @llvm.amdgcn.umed3.i16(i16, i16, i16) #0
+
+attributes #0 = { nounwind readnone }
+attributes #1 = { nounwind }
diff --git a/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.umed3.ll b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.umed3.ll
new file mode 100644
index 0000000000000..e1bec276d1fb6
--- /dev/null
+++ b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.umed3.ll
@@ -0,0 +1,42 @@
+; RUN: llc -mtriple=amdgcn < %s | FileCheck -check-prefix=GCN %s
+; RUN: llc -mtriple=amdgcn -mcpu=tonga < %s | FileCheck -check-prefix=GCN %s
+
+; GCN-LABEL: {{^}}test_umed3:
+; GCN: v_med3_u32 v{{[0-9]+}}, s{{[0-9]+}}, v{{[0-9]+}}, v{{[0-9]+}}
+define amdgpu_kernel void @test_umed3(ptr addrspace(1) %out, i32 %src0, i32 %src1, i32 %src2) #1 {
+ %med3 = call i32 @llvm.amdgcn.umed3.i32(i32 %src0, i32 %src1, i32 %src2)
+ store i32 %med3, ptr addrspace(1) %out
+ ret void
+}
+
+; GCN-LABEL: {{^}}test_umed3_multi_use:...
[truncated]
``````````
</details>
https://github.com/llvm/llvm-project/pull/157748
More information about the llvm-commits
mailing list