================
@@ -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.
----------------
RyanRio wrote:
Empty mask means zero mask [to me] 😂 I'll change it though, fine with me.
https://github.com/llvm/llvm-project/pull/220442
_______________________________________________
cfe-commits mailing list
[email protected]
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
