This is an archive of the discontinued LLVM Phabricator instance.

Differential D64891

AMDGPU: Set inaccessiblememonly on sendmsg intrinsics
ClosedPublic

Authored by arsenm on Jul 17 2019, 3:29 PM.

Details

Reviewers
rampitec
nhaehnle

Diff Detail

Event Timeline

arsenm created this revision.Jul 17 2019, 3:29 PM
Herald added subscribers: t-tye, tpr, dstuttard and 4 others. · View Herald TranscriptJul 17 2019, 3:29 PM
rampitec accepted this revision.Jul 17 2019, 3:30 PM

LGTM

This revision is now accepted and ready to land.Jul 17 2019, 3:30 PM
arsenm closed this revision.Jul 17 2019, 3:41 PM

r366384

Revision Contents

PathSize
include/
llvm/
IR/
4 lines

Diff 210440

include/llvm/IR/IntrinsicsAMDGPU.td

Show First 20 Lines • Show All 193 Lines • ▼ Show 20 Lines
//===----------------------------------------------------------------------===// //===----------------------------------------------------------------------===//
// Instruction Intrinsics // Instruction Intrinsics
//===----------------------------------------------------------------------===// //===----------------------------------------------------------------------===//
// The first parameter is s_sendmsg immediate (i16), // The first parameter is s_sendmsg immediate (i16),
// the second one is copied to m0 // the second one is copied to m0
def int_amdgcn_s_sendmsg : GCCBuiltin<"__builtin_amdgcn_s_sendmsg">, def int_amdgcn_s_sendmsg : GCCBuiltin<"__builtin_amdgcn_s_sendmsg">,
Intrinsic <[], [llvm_i32_ty, llvm_i32_ty], [ImmArg<0>]>; Intrinsic <[], [llvm_i32_ty, llvm_i32_ty], [ImmArg<0>, IntrInaccessibleMemOnly]>;
def int_amdgcn_s_sendmsghalt : GCCBuiltin<"__builtin_amdgcn_s_sendmsghalt">, def int_amdgcn_s_sendmsghalt : GCCBuiltin<"__builtin_amdgcn_s_sendmsghalt">,
Intrinsic <[], [llvm_i32_ty, llvm_i32_ty], [ImmArg<0>]>; Intrinsic <[], [llvm_i32_ty, llvm_i32_ty], [ImmArg<0>, IntrInaccessibleMemOnly]>;
def int_amdgcn_s_barrier : GCCBuiltin<"__builtin_amdgcn_s_barrier">, def int_amdgcn_s_barrier : GCCBuiltin<"__builtin_amdgcn_s_barrier">,
Intrinsic<[], [], [IntrConvergent]>; Intrinsic<[], [], [IntrConvergent]>;
def int_amdgcn_wave_barrier : GCCBuiltin<"__builtin_amdgcn_wave_barrier">, def int_amdgcn_wave_barrier : GCCBuiltin<"__builtin_amdgcn_wave_barrier">,
Intrinsic<[], [], [IntrConvergent]>; Intrinsic<[], [], [IntrConvergent]>;
def int_amdgcn_s_waitcnt : GCCBuiltin<"__builtin_amdgcn_s_waitcnt">, def int_amdgcn_s_waitcnt : GCCBuiltin<"__builtin_amdgcn_s_waitcnt">,
▲ Show 20 Lines • Show All 1,600 Lines • Show Last 20 Lines