[AMDGPU] add LDS f32 intrinsics

[dfukalov]

Summary

added llvm.amdgcn.atomic.{add|min|max}.f32 intrinsics
to allow generate ds_{add|min|max}[_rtn]_f32 instructions
needed for OpenCL float atomics in LDS

[rampitec]

LGTM

[arsenm]

Since it's atomic it needs to be added to isIntrinsicSourceOfDivergence. Also getTgtMemIntrinsic, and getAddrModeArguments (all with associated tests)

include/llvm/IR/IntrinsicsAMDGPU.td

304

Should this have an operand added for the ordering?

305

This is certainly not IntrNoReturn

[dfukalov]

include/llvm/IR/IntrinsicsAMDGPU.td

304

No, these intrinsics are created by request to be able to generate ds_{add|min|max}[_rtn]_f32 in case of OpenCL local memory atomics only. They work only for pointers to floats located in addrspace 3

305

You're right, I interpreted it as intrinsic that has no return value, diff updated

[arsenm]

include/llvm/IR/IntrinsicsAMDGPU.td

304

That doesn't change the ordering. Also needs an operand for volatile

[t-tye]

include/llvm/IR/IntrinsicsAMDGPU.td

304

How are the memory ordering, memory scope and volatile carried through so that those fields can be set in the Machine Memory Operand? All these properties are needed to generate the correct waitcnt in the memory legalizer (see AMDGPUUsage.rst section on memory model).

[arsenm]

include/llvm/IR/IntrinsicsAMDGPU.td

304

Those are supposed to be handled by the intrinsic callbacks I mentioned that need to be implemented (at least for volatile). I'm not sure anything really correctly considers the atomic scope for possibly atomic intrinsics.

The most similar case we have is for amdgcn_atomic_inc/dec.

[t-tye]

include/llvm/IR/IntrinsicsAMDGPU.td

301

Need to add the same fields as for AMDGPUAtomicIncIntrin, namely:

llvm_i32_ty, // ordering
llvm_i32_ty, // scope
llvm_i1_ty], // isVolatile

306

Need to add the same fields as for AMDGPUAtomicIncIntrin, namely:

llvm_i32_ty, // ordering
llvm_i32_ty, // scope
llvm_i1_ty], // isVolatile

lib/Target/AMDGPU/DSInstructions.td

776–785

Also, somehow a MachineMemoryOperand needs to be created with values from the ordering, scope and isVolatile LLVM IR instruction operands.

[arsenm]

This also should have an intermediate node added, so it is easy to add it to the various switches over memory operations (e.g. we look for all memory operations in PerformDAGCombine)

[dfukalov]

Am I right that since we should have almost the same processing as atomic inc intrinsics,
it would be better idea to define ds_add/min/max intrinsics the same way as AMDGPUAtomicIncIntrin (or unify them),
and then update AMDGPU BE to correctly process these ds_ instrisics the same ways as atomic inc/dec?

[dfukalov]

updates as requested by reviewers

[dfukalov]

ping

[arsenm]

include/llvm/IR/IntrinsicsAMDGPU.td

307

These should probably be named fadd.. etc to match the IR operations

lib/Target/AMDGPU/SIISelLowering.cpp

575–793

This will need to be rebased since I just changed these last week

6736–6740

Tests for these with this combine would be nice

test/CodeGen/AMDGPU/lds_atomic_f32.ll

2

Can you also add a pre-gfx9 run line, and check for the m0 initialization

[dfukalov]

diff updated according to latest comments

[arsenm]

LGTM except for the mystery name mangling in the test

test/CodeGen/AMDGPU/lds_atomic_f32.ll

4

These weren't declared as mangled, so why do these have the .f32? I'm kind of surprised this is accepted

include/llvm/IR/IntrinsicsAMDGPU.td

class AMDGPUAtomicIncIntrin : Intrinsic<[llvm_anyint_ty],
   llvm_i1_ty], // isVolatile
   [IntrArgMemOnly, NoCapture<0>], "",
   [SDNPMemOperand]
 >;
 
 def int_amdgcn_atomic_inc : AMDGPUAtomicIncIntrin;
 def int_amdgcn_atomic_dec : AMDGPUAtomicIncIntrin;
 
+class AMDGPUAtomicF32Intrin<string clang_builtin> :
+  GCCBuiltin<clang_builtin>,
+  Intrinsic<[llvm_float_ty],
+    [LLVMAnyPointerType<llvm_float_ty>,
+    llvm_float_ty,
+    llvm_i32_ty, // ordering
+    llvm_i32_ty, // scope
+    llvm_i1_ty], // isVolatile
+    [IntrArgMemOnly, NoCapture<0>]
+>;
+
+def int_amdgcn_atomic_fadd : AMDGPUAtomicF32Intrin<"__builtin_amdgcn_ds_fadd">;
+def int_amdgcn_atomic_fmin : AMDGPUAtomicF32Intrin<"__builtin_amdgcn_ds_fmin">;
+def int_amdgcn_atomic_fmax : AMDGPUAtomicF32Intrin<"__builtin_amdgcn_ds_fmax">;
+
 class AMDGPUImageLoad<bit NoMem = 0> : Intrinsic <
   [llvm_anyfloat_ty], // vdata(VGPR)
   [llvm_anyint_ty,    // vaddr(VGPR)
    llvm_anyint_ty,    // rsrc(SGPR)
    llvm_i32_ty,       // dmask(imm)
    llvm_i1_ty,        // glc(imm)
    llvm_i1_ty,        // slc(imm)
    llvm_i1_ty,        // lwe(imm)

lib/Target/AMDGPU/AMDGPUISelDAGToDAG.cpp

 void AMDGPUDAGToDAGISel::Select(SDNode *N) {
   unsigned int Opc = N->getOpcode();
   if (N->isMachineOpcode()) {
     N->setNodeId(-1);
     return;   // Already selected.
   }
 
   if (isa<AtomicSDNode>(N) ||
-      (Opc == AMDGPUISD::ATOMIC_INC || Opc == AMDGPUISD::ATOMIC_DEC))
+      (Opc == AMDGPUISD::ATOMIC_INC || Opc == AMDGPUISD::ATOMIC_DEC ||
+       Opc == AMDGPUISD::ATOMIC_LOAD_FADD ||
+       Opc == AMDGPUISD::ATOMIC_LOAD_FMIN ||
+       Opc == AMDGPUISD::ATOMIC_LOAD_FMAX))
     N = glueCopyToM0(N);
 
   switch (Opc) {
   default:
     break;
   // We are selecting i64 ADD here instead of custom lower it during
   // DAG legalization, so we can fold some i64 ADDs used for address
   // calculation into the LOAD and STORE instructions.

lib/Target/AMDGPU/AMDGPUISelLowering.h

enum NodeType : unsigned {
   STORE_MSKOR,
   LOAD_CONSTANT,
   TBUFFER_STORE_FORMAT,
   TBUFFER_STORE_FORMAT_X3,
   TBUFFER_LOAD_FORMAT,
   ATOMIC_CMP_SWAP,
   ATOMIC_INC,
   ATOMIC_DEC,
+  ATOMIC_LOAD_FADD,
+  ATOMIC_LOAD_FMIN,
+  ATOMIC_LOAD_FMAX,
   BUFFER_LOAD,
   BUFFER_LOAD_FORMAT,
   BUFFER_STORE,
   BUFFER_STORE_FORMAT,
   BUFFER_ATOMIC_SWAP,
   BUFFER_ATOMIC_ADD,
   BUFFER_ATOMIC_SUB,
   BUFFER_ATOMIC_SMIN,

lib/Target/AMDGPU/AMDGPUISelLowering.cpp

const char* AMDGPUTargetLowering::getTargetNodeName(unsigned Opcode) const {
   NODE_NAME_CASE(STORE_MSKOR)
   NODE_NAME_CASE(LOAD_CONSTANT)
   NODE_NAME_CASE(TBUFFER_STORE_FORMAT)
   NODE_NAME_CASE(TBUFFER_STORE_FORMAT_X3)
   NODE_NAME_CASE(TBUFFER_LOAD_FORMAT)
   NODE_NAME_CASE(ATOMIC_CMP_SWAP)
   NODE_NAME_CASE(ATOMIC_INC)
   NODE_NAME_CASE(ATOMIC_DEC)
+  NODE_NAME_CASE(ATOMIC_LOAD_FADD)
+  NODE_NAME_CASE(ATOMIC_LOAD_FMIN)
+  NODE_NAME_CASE(ATOMIC_LOAD_FMAX)
   NODE_NAME_CASE(BUFFER_LOAD)
   NODE_NAME_CASE(BUFFER_LOAD_FORMAT)
   NODE_NAME_CASE(BUFFER_STORE)
   NODE_NAME_CASE(BUFFER_STORE_FORMAT)
   NODE_NAME_CASE(BUFFER_ATOMIC_SWAP)
   NODE_NAME_CASE(BUFFER_ATOMIC_ADD)
   NODE_NAME_CASE(BUFFER_ATOMIC_SUB)
   NODE_NAME_CASE(BUFFER_ATOMIC_SMIN)

lib/Target/AMDGPU/AMDGPUTargetTransformInfo.cpp

static bool isIntrinsicSourceOfDivergence(const IntrinsicInst *I) {
   case Intrinsic::amdgcn_interp_p2:
   case Intrinsic::amdgcn_mbcnt_hi:
   case Intrinsic::amdgcn_mbcnt_lo:
   case Intrinsic::r600_read_tidig_x:
   case Intrinsic::r600_read_tidig_y:
   case Intrinsic::r600_read_tidig_z:
   case Intrinsic::amdgcn_atomic_inc:
   case Intrinsic::amdgcn_atomic_dec:
+  case Intrinsic::amdgcn_atomic_fadd:
+  case Intrinsic::amdgcn_atomic_fmin:
+  case Intrinsic::amdgcn_atomic_fmax:
   case Intrinsic::amdgcn_image_atomic_swap:
   case Intrinsic::amdgcn_image_atomic_add:
   case Intrinsic::amdgcn_image_atomic_sub:
   case Intrinsic::amdgcn_image_atomic_smin:
   case Intrinsic::amdgcn_image_atomic_umin:
   case Intrinsic::amdgcn_image_atomic_smax:
   case Intrinsic::amdgcn_image_atomic_umax:
   case Intrinsic::amdgcn_image_atomic_and:

lib/Target/AMDGPU/DSInstructions.td

 defm DS_MIN_RTN_U32   : DS_1A1D_RET_mc<"ds_min_rtn_u32", VGPR_32, "ds_min_u32">;
 defm DS_MAX_RTN_U32   : DS_1A1D_RET_mc<"ds_max_rtn_u32", VGPR_32, "ds_max_u32">;
 defm DS_AND_RTN_B32   : DS_1A1D_RET_mc<"ds_and_rtn_b32", VGPR_32, "ds_and_b32">;
 defm DS_OR_RTN_B32    : DS_1A1D_RET_mc<"ds_or_rtn_b32", VGPR_32, "ds_or_b32">;
 defm DS_XOR_RTN_B32   : DS_1A1D_RET_mc<"ds_xor_rtn_b32", VGPR_32, "ds_xor_b32">;
 defm DS_MSKOR_RTN_B32 : DS_1A2D_RET_mc<"ds_mskor_rtn_b32", VGPR_32, "ds_mskor_b32">;
 defm DS_CMPST_RTN_B32 : DS_1A2D_RET_mc<"ds_cmpst_rtn_b32", VGPR_32, "ds_cmpst_b32">;
 defm DS_CMPST_RTN_F32 : DS_1A2D_RET_mc<"ds_cmpst_rtn_f32", VGPR_32, "ds_cmpst_f32">;
-defm DS_MIN_RTN_F32   : DS_1A1D_RET_mc <"ds_min_rtn_f32", VGPR_32, "ds_min_f32">;
+defm DS_MIN_RTN_F32   : DS_1A1D_RET_mc<"ds_min_rtn_f32", VGPR_32, "ds_min_f32">;
 defm DS_MAX_RTN_F32   : DS_1A1D_RET_mc<"ds_max_rtn_f32", VGPR_32, "ds_max_f32">;
 
 defm DS_WRXCHG_RTN_B32 : DS_1A1D_RET_mc<"ds_wrxchg_rtn_b32">;
 defm DS_WRXCHG2_RTN_B32 : DS_1A2D_Off8_RET_mc<"ds_wrxchg2_rtn_b32", VReg_64, VGPR_32>;
 defm DS_WRXCHG2ST64_RTN_B32 : DS_1A2D_Off8_RET_mc<"ds_wrxchg2st64_rtn_b32", VReg_64, VGPR_32>;
 
 defm DS_ADD_RTN_U64  : DS_1A1D_RET_mc<"ds_add_rtn_u64", VReg_64, "ds_add_u64">;
 defm DS_SUB_RTN_U64  : DS_1A1D_RET_mc<"ds_sub_rtn_u64", VReg_64, "ds_sub_u64">;

 defm : DSAtomicRetPat_mc<DS_AND_RTN_B32, i32, "atomic_load_and_local">;
 defm : DSAtomicRetPat_mc<DS_OR_RTN_B32, i32, "atomic_load_or_local">;
 defm : DSAtomicRetPat_mc<DS_XOR_RTN_B32, i32, "atomic_load_xor_local">;
 defm : DSAtomicRetPat_mc<DS_MIN_RTN_I32, i32, "atomic_load_min_local">;
 defm : DSAtomicRetPat_mc<DS_MAX_RTN_I32, i32, "atomic_load_max_local">;
 defm : DSAtomicRetPat_mc<DS_MIN_RTN_U32, i32, "atomic_load_umin_local">;
 defm : DSAtomicRetPat_mc<DS_MAX_RTN_U32, i32, "atomic_load_umax_local">;
 defm : DSAtomicCmpXChg_mc<DS_CMPST_RTN_B32, i32, "atomic_cmp_swap_local">;
+defm : DSAtomicRetPat_mc<DS_MIN_RTN_F32, f32, "atomic_load_fmin_local">;
+defm : DSAtomicRetPat_mc<DS_MAX_RTN_F32, f32, "atomic_load_fmax_local">;
+defm : DSAtomicRetPat_mc<DS_ADD_RTN_F32, f32, "atomic_load_fadd_local">;
 
 // 64-bit atomics.
 defm : DSAtomicRetPat_mc<DS_WRXCHG_RTN_B64, i64, "atomic_swap_local">;
 defm : DSAtomicRetPat_mc<DS_ADD_RTN_U64, i64, "atomic_load_add_local">;
 defm : DSAtomicRetPat_mc<DS_SUB_RTN_U64, i64, "atomic_load_sub_local">;
 defm : DSAtomicRetPat_mc<DS_INC_RTN_U64, i64, "atomic_inc_local">;
 defm : DSAtomicRetPat_mc<DS_DEC_RTN_U64, i64, "atomic_dec_local">;
 defm : DSAtomicRetPat_mc<DS_AND_RTN_B64, i64, "atomic_load_and_local">;
 defm : DSAtomicRetPat_mc<DS_OR_RTN_B64, i64, "atomic_load_or_local">;
 defm : DSAtomicRetPat_mc<DS_XOR_RTN_B64, i64, "atomic_load_xor_local">;
 defm : DSAtomicRetPat_mc<DS_MIN_RTN_I64, i64, "atomic_load_min_local">;
 defm : DSAtomicRetPat_mc<DS_MAX_RTN_I64, i64, "atomic_load_max_local">;
 defm : DSAtomicRetPat_mc<DS_MIN_RTN_U64, i64, "atomic_load_umin_local">;
 defm : DSAtomicRetPat_mc<DS_MAX_RTN_U64, i64, "atomic_load_umax_local">;
 
 defm : DSAtomicCmpXChg_mc<DS_CMPST_RTN_B64, i64, "atomic_cmp_swap_local">;
 
 //===----------------------------------------------------------------------===//
 // Real instructions

lib/Target/AMDGPU/SIISelLowering.cpp

 }
 
 bool SITargetLowering::getTgtMemIntrinsic(IntrinsicInfo &Info,
                                           const CallInst &CI,
                                           MachineFunction &MF,
                                           unsigned IntrID) const {
   switch (IntrID) {
   case Intrinsic::amdgcn_atomic_inc:
-  case Intrinsic::amdgcn_atomic_dec: {
+  case Intrinsic::amdgcn_atomic_dec:
+  case Intrinsic::amdgcn_atomic_fadd:
+  case Intrinsic::amdgcn_atomic_fmin:
+  case Intrinsic::amdgcn_atomic_fmax: {
     Info.opc = ISD::INTRINSIC_W_CHAIN;
     Info.memVT = MVT::getVT(CI.getType());
     Info.ptrVal = CI.getOperand(0);
     Info.align = 0;
     Info.flags = MachineMemOperand::MOLoad | MachineMemOperand::MOStore;
 
     const ConstantInt *Vol = dyn_cast<ConstantInt>(CI.getOperand(4));
     if (!Vol || !Vol->isZero())
       Info.flags |= MachineMemOperand::MOVolatile;
 
     return true;
   }
 
   // Image load.
   case Intrinsic::amdgcn_image_load:
   case Intrinsic::amdgcn_image_load_mip:
 
   // Sample.
   case Intrinsic::amdgcn_image_sample:
   case Intrinsic::amdgcn_image_sample_cl:
   case Intrinsic::amdgcn_image_sample_d:
   case Intrinsic::amdgcn_image_sample_d_cl:
   case Intrinsic::amdgcn_image_sample_l:
   case Intrinsic::amdgcn_image_sample_b:
   case Intrinsic::amdgcn_image_sample_b_cl:
   case Intrinsic::amdgcn_image_sample_lz:
   case Intrinsic::amdgcn_image_sample_cd:
   case Intrinsic::amdgcn_image_sample_cd_cl:
 
     // Sample with comparison.
   case Intrinsic::amdgcn_image_sample_c:
   case Intrinsic::amdgcn_image_sample_c_cl:
   case Intrinsic::amdgcn_image_sample_c_d:
   case Intrinsic::amdgcn_image_sample_c_d_cl:
   case Intrinsic::amdgcn_image_sample_c_l:
   case Intrinsic::amdgcn_image_sample_c_b:
   case Intrinsic::amdgcn_image_sample_c_b_cl:
   case Intrinsic::amdgcn_image_sample_c_lz:
   case Intrinsic::amdgcn_image_sample_c_cd:
   case Intrinsic::amdgcn_image_sample_c_cd_cl:
 
     // Sample with offsets.
   case Intrinsic::amdgcn_image_sample_o:
   case Intrinsic::amdgcn_image_sample_cl_o:
   case Intrinsic::amdgcn_image_sample_d_o:
   case Intrinsic::amdgcn_image_sample_d_cl_o:
   case Intrinsic::amdgcn_image_sample_l_o:
   case Intrinsic::amdgcn_image_sample_b_o:
   case Intrinsic::amdgcn_image_sample_b_cl_o:
   case Intrinsic::amdgcn_image_sample_lz_o:
   case Intrinsic::amdgcn_image_sample_cd_o:
   case Intrinsic::amdgcn_image_sample_cd_cl_o:
 
     // Sample with comparison and offsets.
   case Intrinsic::amdgcn_image_sample_c_o:
   case Intrinsic::amdgcn_image_sample_c_cl_o:
   case Intrinsic::amdgcn_image_sample_c_d_o:
   case Intrinsic::amdgcn_image_sample_c_d_cl_o:
   case Intrinsic::amdgcn_image_sample_c_l_o:
   case Intrinsic::amdgcn_image_sample_c_b_o:
   case Intrinsic::amdgcn_image_sample_c_b_cl_o:
   case Intrinsic::amdgcn_image_sample_c_lz_o:
   case Intrinsic::amdgcn_image_sample_c_cd_o:
   case Intrinsic::amdgcn_image_sample_c_cd_cl_o:
 
     // Basic gather4
   case Intrinsic::amdgcn_image_gather4:
   case Intrinsic::amdgcn_image_gather4_cl:
   case Intrinsic::amdgcn_image_gather4_l:
   case Intrinsic::amdgcn_image_gather4_b:
   case Intrinsic::amdgcn_image_gather4_b_cl:
   case Intrinsic::amdgcn_image_gather4_lz:
 
     // Gather4 with comparison
   case Intrinsic::amdgcn_image_gather4_c:
   case Intrinsic::amdgcn_image_gather4_c_cl:
   case Intrinsic::amdgcn_image_gather4_c_l:
   case Intrinsic::amdgcn_image_gather4_c_b:
   case Intrinsic::amdgcn_image_gather4_c_b_cl:
   case Intrinsic::amdgcn_image_gather4_c_lz:
 
     // Gather4 with offsets
   case Intrinsic::amdgcn_image_gather4_o:
   case Intrinsic::amdgcn_image_gather4_cl_o:
   case Intrinsic::amdgcn_image_gather4_l_o:
   case Intrinsic::amdgcn_image_gather4_b_o:
   case Intrinsic::amdgcn_image_gather4_b_cl_o:
   case Intrinsic::amdgcn_image_gather4_lz_o:
 
     // Gather4 with comparison and offsets
   case Intrinsic::amdgcn_image_gather4_c_o:
   case Intrinsic::amdgcn_image_gather4_c_cl_o:
   case Intrinsic::amdgcn_image_gather4_c_l_o:
   case Intrinsic::amdgcn_image_gather4_c_b_o:
   case Intrinsic::amdgcn_image_gather4_c_b_cl_o:
   case Intrinsic::amdgcn_image_gather4_c_lz_o: {
     SIMachineFunctionInfo *MFI = MF.getInfo<SIMachineFunctionInfo>();
     Info.opc = ISD::INTRINSIC_W_CHAIN;
     Info.memVT = MVT::getVT(CI.getType());
     Info.ptrVal = MFI->getImagePSV(
       *MF.getSubtarget<SISubtarget>().getInstrInfo(),
       CI.getArgOperand(1));
     Info.align = 0;
     Info.flags = MachineMemOperand::MOLoad |
                  MachineMemOperand::MODereferenceable;
     return true;
   }
   case Intrinsic::amdgcn_image_store:
   case Intrinsic::amdgcn_image_store_mip: {
     SIMachineFunctionInfo *MFI = MF.getInfo<SIMachineFunctionInfo>();
     Info.opc = ISD::INTRINSIC_VOID;
     Info.memVT = MVT::getVT(CI.getArgOperand(0)->getType());
     Info.ptrVal = MFI->getImagePSV(
       *MF.getSubtarget<SISubtarget>().getInstrInfo(),
       CI.getArgOperand(2));
     Info.flags = MachineMemOperand::MOStore |
                  MachineMemOperand::MODereferenceable;
     Info.align = 0;
     return true;
   }
   case Intrinsic::amdgcn_image_atomic_swap:
   case Intrinsic::amdgcn_image_atomic_add:
   case Intrinsic::amdgcn_image_atomic_sub:
   case Intrinsic::amdgcn_image_atomic_smin:
   case Intrinsic::amdgcn_image_atomic_umin:
   case Intrinsic::amdgcn_image_atomic_smax:
   case Intrinsic::amdgcn_image_atomic_umax:
   case Intrinsic::amdgcn_image_atomic_and:
   case Intrinsic::amdgcn_image_atomic_or:
   case Intrinsic::amdgcn_image_atomic_xor:
   case Intrinsic::amdgcn_image_atomic_inc:
   case Intrinsic::amdgcn_image_atomic_dec: {
     SIMachineFunctionInfo *MFI = MF.getInfo<SIMachineFunctionInfo>();
     Info.opc = ISD::INTRINSIC_W_CHAIN;
     Info.memVT = MVT::getVT(CI.getType());
     Info.ptrVal = MFI->getImagePSV(
       *MF.getSubtarget<SISubtarget>().getInstrInfo(),
       CI.getArgOperand(2));
 
     Info.flags = MachineMemOperand::MOLoad |
                  MachineMemOperand::MOStore |
                  MachineMemOperand::MODereferenceable;
 
     // XXX - Should this be volatile without known ordering?
     Info.flags |= MachineMemOperand::MOVolatile;
     return true;
   }
   case Intrinsic::amdgcn_image_atomic_cmpswap: {
     SIMachineFunctionInfo *MFI = MF.getInfo<SIMachineFunctionInfo>();
     Info.opc = ISD::INTRINSIC_W_CHAIN;
     Info.memVT = MVT::getVT(CI.getType());
     Info.ptrVal = MFI->getImagePSV(
       *MF.getSubtarget<SISubtarget>().getInstrInfo(),
       CI.getArgOperand(3));
 
     Info.flags = MachineMemOperand::MOLoad |
                  MachineMemOperand::MOStore |
                  MachineMemOperand::MODereferenceable;
 
     // XXX - Should this be volatile without known ordering?
     Info.flags |= MachineMemOperand::MOVolatile;
     return true;
   }
   case Intrinsic::amdgcn_tbuffer_load:
   case Intrinsic::amdgcn_buffer_load:
   case Intrinsic::amdgcn_buffer_load_format: {
     SIMachineFunctionInfo *MFI = MF.getInfo<SIMachineFunctionInfo>();
     Info.opc = ISD::INTRINSIC_W_CHAIN;
     Info.ptrVal = MFI->getBufferPSV(
       *MF.getSubtarget<SISubtarget>().getInstrInfo(),
       CI.getArgOperand(0));
     Info.memVT = MVT::getVT(CI.getType());
     Info.flags = MachineMemOperand::MOLoad |
                  MachineMemOperand::MODereferenceable;
 
     // There is a constant offset component, but there are additional register
     // offsets which could break AA if we set the offset to anything non-0.
     return true;
   }
   case Intrinsic::amdgcn_tbuffer_store:
   case Intrinsic::amdgcn_buffer_store:
   case Intrinsic::amdgcn_buffer_store_format: {
     SIMachineFunctionInfo *MFI = MF.getInfo<SIMachineFunctionInfo>();
     Info.opc = ISD::INTRINSIC_VOID;
     Info.ptrVal = MFI->getBufferPSV(
       *MF.getSubtarget<SISubtarget>().getInstrInfo(),
       CI.getArgOperand(1));
     Info.memVT = MVT::getVT(CI.getArgOperand(0)->getType());
     Info.flags = MachineMemOperand::MOStore |
                  MachineMemOperand::MODereferenceable;
     return true;
   }
   case Intrinsic::amdgcn_buffer_atomic_swap:
   case Intrinsic::amdgcn_buffer_atomic_add:
   case Intrinsic::amdgcn_buffer_atomic_sub:
   case Intrinsic::amdgcn_buffer_atomic_smin:
   case Intrinsic::amdgcn_buffer_atomic_umin:
   case Intrinsic::amdgcn_buffer_atomic_smax:
   case Intrinsic::amdgcn_buffer_atomic_umax:
   case Intrinsic::amdgcn_buffer_atomic_and:
   case Intrinsic::amdgcn_buffer_atomic_or:
   case Intrinsic::amdgcn_buffer_atomic_xor: {
     SIMachineFunctionInfo *MFI = MF.getInfo<SIMachineFunctionInfo>();
     Info.opc = ISD::INTRINSIC_W_CHAIN;
     Info.ptrVal = MFI->getBufferPSV(
       *MF.getSubtarget<SISubtarget>().getInstrInfo(),
       CI.getArgOperand(1));
     Info.memVT = MVT::getVT(CI.getType());
     Info.flags = MachineMemOperand::MOLoad |
                  MachineMemOperand::MOStore |
                  MachineMemOperand::MODereferenceable |
                  MachineMemOperand::MOVolatile;
     return true;
   }
   case Intrinsic::amdgcn_buffer_atomic_cmpswap: {
     SIMachineFunctionInfo *MFI = MF.getInfo<SIMachineFunctionInfo>();
     Info.opc = ISD::INTRINSIC_W_CHAIN;
     Info.ptrVal = MFI->getBufferPSV(
       *MF.getSubtarget<SISubtarget>().getInstrInfo(),
       CI.getArgOperand(2));
     Info.memVT = MVT::getVT(CI.getType());
     Info.flags = MachineMemOperand::MOLoad |
                  MachineMemOperand::MOStore |
                  MachineMemOperand::MODereferenceable |
                  MachineMemOperand::MOVolatile;
     return true;
   }
   default:
     return false;
   }
 }
 
 bool SITargetLowering::getAddrModeArguments(IntrinsicInst *II,
                                             SmallVectorImpl<Value*> &Ops,
                                             Type *&AccessTy) const {
   switch (II->getIntrinsicID()) {
   case Intrinsic::amdgcn_atomic_inc:
-  case Intrinsic::amdgcn_atomic_dec: {
+  case Intrinsic::amdgcn_atomic_dec:
+  case Intrinsic::amdgcn_atomic_fadd:
+  case Intrinsic::amdgcn_atomic_fmin:
+  case Intrinsic::amdgcn_atomic_fmax: {
     Value *Ptr = II->getArgOperand(0);
     AccessTy = II->getType();
     Ops.push_back(Ptr);
     return true;
   }
   default:
     return false;
   }

 
 SDValue SITargetLowering::LowerINTRINSIC_W_CHAIN(SDValue Op,
                                                  SelectionDAG &DAG) const {
   unsigned IntrID = cast<ConstantSDNode>(Op.getOperand(1))->getZExtValue();
   SDLoc DL(Op);
 
   switch (IntrID) {
   case Intrinsic::amdgcn_atomic_inc:
-  case Intrinsic::amdgcn_atomic_dec: {
+  case Intrinsic::amdgcn_atomic_dec:
+  case Intrinsic::amdgcn_atomic_fadd:
+  case Intrinsic::amdgcn_atomic_fmin:
+  case Intrinsic::amdgcn_atomic_fmax: {
     MemSDNode *M = cast<MemSDNode>(Op);
-    unsigned Opc = (IntrID == Intrinsic::amdgcn_atomic_inc) ?
-      AMDGPUISD::ATOMIC_INC : AMDGPUISD::ATOMIC_DEC;
+    unsigned Opc;
+    switch (IntrID) {
+    case Intrinsic::amdgcn_atomic_inc:
+      Opc = AMDGPUISD::ATOMIC_INC;
+      break;
+    case Intrinsic::amdgcn_atomic_dec:
+      Opc = AMDGPUISD::ATOMIC_DEC;
+      break;
+    case Intrinsic::amdgcn_atomic_fadd:
+      Opc = AMDGPUISD::ATOMIC_LOAD_FADD;
+      break;
+    case Intrinsic::amdgcn_atomic_fmin:
+      Opc = AMDGPUISD::ATOMIC_LOAD_FMIN;
+      break;
+    case Intrinsic::amdgcn_atomic_fmax:
+      Opc = AMDGPUISD::ATOMIC_LOAD_FMAX;
+      break;
+    default:
+      llvm_unreachable("Unknown intrinsic!");
+    }
     SDValue Ops[] = {
       M->getOperand(0), // Chain
       M->getOperand(2), // Ptr
       M->getOperand(3)  // Value
     };
 
     return DAG.getMemIntrinsicNode(Opc, SDLoc(Op), M->getVTList(), Ops,
                                    M->getMemoryVT(), M->getMemOperand());
SDValue SITargetLowering::PerformDAGCombine(SDNode *N,
   case ISD::ATOMIC_LOAD_OR:
   case ISD::ATOMIC_LOAD_XOR:
   case ISD::ATOMIC_LOAD_NAND:
   case ISD::ATOMIC_LOAD_MIN:
   case ISD::ATOMIC_LOAD_MAX:
   case ISD::ATOMIC_LOAD_UMIN:
   case ISD::ATOMIC_LOAD_UMAX:
   case AMDGPUISD::ATOMIC_INC:
-  case AMDGPUISD::ATOMIC_DEC: // TODO: Target mem intrinsics.
+  case AMDGPUISD::ATOMIC_DEC:
+  case AMDGPUISD::ATOMIC_LOAD_FADD:
+  case AMDGPUISD::ATOMIC_LOAD_FMIN:
+  case AMDGPUISD::ATOMIC_LOAD_FMAX:  // TODO: Target mem intrinsics.
     if (DCI.isBeforeLegalize())
       break;
     return performMemSDNodeCombine(cast<MemSDNode>(N), DCI);
   case ISD::AND:
     return performAndCombine(N, DCI);
   case ISD::OR:
     return performOrCombine(N, DCI);
   case ISD::XOR:
     return performXorCombine(N, DCI);

lib/Target/AMDGPU/SIInstrInfo.td

 def SIatomic_inc : SDNode<"AMDGPUISD::ATOMIC_INC", SDTAtomic2,
   [SDNPMayLoad, SDNPMayStore, SDNPMemOperand, SDNPHasChain]
 >;
 
 def SIatomic_dec : SDNode<"AMDGPUISD::ATOMIC_DEC", SDTAtomic2,
   [SDNPMayLoad, SDNPMayStore, SDNPMemOperand, SDNPHasChain]
 >;
 
+def SDTAtomic2_f32 : SDTypeProfile<1, 2, [
+  SDTCisSameAs<0,2>, SDTCisFP<0>, SDTCisPtrTy<1>
+]>;
+
+def SIatomic_fadd : SDNode<"AMDGPUISD::ATOMIC_LOAD_FADD", SDTAtomic2_f32,
+  [SDNPMayLoad, SDNPMayStore, SDNPMemOperand, SDNPHasChain]
+>;
+
+def SIatomic_fmin : SDNode<"AMDGPUISD::ATOMIC_LOAD_FMIN", SDTAtomic2_f32,
+  [SDNPMayLoad, SDNPMayStore, SDNPMemOperand, SDNPHasChain]
+>;
+
+def SIatomic_fmax : SDNode<"AMDGPUISD::ATOMIC_LOAD_FMAX", SDTAtomic2_f32,
+  [SDNPMayLoad, SDNPMayStore, SDNPMemOperand, SDNPHasChain]
+>;
+
 def SItbuffer_load : SDNode<"AMDGPUISD::TBUFFER_LOAD_FORMAT",
   SDTypeProfile<1, 9,
     [                     // vdata
      SDTCisVT<1, v4i32>,  // rsrc
      SDTCisVT<2, i32>,    // vindex(VGPR)
      SDTCisVT<3, i32>,    // voffset(VGPR)
      SDTCisVT<4, i32>,    // soffset(SGPR)
      SDTCisVT<5, i32>,    // offset(imm)

 // PatFrags for global memory operations
 //===----------------------------------------------------------------------===//
 
 defm atomic_inc_global : global_binary_atomic_op<SIatomic_inc>;
 defm atomic_dec_global : global_binary_atomic_op<SIatomic_dec>;
 
 def atomic_inc_local : local_binary_atomic_op<SIatomic_inc>;
 def atomic_dec_local : local_binary_atomic_op<SIatomic_dec>;
+def atomic_load_fadd_local : local_binary_atomic_op<SIatomic_fadd>;
+def atomic_load_fmin_local : local_binary_atomic_op<SIatomic_fmin>;
+def atomic_load_fmax_local : local_binary_atomic_op<SIatomic_fmax>;
 
 //===----------------------------------------------------------------------===//
 // SDNodes PatFrags for loads/stores with a glue input.
 // This is for SDNodes and PatFrag for local loads and stores to
 // enable s_mov_b32 m0, -1 to be glued to the memory instructions.
 //
 // These mirror the regular load/store PatFrags and rely on special
 // processing during Select() to add the glued copy.
def ashr_rev : PatFrag <
   (sra $src0, $src1)
 >;
 
 def lshl_rev : PatFrag <
   (ops node:$src1, node:$src0),
   (shl $src0, $src1)
 >;
 
-multiclass SIAtomicM0Glue2 <string op_name, bit is_amdgpu = 0> {
+multiclass SIAtomicM0Glue2 <string op_name, bit is_amdgpu = 0,
+                            SDTypeProfile tc = SDTAtomic2> {
 
   def _glue : SDNode <
-    !if(is_amdgpu, "AMDGPUISD", "ISD")#"::ATOMIC_"#op_name, SDTAtomic2,
+    !if(is_amdgpu, "AMDGPUISD", "ISD")#"::ATOMIC_"#op_name, tc,
     [SDNPHasChain, SDNPMayStore, SDNPMayLoad, SDNPMemOperand, SDNPInGlue]
   >;
 
   def _local_m0 : local_binary_atomic_op <!cast<SDNode>(NAME#"_glue")>;
 }
 
 defm atomic_load_add : SIAtomicM0Glue2 <"LOAD_ADD">;
 defm atomic_load_sub : SIAtomicM0Glue2 <"LOAD_SUB">;
 defm atomic_inc : SIAtomicM0Glue2 <"INC", 1>;
 defm atomic_dec : SIAtomicM0Glue2 <"DEC", 1>;
 defm atomic_load_and : SIAtomicM0Glue2 <"LOAD_AND">;
 defm atomic_load_min : SIAtomicM0Glue2 <"LOAD_MIN">;
 defm atomic_load_max : SIAtomicM0Glue2 <"LOAD_MAX">;
 defm atomic_load_or : SIAtomicM0Glue2 <"LOAD_OR">;
 defm atomic_load_xor : SIAtomicM0Glue2 <"LOAD_XOR">;
 defm atomic_load_umin : SIAtomicM0Glue2 <"LOAD_UMIN">;
 defm atomic_load_umax : SIAtomicM0Glue2 <"LOAD_UMAX">;
 defm atomic_swap : SIAtomicM0Glue2 <"SWAP">;
+defm atomic_load_fadd : SIAtomicM0Glue2 <"LOAD_FADD", 1, SDTAtomic2_f32>;
+defm atomic_load_fmin : SIAtomicM0Glue2 <"LOAD_FMIN", 1, SDTAtomic2_f32>;
+defm atomic_load_fmax : SIAtomicM0Glue2 <"LOAD_FMAX", 1, SDTAtomic2_f32>;
 
 def atomic_cmp_swap_glue : SDNode <"ISD::ATOMIC_CMP_SWAP", SDTAtomic3,
   [SDNPHasChain, SDNPMayStore, SDNPMayLoad, SDNPMemOperand, SDNPInGlue]
 >;
 
 def atomic_cmp_swap_local_m0 : AtomicCmpSwapLocal<atomic_cmp_swap_glue>;
 
 

lib/Transforms/Scalar/InferAddressSpaces.cpp

 // TODO: Move logic to TTI?
 bool InferAddressSpaces::rewriteIntrinsicOperands(IntrinsicInst *II,
                                                   Value *OldV,
                                                   Value *NewV) const {
   Module *M = II->getParent()->getParent()->getParent();
 
   switch (II->getIntrinsicID()) {
   case Intrinsic::amdgcn_atomic_inc:
-  case Intrinsic::amdgcn_atomic_dec:{
+  case Intrinsic::amdgcn_atomic_dec:
+  case Intrinsic::amdgcn_atomic_fadd:
+  case Intrinsic::amdgcn_atomic_fmin:
+  case Intrinsic::amdgcn_atomic_fmax: {
     const ConstantInt *IsVolatile = dyn_cast<ConstantInt>(II->getArgOperand(4));
     if (!IsVolatile || !IsVolatile->isZero())
       return false;
 
     LLVM_FALLTHROUGH;
   }
   case Intrinsic::objectsize: {
     Type *DestTy = II->getType();

 // TODO: Move logic to TTI?
 void InferAddressSpaces::collectRewritableIntrinsicOperands(
     IntrinsicInst *II, std::vector<std::pair<Value *, bool>> &PostorderStack,
     DenseSet<Value *> &Visited) const {
   switch (II->getIntrinsicID()) {
   case Intrinsic::objectsize:
   case Intrinsic::amdgcn_atomic_inc:
   case Intrinsic::amdgcn_atomic_dec:
+  case Intrinsic::amdgcn_atomic_fadd:
+  case Intrinsic::amdgcn_atomic_fmin:
+  case Intrinsic::amdgcn_atomic_fmax:
     appendsFlatAddressExpressionToPostorderStack(II->getArgOperand(0),
                                                  PostorderStack, Visited);
     break;
   default:
     break;
   }
 }
 

test/CodeGen/AMDGPU/lds_atomic_f32.ll

+; RUN: llc -march=amdgcn -mcpu=tonga -verify-machineinstrs < %s | FileCheck -enable-var-scope -check-prefixes=GCN,VI %s
+; RUN: llc -march=amdgcn -mcpu=gfx900 -verify-machineinstrs < %s | FileCheck -enable-var-scope -check-prefixes=GCN,GFX9 %s
+
+declare float @llvm.amdgcn.atomic.fadd.f32(float addrspace(3)* nocapture, float, i32, i32, i1)
+declare float @llvm.amdgcn.atomic.fmin.f32(float addrspace(3)* nocapture, float, i32, i32, i1)
+declare float @llvm.amdgcn.atomic.fmax.f32(float addrspace(3)* nocapture, float, i32, i32, i1)
+
+; GCN-LABEL: {{^}}lds_atomic_fadd_f32:
+; VI-DAG: s_mov_b32 m0
+; GFX9-NOT: m0
+; GCN-DAG: v_mov_b32_e32 [[V0:v[0-9]+]], 0x42280000
+; GCN: ds_add_rtn_f32 [[V2:v[0-9]+]], [[V1:v[0-9]+]], [[V0]] offset:32
+; GCN: ds_add_f32 [[V3:v[0-9]+]], [[V0]] offset:64
+; GCN: s_waitcnt lgkmcnt(1)
+; GCN: ds_add_rtn_f32 {{v[0-9]+}}, {{v[0-9]+}}, [[V2]]
+define amdgpu_kernel void @lds_atomic_fadd_f32(float addrspace(1)* %out, float addrspace(3)* %ptrf, i32 %idx) {
+  %idx.add = add nuw i32 %idx, 4
+  %shl0 = shl i32 %idx.add, 3
+  %shl1 = shl i32 %idx.add, 4
+  %ptr0 = inttoptr i32 %shl0 to float addrspace(3)*
+  %ptr1 = inttoptr i32 %shl1 to float addrspace(3)*
+  %a1 = call float @llvm.amdgcn.atomic.fadd.f32(float addrspace(3)* %ptr0, float 4.2e+1, i32 0, i32 0, i1 false)
+  %a2 = call float @llvm.amdgcn.atomic.fadd.f32(float addrspace(3)* %ptr1, float 4.2e+1, i32 0, i32 0, i1 false)
+  %a3 = call float @llvm.amdgcn.atomic.fadd.f32(float addrspace(3)* %ptrf, float %a1, i32 0, i32 0, i1 false)
+  store float %a3, float addrspace(1)* %out
+  ret void
+}
+
+; GCN-LABEL: {{^}}lds_atomic_fmin_f32:
+; VI-DAG: s_mov_b32 m0
+; GFX9-NOT: m0
+; GCN-DAG: v_mov_b32_e32 [[V0:v[0-9]+]], 0x42280000
+; GCN: ds_min_rtn_f32 [[V2:v[0-9]+]], [[V1:v[0-9]+]], [[V0]] offset:32
+; GCN: ds_min_f32 [[V3:v[0-9]+]], [[V0]] offset:64
+; GCN: s_waitcnt lgkmcnt(1)
+; GCN: ds_min_rtn_f32 {{v[0-9]+}}, {{v[0-9]+}}, [[V2]]
+define amdgpu_kernel void @lds_atomic_fmin_f32(float addrspace(1)* %out, float addrspace(3)* %ptrf, i32 %idx) {
+  %idx.add = add nuw i32 %idx, 4
+  %shl0 = shl i32 %idx.add, 3
+  %shl1 = shl i32 %idx.add, 4
+  %ptr0 = inttoptr i32 %shl0 to float addrspace(3)*
+  %ptr1 = inttoptr i32 %shl1 to float addrspace(3)*
+  %a1 = call float @llvm.amdgcn.atomic.fmin.f32(float addrspace(3)* %ptr0, float 4.2e+1, i32 0, i32 0, i1 false)
+  %a2 = call float @llvm.amdgcn.atomic.fmin.f32(float addrspace(3)* %ptr1, float 4.2e+1, i32 0, i32 0, i1 false)
+  %a3 = call float @llvm.amdgcn.atomic.fmin.f32(float addrspace(3)* %ptrf, float %a1, i32 0, i32 0, i1 false)
+  store float %a3, float addrspace(1)* %out
+  ret void
+}
+
+; GCN-LABEL: {{^}}lds_atomic_fmax_f32:
+; VI-DAG: s_mov_b32 m0
+; GFX9-NOT: m0
+; GCN-DAG: v_mov_b32_e32 [[V0:v[0-9]+]], 0x42280000
+; GCN: ds_max_rtn_f32 [[V2:v[0-9]+]], [[V1:v[0-9]+]], [[V0]] offset:32
+; GCN: ds_max_f32 [[V3:v[0-9]+]], [[V0]] offset:64
+; GCN: s_waitcnt lgkmcnt(1)
+; GCN: ds_max_rtn_f32 {{v[0-9]+}}, {{v[0-9]+}}, [[V2]]
+define amdgpu_kernel void @lds_atomic_fmax_f32(float addrspace(1)* %out, float addrspace(3)* %ptrf, i32 %idx) {
+  %idx.add = add nuw i32 %idx, 4
+  %shl0 = shl i32 %idx.add, 3
+  %shl1 = shl i32 %idx.add, 4
+  %ptr0 = inttoptr i32 %shl0 to float addrspace(3)*
+  %ptr1 = inttoptr i32 %shl1 to float addrspace(3)*
+  %a1 = call float @llvm.amdgcn.atomic.fmax.f32(float addrspace(3)* %ptr0, float 4.2e+1, i32 0, i32 0, i1 false)
+  %a2 = call float @llvm.amdgcn.atomic.fmax.f32(float addrspace(3)* %ptr1, float 4.2e+1, i32 0, i32 0, i1 false)
+  %a3 = call float @llvm.amdgcn.atomic.fmax.f32(float addrspace(3)* %ptrf, float %a1, i32 0, i32 0, i1 false)
+  store float %a3, float addrspace(1)* %out
+  ret void
+}
+