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

Stanislav Mekhanoshin via Phabricator via cfe-commits cfe-commits at lists.llvm.org
Fri Apr 29 15:08:49 PDT 2022

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?

  rG LLVM Github Monorepo



More information about the cfe-commits mailing list