AMDGPU: Add llvm.amdgcn.exp2 intrinsic

[arsenm]

Summary

Provide direct access to v_exp_f32 and v_exp_f16, so we can start
correctly lowering the generic exp intrinsics.

Unfortunately have to break from the usual naming convention of
matching the instruction name and stripping the v_ prefix. exp is
already taken by the export intrinsic. On the clang builtin side, we
have a choice of maintaining the convention to the instruction name,
or following the intrinsic name.

[arsenm]

Fix patch split

[foad]

llvm/lib/Target/AMDGPU/AMDGPUISelLowering.h

463

Is this used anywhere?

[arsenm]

llvm/lib/Target/AMDGPU/AMDGPUISelLowering.h

463

Will be in the next patch

[arsenm]

Release notes

[arsenm]

28f3edd2be2a27b9dc7739d137147007c5fd9e41

clang/include/clang/Basic/BuiltinsAMDGPU.def

 BUILTIN(__builtin_amdgcn_sqrtf, "ff", "nc")
 BUILTIN(__builtin_amdgcn_rsq, "dd", "nc")
 BUILTIN(__builtin_amdgcn_rsqf, "ff", "nc")
 BUILTIN(__builtin_amdgcn_rsq_clamp, "dd", "nc")
 BUILTIN(__builtin_amdgcn_rsq_clampf, "ff", "nc")
 BUILTIN(__builtin_amdgcn_sinf, "ff", "nc")
 BUILTIN(__builtin_amdgcn_cosf, "ff", "nc")
 BUILTIN(__builtin_amdgcn_logf, "ff", "nc")
+BUILTIN(__builtin_amdgcn_exp2f, "ff", "nc")
 BUILTIN(__builtin_amdgcn_log_clampf, "ff", "nc")
 BUILTIN(__builtin_amdgcn_ldexp, "ddi", "nc")
 BUILTIN(__builtin_amdgcn_ldexpf, "ffi", "nc")
 BUILTIN(__builtin_amdgcn_frexp_mant, "dd", "nc")
 BUILTIN(__builtin_amdgcn_frexp_mantf, "ff", "nc")
 BUILTIN(__builtin_amdgcn_frexp_exp, "id", "nc")
 BUILTIN(__builtin_amdgcn_frexp_expf, "if", "nc")
 BUILTIN(__builtin_amdgcn_fract, "dd", "nc")

clang/lib/CodeGen/CGBuiltin.cpp

case AMDGPU::BI__builtin_amdgcn_sinh:
     return emitUnaryBuiltin(*this, E, Intrinsic::amdgcn_sin);
   case AMDGPU::BI__builtin_amdgcn_cosf:
   case AMDGPU::BI__builtin_amdgcn_cosh:
     return emitUnaryBuiltin(*this, E, Intrinsic::amdgcn_cos);
   case AMDGPU::BI__builtin_amdgcn_dispatch_ptr:
     return EmitAMDGPUDispatchPtr(*this, E);
   case AMDGPU::BI__builtin_amdgcn_logf:
     return emitUnaryBuiltin(*this, E, Intrinsic::amdgcn_log);
+  case AMDGPU::BI__builtin_amdgcn_exp2f:
+    return emitUnaryBuiltin(*this, E, Intrinsic::amdgcn_exp2);
   case AMDGPU::BI__builtin_amdgcn_log_clampf:
     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: {
     llvm::Value *Src0 = EmitScalarExpr(E->getArg(0));
     llvm::Value *Src1 = EmitScalarExpr(E->getArg(1));
     llvm::Function *F =

clang/test/CodeGenOpenCL/builtins-amdgcn.cl

 
 // CHECK-LABEL: @test_log_f32
 // CHECK: call float @llvm.amdgcn.log.f32
 void test_log_f32(global float* out, float a)
 {
   *out = __builtin_amdgcn_logf(a);
 }
 
+// CHECK-LABEL: @test_exp2_f32
+// CHECK: call float @llvm.amdgcn.exp2.f32
+void test_exp2_f32(global float* out, float a)
+{
+  *out = __builtin_amdgcn_exp2f(a);
+}
+
 // CHECK-LABEL: @test_log_clamp_f32
 // CHECK: call float @llvm.amdgcn.log.clamp.f32
 void test_log_clamp_f32(global float* out, float a)
 {
   *out = __builtin_amdgcn_log_clampf(a);
 }
 
 // CHECK-LABEL: @test_ldexp_f32

llvm/docs/AMDGPUUsage.rst

.. table:: AMDGPU LLVM IR Intrinsics
 
   :ref:`llvm.log <int_log>`                  Implemented for float and half (and vectors).
 
   :ref:`llvm.log10 <int_log10>`              Implemented for float and half (and vectors).
 
   llvm.amdgcn.log                            Provides direct access to v_log_f32 and v_log_f16
                                              (on targets with half support). Peforms log2 function.
 
+  llvm.amdgcn.exp2                           Provides direct access to v_exp_f32 and v_exp_f16
+                                             (on targets with half support). Performs exp2 function.
+
   =========================================  ==========================================================
 
 
 .. TODO::
 
    List AMDGPU intrinsics.
 
 LLVM IR Attributes

llvm/docs/ReleaseNotes.rst

 
 * Added llvm.amdgcn.log.f32 intrinsic. This provides direct access to
   v_log_f32.
 
 * llvm.log2.f32, llvm.log10.f32, and llvm.log.f32 are now lowered
   accurately. Use llvm.amdgcn.log.f32 to access the old behavior for
   llvm.log2.f32.
 
+* Added llvm.amdgcn.exp2.f32 intrinsic. This provides direct access to
+  v_exp_f32.
+
+
 Changes to the ARM Backend
 --------------------------
 
 - The hard-float ABI is now available in Armv8.1-M configurations that
   have integer MVE instructions (and therefore have FP registers) but
   no scalar or vector floating point computation.
 
 - The ``.arm`` directive now aligns code to the next 4-byte boundary, and

llvm/include/llvm/IR/IntrinsicsAMDGPU.td

 // v_log_{f16|f32}, performs log2. f32 version does not handle
 // denormals. There is no reason to use this for f16 as it does
 // support denormals, and the generic log2 intrinsic should be
 // preferred.
 def int_amdgcn_log : DefaultAttrsIntrinsic<
   [llvm_anyfloat_ty], [LLVMMatchType<0>], [IntrNoMem, IntrSpeculatable]
 >;
 
+// v_exp_{f16|f32} (int_amdgcn_exp was taken by export
+// already). Performs exp2. f32 version does not handle
+// denormals. There is no reason to use this for f16 as it does
+// support denormals, and the generic exp2 intrinsic should be
+// preferred.
+def int_amdgcn_exp2 : DefaultAttrsIntrinsic<
+  [llvm_anyfloat_ty], [LLVMMatchType<0>], [IntrNoMem, IntrSpeculatable]
+>;
+
 def int_amdgcn_log_clamp : DefaultAttrsIntrinsic<
   [llvm_anyfloat_ty], [LLVMMatchType<0>], [IntrNoMem, IntrSpeculatable]
 >;
 
 def int_amdgcn_fmul_legacy : ClangBuiltin<"__builtin_amdgcn_fmul_legacy">,
   DefaultAttrsIntrinsic<[llvm_float_ty], [llvm_float_ty, llvm_float_ty],
   [IntrNoMem, IntrSpeculatable, Commutative]
 >;

llvm/lib/Target/AMDGPU/AMDGPUISelLowering.h

enum NodeType : unsigned {
   RCP,
   RSQ,
   RCP_LEGACY,
   RCP_IFLAG,
 
   // log2, no denormal handling for f32.
   LOG,
 
+  // exp2, no denormal handling for f32.
+  EXP,
+
   FMUL_LEGACY,
   RSQ_CLAMP,
   FP_CLASS,
   DOT4,
   CARRY,
   BORROW,
   BFE_U32,  // Extract range of bits with zero extension to 32-bits.
   BFE_I32,  // Extract range of bits with sign extension to 32-bits.

llvm/lib/Target/AMDGPU/AMDGPUISelLowering.cpp

const char* AMDGPUTargetLowering::getTargetNodeName(unsigned Opcode) const {
   NODE_NAME_CASE(DIV_FMAS)
   NODE_NAME_CASE(DIV_FIXUP)
   NODE_NAME_CASE(FMAD_FTZ)
   NODE_NAME_CASE(RCP)
   NODE_NAME_CASE(RSQ)
   NODE_NAME_CASE(RCP_LEGACY)
   NODE_NAME_CASE(RCP_IFLAG)
   NODE_NAME_CASE(LOG)
+  NODE_NAME_CASE(EXP)
   NODE_NAME_CASE(FMUL_LEGACY)
   NODE_NAME_CASE(RSQ_CLAMP)
   NODE_NAME_CASE(FP_CLASS)
   NODE_NAME_CASE(DOT4)
   NODE_NAME_CASE(CARRY)
   NODE_NAME_CASE(BORROW)
   NODE_NAME_CASE(BFE_U32)
   NODE_NAME_CASE(BFE_I32)

llvm/lib/Target/AMDGPU/AMDGPUInstrInfo.td

 def AMDGPUfract_impl : SDNode<"AMDGPUISD::FRACT", SDTFPUnaryOp>;
 
 // out = 1.0 / a
 def AMDGPUrcp_impl : SDNode<"AMDGPUISD::RCP", SDTFPUnaryOp>;
 
 // v_log_f32, which is log2
 def AMDGPUlog_impl : SDNode<"AMDGPUISD::LOG", SDTFPUnaryOp>;
 
+// v_exp_f32, which is exp2
+def AMDGPUexp_impl : SDNode<"AMDGPUISD::EXP", SDTFPUnaryOp>;
+
 // out = 1.0 / sqrt(a)
 def AMDGPUrsq_impl : SDNode<"AMDGPUISD::RSQ", SDTFPUnaryOp>;
 
 def AMDGPUrcp_legacy_impl : SDNode<"AMDGPUISD::RCP_LEGACY", SDTFPUnaryOp>;
 
 def AMDGPUrcp_iflag : SDNode<"AMDGPUISD::RCP_IFLAG", SDTFPUnaryOp>;
 
 // out = 1.0 / sqrt(a) result clamped to +/- max_float.
def AMDGPUcos : PatFrags<(ops node:$src), [(int_amdgcn_cos node:$src),
                                            (AMDGPUcos_impl node:$src)]>;
 def AMDGPUfract : PatFrags<(ops node:$src), [(int_amdgcn_fract node:$src),
                                              (AMDGPUfract_impl node:$src)]>;
 def AMDGPUlog : PatFrags<(ops node:$src), [(int_amdgcn_log node:$src),
                                            (AMDGPUlog_impl node:$src)]>;
 def AMDGPUlogf16 : PatFrags<(ops node:$src), [(int_amdgcn_log node:$src),
                                               (flog2 node:$src)]>;
 
+def AMDGPUexp : PatFrags<(ops node:$src), [(int_amdgcn_exp2 node:$src),
+                                           (AMDGPUexp_impl node:$src),
+                                           (fexp2 node:$src)]>; // FIXME: Remove me
+def AMDGPUexpf16 : PatFrags<(ops node:$src), [(int_amdgcn_exp2 node:$src),
+                                              (fexp2 node:$src)]>;
 
 def AMDGPUfp_class : PatFrags<(ops node:$src0, node:$src1),
   [(int_amdgcn_class node:$src0, node:$src1),
    (AMDGPUfp_class_impl node:$src0, node:$src1)]>;
 
 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)]>;

llvm/lib/Target/AMDGPU/AMDGPURegisterBankInfo.cpp

default:
       return getInvalidInstructionMapping();
     case Intrinsic::amdgcn_div_fmas:
     case Intrinsic::amdgcn_div_fixup:
     case Intrinsic::amdgcn_trig_preop:
     case Intrinsic::amdgcn_sin:
     case Intrinsic::amdgcn_cos:
     case Intrinsic::amdgcn_log_clamp:
     case Intrinsic::amdgcn_log:
+    case Intrinsic::amdgcn_exp2:
     case Intrinsic::amdgcn_rcp:
     case Intrinsic::amdgcn_rcp_legacy:
     case Intrinsic::amdgcn_sqrt:
     case Intrinsic::amdgcn_rsq:
     case Intrinsic::amdgcn_rsq_legacy:
     case Intrinsic::amdgcn_rsq_clamp:
     case Intrinsic::amdgcn_fmul_legacy:
     case Intrinsic::amdgcn_fma_legacy:

llvm/lib/Target/AMDGPU/R600Instructions.td

> {
     let UseNamedOperandTable = 1;
   }
 
   def _real : R600_2OP <inst, "CUBE", []>;
 }
 } // End mayLoad = 0, mayStore = 0, hasSideEffects = 0
 
 class EXP_IEEE_Common <bits<11> inst> : R600_1OP_Helper <
-  inst, "EXP_IEEE", fexp2
+  inst, "EXP_IEEE", AMDGPUexp
 > {
   let Itinerary = TransALU;
 }
 
 class FLT_TO_INT_Common <bits<11> inst> : R600_1OP_Helper <
   inst, "FLT_TO_INT", fp_to_sint
 > {
   let Itinerary = TransALU;

llvm/lib/Target/AMDGPU/VOP1Instructions.td

 
 defm V_FRACT_F32 : VOP1Inst <"v_fract_f32", VOP_F32_F32, AMDGPUfract>;
 defm V_TRUNC_F32 : VOP1Inst <"v_trunc_f32", VOP_F32_F32, ftrunc>;
 defm V_CEIL_F32 : VOP1Inst <"v_ceil_f32", VOP_F32_F32, fceil>;
 defm V_RNDNE_F32 : VOP1Inst <"v_rndne_f32", VOP_F32_F32, frint>;
 defm V_FLOOR_F32 : VOP1Inst <"v_floor_f32", VOP_F32_F32, ffloor>;
 
 let TRANS = 1, SchedRW = [WriteTrans32] in {
-defm V_EXP_F32 : VOP1Inst <"v_exp_f32", VOP_F32_F32, fexp2>;
+defm V_EXP_F32 : VOP1Inst <"v_exp_f32", VOP_F32_F32, AMDGPUexp>;
 defm V_LOG_F32 : VOP1Inst <"v_log_f32", VOP_F32_F32, AMDGPUlog>;
 defm V_RCP_F32 : VOP1Inst <"v_rcp_f32", VOP_F32_F32, AMDGPUrcp>;
 defm V_RCP_IFLAG_F32 : VOP1Inst <"v_rcp_iflag_f32", VOP_F32_F32, AMDGPUrcp_iflag>;
 defm V_RSQ_F32 : VOP1Inst <"v_rsq_f32", VOP_F32_F32, AMDGPUrsq>;
 defm V_SQRT_F32 : VOP1Inst <"v_sqrt_f32", VOP_F32_F32, any_amdgcn_sqrt>;
 } // End TRANS = 1, SchedRW = [WriteTrans32]
 
 let TRANS = 1, SchedRW = [WriteTrans64] in {

 defm V_CVT_U16_F16_t16 : VOP1Inst <"v_cvt_u16_f16_t16", VOP_I16_F16_SPECIAL_OMOD_t16, fp_to_uint>;
 defm V_CVT_I16_F16_t16 : VOP1Inst <"v_cvt_i16_f16_t16", VOP_I16_F16_SPECIAL_OMOD_t16, fp_to_sint>;
 }
 let TRANS = 1, SchedRW = [WriteTrans32] in {
 defm V_RCP_F16 : VOP1Inst_t16 <"v_rcp_f16", VOP_F16_F16, AMDGPUrcp>;
 defm V_SQRT_F16 : VOP1Inst_t16 <"v_sqrt_f16", VOP_F16_F16, any_amdgcn_sqrt>;
 defm V_RSQ_F16 : VOP1Inst_t16 <"v_rsq_f16", VOP_F16_F16, AMDGPUrsq>;
 defm V_LOG_F16 : VOP1Inst_t16 <"v_log_f16", VOP_F16_F16, AMDGPUlogf16>;
-defm V_EXP_F16 : VOP1Inst_t16 <"v_exp_f16", VOP_F16_F16, fexp2>;
+defm V_EXP_F16 : VOP1Inst_t16 <"v_exp_f16", VOP_F16_F16, AMDGPUexpf16>;
 defm V_SIN_F16 : VOP1Inst_t16 <"v_sin_f16", VOP_F16_F16, AMDGPUsin>;
 defm V_COS_F16 : VOP1Inst_t16 <"v_cos_f16", VOP_F16_F16, AMDGPUcos>;
 } // End TRANS = 1, SchedRW = [WriteTrans32]
 defm V_FREXP_MANT_F16 : VOP1Inst_t16 <"v_frexp_mant_f16", VOP_F16_F16, int_amdgcn_frexp_mant>;
 let OtherPredicates = [Has16BitInsts, NotHasTrue16BitInsts] in {
 defm V_FREXP_EXP_I16_F16 : VOP1Inst <"v_frexp_exp_i16_f16", VOP_I16_F16_SPECIAL_OMOD, int_amdgcn_frexp_exp>;
 }
 let OtherPredicates = [HasTrue16BitInsts] in {

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

+; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 2
+; RUN: llc -global-isel=0 -mtriple=amdgcn-amd-amdhsa -mcpu=fiji < %s | FileCheck -check-prefixes=GCN,SDAG %s
+; RUN: llc -global-isel=1 -mtriple=amdgcn-amd-amdhsa -mcpu=fiji < %s | FileCheck -check-prefixes=GCN,GISEL %s
+
+define float @v_exp2_f32(float %src)  {
+; GCN-LABEL: v_exp2_f32:
+; GCN:       ; %bb.0:
+; GCN-NEXT:    s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
+; GCN-NEXT:    v_exp_f32_e32 v0, v0
+; GCN-NEXT:    s_setpc_b64 s[30:31]
+  %exp2 = call float @llvm.amdgcn.exp2.f32(float %src)
+  ret float %exp2
+}
+
+define float @v_fabs_exp2_f32(float %src)  {
+; GCN-LABEL: v_fabs_exp2_f32:
+; GCN:       ; %bb.0:
+; GCN-NEXT:    s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
+; GCN-NEXT:    v_exp_f32_e64 v0, |v0|
+; GCN-NEXT:    s_setpc_b64 s[30:31]
+  %fabs.src = call float @llvm.fabs.f32(float %src)
+  %exp2 = call float @llvm.amdgcn.exp2.f32(float %fabs.src)
+  ret float %exp2
+}
+
+define float @v_fneg_fabs_exp2_f32(float %src)  {
+; GCN-LABEL: v_fneg_fabs_exp2_f32:
+; GCN:       ; %bb.0:
+; GCN-NEXT:    s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
+; GCN-NEXT:    v_exp_f32_e64 v0, -|v0|
+; GCN-NEXT:    s_setpc_b64 s[30:31]
+  %fabs.src = call float @llvm.fabs.f32(float %src)
+  %neg.fabs.src = fneg float %fabs.src
+  %exp2 = call float @llvm.amdgcn.exp2.f32(float %neg.fabs.src)
+  ret float %exp2
+}
+
+define half @v_exp2_f16(half %src)  {
+; GCN-LABEL: v_exp2_f16:
+; GCN:       ; %bb.0:
+; GCN-NEXT:    s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
+; GCN-NEXT:    v_exp_f16_e32 v0, v0
+; GCN-NEXT:    s_setpc_b64 s[30:31]
+  %exp2 = call half @llvm.amdgcn.exp2.f16(half %src)
+  ret half %exp2
+}
+
+define half @v_fabs_exp2_f16(half %src)  {
+; GCN-LABEL: v_fabs_exp2_f16:
+; GCN:       ; %bb.0:
+; GCN-NEXT:    s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
+; GCN-NEXT:    v_exp_f16_e64 v0, |v0|
+; GCN-NEXT:    s_setpc_b64 s[30:31]
+  %fabs.src = call half @llvm.fabs.f16(half %src)
+  %exp2 = call half @llvm.amdgcn.exp2.f16(half %fabs.src)
+  ret half %exp2
+}
+
+define half @v_fneg_fabs_exp2_f16(half %src)  {
+; GCN-LABEL: v_fneg_fabs_exp2_f16:
+; GCN:       ; %bb.0:
+; GCN-NEXT:    s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
+; GCN-NEXT:    v_exp_f16_e64 v0, -|v0|
+; GCN-NEXT:    s_setpc_b64 s[30:31]
+  %fabs.src = call half @llvm.fabs.f16(half %src)
+  %neg.fabs.src = fneg half %fabs.src
+  %exp2 = call half @llvm.amdgcn.exp2.f16(half %neg.fabs.src)
+  ret half %exp2
+}
+
+declare half @llvm.amdgcn.exp2.f16(half) #0
+declare float @llvm.amdgcn.exp2.f32(float) #0
+declare float @llvm.fabs.f32(float) #0
+declare half @llvm.fabs.f16(half) #0
+
+attributes #0 = { nounwind readnone speculatable willreturn }
+;; NOTE: These prefixes are unused and the list is autogenerated. Do not add tests below this line:
+; GISEL: {{.*}}
+; SDAG: {{.*}}