This is an archive of the discontinued LLVM Phabricator instance.

Differential D17318

[NVPTX] Annotate convergent intrinsics as convergent.
ClosedPublic

Authored by jlebar on Feb 16 2016, 5:03 PM.

Details

Reviewers
jingyue
hfinkel
Commits
rGf9b5add6adb7: [NVPTX] Annotate convergent intrinsics as convergent.
rL261115: [NVPTX] Annotate convergent intrinsics as convergent.
Summary

Previously the machine instructions for bar.sync &co. were not marked as
convergent. This resulted in some MI passes (such as TailDuplication,
fixed in an upcoming patch) doing unsafe things to these instructions.

Diff Detail

Event Timeline

jlebar updated this revision to Diff 48133.Feb 16 2016, 5:03 PM
jlebar retitled this revision from to [NVPTX] Annotate convergent intrinsics as convergent..
jlebar updated this object.
jlebar added reviewers: jingyue, hfinkel.
jlebar added subscribers: tra, llvm-commits.
Herald added a subscriber: jholewinski. · View Herald TranscriptFeb 16 2016, 5:03 PM
jingyue accepted this revision.Feb 16 2016, 9:09 PM
jingyue edited edge metadata.
This revision is now accepted and ready to land.Feb 16 2016, 9:09 PM
Closed by commit rL261115: [NVPTX] Annotate convergent intrinsics as convergent. (authored by jlebar). · Explain WhyFeb 17 2016, 9:51 AM
This revision was automatically updated to reflect the committed changes.

Revision Contents

PathSize
lib/
Target/
NVPTX/
2 lines

Diff 48133

lib/Target/NVPTX/NVPTXIntrinsics.td

  • This file is larger than 256 KB, so syntax highlighting is disabled by default.
Show All 26 Linesdef immDouble1 : PatLeaf<(fpimm), [{
return (d==1.0); return (d==1.0);
}]>; }]>;
//----------------------------------- //-----------------------------------
// Synchronization Functions // Synchronization Functions
//----------------------------------- //-----------------------------------
let isConvergent = 1 in {
def INT_CUDA_SYNCTHREADS : NVPTXInst<(outs), (ins), def INT_CUDA_SYNCTHREADS : NVPTXInst<(outs), (ins),
"bar.sync \t0;", "bar.sync \t0;",
[(int_cuda_syncthreads)]>; [(int_cuda_syncthreads)]>;
def INT_BARRIER0 : NVPTXInst<(outs), (ins), def INT_BARRIER0 : NVPTXInst<(outs), (ins),
"bar.sync \t0;", "bar.sync \t0;",
[(int_nvvm_barrier0)]>; [(int_nvvm_barrier0)]>;
def INT_BARRIER0_POPC : NVPTXInst<(outs Int32Regs:$dst), (ins Int32Regs:$pred), def INT_BARRIER0_POPC : NVPTXInst<(outs Int32Regs:$dst), (ins Int32Regs:$pred),
!strconcat("{{ \n\t", !strconcat("{{ \n\t",
Show All 15 Linesdef INT_BARRIER0_OR : NVPTXInst<(outs Int32Regs:$dst), (ins Int32Regs:$pred),
!strconcat("{{ \n\t", !strconcat("{{ \n\t",
!strconcat(".reg .pred \t%p1; \n\t", !strconcat(".reg .pred \t%p1; \n\t",
!strconcat(".reg .pred \t%p2; \n\t", !strconcat(".reg .pred \t%p2; \n\t",
!strconcat("setp.ne.u32 \t%p1, $pred, 0; \n\t", !strconcat("setp.ne.u32 \t%p1, $pred, 0; \n\t",
!strconcat("bar.red.or.pred \t%p2, 0, %p1; \n\t", !strconcat("bar.red.or.pred \t%p2, 0, %p1; \n\t",
!strconcat("selp.u32 \t$dst, 1, 0, %p2; \n\t", !strconcat("selp.u32 \t$dst, 1, 0, %p2; \n\t",
!strconcat("}}", ""))))))), !strconcat("}}", ""))))))),
[(set Int32Regs:$dst, (int_nvvm_barrier0_or Int32Regs:$pred))]>; [(set Int32Regs:$dst, (int_nvvm_barrier0_or Int32Regs:$pred))]>;
} // isConvergent = 1
//----------------------------------- //-----------------------------------
// Explicit Memory Fence Functions // Explicit Memory Fence Functions
//----------------------------------- //-----------------------------------
class MEMBAR<string StrOp, Intrinsic IntOP> : class MEMBAR<string StrOp, Intrinsic IntOP> :
NVPTXInst<(outs), (ins), NVPTXInst<(outs), (ins),
StrOp, [(IntOP)]>; StrOp, [(IntOP)]>;
▲ Show 20 Lines • Show All 6,974 Lines • Show Last 20 Lines