AMDGPU: Make amdgcn.s.get.waveid.in.workgroup inaccessiblememonly

This should probably be readnone, even though the instruction looks
like a load.

git-svn-id: https://llvm.org/svn/llvm-project/llvm/trunk@364304 91177308-0d34-0410-b5e6-96231b3b80d8

diff --git a/include/llvm/IR/IntrinsicsAMDGPU.td b/include/llvm/IR/IntrinsicsAMDGPU.td
index b646593344cc444e9fcd9dbea0269ac9fbeb815e..4ac8d6ab8bef29cd6d6b4ebe9e6e6477b7187b50 100644 (file)
--- a/include/llvm/IR/IntrinsicsAMDGPU.td
+++ b/include/llvm/IR/IntrinsicsAMDGPU.td
@@ -1524,7 +1524,7 @@ def int_amdgcn_mov_dpp8 :
 
 def int_amdgcn_s_get_waveid_in_workgroup :
   GCCBuiltin<"__builtin_amdgcn_s_get_waveid_in_workgroup">,
-  Intrinsic<[llvm_i32_ty], [], [IntrReadMem]>;
+  Intrinsic<[llvm_i32_ty], [], [IntrReadMem, IntrInaccessibleMemOnly]>;
 
 //===----------------------------------------------------------------------===//
 // Deep learning intrinsics.