From 8f42855556f8d2675b2ada33c59d8b3baf5ed391 Mon Sep 17 00:00:00 2001
From: Matt Arsenault <Matthew.Arsenault@amd.com>
Date: Tue, 25 Jun 2019 13:03:06 +0000
Subject: [PATCH] AMDGPU: Make amdgcn.s.get.waveid.in.workgroup
 inaccessiblememonly

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

llvm-svn: 364304
---
 include/llvm/IR/IntrinsicsAMDGPU.td | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/include/llvm/IR/IntrinsicsAMDGPU.td b/include/llvm/IR/IntrinsicsAMDGPU.td
index b646593344c..4ac8d6ab8be 100644
--- 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.