[Mlir-commits] [llvm] [mlir] [LLVM][NVPTX] Add Rubin extensions for G2S Tensor intrinsics (PR #220029)

[Rajat Bajpai]

================
@@ -677,58 +677,94 @@ def CTAGroupFlags : Operand<i32> {
 def tma_cta_group_imm0 : TImmLeaf<i32, [{return Imm == 0;}]>;
 def tma_cta_group_imm_any : TImmLeaf<i32, [{return Imm >= 0;}]>;
 
+def TMAValidateDataFlags : Operand<i32> {
+  let PrintMethod = "printTMAValidateDataFlags";
+}
+
+// Matches validate pattern values 0-5. The 1-5 values are Rubin-only
+// validate patterns. 0 denotes disabled validate pattern (default).
+def tma_validate_data_imm : TImmLeaf<i32, [{
+  return Imm == 0 || (Imm >= 1 && Imm <= 5 &&
+                      Subtarget->hasFeature(NVPTX::SM107f));
+}]>;
+
 multiclass TMA_TENSOR_G2S_INTR<int dim, string mode, list<Predicate> pred,
                                TImmLeaf cta_group_type = tma_cta_group_imm_any> {
   defvar dims_dag = TMA_DIMS_DAG_UTIL<dim>.ins_dag;
   defvar dims_str = TMA_DIMS_DAG_UTIL<dim>.base_str;
-  defvar asm_str_base = "$cg [$dst], [$tmap, {{" # dims_str # "}}], [$mbar]";
-
+  defvar asm_str_base = " [$dst], [$tmap, {{" # dims_str # "}}], [$mbar]";
   defvar im2col_dag = TMA_IM2COL_UTIL<dim, mode>.ins_dag;
   defvar im2col_str = TMA_IM2COL_UTIL<dim, mode>.base_str;
-  defvar asm_str = !if(!empty(im2col_str),
-                       asm_str_base,
-                       asm_str_base # ", {{" # im2col_str # "}}");
-
+  defvar asm_str = !if(!empty(im2col_str), asm_str_base, asm_str_base # ", {{" # im2col_str # "}}");
   defvar dim_val = TMA_DIMS_UTIL<dim, mode>.num_dims;
+
   defvar inst_name = "cp.async.bulk.tensor"
                      # "." # dim_val # "d"
                      # "." # "shared::cluster.global"
                      # "." # !subst("_", "::", mode)
-                     # "." # "mbarrier::complete_tx::bytes";
+                     # "." # "mbarrier::complete_tx::bytes"
+                     # "$validate" # "$cg";
+
+  defvar ins_dag_base = !con(
+                          (ins ADDR:$dst, ADDR:$mbar, B64:$tmap),
+                          dims_dag, im2col_dag);
+  defvar ins_tail = (ins TMAValidateDataFlags:$validate, CTAGroupFlags:$cg);
----------------
rajatbajpai wrote:

sure, fine with me.

https://github.com/llvm/llvm-project/pull/220029