[PATCH] D67218: [AMDGPU] Mark s_barrier as having side effects but not accessing memory.

[Matt Arsenault via Phabricator via llvm-commits]

arsenm accepted this revision.
arsenm added a comment.
This revision is now accepted and ready to land.

LGTM



================
Comment at: llvm/include/llvm/IR/IntrinsicsAMDGPU.td:209
 def int_amdgcn_s_barrier : GCCBuiltin<"__builtin_amdgcn_s_barrier">,
-  Intrinsic<[], [], [IntrConvergent]>;
+  Intrinsic<[], [], [IntrNoMem, IntrHasSideEffects, IntrConvergent]>;
 
----------------
This does end up not adding readnone to the intrinsic declaration, correct?


================
Comment at: llvm/test/CodeGen/AMDGPU/schedule-barrier.mir:44
+    %23:vgpr_32 = V_ADD_U32_e32 %39, %40, implicit $exec
+    GLOBAL_STORE_DWORD %38, %23, 0, 0, 0, 0, implicit $exec :: (store 4 into `i32 addrspace(1)* undef`, addrspace 1)
+    S_ENDPGM 0
----------------
You can remove the into and IR fragment 


Repository:
  rG LLVM Github Monorepo

CHANGES SINCE LAST ACTION
  https://reviews.llvm.org/D67218/new/

https://reviews.llvm.org/D67218