RPCS3/llvm
1
0
Fork 0
mirror of https://github.com/RPCS3/llvm.git synced 2025-03-04 00:29:28 +00:00
Code Issues Actions Packages Projects Releases Wiki Activity

[NVPTX] Annotate convergent intrinsics as convergent.

Browse Source 
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.

Reviewers: jingyue

Subscribers: llvm-commits, tra, jholewinski, hfinkel

Differential Revision: http://reviews.llvm.org/D17318

git-svn-id: https://llvm.org/svn/llvm-project/llvm/trunk@261115 91177308-0d34-0410-b5e6-96231b3b80d8
This commit is contained in:
Justin Lebar 2016-02-17 17:46:54 +00:00
parent 97cc80a0e8
commit 1d3171f22a
1 changed files with 2 additions and 0 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

2
lib/Target/NVPTX/NVPTXIntrinsics.td
View File

@ -32,6 +32,7 @@ def immDouble1 : PatLeaf<(fpimm), [{
//-----------------------------------
// Synchronization Functions
//-----------------------------------
let isConvergent = 1 in {
def INT_CUDA_SYNCTHREADS : NVPTXInst<(outs), (ins),
"bar.sync \t0;",
[(int_cuda_syncthreads)]>;
@ -63,6 +64,7 @@ def INT_BARRIER0_OR : NVPTXInst<(outs Int32Regs:$dst), (ins Int32Regs:$pred),
!strconcat("selp.u32 \t$dst, 1, 0, %p2; \n\t",
!strconcat("}}", ""))))))),
[(set Int32Regs:$dst, (int_nvvm_barrier0_or Int32Regs:$pred))]>;
} // isConvergent = 1
//-----------------------------------