From: Matt Arsenault <Matthew.Arsenault@amd.com>
Date: Fri, 8 Dec 2017 20:01:02 +0000 (+0000)
Subject: AMDGPU: Set IntrReadMem on memtime intrinsics
X-Git-Url: https://granicus.if.org/sourcecode?a=commitdiff_plain;h=46bfd217ebd9465c9d3575bcf4d3b3bd11bc438b;p=llvm

AMDGPU: Set IntrReadMem on memtime intrinsics

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

diff --git a/include/llvm/IR/IntrinsicsAMDGPU.td b/include/llvm/IR/IntrinsicsAMDGPU.td
index d72a8c6798a..d7999cd3323 100644
--- a/include/llvm/IR/IntrinsicsAMDGPU.td
+++ b/include/llvm/IR/IntrinsicsAMDGPU.td
@@ -570,7 +570,7 @@ def int_amdgcn_s_dcache_inv :
 
 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_dcache_wb_vol :
 
 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 :
diff --git a/lib/Target/AMDGPU/SMInstructions.td b/lib/Target/AMDGPU/SMInstructions.td
index 5e72a2e8828..8f347986eb8 100644
--- a/lib/Target/AMDGPU/SMInstructions.td
+++ b/lib/Target/AMDGPU/SMInstructions.td
@@ -129,11 +129,8 @@ class SM_Time_Pseudo<string opName, SDPatternOperator node> : SM_Pseudo<
   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;
 }