[Mlir-commits] [mlir] [MLIR][NVVM] Update TMA Load Op (PR #156347)
Guray Ozen
llvmlistbot at llvm.org
Mon Sep 1 09:10:49 PDT 2025
================
@@ -2367,6 +2367,23 @@ def TMAStoreModeAttr : EnumAttr<NVVM_Dialect, TMAStoreMode, "tma_store_mode"> {
let assemblyFormat = "`<` $value `>`";
}
+// Num CTAs in a group participating in the TCGEN05 operation.
+// This corresponds to the "cta_group::1", "cta_group::2"
+// modifiers in the PTX instructions.
+def Tcgen05GroupCTA_1 : I32EnumAttrCase<"CTA_1", 0, "cta_1">;
+def Tcgen05GroupCTA_2 : I32EnumAttrCase<"CTA_2", 1, "cta_2">;
+
+def Tcgen05GroupKind : I32EnumAttr<"Tcgen05GroupKind",
+ "NVVM Tcgen05 group kind",
+ [Tcgen05GroupCTA_1, Tcgen05GroupCTA_2]> {
+ let genSpecializedAttr = 0;
+ let cppNamespace = "::mlir::NVVM";
+}
+def Tcgen05GroupKindAttr :
----------------
grypp wrote:
Can we rename this to `CTAGroupKindAttr`. I think TMA lives beyond tcgen05.
https://github.com/llvm/llvm-project/pull/156347
More information about the Mlir-commits
mailing list