summaryrefslogtreecommitdiffstats
path: root/llvm
diff options
context:
space:
mode:
authorMatt Arsenault <Matthew.Arsenault@amd.com>2019-07-17 22:41:53 +0000
committerMatt Arsenault <Matthew.Arsenault@amd.com>2019-07-17 22:41:53 +0000
commited9a91ce7ec6f8154ae5ded8c54f1521731622a6 (patch)
tree0452abd76537ef5f31dbb2c9d15df7bc3c4f57d0 /llvm
parent3628a8fae9f30ec024a5781aaf607c96c4ded05a (diff)
downloadbcm5719-llvm-ed9a91ce7ec6f8154ae5ded8c54f1521731622a6.tar.gz
bcm5719-llvm-ed9a91ce7ec6f8154ae5ded8c54f1521731622a6.zip
AMDGPU: Set inaccessiblememonly on sendmsg intrinsics
llvm-svn: 366384
Diffstat (limited to 'llvm')
-rw-r--r--llvm/include/llvm/IR/IntrinsicsAMDGPU.td4
1 files changed, 2 insertions, 2 deletions
diff --git a/llvm/include/llvm/IR/IntrinsicsAMDGPU.td b/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
index 1f835171386..3982444b540 100644
--- a/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
+++ b/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
@@ -199,9 +199,9 @@ def int_amdgcn_wavefrontsize :
// The first parameter is s_sendmsg immediate (i16),
// the second one is copied to m0
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">,
- 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">,
Intrinsic<[], [], [IntrConvergent]>;
OpenPOWER on IntegriCloud