[llvm-branch-commits] [llvm-branch] r351443 - Merging r351351:
Hans Wennborg via llvm-branch-commits
llvm-branch-commits at lists.llvm.org
Thu Jan 17 05:41:27 PST 2019
Author: hans
Date: Thu Jan 17 05:41:26 2019
New Revision: 351443
URL: http://llvm.org/viewvc/llvm-project?rev=351443&view=rev
Log:
Merging r351351:
------------------------------------------------------------------------
r351351 | mareko | 2019-01-16 16:43:53 +0100 (Wed, 16 Jan 2019) | 7 lines
AMDGPU: Add llvm.amdgcn.ds.ordered.add & swap
Reviewers: arsenm, nhaehnle
Subscribers: kzhuravl, jvesely, wdng, yaxunl, dstuttard, tpr, t-tye, llvm-commits
Differential Revision: https://reviews.llvm.org/D52944
------------------------------------------------------------------------
Added:
llvm/branches/release_80/test/CodeGen/AMDGPU/llvm.amdgcn.ds.ordered.add.ll
- copied unchanged from r351351, llvm/trunk/test/CodeGen/AMDGPU/llvm.amdgcn.ds.ordered.add.ll
llvm/branches/release_80/test/CodeGen/AMDGPU/llvm.amdgcn.ds.ordered.swap.ll
- copied unchanged from r351351, llvm/trunk/test/CodeGen/AMDGPU/llvm.amdgcn.ds.ordered.swap.ll
Modified:
llvm/branches/release_80/ (props changed)
llvm/branches/release_80/include/llvm/IR/IntrinsicsAMDGPU.td
llvm/branches/release_80/lib/Target/AMDGPU/AMDGPU.h
llvm/branches/release_80/lib/Target/AMDGPU/AMDGPUISelLowering.cpp
llvm/branches/release_80/lib/Target/AMDGPU/AMDGPUISelLowering.h
llvm/branches/release_80/lib/Target/AMDGPU/AMDGPUSearchableTables.td
llvm/branches/release_80/lib/Target/AMDGPU/AMDGPUTargetTransformInfo.cpp
llvm/branches/release_80/lib/Target/AMDGPU/DSInstructions.td
llvm/branches/release_80/lib/Target/AMDGPU/GCNHazardRecognizer.cpp
llvm/branches/release_80/lib/Target/AMDGPU/SIISelLowering.cpp
llvm/branches/release_80/lib/Target/AMDGPU/SIInsertWaitcnts.cpp
llvm/branches/release_80/lib/Target/AMDGPU/SIInstrInfo.cpp
llvm/branches/release_80/lib/Target/AMDGPU/SIInstrInfo.h
llvm/branches/release_80/lib/Target/AMDGPU/SIInstrInfo.td
Propchange: llvm/branches/release_80/
------------------------------------------------------------------------------
--- svn:mergeinfo (original)
+++ svn:mergeinfo Thu Jan 17 05:41:26 2019
@@ -1,3 +1,3 @@
/llvm/branches/Apple/Pertwee:110850,110961
/llvm/branches/type-system-rewrite:133420-134817
-/llvm/trunk:155241,351344-351345,351349,351436
+/llvm/trunk:155241,351344-351345,351349,351351,351436
Modified: llvm/branches/release_80/include/llvm/IR/IntrinsicsAMDGPU.td
URL: http://llvm.org/viewvc/llvm-project/llvm/branches/release_80/include/llvm/IR/IntrinsicsAMDGPU.td?rev=351443&r1=351442&r2=351443&view=diff
==============================================================================
--- llvm/branches/release_80/include/llvm/IR/IntrinsicsAMDGPU.td (original)
+++ llvm/branches/release_80/include/llvm/IR/IntrinsicsAMDGPU.td Thu Jan 17 05:41:26 2019
@@ -392,6 +392,24 @@ class AMDGPULDSF32Intrin<string clang_bu
[IntrArgMemOnly, NoCapture<0>]
>;
+class AMDGPUDSOrderedIntrinsic : Intrinsic<
+ [llvm_i32_ty],
+ // M0 = {hi16:address, lo16:waveID}. Allow passing M0 as a pointer, so that
+ // the bit packing can be optimized at the IR level.
+ [LLVMQualPointerType<llvm_i32_ty, 2>, // IntToPtr(M0)
+ llvm_i32_ty, // value to add or swap
+ llvm_i32_ty, // ordering
+ llvm_i32_ty, // scope
+ llvm_i1_ty, // isVolatile
+ llvm_i32_ty, // ordered count index (OA index), also added to the address
+ llvm_i1_ty, // wave release, usually set to 1
+ llvm_i1_ty], // wave done, set to 1 for the last ordered instruction
+ [NoCapture<0>]
+>;
+
+def int_amdgcn_ds_ordered_add : AMDGPUDSOrderedIntrinsic;
+def int_amdgcn_ds_ordered_swap : AMDGPUDSOrderedIntrinsic;
+
def int_amdgcn_ds_fadd : AMDGPULDSF32Intrin<"__builtin_amdgcn_ds_faddf">;
def int_amdgcn_ds_fmin : AMDGPULDSF32Intrin<"__builtin_amdgcn_ds_fminf">;
def int_amdgcn_ds_fmax : AMDGPULDSF32Intrin<"__builtin_amdgcn_ds_fmaxf">;
Modified: llvm/branches/release_80/lib/Target/AMDGPU/AMDGPU.h
URL: http://llvm.org/viewvc/llvm-project/llvm/branches/release_80/lib/Target/AMDGPU/AMDGPU.h?rev=351443&r1=351442&r2=351443&view=diff
==============================================================================
--- llvm/branches/release_80/lib/Target/AMDGPU/AMDGPU.h (original)
+++ llvm/branches/release_80/lib/Target/AMDGPU/AMDGPU.h Thu Jan 17 05:41:26 2019
@@ -254,7 +254,7 @@ namespace AMDGPUAS {
FLAT_ADDRESS = 0, ///< Address space for flat memory.
GLOBAL_ADDRESS = 1, ///< Address space for global memory (RAT0, VTX0).
- REGION_ADDRESS = 2, ///< Address space for region memory.
+ REGION_ADDRESS = 2, ///< Address space for region memory. (GDS)
CONSTANT_ADDRESS = 4, ///< Address space for constant memory (VTX2)
LOCAL_ADDRESS = 3, ///< Address space for local memory.
Modified: llvm/branches/release_80/lib/Target/AMDGPU/AMDGPUISelLowering.cpp
URL: http://llvm.org/viewvc/llvm-project/llvm/branches/release_80/lib/Target/AMDGPU/AMDGPUISelLowering.cpp?rev=351443&r1=351442&r2=351443&view=diff
==============================================================================
--- llvm/branches/release_80/lib/Target/AMDGPU/AMDGPUISelLowering.cpp (original)
+++ llvm/branches/release_80/lib/Target/AMDGPU/AMDGPUISelLowering.cpp Thu Jan 17 05:41:26 2019
@@ -4192,6 +4192,7 @@ const char* AMDGPUTargetLowering::getTar
NODE_NAME_CASE(TBUFFER_STORE_FORMAT_D16)
NODE_NAME_CASE(TBUFFER_LOAD_FORMAT)
NODE_NAME_CASE(TBUFFER_LOAD_FORMAT_D16)
+ NODE_NAME_CASE(DS_ORDERED_COUNT)
NODE_NAME_CASE(ATOMIC_CMP_SWAP)
NODE_NAME_CASE(ATOMIC_INC)
NODE_NAME_CASE(ATOMIC_DEC)
Modified: llvm/branches/release_80/lib/Target/AMDGPU/AMDGPUISelLowering.h
URL: http://llvm.org/viewvc/llvm-project/llvm/branches/release_80/lib/Target/AMDGPU/AMDGPUISelLowering.h?rev=351443&r1=351442&r2=351443&view=diff
==============================================================================
--- llvm/branches/release_80/lib/Target/AMDGPU/AMDGPUISelLowering.h (original)
+++ llvm/branches/release_80/lib/Target/AMDGPU/AMDGPUISelLowering.h Thu Jan 17 05:41:26 2019
@@ -474,6 +474,7 @@ enum NodeType : unsigned {
TBUFFER_STORE_FORMAT_D16,
TBUFFER_LOAD_FORMAT,
TBUFFER_LOAD_FORMAT_D16,
+ DS_ORDERED_COUNT,
ATOMIC_CMP_SWAP,
ATOMIC_INC,
ATOMIC_DEC,
Modified: llvm/branches/release_80/lib/Target/AMDGPU/AMDGPUSearchableTables.td
URL: http://llvm.org/viewvc/llvm-project/llvm/branches/release_80/lib/Target/AMDGPU/AMDGPUSearchableTables.td?rev=351443&r1=351442&r2=351443&view=diff
==============================================================================
--- llvm/branches/release_80/lib/Target/AMDGPU/AMDGPUSearchableTables.td (original)
+++ llvm/branches/release_80/lib/Target/AMDGPU/AMDGPUSearchableTables.td Thu Jan 17 05:41:26 2019
@@ -72,6 +72,8 @@ def : SourceOfDivergence<int_amdgcn_buff
def : SourceOfDivergence<int_amdgcn_buffer_atomic_cmpswap>;
def : SourceOfDivergence<int_amdgcn_ps_live>;
def : SourceOfDivergence<int_amdgcn_ds_swizzle>;
+def : SourceOfDivergence<int_amdgcn_ds_ordered_add>;
+def : SourceOfDivergence<int_amdgcn_ds_ordered_swap>;
foreach intr = AMDGPUImageDimAtomicIntrinsics in
def : SourceOfDivergence<intr>;
Modified: llvm/branches/release_80/lib/Target/AMDGPU/AMDGPUTargetTransformInfo.cpp
URL: http://llvm.org/viewvc/llvm-project/llvm/branches/release_80/lib/Target/AMDGPU/AMDGPUTargetTransformInfo.cpp?rev=351443&r1=351442&r2=351443&view=diff
==============================================================================
--- llvm/branches/release_80/lib/Target/AMDGPU/AMDGPUTargetTransformInfo.cpp (original)
+++ llvm/branches/release_80/lib/Target/AMDGPU/AMDGPUTargetTransformInfo.cpp Thu Jan 17 05:41:26 2019
@@ -308,6 +308,8 @@ bool GCNTTIImpl::getTgtMemIntrinsic(Intr
switch (Inst->getIntrinsicID()) {
case Intrinsic::amdgcn_atomic_inc:
case Intrinsic::amdgcn_atomic_dec:
+ case Intrinsic::amdgcn_ds_ordered_add:
+ case Intrinsic::amdgcn_ds_ordered_swap:
case Intrinsic::amdgcn_ds_fadd:
case Intrinsic::amdgcn_ds_fmin:
case Intrinsic::amdgcn_ds_fmax: {
Modified: llvm/branches/release_80/lib/Target/AMDGPU/DSInstructions.td
URL: http://llvm.org/viewvc/llvm-project/llvm/branches/release_80/lib/Target/AMDGPU/DSInstructions.td?rev=351443&r1=351442&r2=351443&view=diff
==============================================================================
--- llvm/branches/release_80/lib/Target/AMDGPU/DSInstructions.td (original)
+++ llvm/branches/release_80/lib/Target/AMDGPU/DSInstructions.td Thu Jan 17 05:41:26 2019
@@ -817,6 +817,11 @@ defm : DSAtomicRetPat_mc<DS_MAX_RTN_U64,
defm : DSAtomicCmpXChg_mc<DS_CMPST_RTN_B64, i64, "atomic_cmp_swap_local">;
+def : Pat <
+ (SIds_ordered_count i32:$value, i16:$offset),
+ (DS_ORDERED_COUNT $value, (as_i16imm $offset))
+>;
+
//===----------------------------------------------------------------------===//
// Real instructions
//===----------------------------------------------------------------------===//
Modified: llvm/branches/release_80/lib/Target/AMDGPU/GCNHazardRecognizer.cpp
URL: http://llvm.org/viewvc/llvm-project/llvm/branches/release_80/lib/Target/AMDGPU/GCNHazardRecognizer.cpp?rev=351443&r1=351442&r2=351443&view=diff
==============================================================================
--- llvm/branches/release_80/lib/Target/AMDGPU/GCNHazardRecognizer.cpp (original)
+++ llvm/branches/release_80/lib/Target/AMDGPU/GCNHazardRecognizer.cpp Thu Jan 17 05:41:26 2019
@@ -88,14 +88,28 @@ static bool isSMovRel(unsigned Opcode) {
}
}
-static bool isSendMsgTraceDataOrGDS(const MachineInstr &MI) {
+static bool isSendMsgTraceDataOrGDS(const SIInstrInfo &TII,
+ const MachineInstr &MI) {
+ if (TII.isAlwaysGDS(MI.getOpcode()))
+ return true;
+
switch (MI.getOpcode()) {
case AMDGPU::S_SENDMSG:
case AMDGPU::S_SENDMSGHALT:
case AMDGPU::S_TTRACEDATA:
return true;
+ // These DS opcodes don't support GDS.
+ case AMDGPU::DS_NOP:
+ case AMDGPU::DS_PERMUTE_B32:
+ case AMDGPU::DS_BPERMUTE_B32:
+ return false;
default:
- // TODO: GDS
+ if (TII.isDS(MI.getOpcode())) {
+ int GDS = AMDGPU::getNamedOperandIdx(MI.getOpcode(),
+ AMDGPU::OpName::gds);
+ if (MI.getOperand(GDS).getImm())
+ return true;
+ }
return false;
}
}
@@ -145,7 +159,7 @@ GCNHazardRecognizer::getHazardType(SUnit
checkReadM0Hazards(MI) > 0)
return NoopHazard;
- if (ST.hasReadM0SendMsgHazard() && isSendMsgTraceDataOrGDS(*MI) &&
+ if (ST.hasReadM0SendMsgHazard() && isSendMsgTraceDataOrGDS(TII, *MI) &&
checkReadM0Hazards(MI) > 0)
return NoopHazard;
@@ -199,7 +213,7 @@ unsigned GCNHazardRecognizer::PreEmitNoo
isSMovRel(MI->getOpcode())))
return std::max(WaitStates, checkReadM0Hazards(MI));
- if (ST.hasReadM0SendMsgHazard() && isSendMsgTraceDataOrGDS(*MI))
+ if (ST.hasReadM0SendMsgHazard() && isSendMsgTraceDataOrGDS(TII, *MI))
return std::max(WaitStates, checkReadM0Hazards(MI));
return WaitStates;
Modified: llvm/branches/release_80/lib/Target/AMDGPU/SIISelLowering.cpp
URL: http://llvm.org/viewvc/llvm-project/llvm/branches/release_80/lib/Target/AMDGPU/SIISelLowering.cpp?rev=351443&r1=351442&r2=351443&view=diff
==============================================================================
--- llvm/branches/release_80/lib/Target/AMDGPU/SIISelLowering.cpp (original)
+++ llvm/branches/release_80/lib/Target/AMDGPU/SIISelLowering.cpp Thu Jan 17 05:41:26 2019
@@ -910,6 +910,8 @@ bool SITargetLowering::getTgtMemIntrinsi
switch (IntrID) {
case Intrinsic::amdgcn_atomic_inc:
case Intrinsic::amdgcn_atomic_dec:
+ case Intrinsic::amdgcn_ds_ordered_add:
+ case Intrinsic::amdgcn_ds_ordered_swap:
case Intrinsic::amdgcn_ds_fadd:
case Intrinsic::amdgcn_ds_fmin:
case Intrinsic::amdgcn_ds_fmax: {
@@ -937,6 +939,8 @@ bool SITargetLowering::getAddrModeArgume
switch (II->getIntrinsicID()) {
case Intrinsic::amdgcn_atomic_inc:
case Intrinsic::amdgcn_atomic_dec:
+ case Intrinsic::amdgcn_ds_ordered_add:
+ case Intrinsic::amdgcn_ds_ordered_swap:
case Intrinsic::amdgcn_ds_fadd:
case Intrinsic::amdgcn_ds_fmin:
case Intrinsic::amdgcn_ds_fmax: {
@@ -5438,6 +5442,63 @@ SDValue SITargetLowering::LowerINTRINSIC
SDLoc DL(Op);
switch (IntrID) {
+ case Intrinsic::amdgcn_ds_ordered_add:
+ case Intrinsic::amdgcn_ds_ordered_swap: {
+ MemSDNode *M = cast<MemSDNode>(Op);
+ SDValue Chain = M->getOperand(0);
+ SDValue M0 = M->getOperand(2);
+ SDValue Value = M->getOperand(3);
+ unsigned OrderedCountIndex = M->getConstantOperandVal(7);
+ unsigned WaveRelease = M->getConstantOperandVal(8);
+ unsigned WaveDone = M->getConstantOperandVal(9);
+ unsigned ShaderType;
+ unsigned Instruction;
+
+ switch (IntrID) {
+ case Intrinsic::amdgcn_ds_ordered_add:
+ Instruction = 0;
+ break;
+ case Intrinsic::amdgcn_ds_ordered_swap:
+ Instruction = 1;
+ break;
+ }
+
+ if (WaveDone && !WaveRelease)
+ report_fatal_error("ds_ordered_count: wave_done requires wave_release");
+
+ switch (DAG.getMachineFunction().getFunction().getCallingConv()) {
+ case CallingConv::AMDGPU_CS:
+ case CallingConv::AMDGPU_KERNEL:
+ ShaderType = 0;
+ break;
+ case CallingConv::AMDGPU_PS:
+ ShaderType = 1;
+ break;
+ case CallingConv::AMDGPU_VS:
+ ShaderType = 2;
+ break;
+ case CallingConv::AMDGPU_GS:
+ ShaderType = 3;
+ break;
+ default:
+ report_fatal_error("ds_ordered_count unsupported for this calling conv");
+ }
+
+ unsigned Offset0 = OrderedCountIndex << 2;
+ unsigned Offset1 = WaveRelease | (WaveDone << 1) | (ShaderType << 2) |
+ (Instruction << 4);
+ unsigned Offset = Offset0 | (Offset1 << 8);
+
+ SDValue Ops[] = {
+ Chain,
+ Value,
+ DAG.getTargetConstant(Offset, DL, MVT::i16),
+ copyToM0(DAG, Chain, DL, M0).getValue(1), // Glue
+ };
+ return DAG.getMemIntrinsicNode(AMDGPUISD::DS_ORDERED_COUNT, DL,
+ M->getVTList(), Ops, M->getMemoryVT(),
+ M->getMemOperand());
+ }
case Intrinsic::amdgcn_atomic_inc:
case Intrinsic::amdgcn_atomic_dec:
case Intrinsic::amdgcn_ds_fadd:
Modified: llvm/branches/release_80/lib/Target/AMDGPU/SIInsertWaitcnts.cpp
URL: http://llvm.org/viewvc/llvm-project/llvm/branches/release_80/lib/Target/AMDGPU/SIInsertWaitcnts.cpp?rev=351443&r1=351442&r2=351443&view=diff
==============================================================================
--- llvm/branches/release_80/lib/Target/AMDGPU/SIInsertWaitcnts.cpp (original)
+++ llvm/branches/release_80/lib/Target/AMDGPU/SIInsertWaitcnts.cpp Thu Jan 17 05:41:26 2019
@@ -536,10 +536,13 @@ void WaitcntBrackets::updateByEvent(cons
CurrScore);
}
if (Inst.mayStore()) {
- setExpScore(
- &Inst, TII, TRI, MRI,
- AMDGPU::getNamedOperandIdx(Inst.getOpcode(), AMDGPU::OpName::data0),
- CurrScore);
+ if (AMDGPU::getNamedOperandIdx(Inst.getOpcode(),
+ AMDGPU::OpName::data0) != -1) {
+ setExpScore(
+ &Inst, TII, TRI, MRI,
+ AMDGPU::getNamedOperandIdx(Inst.getOpcode(), AMDGPU::OpName::data0),
+ CurrScore);
+ }
if (AMDGPU::getNamedOperandIdx(Inst.getOpcode(),
AMDGPU::OpName::data1) != -1) {
setExpScore(&Inst, TII, TRI, MRI,
@@ -1093,7 +1096,8 @@ void SIInsertWaitcnts::updateEventWaitcn
// bracket and the destination operand scores.
// TODO: Use the (TSFlags & SIInstrFlags::LGKM_CNT) property everywhere.
if (TII->isDS(Inst) && TII->usesLGKM_CNT(Inst)) {
- if (TII->hasModifiersSet(Inst, AMDGPU::OpName::gds)) {
+ if (TII->isAlwaysGDS(Inst.getOpcode()) ||
+ TII->hasModifiersSet(Inst, AMDGPU::OpName::gds)) {
ScoreBrackets->updateByEvent(TII, TRI, MRI, GDS_ACCESS, Inst);
ScoreBrackets->updateByEvent(TII, TRI, MRI, GDS_GPR_LOCK, Inst);
} else {
Modified: llvm/branches/release_80/lib/Target/AMDGPU/SIInstrInfo.cpp
URL: http://llvm.org/viewvc/llvm-project/llvm/branches/release_80/lib/Target/AMDGPU/SIInstrInfo.cpp?rev=351443&r1=351442&r2=351443&view=diff
==============================================================================
--- llvm/branches/release_80/lib/Target/AMDGPU/SIInstrInfo.cpp (original)
+++ llvm/branches/release_80/lib/Target/AMDGPU/SIInstrInfo.cpp Thu Jan 17 05:41:26 2019
@@ -2390,6 +2390,16 @@ bool SIInstrInfo::isSchedulingBoundary(c
changesVGPRIndexingMode(MI);
}
+bool SIInstrInfo::isAlwaysGDS(uint16_t Opcode) const {
+ return Opcode == AMDGPU::DS_ORDERED_COUNT ||
+ Opcode == AMDGPU::DS_GWS_INIT ||
+ Opcode == AMDGPU::DS_GWS_SEMA_V ||
+ Opcode == AMDGPU::DS_GWS_SEMA_BR ||
+ Opcode == AMDGPU::DS_GWS_SEMA_P ||
+ Opcode == AMDGPU::DS_GWS_SEMA_RELEASE_ALL ||
+ Opcode == AMDGPU::DS_GWS_BARRIER;
+}
+
bool SIInstrInfo::hasUnwantedEffectsWhenEXECEmpty(const MachineInstr &MI) const {
unsigned Opcode = MI.getOpcode();
@@ -2403,7 +2413,8 @@ bool SIInstrInfo::hasUnwantedEffectsWhen
// EXEC = 0, but checking for that case here seems not worth it
// given the typical code patterns.
if (Opcode == AMDGPU::S_SENDMSG || Opcode == AMDGPU::S_SENDMSGHALT ||
- Opcode == AMDGPU::EXP || Opcode == AMDGPU::EXP_DONE)
+ Opcode == AMDGPU::EXP || Opcode == AMDGPU::EXP_DONE ||
+ Opcode == AMDGPU::DS_ORDERED_COUNT)
return true;
if (MI.isInlineAsm())
Modified: llvm/branches/release_80/lib/Target/AMDGPU/SIInstrInfo.h
URL: http://llvm.org/viewvc/llvm-project/llvm/branches/release_80/lib/Target/AMDGPU/SIInstrInfo.h?rev=351443&r1=351442&r2=351443&view=diff
==============================================================================
--- llvm/branches/release_80/lib/Target/AMDGPU/SIInstrInfo.h (original)
+++ llvm/branches/release_80/lib/Target/AMDGPU/SIInstrInfo.h Thu Jan 17 05:41:26 2019
@@ -450,6 +450,8 @@ public:
return get(Opcode).TSFlags & SIInstrFlags::DS;
}
+ bool isAlwaysGDS(uint16_t Opcode) const;
+
static bool isMIMG(const MachineInstr &MI) {
return MI.getDesc().TSFlags & SIInstrFlags::MIMG;
}
Modified: llvm/branches/release_80/lib/Target/AMDGPU/SIInstrInfo.td
URL: http://llvm.org/viewvc/llvm-project/llvm/branches/release_80/lib/Target/AMDGPU/SIInstrInfo.td?rev=351443&r1=351442&r2=351443&view=diff
==============================================================================
--- llvm/branches/release_80/lib/Target/AMDGPU/SIInstrInfo.td (original)
+++ llvm/branches/release_80/lib/Target/AMDGPU/SIInstrInfo.td Thu Jan 17 05:41:26 2019
@@ -45,6 +45,11 @@ def SIsbuffer_load : SDNode<"AMDGPUISD::
[SDNPMayLoad, SDNPMemOperand]
>;
+def SIds_ordered_count : SDNode<"AMDGPUISD::DS_ORDERED_COUNT",
+ SDTypeProfile<1, 2, [SDTCisVT<0, i32>, SDTCisVT<1, i32>, SDTCisVT<2, i16>]>,
+ [SDNPMayLoad, SDNPMayStore, SDNPMemOperand, SDNPHasChain, SDNPInGlue]
+>;
+
def SIatomic_inc : SDNode<"AMDGPUISD::ATOMIC_INC", SDTAtomic2,
[SDNPMayLoad, SDNPMayStore, SDNPMemOperand, SDNPHasChain]
>;
More information about the llvm-branch-commits
mailing list