Skip to content

Commit 1159c67

Browse files
arsenmarsenm
committed
AMDGPU: Drop and auto-upgrade llvm.amdgcn.ldexp to llvm.ldexp
1 parent ab85553 commit 1159c67

File tree

9 files changed

+40
-319
lines changed

9 files changed

+40
-319
lines changed

‎llvm/include/llvm/IR/IntrinsicsAMDGPU.td

Lines changed: 0 additions & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -362,12 +362,6 @@ def int_amdgcn_rsq_legacy : ClangBuiltin<"__builtin_amdgcn_rsq_legacy">,
362362
def int_amdgcn_rsq_clamp : DefaultAttrsIntrinsic<
363363
[llvm_anyfloat_ty], [LLVMMatchType<0>], [IntrNoMem, IntrSpeculatable]>;
364364

365-
// For int_amdgcn_ldexp_f16, only the low 16 bits of the i32 src1 operand will used.
366-
def int_amdgcn_ldexp : DefaultAttrsIntrinsic<
367-
[llvm_anyfloat_ty], [LLVMMatchType<0>, llvm_i32_ty],
368-
[IntrNoMem, IntrSpeculatable]
369-
>;
370-
371365
def int_amdgcn_frexp_mant : DefaultAttrsIntrinsic<
372366
[llvm_anyfloat_ty], [LLVMMatchType<0>], [IntrNoMem, IntrSpeculatable]
373367
>;

‎llvm/lib/IR/AutoUpgrade.cpp

Lines changed: 7 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -836,6 +836,13 @@ static bool UpgradeIntrinsicFunction1(Function *F, Function *&NewFn) {
836836
{F->getReturnType()});
837837
return true;
838838
}
839+
if (Name.startswith("amdgcn.ldexp")) {
840+
// Target specific intrinsic became redundant
841+
NewFn = Intrinsic::getDeclaration(
842+
F->getParent(), Intrinsic::ldexp,
843+
{F->getReturnType(), F->getArg(1)->getType()});
844+
return true;
845+
}
839846

840847
break;
841848
}

‎llvm/lib/Target/AMDGPU/AMDGPUISelLowering.cpp

Lines changed: 1 addition & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -3220,8 +3220,7 @@ SDValue AMDGPUTargetLowering::performIntrinsicWOChainCombine(
32203220
case Intrinsic::amdgcn_rsq:
32213221
case Intrinsic::amdgcn_rcp_legacy:
32223222
case Intrinsic::amdgcn_rsq_legacy:
3223-
case Intrinsic::amdgcn_rsq_clamp:
3224-
case Intrinsic::amdgcn_ldexp: {
3223+
case Intrinsic::amdgcn_rsq_clamp: {
32253224
// FIXME: This is probably wrong. If src is an sNaN, it won't be quieted
32263225
SDValue Src = N->getOperand(1);
32273226
return Src.isUndef() ? Src : SDValue();

‎llvm/lib/Target/AMDGPU/AMDGPUInstCombineIntrinsic.cpp

Lines changed: 0 additions & 44 deletions
Original file line numberDiff line numberDiff line change
@@ -996,50 +996,6 @@ GCNTTIImpl::instCombineIntrinsic(InstCombiner &IC, IntrinsicInst &II) const {
996996

997997
break;
998998
}
999-
case Intrinsic::amdgcn_ldexp: {
1000-
// FIXME: This doesn't introduce new instructions and belongs in
1001-
// InstructionSimplify.
1002-
Type *Ty = II.getType();
1003-
Value *Op0 = II.getArgOperand(0);
1004-
Value *Op1 = II.getArgOperand(1);
1005-
1006-
// Folding undef to qnan is safe regardless of the FP mode.
1007-
if (isa<UndefValue>(Op0)) {
1008-
auto *QNaN = ConstantFP::get(Ty, APFloat::getQNaN(Ty->getFltSemantics()));
1009-
return IC.replaceInstUsesWith(II, QNaN);
1010-
}
1011-
1012-
const APFloat *C = nullptr;
1013-
match(Op0, PatternMatch::m_APFloat(C));
1014-
1015-
// FIXME: Should flush denorms depending on FP mode, but that's ignored
1016-
// everywhere else.
1017-
//
1018-
// These cases should be safe, even with strictfp.
1019-
// ldexp(0.0, x) -> 0.0
1020-
// ldexp(-0.0, x) -> -0.0
1021-
// ldexp(inf, x) -> inf
1022-
// ldexp(-inf, x) -> -inf
1023-
if (C && (C->isZero() || C->isInfinity())) {
1024-
return IC.replaceInstUsesWith(II, Op0);
1025-
}
1026-
1027-
// With strictfp, be more careful about possibly needing to flush denormals
1028-
// or not, and snan behavior depends on ieee_mode.
1029-
if (II.isStrictFP())
1030-
break;
1031-
1032-
if (C && C->isNaN())
1033-
return IC.replaceInstUsesWith(II, ConstantFP::get(Ty, C->makeQuiet()));
1034-
1035-
// ldexp(x, 0) -> x
1036-
// ldexp(x, undef) -> x
1037-
if (isa<UndefValue>(Op1) || match(Op1, PatternMatch::m_ZeroInt())) {
1038-
return IC.replaceInstUsesWith(II, Op0);
1039-
}
1040-
1041-
break;
1042-
}
1043999
case Intrinsic::amdgcn_fmul_legacy: {
10441000
Value *Op0 = II.getArgOperand(0);
10451001
Value *Op1 = II.getArgOperand(1);

‎llvm/lib/Target/AMDGPU/SIISelLowering.cpp

Lines changed: 0 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -7223,9 +7223,6 @@ SDValue SITargetLowering::LowerINTRINSIC_WO_CHAIN(SDValue Op,
72237223

72247224
return emitRemovedIntrinsicError(DAG, DL, VT);
72257225
}
7226-
case Intrinsic::amdgcn_ldexp:
7227-
return DAG.getNode(ISD::FLDEXP, DL, VT, Op.getOperand(1), Op.getOperand(2));
7228-
72297226
case Intrinsic::amdgcn_fract:
72307227
return DAG.getNode(AMDGPUISD::FRACT, DL, VT, Op.getOperand(1));
72317228

@@ -10672,7 +10669,6 @@ bool SITargetLowering::isCanonicalized(Register Reg, MachineFunction &MF,
1067210669
case Intrinsic::amdgcn_div_fmas:
1067310670
case Intrinsic::amdgcn_div_fixup:
1067410671
case Intrinsic::amdgcn_fract:
10675-
case Intrinsic::amdgcn_ldexp:
1067610672
case Intrinsic::amdgcn_cvt_pkrtz:
1067710673
case Intrinsic::amdgcn_cubeid:
1067810674
case Intrinsic::amdgcn_cubema:

‎llvm/test/Bitcode/amdgcn-ldexp.ll

Lines changed: 30 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,30 @@
1+
; RUN: llvm-as < %s | llvm-dis | FileCheck %s
2+
3+
define float @f32(float %a, i32 %b) {
4+
; CHECK: %call = call float @llvm.ldexp.f32.i32(float %a, i32 %b)
5+
; CHECK-NOT: amdgcn.ldexp
6+
%call = call float @llvm.amdgcn.ldexp.f32(float %a, i32 %b)
7+
ret float %call
8+
}
9+
10+
define double @f64(double %a, i32 %b) {
11+
; CHECK: %call = call double @llvm.ldexp.f64.i32(double %a, i32 %b)
12+
; CHECK-NOT: amdgcn.ldexp
13+
%call = call double @llvm.amdgcn.ldexp.f64(double %a, i32 %b)
14+
ret double %call
15+
}
16+
17+
define half @f16(half %a, i32 %b) {
18+
; CHECK: %call = call half @llvm.ldexp.f16.i32(half %a, i32 %b)
19+
; CHECK-NOT: amdgcn.ldexp
20+
%call = call half @llvm.amdgcn.ldexp.f16(half %a, i32 %b)
21+
ret half %call
22+
}
23+
24+
declare half @llvm.amdgcn.ldexp.f16(half, i32)
25+
declare float @llvm.amdgcn.ldexp.f32(float, i32)
26+
declare double @llvm.amdgcn.ldexp.f64(double, i32)
27+
; CHECK: declare half @llvm.ldexp.f16.i32(half, i32)
28+
; CHECK: declare float @llvm.ldexp.f32.i32(float, i32)
29+
; CHECK: declare double @llvm.ldexp.f64.i32(double, i32)
30+
; CHECK-NOT: amdgcn.ldexp

‎llvm/test/CodeGen/AMDGPU/known-never-snan.ll

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -516,7 +516,7 @@ define float @v_test_known_not_snan_ldexp_input_fmed3_r_i_i_f32(float %a, i32 %b
516516
; GCN-NEXT: v_ldexp_f32 v0, v0, v1
517517
; GCN-NEXT: v_med3_f32 v0, v0, 2.0, 4.0
518518
; GCN-NEXT: s_setpc_b64 s[30:31]
519-
%known.not.snan = call float @llvm.amdgcn.ldexp.f32(float %a, i32 %b)
519+
%known.not.snan = call float @llvm.ldexp.f32.i32(float %a, i32 %b)
520520
%max = call float @llvm.maxnum.f32(float %known.not.snan, float 2.0)
521521
%med = call float @llvm.minnum.f32(float %max, float 4.0)
522522
ret float %med
@@ -658,7 +658,7 @@ declare float @llvm.maxnum.f32(float, float) #1
658658
declare float @llvm.copysign.f32(float, float) #1
659659
declare float @llvm.fma.f32(float, float, float) #1
660660
declare float @llvm.fmuladd.f32(float, float, float) #1
661-
declare float @llvm.amdgcn.ldexp.f32(float, i32) #1
661+
declare float @llvm.ldexp.f32.i32(float, i32) #1
662662
declare float @llvm.amdgcn.fmul.legacy(float, float) #1
663663
declare float @llvm.amdgcn.fmed3.f32(float, float, float) #1
664664
declare float @llvm.amdgcn.frexp.mant.f32(float) #1

‎llvm/test/CodeGen/AMDGPU/llvm.amdgcn.ldexp.f16.ll

Lines changed: 0 additions & 230 deletions
This file was deleted.

0 commit comments

Comments
 (0)