This is an archive of the discontinued LLVM Phabricator instance.

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

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

Reviewers

tra

Commits

Summary

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

jlebar retitled this revision from to [NVPTX] Mark bar.sync intrinsic as convergent..

jlebar updated this object.

jlebar added a reviewer: tra.

Looks good to me!

LGTM

This revision is now accepted and ready to land.Jun 9 2016, 11:02 AM

This revision was automatically updated to reflect the committed changes.

Path

Size

include/

llvm/

IR/

2 lines

Show First 20 Lines • Show All 3,738 Lines • ▼ Show 20 Lines	def 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">;

//		//
// SHUFFLE		// SHUFFLE
//		//

// shfl.down.b32 dest, val, offset, mask_and_clamp		// shfl.down.b32 dest, val, offset, mask_and_clamp
def int_ptx_shfl_down_i32 :		def int_ptx_shfl_down_i32 :
Show All 37 Lines