IR: Add llvm.ldexp and llvm.experimental.constrained.ldexp intrinsics

AMDGPU has native instructions and target intrinsics for this, but these really should be subject to legalization and generic optimizations. This will enable legalization of f16->f32 on targets without f16 support. Implement a somewhat horrible inline expansion for targets without libcall support. This could be better if we could introduce control flow (GlobalISel version not yet implemented). Support for strictfp legalization is less complete but works for the simple cases.

Author: arsenm

clang/lib/CodeGen/CGBuiltin.cpp

@@ -17146,8 +17146,13 @@ Value *CodeGenFunction::EmitAMDGPUBuiltinExpr(unsigned BuiltinID,
  return emitUnaryBuiltin(*this, E, Intrinsic::amdgcn_log_clamp);
  case AMDGPU::BI__builtin_amdgcn_ldexp:
  case AMDGPU::BI__builtin_amdgcn_ldexpf:
- case AMDGPU::BI__builtin_amdgcn_ldexph:
- return emitFPIntBuiltin(*this, E, Intrinsic::amdgcn_ldexp);
+ case AMDGPU::BI__builtin_amdgcn_ldexph: {
+ llvm::Value *Src0 = EmitScalarExpr(E->getArg(0));
+ llvm::Value *Src1 = EmitScalarExpr(E->getArg(1));
+ llvm::Function *F =
+ CGM.getIntrinsic(Intrinsic::ldexp, {Src0->getType(), Src1->getType()});
+ return Builder.CreateCall(F, {Src0, Src1});
+ }
  case AMDGPU::BI__builtin_amdgcn_frexp_mant:
  case AMDGPU::BI__builtin_amdgcn_frexp_mantf:
  case AMDGPU::BI__builtin_amdgcn_frexp_manth:

clang/test/CodeGenOpenCL/builtins-amdgcn-vi.cl

@@ -52,7 +52,7 @@ void test_cos_f16(global half* out, half a)
 }
 
 // CHECK-LABEL: @test_ldexp_f16
-// CHECK: call half @llvm.amdgcn.ldexp.f16
+// CHECK: call half @llvm.ldexp.f16.i32
 void test_ldexp_f16(global half* out, half a, int b)
 {
  *out = __builtin_amdgcn_ldexph(a, b);

clang/test/CodeGenOpenCL/builtins-amdgcn.cl

@@ -180,14 +180,14 @@ void test_log_clamp_f32(global float* out, float a)
 }
 
 // CHECK-LABEL: @test_ldexp_f32
-// CHECK: call float @llvm.amdgcn.ldexp.f32
+// CHECK: call float @llvm.ldexp.f32.i32
 void test_ldexp_f32(global float* out, float a, int b)
 {
  *out = __builtin_amdgcn_ldexpf(a, b);
 }
 
 // CHECK-LABEL: @test_ldexp_f64
-// CHECK: call double @llvm.amdgcn.ldexp.f64
+// CHECK: call double @llvm.ldexp.f64.i32
 void test_ldexp_f64(global double* out, double a, int b)
 {
  *out = __builtin_amdgcn_ldexp(a, b);

llvm/docs/LangRef.rst

@@ -14713,6 +14713,47 @@ trapping or setting ``errno``.
 When specified with the fast-math-flag 'afn', the result may be approximated
 using a less accurate calculation.
 
+'``llvm.ldexp.*``' Intrinsic
+^^^^^^^^^^^^^^^^^^^^^^^^^^^
+
+Syntax:
+"""""""
+
+This is an overloaded intrinsic. You can use ``llvm.ldexp`` on any
+floating point or vector of floating point type. Not all targets support
+all types however.
+
+::
+
+ declare float @llvm.ldexp.f32.i32(float %Val, i32 %Exp)
+ declare double @llvm.ldexp.f64.i32(double %Val, i32 %Exp)
+ declare x86_fp80 @llvm.ldexp.f80.i32(x86_fp80 %Val, i32 %Exp)
+ declare fp128 @llvm.ldexp.f128.i32(fp128 %Val, i32 %Exp)
+ declare ppc_fp128 @llvm.ldexp.ppcf128.i32(ppc_fp128 %Val, i32 %Exp)
+ declare <2 x float> @llvm.ldexp.v2f32.v2i32(<2 x float> %Val, <2 x i32> %Exp)
+
+Overview:
+"""""""""
+
+The '``llvm.ldexp.*``' intrinsics perform the ldexp function.
+
+Arguments:
+""""""""""
+
+The first argument and the return value are :ref:`floating-point
+<t_floating>` or :ref:`vector <t_vector>` of floating-point values of
+the same type. The second argument is an integer with the same number
+of elements.
+
+Semantics:
+""""""""""
+
+This function multiplies the first argument by 2 raised to the second
+argument's power. If the first argument is NaN or infinite, the same
+value is returned. If the result underflows a zero with the same sign
+is returned. If the result overflows, the result is an infinity with
+the same sign.
+
 '``llvm.log.*``' Intrinsic
 ^^^^^^^^^^^^^^^^^^^^^^^^^^
 
@@ -24306,6 +24347,47 @@ This function returns the first value raised to the second power with an
 unspecified sequence of rounding operations.
 
 
+'``llvm.experimental.constrained.ldexp``' Intrinsic
+^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
+
+Syntax:
+"""""""
+
+::
+
+ declare <type0>
+ @llvm.experimental.constrained.ldexp(<type0> <op1>, <type1> <op2>,
+ metadata <rounding mode>,
+ metadata <exception behavior>)
+
+Overview:
+"""""""""
+
+The '``llvm.experimental.constrained.ldexp``' performs the ldexp function.
+
+
+Arguments:
+""""""""""
+
+The first argument and the return value are :ref:`floating-point
+<t_floating>` or :ref:`vector <t_vector>` of floating-point values of
+the same type. The second argument is an integer with the same number
+of elements.
+
+
+The third and fourth arguments specify the rounding mode and exception
+behavior as described above.
+
+Semantics:
+""""""""""
+
+This function multiplies the first argument by 2 raised to the second
+argument's power. If the first argument is NaN or infinite, the same
+value is returned. If the result underflows a zero with the same sign
+is returned. If the result overflows, the result is an infinity with
+the same sign.
+
+
 '``llvm.experimental.constrained.sin``' Intrinsic
 ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
 

llvm/docs/ReleaseNotes.rst

@@ -61,6 +61,8 @@ Changes to the LLVM IR
 * The ``nofpclass`` attribute was introduced. This allows more
  optimizations around special floating point value comparisons.
 
+* Introduced new ``llvm.ldexp`` and ``llvm.experimental.constrained.ldexp`` intrinsics.
+
 * The constant expression variants of the following instructions have been
  removed:
 

llvm/include/llvm/Analysis/TargetLibraryInfo.h

@@ -378,6 +378,7 @@ class TargetLibraryInfo {
  case LibFunc_trunc: case LibFunc_truncf: case LibFunc_truncl:
  case LibFunc_log2: case LibFunc_log2f: case LibFunc_log2l:
  case LibFunc_exp2: case LibFunc_exp2f: case LibFunc_exp2l:
+ case LibFunc_ldexp: case LibFunc_ldexpf: case LibFunc_ldexpl:
  case LibFunc_memcpy: case LibFunc_memset: case LibFunc_memmove:
  case LibFunc_memcmp: case LibFunc_bcmp: case LibFunc_strcmp:
  case LibFunc_strcpy: case LibFunc_stpcpy: case LibFunc_strlen:

llvm/include/llvm/CodeGen/GlobalISel/LegalizerHelper.h

@@ -357,6 +357,7 @@ class LegalizerHelper {
  LegalizeResult narrowScalarCTLZ(MachineInstr &MI, unsigned TypeIdx, LLT Ty);
  LegalizeResult narrowScalarCTTZ(MachineInstr &MI, unsigned TypeIdx, LLT Ty);
  LegalizeResult narrowScalarCTPOP(MachineInstr &MI, unsigned TypeIdx, LLT Ty);
+ LegalizeResult narrowScalarFLDEXP(MachineInstr &MI, unsigned TypeIdx, LLT Ty);
 
  /// Perform Bitcast legalize action on G_EXTRACT_VECTOR_ELT.
  LegalizeResult bitcastExtractVectorElt(MachineInstr &MI, unsigned TypeIdx,

llvm/include/llvm/CodeGen/GlobalISel/MachineIRBuilder.h

@@ -1825,6 +1825,13 @@ class MachineIRBuilder {
  return buildInstr(TargetOpcode::G_FPOW, {Dst}, {Src0, Src1}, Flags);
  }
 
+ /// Build and insert \p Dst = G_FLDEXP \p Src0, \p Src1
+ MachineInstrBuilder
+ buildFLdexp(const DstOp &Dst, const SrcOp &Src0, const SrcOp &Src1,
+ std::optional<unsigned> Flags = std::nullopt) {
+ return buildInstr(TargetOpcode::G_FLDEXP, {Dst}, {Src0, Src1}, Flags);
+ }
+
  /// Build and insert \p Res = G_FCOPYSIGN \p Op0, \p Op1
  MachineInstrBuilder buildFCopysign(const DstOp &Dst, const SrcOp &Src0,
  const SrcOp &Src1) {

llvm/include/llvm/CodeGen/ISDOpcodes.h

@@ -411,6 +411,7 @@ enum NodeType {
  STRICT_FSQRT,
  STRICT_FPOW,
  STRICT_FPOWI,
+ STRICT_FLDEXP,
  STRICT_FSIN,
  STRICT_FCOS,
  STRICT_FEXP,
@@ -926,8 +927,10 @@ enum NodeType {
  FCBRT,
  FSIN,
  FCOS,
- FPOWI,
  FPOW,
+ FPOWI,
+ /// FLDEXP - ldexp, inspired by libm (op0 * 2**op1).
+ FLDEXP,
  FLOG,
  FLOG2,
  FLOG10,

llvm/include/llvm/CodeGen/RuntimeLibcalls.h

@@ -70,6 +70,10 @@ namespace RTLIB {
  /// UNKNOWN_LIBCALL if there is none.
  Libcall getPOWI(EVT RetVT);
 
+ /// getLDEXP - Return the LDEXP_* value for the given types, or
+ /// UNKNOWN_LIBCALL if there is none.
+ Libcall getLDEXP(EVT RetVT);
+
  /// Return the SYNC_FETCH_AND_* value for the given opcode and type, or
  /// UNKNOWN_LIBCALL if there is none.
  Libcall getSYNC(unsigned Opc, MVT VT);