================
@@ -214,6 +215,22 @@ bool SemaAMDGPU::CheckAMDGCNBuiltinFunctionCall(const
TargetInfo &TI,
case AMDGPU::BI__builtin_amdgcn_cvt_scale_pk16_f32_fp6:
case AMDGPU::BI__builtin_amdgcn_cvt_scale_pk16_f32_bf6:
return SemaRef.BuiltinConstantArgRange(TheCall, 2, 0, 15);
+ case AMDGPU::BI__builtin_amdgcn_asyncmark:
+ case AMDGPU::BI__builtin_amdgcn_wait_asyncmark: {
+ bool IsMark = BuiltinID == AMDGPU::BI__builtin_amdgcn_asyncmark;
+
+ // Check the sequence length limit is a constant.
+ llvm::APSInt NumMarks;
+ if (!IsMark && SemaRef.BuiltinConstantArg(TheCall, 0, NumMarks))
+ return true;
+
+ // The stage mask names the stages to act on, so every combination of known
+ // stage bits is meaningful, including none of them: the empty mask names
+ // every stage.
----------------
ssahasra wrote:
> the empty mask names every stage
We say this in multiple places, but actually there is no way to pass an empty
mask in the current change. There is a zero mask, which is not the mask being
absent. Do you plan to keep the old intrinsic name with the actual empty
parameter list, and create a new name which takes an explicit mask argument?
https://github.com/llvm/llvm-project/pull/220442
_______________________________________________
cfe-commits mailing list
[email protected]
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
