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

[Stanislav Mekhanoshin via Phabricator via cfe-commits]

rampitec added a comment.

You do not handle masks other than 0 yet?



================
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.
----------------
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.


================
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]>;
----------------
Why not full i32? This is immediate anyway but you will have more bits for the future.


================
Comment at: llvm/lib/Target/AMDGPU/AMDGPUMCInstLower.cpp:213
+        OutStreamer->emitRawComment(" sched_barrier mask(" +
+                                    Twine(MI->getOperand(0).getImm()) + ")");
+      }
----------------
Use hex?


Repository:
  rG LLVM Github Monorepo

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

https://reviews.llvm.org/D124700