This is an archive of the discontinued LLVM Phabricator instance.

Differential D21161

[NVPTX] Mark bar.sync intrinsic as convergent.
ClosedPublic

Authored by jlebar on Jun 8 2016, 5:22 PM.

Details

Reviewers
tra
Commits
rG17b4701070e0: [NVPTX] Mark bar.sync intrinsic as convergent.
rL272295: [NVPTX] Mark bar.sync intrinsic as convergent.
Summary

__syncthreads, which corresponds to bar.sync 0, is already convergent.
This makes the more general bar.sync n likewise convergent.

Diff Detail

Repository
rL LLVM

Event Timeline

jlebar updated this revision to Diff 60124.Jun 8 2016, 5:22 PM
jlebar retitled this revision from to [NVPTX] Mark bar.sync intrinsic as convergent..
jlebar updated this object.
jlebar added a reviewer: tra.
jlebar added subscribers: jholewinski, llvm-commits.
jholewinski added a comment.Jun 9 2016, 5:00 AM

Looks good to me!

tra accepted this revision.Jun 9 2016, 11:02 AM
tra edited edge metadata.

LGTM

This revision is now accepted and ready to land.Jun 9 2016, 11:02 AM
Closed by commit rL272295: [NVPTX] Mark bar.sync intrinsic as convergent. (authored by jlebar). · Explain WhyJun 9 2016, 12:56 PM
This revision was automatically updated to reflect the committed changes.

Revision Contents

PathSize
llvm/
trunk/
include/
llvm/
IR/
2 lines

Diff 60225

llvm/trunk/include/llvm/IR/IntrinsicsNVVM.td

Show First 20 Lines • Show All 3,738 Lines • ▼ Show 20 Linesdef int_ptx_read_pm0 : PTXReadSpecialRegisterIntrinsic_r32
<"__builtin_ptx_read_pm0">; <"__builtin_ptx_read_pm0">;
def int_ptx_read_pm1 : PTXReadSpecialRegisterIntrinsic_r32 def int_ptx_read_pm1 : PTXReadSpecialRegisterIntrinsic_r32
<"__builtin_ptx_read_pm1">; <"__builtin_ptx_read_pm1">;
def int_ptx_read_pm2 : PTXReadSpecialRegisterIntrinsic_r32 def int_ptx_read_pm2 : PTXReadSpecialRegisterIntrinsic_r32
<"__builtin_ptx_read_pm2">; <"__builtin_ptx_read_pm2">;
def int_ptx_read_pm3 : PTXReadSpecialRegisterIntrinsic_r32 def int_ptx_read_pm3 : PTXReadSpecialRegisterIntrinsic_r32
<"__builtin_ptx_read_pm3">; <"__builtin_ptx_read_pm3">;
def int_ptx_bar_sync : Intrinsic<[], [llvm_i32_ty], []>, def int_ptx_bar_sync : Intrinsic<[], [llvm_i32_ty], [IntrConvergent]>,
GCCBuiltin<"__builtin_ptx_bar_sync">; GCCBuiltin<"__builtin_ptx_bar_sync">;