[PATCH] D124700: [AMDGPU] Add llvm.amdgcn.sched.barrier intrinsic

[Austin Kerbow via Phabricator via cfe-commits]

kerbowa added a comment.

In D124700#3483556 <https://reviews.llvm.org/D124700#3483556>, @rampitec wrote:

> You do not handle masks other than 0 yet?

We handle 0 and 1 only.



================
Comment at: llvm/include/llvm/IR/IntrinsicsAMDGPU.td:219
+//     MASK = 0: No instructions may be scheduled across SCHED_BARRIER.
+//     MASK = 1: Non-memory, non-side-effect producing instructions may be
+//               scheduled across SCHED_BARRIER, i.e. allow ALU instructions to pass.
----------------
rampitec wrote:
> Since you are going to extend it I'd suggest this is -1. Then you will start carving bits outs of it. That way if someone start to use it it will still work after update.
Since the most common use case will be to block all instruction types I thought having that be MASK = 0 made the most sense. After that, we carve out bits for types of instructions that should be scheduled across it.

There may be modes where we restrict certain types of memops, so we cannot have MASK = 1 above changed to -1. Since this (MASK = 1) is allowing all ALU across we could define which bits mean VALU/SALU/MFMA etc and use that mask if you think it's better. I'm worried we won't be able to anticipate all the types that we could want to be maskable. It might be better to just have a single bit that can mean all ALU, or all MemOps, and so on to avoid this problem.


================
Comment at: llvm/include/llvm/IR/IntrinsicsAMDGPU.td:222
+def int_amdgcn_sched_barrier : GCCBuiltin<"__builtin_amdgcn_sched_barrier">,
+  Intrinsic<[], [llvm_i16_ty], [ImmArg<ArgIndex<0>>, IntrNoMem,
+                                IntrHasSideEffects, IntrConvergent, IntrWillReturn]>;
----------------
rampitec wrote:
> Why not full i32? This is immediate anyway but you will have more bits for the future.
Good point thanks.


Repository:
  rG LLVM Github Monorepo

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

https://reviews.llvm.org/D124700