[clang] [llvm] [AMDGPU] Extend __builtin_amdgcn_sched_group_barrier to support rules. (PR #85304)
Matt Arsenault via cfe-commits
cfe-commits at lists.llvm.org
Wed Apr 24 00:56:41 PDT 2024
================
@@ -18763,19 +18763,28 @@ Value *CodeGenFunction::EmitAMDGPUBuiltinExpr(unsigned BuiltinID,
// scheduling builtins
case AMDGPU::BI__builtin_amdgcn_sched_group_barrier: {
- return E->getNumArgs() == 3
- ? Builder.CreateCall(
- CGM.getIntrinsic(Intrinsic::amdgcn_sched_group_barrier),
- {EmitScalarExpr(E->getArg(0)),
- EmitScalarExpr(E->getArg(1)),
- EmitScalarExpr(E->getArg(2))})
- : Builder.CreateCall(
- CGM.getIntrinsic(
- Intrinsic::amdgcn_sched_group_barrier_rule),
- {EmitScalarExpr(E->getArg(0)),
- EmitScalarExpr(E->getArg(1)),
- EmitScalarExpr(E->getArg(2)),
- EmitScalarExpr(E->getArg(3))});
+ if (E->getNumArgs() == 3)
+ return Builder.CreateCall(
+ CGM.getIntrinsic(Intrinsic::amdgcn_sched_group_barrier),
+ {EmitScalarExpr(E->getArg(0)), EmitScalarExpr(E->getArg(1)),
+ EmitScalarExpr(E->getArg(2))});
+
+ uint64_t Mask = 0;
+ for (unsigned I = 3; I < E->getNumArgs(); I++) {
+ auto NextArg = EmitScalarExpr(E->getArg(I));
+ auto ArgLiteral = cast<ConstantInt>(NextArg)->getZExtValue();
+ if (ArgLiteral > 63) {
+ CGM.Error(E->getExprLoc(),
+ getContext().BuiltinInfo.getName(BuiltinID).str() +
+ " RuleID must be within [0,63].");
----------------
arsenm wrote:
Should such checks go in Sema instead?
https://github.com/llvm/llvm-project/pull/85304
More information about the cfe-commits
mailing list