AMDGPU: Remove IntrReadMem from memtime/memrealtime intrinsics

[arsenm]

Summary

EarlyCSE with MemorySSA was able to use this to merge multiple calls
with no intervening store.

[rampitec]

1. These intrinsics read internal state. If not IntrReadMem something needs to be added to model it.
2. Patch does more than stated in the description with no justification.

lib/Target/AMDGPU/SMInstructions.td

160

It was zero before this patch.

803

It does not belong to the patch.

[arsenm]

In D58635#1409379, @rampitec wrote:

1. These intrinsics read internal state. If not IntrReadMem something needs to be added to model it.

IntrReadMem doesn't mean it reads memory, it means it only reads memory. Reading and writing memory is the default assumption. The negated behavior of the intrinsic memory attributes is part of why the tablegen inference is so broken

2. Patch does more than stated in the description with no justification.

lib/Target/AMDGPU/SMInstructions.td

160

This is to defeat tablegen's completely broken inference of instruction flags based on intrinsic properties. This conflicts with the assumed mayStore since the intrinsic isn't IntrReadMem. The assumption the flags directly correspond to the intrinsic attributes is broken because the weird memory dependency is tracked here with hasSideEffects. It's even more broken because this was working with readcyclecounter already, since it isn't a target intrinsic.

803

This was part of trying to avoid not explicitly setting the mayStore flag, but that ultimately didn't work

[arsenm]

Remove leftover attempt to defeat property inference

[rampitec]

lib/Target/AMDGPU/SMInstructions.td

160

At the very least it needs a comment explaining why is it '?'.

[arsenm]

Add comment

[rampitec]

LGTM

[arsenm]

r354814

include/llvm/IR/IntrinsicsAMDGPU.td

def int_amdgcn_buffer_wbinvl1 :
   Intrinsic<[], [], []>;
 
 def int_amdgcn_s_dcache_inv :
   GCCBuiltin<"__builtin_amdgcn_s_dcache_inv">,
   Intrinsic<[], [], []>;
 
 def int_amdgcn_s_memtime :
   GCCBuiltin<"__builtin_amdgcn_s_memtime">,
-  Intrinsic<[llvm_i64_ty], [], [IntrReadMem]>;
+  Intrinsic<[llvm_i64_ty], []>;
 
 def int_amdgcn_s_sleep :
   GCCBuiltin<"__builtin_amdgcn_s_sleep">,
   Intrinsic<[], [llvm_i32_ty], []> {
 }
 
 def int_amdgcn_s_incperflevel :
   GCCBuiltin<"__builtin_amdgcn_s_incperflevel">,
def int_amdgcn_s_dcache_wb :
   Intrinsic<[], [], []>;
 
 def int_amdgcn_s_dcache_wb_vol :
   GCCBuiltin<"__builtin_amdgcn_s_dcache_wb_vol">,
   Intrinsic<[], [], []>;
 
 def int_amdgcn_s_memrealtime :
   GCCBuiltin<"__builtin_amdgcn_s_memrealtime">,
-  Intrinsic<[llvm_i64_ty], [], [IntrReadMem]>;
+  Intrinsic<[llvm_i64_ty]>;
 
 // llvm.amdgcn.ds.permute <index> <src>
 def int_amdgcn_ds_permute :
   GCCBuiltin<"__builtin_amdgcn_ds_permute">,
   Intrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty], [IntrNoMem, IntrConvergent]>;
 
 // llvm.amdgcn.ds.bpermute <index> <src>
 def int_amdgcn_ds_bpermute :

lib/Target/AMDGPU/SMInstructions.td

multiclass SM_Pseudo_Stores<string opName,
   }
 }
 
 multiclass SM_Pseudo_Discards<string opName> {
   def _IMM  : SM_Discard_Pseudo <opName, (ins SReg_64:$sbase, smrd_offset_20:$offset), 1>;
   def _SGPR : SM_Discard_Pseudo <opName, (ins SReg_64:$sbase, SReg_32:$offset), 0>;
 }
 
-class SM_Time_Pseudo<string opName, SDPatternOperator node> : SM_Pseudo<
+class SM_Time_Pseudo<string opName, SDPatternOperator node = null_frag> : SM_Pseudo<
   opName, (outs SReg_64_XEXEC:$sdst), (ins),
   " $sdst", [(set i64:$sdst, (node))]> {
   let hasSideEffects = 1;
-  let mayStore = 0;
+
+  // FIXME: This should be definitively mayStore = 0. TableGen
+  // brokenly tries to infer these based on the intrinsic properties
+  // corresponding to the IR attributes. The target intrinsics are
+  // considered as writing to memory for IR dependency purposes, but
+  // those can be modeled with hasSideEffects here. These also end up
+  // inferring differently for llvm.readcyclecounter and the amdgcn
+  // intrinsics.
+  let mayStore = ?;
   let mayLoad = 1;
   let has_sbase = 0;
   let has_offset = 0;
 }
 
 class SM_Inval_Pseudo <string opName, SDPatternOperator node> : SM_Pseudo<
   opName, (outs), (ins), "", [(node)]> {
   let hasSideEffects = 1;

 
 let OtherPredicates = [isVI] in {
 
 def : GCNPat <
   (i64 (readcyclecounter)),
   (S_MEMREALTIME)
 >;
 
 } // let OtherPredicates = [isVI]

test/Transforms/EarlyCSE/AMDGPU/lit.local.cfg

+config.suffixes = ['.ll']
+
+targets = set(config.root.targets_to_build.split())
+if not 'AMDGPU' in targets:
+    config.unsupported = True

test/Transforms/EarlyCSE/AMDGPU/memrealtime.ll

+; RUN: opt -S -mtriple=amdgcn-amd-amdhsa -early-cse-memssa < %s | FileCheck %s
+target datalayout = "e-p:64:64-p1:64:64-p2:32:32-p3:32:32-p4:64:64-p5:32:32-p6:32:32-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024-v2048:2048-n32:64-S32-A5"
+
+; CHECK-LABEL: @memrealtime(
+; CHECK: call i64 @llvm.amdgcn.s.memrealtime()
+; CHECK: call i64 @llvm.amdgcn.s.memrealtime()
+define amdgpu_kernel void @memrealtime(i64 %cycles) #0 {
+entry:
+  %0 = tail call i64 @llvm.amdgcn.s.memrealtime()
+  %cmp3 = icmp sgt i64 %cycles, 0
+  br i1 %cmp3, label %while.body, label %while.end
+
+while.body:
+  %1 = tail call i64 @llvm.amdgcn.s.memrealtime()
+  %sub = sub nsw i64 %1, %0
+  %cmp = icmp slt i64 %sub, %cycles
+  br i1 %cmp, label %while.body, label %while.end
+
+while.end:
+  ret void
+}
+
+; CHECK-LABEL: @memtime(
+; CHECK: call i64 @llvm.amdgcn.s.memtime()
+; CHECK: call i64 @llvm.amdgcn.s.memtime()
+define amdgpu_kernel void @memtime(i64 %cycles) #0 {
+entry:
+  %0 = tail call i64 @llvm.amdgcn.s.memtime()
+  %cmp3 = icmp sgt i64 %cycles, 0
+  br i1 %cmp3, label %while.body, label %while.end
+
+while.body:
+  %1 = tail call i64 @llvm.amdgcn.s.memtime()
+  %sub = sub nsw i64 %1, %0
+  %cmp = icmp slt i64 %sub, %cycles
+  br i1 %cmp, label %while.body, label %while.end
+
+while.end:
+  ret void
+}
+
+declare i64 @llvm.amdgcn.s.memrealtime()
+declare i64 @llvm.amdgcn.s.memtime()