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
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
