AMDGPU: Drop and auto-upgrade llvm.amdgcn.ldexp to llvm.ldexp

[arsenm]

Missed test update

[Joe_Nash]

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

6

Should we not update these tests to use the generic intrinsic rather than delete them? I am working on Codegen for the new version of this instruction where src1 is 16-bit, and these tests are covering useful output.

[arsenm]

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

6

New ones are added in D14327

[Joe_Nash]

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

661

Why did this case need a manual update rather than use the autoupgrade?

[arsenm]

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

661

The auto-upgrade would work functionality wise. But there are 2 issues. For the optimization purposes, the old form wasn't recognized in isCanonicalized so the optimization broke. Second I think it's good practice to upgrade tests to the new intrinsics and only test the old forms in upgrade tests

[Joe_Nash]

LGTM

[arsenm]

1159c670d40e3ef302264c681fe7e0268a550874

[foad]

In D149589#4429514, @arsenm wrote:

1159c670d40e3ef302264c681fe7e0268a550874

That commit was reverted with d9333e360a7c52587ab6e4328e7493b357fb2cf3

[arsenm]

edecb6048148b7c22cfe567ccbf2dba5e1000bdc

llvm/include/llvm/IR/IntrinsicsAMDGPU.td

def int_amdgcn_rsq_legacy :  ClangBuiltin<"__builtin_amdgcn_rsq_legacy">,
   DefaultAttrsIntrinsic<
   [llvm_float_ty], [llvm_float_ty], [IntrNoMem, IntrSpeculatable]
 >;
 
 // out = 1.0 / sqrt(a) result clamped to +/- max_float.
 def int_amdgcn_rsq_clamp : DefaultAttrsIntrinsic<
   [llvm_anyfloat_ty], [LLVMMatchType<0>], [IntrNoMem, IntrSpeculatable]>;
 
-// For int_amdgcn_ldexp_f16, only the low 16 bits of the i32 src1 operand will used.
-def int_amdgcn_ldexp : DefaultAttrsIntrinsic<
-  [llvm_anyfloat_ty], [LLVMMatchType<0>, llvm_i32_ty],
-  [IntrNoMem, IntrSpeculatable]
->;
-
 def int_amdgcn_frexp_mant : DefaultAttrsIntrinsic<
   [llvm_anyfloat_ty], [LLVMMatchType<0>], [IntrNoMem, IntrSpeculatable]
 >;
 
 def int_amdgcn_frexp_exp : DefaultAttrsIntrinsic<
   [llvm_anyint_ty], [llvm_anyfloat_ty], [IntrNoMem, IntrSpeculatable]
 >;
 

llvm/lib/IR/AutoUpgrade.cpp

if (Name == "arm.mve.mull.int.predicated.v2i64.v4i32.v4i1" ||
       return true;
 
     if (Name == "amdgcn.alignbit") {
       // Target specific intrinsic became redundant
       NewFn = Intrinsic::getDeclaration(F->getParent(), Intrinsic::fshr,
                                         {F->getReturnType()});
       return true;
     }
+    if (Name.startswith("amdgcn.ldexp")) {
+      // Target specific intrinsic became redundant
+      NewFn = Intrinsic::getDeclaration(
+          F->getParent(), Intrinsic::ldexp,
+          {F->getReturnType(), F->getArg(1)->getType()});
+      return true;
+    }
 
     break;
   }
 
   case 'c': {
     if (Name.startswith("ctlz.") && F->arg_size() == 1) {
       rename(F);
       NewFn = Intrinsic::getDeclaration(F->getParent(), Intrinsic::ctlz,

llvm/lib/Target/AMDGPU/AMDGPUISelLowering.cpp

SDValue AMDGPUTargetLowering::performIntrinsicWOChainCombine(
   case Intrinsic::amdgcn_mul_u24:
   case Intrinsic::amdgcn_mulhi_i24:
   case Intrinsic::amdgcn_mulhi_u24:
     return simplifyMul24(N, DCI);
   case Intrinsic::amdgcn_fract:
   case Intrinsic::amdgcn_rsq:
   case Intrinsic::amdgcn_rcp_legacy:
   case Intrinsic::amdgcn_rsq_legacy:
-  case Intrinsic::amdgcn_rsq_clamp:
-  case Intrinsic::amdgcn_ldexp: {
+  case Intrinsic::amdgcn_rsq_clamp: {
     // FIXME: This is probably wrong. If src is an sNaN, it won't be quieted
     SDValue Src = N->getOperand(1);
     return Src.isUndef() ? Src : SDValue();
   }
   default:
     return SDValue();
   }
 }

llvm/lib/Target/AMDGPU/AMDGPUInstCombineIntrinsic.cpp

if (IID == Intrinsic::amdgcn_readfirstlane) {
                          PatternMatch::m_Value(),
                          PatternMatch::m_Specific(II.getArgOperand(1))))) {
         return IC.replaceInstUsesWith(II, Src);
       }
     }
 
     break;
   }
-  case Intrinsic::amdgcn_ldexp: {
-    // FIXME: This doesn't introduce new instructions and belongs in
-    // InstructionSimplify.
-    Type *Ty = II.getType();
-    Value *Op0 = II.getArgOperand(0);
-    Value *Op1 = II.getArgOperand(1);
-
-    // Folding undef to qnan is safe regardless of the FP mode.
-    if (isa<UndefValue>(Op0)) {
-      auto *QNaN = ConstantFP::get(Ty, APFloat::getQNaN(Ty->getFltSemantics()));
-      return IC.replaceInstUsesWith(II, QNaN);
-    }
-
-    const APFloat *C = nullptr;
-    match(Op0, PatternMatch::m_APFloat(C));
-
-    // FIXME: Should flush denorms depending on FP mode, but that's ignored
-    // everywhere else.
-    //
-    // These cases should be safe, even with strictfp.
-    // ldexp(0.0, x) -> 0.0
-    // ldexp(-0.0, x) -> -0.0
-    // ldexp(inf, x) -> inf
-    // ldexp(-inf, x) -> -inf
-    if (C && (C->isZero() || C->isInfinity())) {
-      return IC.replaceInstUsesWith(II, Op0);
-    }
-
-    // With strictfp, be more careful about possibly needing to flush denormals
-    // or not, and snan behavior depends on ieee_mode.
-    if (II.isStrictFP())
-      break;
-
-    if (C && C->isNaN())
-      return IC.replaceInstUsesWith(II, ConstantFP::get(Ty, C->makeQuiet()));
-
-    // ldexp(x, 0) -> x
-    // ldexp(x, undef) -> x
-    if (isa<UndefValue>(Op1) || match(Op1, PatternMatch::m_ZeroInt())) {
-      return IC.replaceInstUsesWith(II, Op0);
-    }
-
-    break;
-  }
   case Intrinsic::amdgcn_fmul_legacy: {
     Value *Op0 = II.getArgOperand(0);
     Value *Op1 = II.getArgOperand(1);
 
     // The legacy behaviour is that multiplying +/-0.0 by anything, even NaN or
     // infinity, gives +0.0.
     // TODO: Move to InstSimplify?
     if (match(Op0, PatternMatch::m_AnyZeroFP()) ||

llvm/lib/Target/AMDGPU/SIISelLowering.cpp

case Intrinsic::amdgcn_mul_i24:
     return DAG.getNode(AMDGPUISD::MUL_I24, DL, VT, Op.getOperand(1), Op.getOperand(2));
 
   case Intrinsic::amdgcn_log_clamp: {
     if (Subtarget->getGeneration() < AMDGPUSubtarget::VOLCANIC_ISLANDS)
       return SDValue();
 
     return emitRemovedIntrinsicError(DAG, DL, VT);
   }
-  case Intrinsic::amdgcn_ldexp:
-    return DAG.getNode(ISD::FLDEXP, DL, VT, Op.getOperand(1), Op.getOperand(2));
-
   case Intrinsic::amdgcn_fract:
     return DAG.getNode(AMDGPUISD::FRACT, DL, VT, Op.getOperand(1));
 
   case Intrinsic::amdgcn_class:
     return DAG.getNode(AMDGPUISD::FP_CLASS, DL, VT,
                        Op.getOperand(1), Op.getOperand(2));
   case Intrinsic::amdgcn_div_fmas:
     return DAG.getNode(AMDGPUISD::DIV_FMAS, DL, VT,
case AMDGPU::G_INTRINSIC:
     case Intrinsic::amdgcn_rcp_legacy:
     case Intrinsic::amdgcn_rsq:
     case Intrinsic::amdgcn_rsq_clamp:
     case Intrinsic::amdgcn_rsq_legacy:
     case Intrinsic::amdgcn_div_scale:
     case Intrinsic::amdgcn_div_fmas:
     case Intrinsic::amdgcn_div_fixup:
     case Intrinsic::amdgcn_fract:
-    case Intrinsic::amdgcn_ldexp:
     case Intrinsic::amdgcn_cvt_pkrtz:
     case Intrinsic::amdgcn_cubeid:
     case Intrinsic::amdgcn_cubema:
     case Intrinsic::amdgcn_cubesc:
     case Intrinsic::amdgcn_cubetc:
     case Intrinsic::amdgcn_frexp_mant:
     case Intrinsic::amdgcn_fdot2:
     case Intrinsic::amdgcn_trig_preop:

llvm/test/Bitcode/amdgcn-ldexp.ll

+; RUN: llvm-as < %s | llvm-dis | FileCheck %s
+
+define float @f32(float %a, i32 %b) {
+  ; CHECK: %call = call float @llvm.ldexp.f32.i32(float %a, i32 %b)
+  ; CHECK-NOT: amdgcn.ldexp
+  %call = call float @llvm.amdgcn.ldexp.f32(float %a, i32 %b)
+  ret float %call
+}
+
+define double @f64(double %a, i32 %b) {
+  ; CHECK: %call = call double @llvm.ldexp.f64.i32(double %a, i32 %b)
+  ; CHECK-NOT: amdgcn.ldexp
+  %call = call double @llvm.amdgcn.ldexp.f64(double %a, i32 %b)
+  ret double %call
+}
+
+define half @f16(half %a, i32 %b) {
+  ; CHECK: %call = call half @llvm.ldexp.f16.i32(half %a, i32 %b)
+  ; CHECK-NOT: amdgcn.ldexp
+  %call = call half @llvm.amdgcn.ldexp.f16(half %a, i32 %b)
+  ret half %call
+}
+
+declare half @llvm.amdgcn.ldexp.f16(half, i32)
+declare float @llvm.amdgcn.ldexp.f32(float, i32)
+declare double @llvm.amdgcn.ldexp.f64(double, i32)
+; CHECK: declare half @llvm.ldexp.f16.i32(half, i32)
+; CHECK: declare float @llvm.ldexp.f32.i32(float, i32)
+; CHECK: declare double @llvm.ldexp.f64.i32(double, i32)
+; CHECK-NOT: amdgcn.ldexp

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

 
 define float @v_test_known_not_snan_ldexp_input_fmed3_r_i_i_f32(float %a, i32 %b) #0 {
 ; GCN-LABEL: v_test_known_not_snan_ldexp_input_fmed3_r_i_i_f32:
 ; GCN:       ; %bb.0:
 ; GCN-NEXT:    s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
 ; GCN-NEXT:    v_ldexp_f32 v0, v0, v1
 ; GCN-NEXT:    v_med3_f32 v0, v0, 2.0, 4.0
 ; GCN-NEXT:    s_setpc_b64 s[30:31]
-  %known.not.snan = call float @llvm.amdgcn.ldexp.f32(float %a, i32 %b)
+  %known.not.snan = call float @llvm.ldexp.f32.i32(float %a, i32 %b)
   %max = call float @llvm.maxnum.f32(float %known.not.snan, float 2.0)
   %med = call float @llvm.minnum.f32(float %max, float 4.0)
   ret float %med
 }
 
 define float @v_test_known_not_snan_fmed3_input_fmed3_r_i_i_f32(float %a, float %b, float %c) #0 {
 ; GCN-LABEL: v_test_known_not_snan_fmed3_input_fmed3_r_i_i_f32:
 ; GCN:       ; %bb.0:

 declare float @llvm.rint.f32(float) #1
 declare float @llvm.nearbyint.f32(float) #1
 declare float @llvm.canonicalize.f32(float) #1
 declare float @llvm.minnum.f32(float, float) #1
 declare float @llvm.maxnum.f32(float, float) #1
 declare float @llvm.copysign.f32(float, float) #1
 declare float @llvm.fma.f32(float, float, float) #1
 declare float @llvm.fmuladd.f32(float, float, float) #1
-declare float @llvm.amdgcn.ldexp.f32(float, i32) #1
+declare float @llvm.ldexp.f32.i32(float, i32) #1
 declare float @llvm.amdgcn.fmul.legacy(float, float) #1
 declare float @llvm.amdgcn.fmed3.f32(float, float, float) #1
 declare float @llvm.amdgcn.frexp.mant.f32(float) #1
 declare float @llvm.amdgcn.rcp.f32(float) #1
 declare float @llvm.amdgcn.rsq.f32(float) #1
 declare float @llvm.amdgcn.fract.f32(float) #1
 declare float @llvm.amdgcn.cubeid(float, float, float) #0
 
 attributes #0 = { nounwind "denormal-fp-math-f32"="preserve-sign,preserve-sign" }
 attributes #1 = { nounwind readnone speculatable }
 
 !0 = !{float 2.500000e+00}

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

-; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py
-; RUN:  llc -amdgpu-scalarize-global-loads=false  -march=amdgcn -mcpu=fiji -mattr=-flat-for-global -verify-machineinstrs < %s | FileCheck -check-prefixes=VI %s
-; RUN:  llc -amdgpu-scalarize-global-loads=false  -march=amdgcn -mcpu=gfx1010 -mattr=-flat-for-global -verify-machineinstrs < %s | FileCheck -check-prefixes=GFX10 %s
-; RUN:  llc -amdgpu-scalarize-global-loads=false  -march=amdgcn -mcpu=gfx1100 -mattr=-flat-for-global -verify-machineinstrs < %s | FileCheck -check-prefixes=GFX11 %s
-
-declare half @llvm.amdgcn.ldexp.f16(half %a, i32 %b)
-
-define amdgpu_kernel void @ldexp_f16(
-; VI-LABEL: ldexp_f16:
-; VI:       ; %bb.0:
-; VI-NEXT:    s_load_dwordx4 s[4:7], s[0:1], 0x24
-; VI-NEXT:    s_load_dwordx2 s[8:9], s[0:1], 0x34
-; VI-NEXT:    s_mov_b32 s3, 0xf000
-; VI-NEXT:    s_mov_b32 s2, -1
-; VI-NEXT:    s_mov_b32 s14, s2
-; VI-NEXT:    s_waitcnt lgkmcnt(0)
-; VI-NEXT:    s_mov_b32 s12, s6
-; VI-NEXT:    s_mov_b32 s13, s7
-; VI-NEXT:    s_mov_b32 s15, s3
-; VI-NEXT:    s_mov_b32 s10, s2
-; VI-NEXT:    s_mov_b32 s11, s3
-; VI-NEXT:    buffer_load_ushort v0, off, s[12:15], 0
-; VI-NEXT:    buffer_load_dword v1, off, s[8:11], 0
-; VI-NEXT:    s_mov_b32 s0, s4
-; VI-NEXT:    s_mov_b32 s1, s5
-; VI-NEXT:    s_waitcnt vmcnt(0)
-; VI-NEXT:    v_ldexp_f16_e32 v0, v0, v1
-; VI-NEXT:    buffer_store_short v0, off, s[0:3], 0
-; VI-NEXT:    s_endpgm
-;
-; GFX10-LABEL: ldexp_f16:
-; GFX10:       ; %bb.0:
-; GFX10-NEXT:    s_clause 0x1
-; GFX10-NEXT:    s_load_dwordx4 s[4:7], s[0:1], 0x24
-; GFX10-NEXT:    s_load_dwordx2 s[8:9], s[0:1], 0x34
-; GFX10-NEXT:    s_mov_b32 s2, -1
-; GFX10-NEXT:    s_mov_b32 s3, 0x31016000
-; GFX10-NEXT:    s_mov_b32 s14, s2
-; GFX10-NEXT:    s_mov_b32 s15, s3
-; GFX10-NEXT:    s_mov_b32 s10, s2
-; GFX10-NEXT:    s_mov_b32 s11, s3
-; GFX10-NEXT:    s_waitcnt lgkmcnt(0)
-; GFX10-NEXT:    s_mov_b32 s12, s6
-; GFX10-NEXT:    s_mov_b32 s13, s7
-; GFX10-NEXT:    buffer_load_ushort v0, off, s[12:15], 0
-; GFX10-NEXT:    buffer_load_dword v1, off, s[8:11], 0
-; GFX10-NEXT:    s_mov_b32 s0, s4
-; GFX10-NEXT:    s_mov_b32 s1, s5
-; GFX10-NEXT:    s_waitcnt vmcnt(0)
-; GFX10-NEXT:    v_ldexp_f16_e32 v0, v0, v1
-; GFX10-NEXT:    buffer_store_short v0, off, s[0:3], 0
-; GFX10-NEXT:    s_endpgm
-;
-; GFX11-LABEL: ldexp_f16:
-; GFX11:       ; %bb.0:
-; GFX11-NEXT:    s_clause 0x1
-; GFX11-NEXT:    s_load_b128 s[4:7], s[0:1], 0x24
-; GFX11-NEXT:    s_load_b64 s[0:1], s[0:1], 0x34
-; GFX11-NEXT:    s_mov_b32 s10, -1
-; GFX11-NEXT:    s_mov_b32 s11, 0x31016000
-; GFX11-NEXT:    s_mov_b32 s14, s10
-; GFX11-NEXT:    s_mov_b32 s15, s11
-; GFX11-NEXT:    s_mov_b32 s2, s10
-; GFX11-NEXT:    s_mov_b32 s3, s11
-; GFX11-NEXT:    s_waitcnt lgkmcnt(0)
-; GFX11-NEXT:    s_mov_b32 s12, s6
-; GFX11-NEXT:    s_mov_b32 s13, s7
-; GFX11-NEXT:    buffer_load_u16 v0, off, s[12:15], 0
-; GFX11-NEXT:    buffer_load_b32 v1, off, s[0:3], 0
-; GFX11-NEXT:    s_mov_b32 s8, s4
-; GFX11-NEXT:    s_mov_b32 s9, s5
-; GFX11-NEXT:    s_waitcnt vmcnt(0)
-; GFX11-NEXT:    v_ldexp_f16_e32 v0, v0, v1
-; GFX11-NEXT:    buffer_store_b16 v0, off, s[8:11], 0
-; GFX11-NEXT:    s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
-; GFX11-NEXT:    s_endpgm
-    ptr addrspace(1) %r,
-    ptr addrspace(1) %a,
-    ptr addrspace(1) %b) {
-  %a.val = load half, ptr addrspace(1) %a
-  %b.val = load i32, ptr addrspace(1) %b
-  %r.val = call half @llvm.amdgcn.ldexp.f16(half %a.val, i32 %b.val)
-  store half %r.val, ptr addrspace(1) %r
-  ret void
-}
-
-define amdgpu_kernel void @ldexp_f16_imm_a(
-; VI-LABEL: ldexp_f16_imm_a:
-; VI:       ; %bb.0:
-; VI-NEXT:    s_load_dwordx4 s[0:3], s[0:1], 0x24
-; VI-NEXT:    s_mov_b32 s7, 0xf000
-; VI-NEXT:    s_mov_b32 s6, -1
-; VI-NEXT:    s_mov_b32 s10, s6
-; VI-NEXT:    s_mov_b32 s11, s7
-; VI-NEXT:    s_waitcnt lgkmcnt(0)
-; VI-NEXT:    s_mov_b32 s8, s2
-; VI-NEXT:    s_mov_b32 s9, s3
-; VI-NEXT:    buffer_load_dword v0, off, s[8:11], 0
-; VI-NEXT:    s_mov_b32 s4, s0
-; VI-NEXT:    s_mov_b32 s5, s1
-; VI-NEXT:    s_waitcnt vmcnt(0)
-; VI-NEXT:    v_ldexp_f16_e32 v0, 2.0, v0
-; VI-NEXT:    buffer_store_short v0, off, s[4:7], 0
-; VI-NEXT:    s_endpgm
-;
-; GFX10-LABEL: ldexp_f16_imm_a:
-; GFX10:       ; %bb.0:
-; GFX10-NEXT:    s_load_dwordx4 s[0:3], s[0:1], 0x24
-; GFX10-NEXT:    s_mov_b32 s6, -1
-; GFX10-NEXT:    s_mov_b32 s7, 0x31016000
-; GFX10-NEXT:    s_mov_b32 s10, s6
-; GFX10-NEXT:    s_mov_b32 s11, s7
-; GFX10-NEXT:    s_waitcnt lgkmcnt(0)
-; GFX10-NEXT:    s_mov_b32 s8, s2
-; GFX10-NEXT:    s_mov_b32 s9, s3
-; GFX10-NEXT:    s_mov_b32 s4, s0
-; GFX10-NEXT:    buffer_load_dword v0, off, s[8:11], 0
-; GFX10-NEXT:    s_mov_b32 s5, s1
-; GFX10-NEXT:    s_waitcnt vmcnt(0)
-; GFX10-NEXT:    v_ldexp_f16_e32 v0, 2.0, v0
-; GFX10-NEXT:    buffer_store_short v0, off, s[4:7], 0
-; GFX10-NEXT:    s_endpgm
-;
-; GFX11-LABEL: ldexp_f16_imm_a:
-; GFX11:       ; %bb.0:
-; GFX11-NEXT:    s_load_b128 s[0:3], s[0:1], 0x24
-; GFX11-NEXT:    s_mov_b32 s6, -1
-; GFX11-NEXT:    s_mov_b32 s7, 0x31016000
-; GFX11-NEXT:    s_mov_b32 s10, s6
-; GFX11-NEXT:    s_mov_b32 s11, s7
-; GFX11-NEXT:    s_waitcnt lgkmcnt(0)
-; GFX11-NEXT:    s_mov_b32 s8, s2
-; GFX11-NEXT:    s_mov_b32 s9, s3
-; GFX11-NEXT:    s_mov_b32 s4, s0
-; GFX11-NEXT:    buffer_load_b32 v0, off, s[8:11], 0
-; GFX11-NEXT:    s_mov_b32 s5, s1
-; GFX11-NEXT:    s_waitcnt vmcnt(0)
-; GFX11-NEXT:    v_ldexp_f16_e32 v0, 2.0, v0
-; GFX11-NEXT:    buffer_store_b16 v0, off, s[4:7], 0
-; GFX11-NEXT:    s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
-; GFX11-NEXT:    s_endpgm
-    ptr addrspace(1) %r,
-    ptr addrspace(1) %b) {
-  %b.val = load i32, ptr addrspace(1) %b
-  %r.val = call half @llvm.amdgcn.ldexp.f16(half 2.0, i32 %b.val)
-  store half %r.val, ptr addrspace(1) %r
-  ret void
-}
-
-define amdgpu_kernel void @ldexp_f16_imm_b(
-; VI-LABEL: ldexp_f16_imm_b:
-; VI:       ; %bb.0:
-; VI-NEXT:    s_load_dwordx4 s[0:3], s[0:1], 0x24
-; VI-NEXT:    s_mov_b32 s7, 0xf000
-; VI-NEXT:    s_mov_b32 s6, -1
-; VI-NEXT:    s_mov_b32 s10, s6
-; VI-NEXT:    s_mov_b32 s11, s7
-; VI-NEXT:    s_waitcnt lgkmcnt(0)
-; VI-NEXT:    s_mov_b32 s8, s2
-; VI-NEXT:    s_mov_b32 s9, s3
-; VI-NEXT:    buffer_load_ushort v0, off, s[8:11], 0
-; VI-NEXT:    s_mov_b32 s4, s0
-; VI-NEXT:    s_mov_b32 s5, s1
-; VI-NEXT:    s_waitcnt vmcnt(0)
-; VI-NEXT:    v_ldexp_f16_e64 v0, v0, 2
-; VI-NEXT:    buffer_store_short v0, off, s[4:7], 0
-; VI-NEXT:    s_endpgm
-;
-; GFX10-LABEL: ldexp_f16_imm_b:
-; GFX10:       ; %bb.0:
-; GFX10-NEXT:    s_load_dwordx4 s[0:3], s[0:1], 0x24
-; GFX10-NEXT:    s_mov_b32 s6, -1
-; GFX10-NEXT:    s_mov_b32 s7, 0x31016000
-; GFX10-NEXT:    s_mov_b32 s10, s6
-; GFX10-NEXT:    s_mov_b32 s11, s7
-; GFX10-NEXT:    s_waitcnt lgkmcnt(0)
-; GFX10-NEXT:    s_mov_b32 s8, s2
-; GFX10-NEXT:    s_mov_b32 s9, s3
-; GFX10-NEXT:    s_mov_b32 s4, s0
-; GFX10-NEXT:    buffer_load_ushort v0, off, s[8:11], 0
-; GFX10-NEXT:    s_mov_b32 s5, s1
-; GFX10-NEXT:    s_waitcnt vmcnt(0)
-; GFX10-NEXT:    v_ldexp_f16_e64 v0, v0, 2
-; GFX10-NEXT:    buffer_store_short v0, off, s[4:7], 0
-; GFX10-NEXT:    s_endpgm
-;
-; GFX11-LABEL: ldexp_f16_imm_b:
-; GFX11:       ; %bb.0:
-; GFX11-NEXT:    s_load_b128 s[0:3], s[0:1], 0x24
-; GFX11-NEXT:    s_mov_b32 s6, -1
-; GFX11-NEXT:    s_mov_b32 s7, 0x31016000
-; GFX11-NEXT:    s_mov_b32 s10, s6
-; GFX11-NEXT:    s_mov_b32 s11, s7
-; GFX11-NEXT:    s_waitcnt lgkmcnt(0)
-; GFX11-NEXT:    s_mov_b32 s8, s2
-; GFX11-NEXT:    s_mov_b32 s9, s3
-; GFX11-NEXT:    s_mov_b32 s4, s0
-; GFX11-NEXT:    buffer_load_u16 v0, off, s[8:11], 0
-; GFX11-NEXT:    s_mov_b32 s5, s1
-; GFX11-NEXT:    s_waitcnt vmcnt(0)
-; GFX11-NEXT:    v_ldexp_f16_e64 v0, v0, 2
-; GFX11-NEXT:    buffer_store_b16 v0, off, s[4:7], 0
-; GFX11-NEXT:    s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
-; GFX11-NEXT:    s_endpgm
-    ptr addrspace(1) %r,
-    ptr addrspace(1) %a) {
-  %a.val = load half, ptr addrspace(1) %a
-  %r.val = call half @llvm.amdgcn.ldexp.f16(half %a.val, i32 2)
-  store half %r.val, ptr addrspace(1) %r
-  ret void
-}

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

-; RUN: llc -march=amdgcn -verify-machineinstrs < %s | FileCheck -check-prefix=SI %s
-; RUN: llc -march=amdgcn -mcpu=tonga -verify-machineinstrs < %s | FileCheck -check-prefix=SI %s
-
-declare float @llvm.amdgcn.ldexp.f32(float, i32) nounwind readnone
-declare double @llvm.amdgcn.ldexp.f64(double, i32) nounwind readnone
-
-; SI-LABEL: {{^}}test_ldexp_f32:
-; SI: v_ldexp_f32
-; SI: s_endpgm
-define amdgpu_kernel void @test_ldexp_f32(ptr addrspace(1) %out, float %a, i32 %b) nounwind {
-  %result = call float @llvm.amdgcn.ldexp.f32(float %a, i32 %b) nounwind readnone
-  store float %result, ptr addrspace(1) %out, align 4
-  ret void
-}
-
-; SI-LABEL: {{^}}test_ldexp_f64:
-; SI: v_ldexp_f64
-; SI: s_endpgm
-define amdgpu_kernel void @test_ldexp_f64(ptr addrspace(1) %out, double %a, i32 %b) nounwind {
-  %result = call double @llvm.amdgcn.ldexp.f64(double %a, i32 %b) nounwind readnone
-  store double %result, ptr addrspace(1) %out, align 8
-  ret void
-}
-
-; SI-LABEL: {{^}}test_ldexp_undef_f32:
-; SI-NOT: v_ldexp_f32
-define amdgpu_kernel void @test_ldexp_undef_f32(ptr addrspace(1) %out, i32 %b) nounwind {
-  %result = call float @llvm.amdgcn.ldexp.f32(float undef, i32 %b) nounwind readnone
-  store float %result, ptr addrspace(1) %out, align 4
-  ret void
-}