[PATCH] D41028: AMDGPU: Set IntrReadMem on memtime intrinsics

[Matt Arsenault via Phabricator via llvm-commits]

arsenm created this revision.
Herald added subscribers: t-tye, tpr, dstuttard, yaxunl, nhaehnle, wdng, kzhuravl.

https://reviews.llvm.org/D41028

Files:
  include/llvm/IR/IntrinsicsAMDGPU.td
  lib/Target/AMDGPU/SMInstructions.td


Index: lib/Target/AMDGPU/SMInstructions.td
===================================================================
--- lib/Target/AMDGPU/SMInstructions.td
+++ lib/Target/AMDGPU/SMInstructions.td
@@ -129,11 +129,8 @@
   opName, (outs SReg_64_XEXEC:$sdst), (ins),
   " $sdst", [(set i64:$sdst, (node))]> {
   let hasSideEffects = 1;
-  // FIXME: mayStore = ? is a workaround for tablegen bug for different
-  // inferred mayStore flags for the instruction pattern vs. standalone
-  // Pat. Each considers the other contradictory.
-  let mayStore = ?;
-  let mayLoad = ?;
+  let mayStore = 0;
+  let mayLoad = 1;
   let has_sbase = 0;
   let has_offset = 0;
 }
Index: include/llvm/IR/IntrinsicsAMDGPU.td
===================================================================
--- include/llvm/IR/IntrinsicsAMDGPU.td
+++ include/llvm/IR/IntrinsicsAMDGPU.td
@@ -570,7 +570,7 @@
 
 def int_amdgcn_s_memtime :
   GCCBuiltin<"__builtin_amdgcn_s_memtime">,
-  Intrinsic<[llvm_i64_ty], [], []>;
+  Intrinsic<[llvm_i64_ty], [], [IntrReadMem]>;
 
 def int_amdgcn_s_sleep :
   GCCBuiltin<"__builtin_amdgcn_s_sleep">,
@@ -816,7 +816,7 @@
 
 def int_amdgcn_s_memrealtime :
   GCCBuiltin<"__builtin_amdgcn_s_memrealtime">,
-  Intrinsic<[llvm_i64_ty], [], []>;
+  Intrinsic<[llvm_i64_ty], [], [IntrReadMem]>;
 
 // llvm.amdgcn.ds.permute <index> <src>
 def int_amdgcn_ds_permute :


-------------- next part --------------
A non-text attachment was scrubbed...
Name: D41028.126185.patch
Type: text/x-patch
Size: 1356 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/llvm-commits/attachments/20171208/e711433f/attachment-0001.bin>