[clang] [llvm] [AMDGPU] Implement AsyncMark Stages (PR #220442)

[Sameer Sahasrabuddhe via llvm-commits]

================
@@ -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