[Mlir-commits] [mlir] [python] Emit only dialect `EnumAttr` registrations. (PR #117918)

llvmlistbot at llvm.org llvmlistbot at llvm.org
Wed Nov 27 16:09:16 PST 2024 

llvmbot wrote:


<!--LLVM PR SUMMARY COMMENT-->

@llvm/pr-subscribers-backend-amdgpu

Author: Maksim Levental (makslevental)

<details>
<summary>Changes</summary>

## Problem

Currently we emit `EnumAttr` bindings indiscriminately (i.e., without considering `-bind-dialect`). This leads to two things

1. duplicated `class <OTHERDIALECT>Enum`s for every `<DIALECT>_enums_gen.py` that uses `<OTHERDIALECT>Enum`;
2. similarly, duplicated `@<!-- -->register_attribute_builder("<OTHERDIALECT>EnumAttr")` 

Note, the duplication is due to the fact that tablegen has no "memory" across emitted files so it has no knowledge that it has already emitted `XYZEnumAttr` in both `_<DIALECT>_enums_gen.py` and `_<OTHERDIALECT>_enums_gen.py`.

The first isn't an issue because even if they're duplicated, they're mechanically kept in sync. 

**The second is an issue** because [`attributeBuilderMap`](https://github.com/llvm/llvm-project/blob/7c850867b9ef4427375da6d83c34d0b9c944fcb8/mlir/lib/Bindings/Python/Globals.h#L120) doesn't automatically replace/override registrations (nor should it). For example, we have [`linalg.iterator_type`](https://github.com/llvm/llvm-project/blob/e7d09cecc9123f89ace1712a617e252d78b179e9/mlir/include/mlir/Dialect/Linalg/IR/LinalgBase.td#L78) and IREE has [`iree_gpu.iterator_type`](https://github.com/iree-org/iree/blob/fdc2d6a5987343b688226af11ad8f398f4acdcdb/compiler/src/iree/compiler/Codegen/Dialect/GPU/IR/IREEGPUAttrs.td#L22) and thus if you try to use both `_linalg_enums_gen.py` and `_iree_gpu_enums_gen.py` you will get 

```
RuntimeError: Attribute builder for 'IteratorType' is already registered with func: <function _iteratortype at 0x7f4f678cc5e0>
```

## Solution

The only solution is to cease generation of indiscriminate emission of registration calls. In this PR I do that, i.e., I add `-bind-dialect` to `-gen-python-enum-bindings` and filter which `register_attribute_builder` calls are emitted. The effect is that `builtin` dialect attribute builder registration calls are no longer emitted anywhere. So those have to be written by hand somewhere. Here I put them adjacent (in the `<DIALECT>.py`) to where they were being emitted prior. I also actually use the dialect to prefix the builder lookup key. So 

```
@<!-- -->register_attribute_builder("MatchInterfaceEnum")
```

becomes

```
@<!-- -->register_attribute_builder("builtin.MatchInterfaceEnum")
```

---

Patch is 45.42 KiB, truncated to 20.00 KiB below, full version: https://github.com/llvm/llvm-project/pull/117918.diff


26 Files Affected:

- (modified) mlir/cmake/modules/AddMLIRPython.cmake (+2-2) 
- (modified) mlir/python/CMakeLists.txt (+1-2) 
- (modified) mlir/python/mlir/dialects/_ods_common.py (+1) 
- (modified) mlir/python/mlir/dialects/amdgpu.py (+16) 
- (modified) mlir/python/mlir/dialects/arith.py (+35) 
- (modified) mlir/python/mlir/dialects/bufferization.py (+6) 
- (modified) mlir/python/mlir/dialects/gpu/__init__.py (+56) 
- (modified) mlir/python/mlir/dialects/index.py (+6) 
- (modified) mlir/python/mlir/dialects/linalg/__init__.py (+25) 
- (modified) mlir/python/mlir/dialects/linalg/opdsl/ops/core_named_ops.py (+1) 
- (modified) mlir/python/mlir/dialects/llvm.py (+106-1) 
- (modified) mlir/python/mlir/dialects/nvgpu.py (+26) 
- (modified) mlir/python/mlir/dialects/nvvm.py (+76) 
- (modified) mlir/python/mlir/dialects/sparse_tensor.py (+16) 
- (modified) mlir/python/mlir/dialects/transform/__init__.py (+10) 
- (modified) mlir/python/mlir/dialects/transform/extras/__init__.py (+1) 
- (modified) mlir/python/mlir/dialects/transform/structured.py (+10) 
- (modified) mlir/python/mlir/dialects/transform/vector.py (+21) 
- (modified) mlir/python/mlir/dialects/vector.py (+16) 
- (modified) mlir/python/mlir/ir.py (+54-54) 
- (modified) mlir/test/mlir-tblgen/enums-python-bindings.td (+2-14) 
- (modified) mlir/test/mlir-tblgen/op-python-bindings.td (+2-2) 
- (modified) mlir/test/python/dialects/index_dialect.py (+1-1) 
- (modified) mlir/test/python/dialects/transform_structured_ext.py (+1-1) 
- (modified) mlir/tools/mlir-tblgen/EnumPythonBindingGen.cpp (+16-7) 
- (modified) mlir/tools/mlir-tblgen/OpPythonBindingGen.cpp (+6-3) 


``````````diff
diff --git a/mlir/cmake/modules/AddMLIRPython.cmake b/mlir/cmake/modules/AddMLIRPython.cmake
index 7b91f43e2d57fd..d06fc927ea44d4 100644
--- a/mlir/cmake/modules/AddMLIRPython.cmake
+++ b/mlir/cmake/modules/AddMLIRPython.cmake
@@ -318,7 +318,7 @@ function(declare_mlir_dialect_python_bindings)
         set(LLVM_TARGET_DEFINITIONS ${td_file})
       endif()
       set(enum_filename "${relative_td_directory}/_${ARG_DIALECT_NAME}_enum_gen.py")
-      mlir_tablegen(${enum_filename} -gen-python-enum-bindings)
+      mlir_tablegen(${enum_filename} -gen-python-enum-bindings -bind-dialect=${ARG_DIALECT_NAME})
       list(APPEND _sources ${enum_filename})
     endif()
 
@@ -390,7 +390,7 @@ function(declare_mlir_dialect_extension_python_bindings)
         set(LLVM_TARGET_DEFINITIONS ${td_file})
       endif()
       set(enum_filename "${relative_td_directory}/_${ARG_EXTENSION_NAME}_enum_gen.py")
-      mlir_tablegen(${enum_filename} -gen-python-enum-bindings)
+      mlir_tablegen(${enum_filename} -gen-python-enum-bindings -bind-dialect=${ARG_DIALECT_NAME})
       list(APPEND _sources ${enum_filename})
     endif()
 
diff --git a/mlir/python/CMakeLists.txt b/mlir/python/CMakeLists.txt
index 23187f256455bb..9949743b9bf09c 100644
--- a/mlir/python/CMakeLists.txt
+++ b/mlir/python/CMakeLists.txt
@@ -63,8 +63,7 @@ declare_mlir_dialect_python_bindings(
   TD_FILE dialects/AffineOps.td
   SOURCES
     dialects/affine.py
-  DIALECT_NAME affine
-  GEN_ENUM_BINDINGS)
+  DIALECT_NAME affine)
 
 declare_mlir_dialect_python_bindings(
   ADD_TO_PARENT MLIRPythonSources.Dialects
diff --git a/mlir/python/mlir/dialects/_ods_common.py b/mlir/python/mlir/dialects/_ods_common.py
index d40d936cdc83d6..22fc588e180b1d 100644
--- a/mlir/python/mlir/dialects/_ods_common.py
+++ b/mlir/python/mlir/dialects/_ods_common.py
@@ -143,6 +143,7 @@ def get_op_result_or_op_results(
         else op
     )
 
+
 ResultValueTypeTuple = _cext.ir.Operation, _cext.ir.OpView, _cext.ir.Value
 ResultValueT = _Union[ResultValueTypeTuple]
 VariadicResultValueT = _Union[ResultValueT, _Sequence[ResultValueT]]
diff --git a/mlir/python/mlir/dialects/amdgpu.py b/mlir/python/mlir/dialects/amdgpu.py
index 43d905d0c481cc..9b8beaa5571a26 100644
--- a/mlir/python/mlir/dialects/amdgpu.py
+++ b/mlir/python/mlir/dialects/amdgpu.py
@@ -2,5 +2,21 @@
 #  See https://llvm.org/LICENSE.txt for license information.
 #  SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 
+from ..ir import IntegerAttr, IntegerType, register_attribute_builder
 from ._amdgpu_ops_gen import *
 from ._amdgpu_enum_gen import *
+
+
+ at register_attribute_builder("builtin.AMDGPU_DPPPerm")
+def _amdgpu_dppperm(x, context):
+    return IntegerAttr.get(IntegerType.get_signless(32, context=context), int(x))
+
+
+ at register_attribute_builder("builtin.AMDGPU_MFMAPermB")
+def _amdgpu_mfmapermb(x, context):
+    return IntegerAttr.get(IntegerType.get_signless(32, context=context), int(x))
+
+
+ at register_attribute_builder("builtin.AMDGPU_SchedBarrierOpOpt")
+def _amdgpu_schedbarrieropopt(x, context):
+    return IntegerAttr.get(IntegerType.get_signless(32, context=context), int(x))
diff --git a/mlir/python/mlir/dialects/arith.py b/mlir/python/mlir/dialects/arith.py
index 92da5df9bce665..32ba832260d129 100644
--- a/mlir/python/mlir/dialects/arith.py
+++ b/mlir/python/mlir/dialects/arith.py
@@ -108,3 +108,38 @@ def constant(
     result: Type, value: Union[int, float, Attribute, _array], *, loc=None, ip=None
 ) -> Value:
     return _get_op_result_or_op_results(ConstantOp(result, value, loc=loc, ip=ip))
+
+
+ at register_attribute_builder("builtin.Arith_CmpFPredicateAttr")
+def _arith_cmpfpredicateattr(x, context):
+    return IntegerAttr.get(IntegerType.get_signless(64, context=context), int(x))
+
+
+ at register_attribute_builder("builtin.Arith_CmpIPredicateAttr")
+def _arith_cmpipredicateattr(x, context):
+    return IntegerAttr.get(IntegerType.get_signless(64, context=context), int(x))
+
+
+ at register_attribute_builder("builtin.Arith_DenormalMode")
+def _arith_denormalmode(x, context):
+    return IntegerAttr.get(IntegerType.get_signless(32, context=context), int(x))
+
+
+ at register_attribute_builder("builtin.Arith_IntegerOverflowFlags")
+def _arith_integeroverflowflags(x, context):
+    return IntegerAttr.get(IntegerType.get_signless(32, context=context), int(x))
+
+
+ at register_attribute_builder("builtin.Arith_RoundingModeAttr")
+def _arith_roundingmodeattr(x, context):
+    return IntegerAttr.get(IntegerType.get_signless(32, context=context), int(x))
+
+
+ at register_attribute_builder("builtin.AtomicRMWKindAttr")
+def _atomicrmwkindattr(x, context):
+    return IntegerAttr.get(IntegerType.get_signless(64, context=context), int(x))
+
+
+ at register_attribute_builder("builtin.FastMathFlags")
+def _fastmathflags(x, context):
+    return IntegerAttr.get(IntegerType.get_signless(32, context=context), int(x))
diff --git a/mlir/python/mlir/dialects/bufferization.py b/mlir/python/mlir/dialects/bufferization.py
index 759b6aa24a9ff7..6ad76c729ed2dc 100644
--- a/mlir/python/mlir/dialects/bufferization.py
+++ b/mlir/python/mlir/dialects/bufferization.py
@@ -2,5 +2,11 @@
 #  See https://llvm.org/LICENSE.txt for license information.
 #  SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 
+from ..ir import IntegerAttr, IntegerType, register_attribute_builder
 from ._bufferization_ops_gen import *
 from ._bufferization_enum_gen import *
+
+
+ at register_attribute_builder("builtin.LayoutMapOption")
+def _layoutmapoption(x, context):
+    return IntegerAttr.get(IntegerType.get_signless(32, context=context), int(x))
diff --git a/mlir/python/mlir/dialects/gpu/__init__.py b/mlir/python/mlir/dialects/gpu/__init__.py
index 4cd80aa8b7ca85..e0bb07c5dad8be 100644
--- a/mlir/python/mlir/dialects/gpu/__init__.py
+++ b/mlir/python/mlir/dialects/gpu/__init__.py
@@ -2,6 +2,62 @@
 #  See https://llvm.org/LICENSE.txt for license information.
 #  SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 
+from ...ir import IntegerAttr, IntegerType, register_attribute_builder
 from .._gpu_ops_gen import *
 from .._gpu_enum_gen import *
 from ..._mlir_libs._mlirDialectsGPU import *
+
+
+ at register_attribute_builder("builtin.GPU_AddressSpaceEnum")
+def _gpu_addressspaceenum(x, context):
+    return IntegerAttr.get(IntegerType.get_signless(32, context=context), int(x))
+
+
+ at register_attribute_builder("builtin.GPU_AllReduceOperation")
+def _gpu_allreduceoperation(x, context):
+    return IntegerAttr.get(IntegerType.get_signless(32, context=context), int(x))
+
+
+ at register_attribute_builder("builtin.GPU_CompilationTargetEnum")
+def _gpu_compilationtargetenum(x, context):
+    return IntegerAttr.get(IntegerType.get_signless(32, context=context), int(x))
+
+
+ at register_attribute_builder("builtin.GPU_Dimension")
+def _gpu_dimension(x, context):
+    return IntegerAttr.get(IntegerType.get_signless(32, context=context), int(x))
+
+
+ at register_attribute_builder("builtin.GPU_Prune2To4SpMatFlag")
+def _gpu_prune2to4spmatflag(x, context):
+    return IntegerAttr.get(IntegerType.get_signless(32, context=context), int(x))
+
+
+ at register_attribute_builder("builtin.GPU_ShuffleMode")
+def _gpu_shufflemode(x, context):
+    return IntegerAttr.get(IntegerType.get_signless(32, context=context), int(x))
+
+
+ at register_attribute_builder("builtin.GPU_SpGEMMWorkEstimationOrComputeKind")
+def _gpu_spgemmworkestimationorcomputekind(x, context):
+    return IntegerAttr.get(IntegerType.get_signless(32, context=context), int(x))
+
+
+ at register_attribute_builder("builtin.GPU_TransposeMode")
+def _gpu_transposemode(x, context):
+    return IntegerAttr.get(IntegerType.get_signless(32, context=context), int(x))
+
+
+ at register_attribute_builder("builtin.MMAElementWise")
+def _mmaelementwise(x, context):
+    return IntegerAttr.get(IntegerType.get_signless(32, context=context), int(x))
+
+
+ at register_attribute_builder("builtin.MappingIdEnum")
+def _mappingidenum(x, context):
+    return IntegerAttr.get(IntegerType.get_signless(64, context=context), int(x))
+
+
+ at register_attribute_builder("builtin.ProcessorEnum")
+def _processorenum(x, context):
+    return IntegerAttr.get(IntegerType.get_signless(64, context=context), int(x))
diff --git a/mlir/python/mlir/dialects/index.py b/mlir/python/mlir/dialects/index.py
index 73708c7d71a8c8..f00c397965c97c 100644
--- a/mlir/python/mlir/dialects/index.py
+++ b/mlir/python/mlir/dialects/index.py
@@ -2,5 +2,11 @@
 #  See https://llvm.org/LICENSE.txt for license information.
 #  SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 
+from ..ir import IntegerAttr, IntegerType, register_attribute_builder
 from ._index_ops_gen import *
 from ._index_enum_gen import *
+
+
+ at register_attribute_builder("builtin.IndexCmpPredicate")
+def _indexcmppredicate(x, context):
+    return IntegerAttr.get(IntegerType.get_signless(32, context=context), int(x))
diff --git a/mlir/python/mlir/dialects/linalg/__init__.py b/mlir/python/mlir/dialects/linalg/__init__.py
index 8fb1227ee80ff5..4fe9cc40ee910a 100644
--- a/mlir/python/mlir/dialects/linalg/__init__.py
+++ b/mlir/python/mlir/dialects/linalg/__init__.py
@@ -102,3 +102,28 @@ def broadcast(
     )
     fill_builtin_region(op.operation)
     return op
+
+
+ at register_attribute_builder("builtin.BinaryFn")
+def _binaryfn(x, context):
+    return IntegerAttr.get(IntegerType.get_signless(32, context=context), int(x))
+
+
+ at register_attribute_builder("builtin.IteratorType")
+def _iteratortype(x, context):
+    return IntegerAttr.get(IntegerType.get_signless(32, context=context), int(x))
+
+
+ at register_attribute_builder("builtin.TernaryFn")
+def _ternaryfn(x, context):
+    return IntegerAttr.get(IntegerType.get_signless(32, context=context), int(x))
+
+
+ at register_attribute_builder("builtin.TypeFn")
+def _typefn(x, context):
+    return IntegerAttr.get(IntegerType.get_signless(32, context=context), int(x))
+
+
+ at register_attribute_builder("builtin.UnaryFn")
+def _unaryfn(x, context):
+    return IntegerAttr.get(IntegerType.get_signless(32, context=context), int(x))
diff --git a/mlir/python/mlir/dialects/linalg/opdsl/ops/core_named_ops.py b/mlir/python/mlir/dialects/linalg/opdsl/ops/core_named_ops.py
index c95cd5eecfffca..f87b25e8416023 100644
--- a/mlir/python/mlir/dialects/linalg/opdsl/ops/core_named_ops.py
+++ b/mlir/python/mlir/dialects/linalg/opdsl/ops/core_named_ops.py
@@ -888,6 +888,7 @@ def conv_2d_nchw_fchw_q(
         - TypeFn.cast_signed(U, IZp)
     ) * (TypeFn.cast_signed(U, K[D.f, D.c, D.kh, D.kw]) - TypeFn.cast_signed(U, KZp))
 
+
 @linalg_structured_op
 def conv_2d_nchw_fchw(
     I=TensorDef(T1, S.N, S.C, S.OH * S.SH + S.KH * S.DH, S.OW * S.SW + S.KW * S.DW),
diff --git a/mlir/python/mlir/dialects/llvm.py b/mlir/python/mlir/dialects/llvm.py
index 941a584966dcde..456df39ae650ff 100644
--- a/mlir/python/mlir/dialects/llvm.py
+++ b/mlir/python/mlir/dialects/llvm.py
@@ -5,7 +5,7 @@
 from ._llvm_ops_gen import *
 from ._llvm_enum_gen import *
 from .._mlir_libs._mlirDialectsLLVM import *
-from ..ir import Value
+from ..ir import Value, IntegerAttr, IntegerType, register_attribute_builder
 from ._ods_common import get_op_result_or_op_results as _get_op_result_or_op_results
 
 
@@ -13,3 +13,108 @@ def mlir_constant(value, *, loc=None, ip=None) -> Value:
     return _get_op_result_or_op_results(
         ConstantOp(res=value.type, value=value, loc=loc, ip=ip)
     )
+
+
+ at register_attribute_builder("builtin.AsmATTOrIntel")
+def _asmattorintel(x, context):
+    return IntegerAttr.get(IntegerType.get_signless(64, context=context), int(x))
+
+
+ at register_attribute_builder("builtin.AtomicBinOp")
+def _atomicbinop(x, context):
+    return IntegerAttr.get(IntegerType.get_signless(64, context=context), int(x))
+
+
+ at register_attribute_builder("builtin.AtomicOrdering")
+def _atomicordering(x, context):
+    return IntegerAttr.get(IntegerType.get_signless(64, context=context), int(x))
+
+
+ at register_attribute_builder("builtin.CConvEnum")
+def _cconvenum(x, context):
+    return IntegerAttr.get(IntegerType.get_signless(64, context=context), int(x))
+
+
+ at register_attribute_builder("builtin.Comdat")
+def _comdat(x, context):
+    return IntegerAttr.get(IntegerType.get_signless(64, context=context), int(x))
+
+
+ at register_attribute_builder("builtin.DIFlags")
+def _diflags(x, context):
+    return IntegerAttr.get(IntegerType.get_signless(32, context=context), int(x))
+
+
+ at register_attribute_builder("builtin.DISubprogramFlags")
+def _disubprogramflags(x, context):
+    return IntegerAttr.get(IntegerType.get_signless(32, context=context), int(x))
+
+
+ at register_attribute_builder("builtin.FCmpPredicate")
+def _fcmppredicate(x, context):
+    return IntegerAttr.get(IntegerType.get_signless(64, context=context), int(x))
+
+
+ at register_attribute_builder("builtin.FPExceptionBehaviorAttr")
+def _fpexceptionbehaviorattr(x, context):
+    return IntegerAttr.get(IntegerType.get_signless(64, context=context), int(x))
+
+
+ at register_attribute_builder("builtin.FastmathFlags")
+def _fastmathflags(x, context):
+    return IntegerAttr.get(IntegerType.get_signless(32, context=context), int(x))
+
+
+ at register_attribute_builder("builtin.FramePointerKindEnum")
+def _framepointerkindenum(x, context):
+    return IntegerAttr.get(IntegerType.get_signless(64, context=context), int(x))
+
+
+ at register_attribute_builder("builtin.ICmpPredicate")
+def _icmppredicate(x, context):
+    return IntegerAttr.get(IntegerType.get_signless(64, context=context), int(x))
+
+
+ at register_attribute_builder("builtin.IntegerOverflowFlags")
+def _integeroverflowflags(x, context):
+    return IntegerAttr.get(IntegerType.get_signless(32, context=context), int(x))
+
+
+ at register_attribute_builder("builtin.LLVM_DIEmissionKind")
+def _llvm_diemissionkind(x, context):
+    return IntegerAttr.get(IntegerType.get_signless(64, context=context), int(x))
+
+
+ at register_attribute_builder("builtin.LLVM_DINameTableKind")
+def _llvm_dinametablekind(x, context):
+    return IntegerAttr.get(IntegerType.get_signless(64, context=context), int(x))
+
+
+ at register_attribute_builder("builtin.LinkageEnum")
+def _linkageenum(x, context):
+    return IntegerAttr.get(IntegerType.get_signless(64, context=context), int(x))
+
+
+ at register_attribute_builder("builtin.ModRefInfoEnum")
+def _modrefinfoenum(x, context):
+    return IntegerAttr.get(IntegerType.get_signless(64, context=context), int(x))
+
+
+ at register_attribute_builder("builtin.RoundingModeAttr")
+def _roundingmodeattr(x, context):
+    return IntegerAttr.get(IntegerType.get_signless(64, context=context), int(x))
+
+
+ at register_attribute_builder("builtin.TailCallKindEnum")
+def _tailcallkindenum(x, context):
+    return IntegerAttr.get(IntegerType.get_signless(64, context=context), int(x))
+
+
+ at register_attribute_builder("builtin.UnnamedAddr")
+def _unnamedaddr(x, context):
+    return IntegerAttr.get(IntegerType.get_signless(64, context=context), int(x))
+
+
+ at register_attribute_builder("builtin.Visibility")
+def _visibility(x, context):
+    return IntegerAttr.get(IntegerType.get_signless(64, context=context), int(x))
diff --git a/mlir/python/mlir/dialects/nvgpu.py b/mlir/python/mlir/dialects/nvgpu.py
index d6a54f2772f40d..eea132adb0484e 100644
--- a/mlir/python/mlir/dialects/nvgpu.py
+++ b/mlir/python/mlir/dialects/nvgpu.py
@@ -2,6 +2,32 @@
 #  See https://llvm.org/LICENSE.txt for license information.
 #  SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 
+from ..ir import IntegerAttr, IntegerType, register_attribute_builder
 from ._nvgpu_ops_gen import *
 from ._nvgpu_enum_gen import *
 from .._mlir_libs._mlirDialectsNVGPU import *
+
+
+ at register_attribute_builder("builtin.RcpRoundingMode")
+def _rcproundingmode(x, context):
+    return IntegerAttr.get(IntegerType.get_signless(32, context=context), int(x))
+
+
+ at register_attribute_builder("builtin.TensorMapInterleaveKind")
+def _tensormapinterleavekind(x, context):
+    return IntegerAttr.get(IntegerType.get_signless(32, context=context), int(x))
+
+
+ at register_attribute_builder("builtin.TensorMapL2PromoKind")
+def _tensormapl2promokind(x, context):
+    return IntegerAttr.get(IntegerType.get_signless(32, context=context), int(x))
+
+
+ at register_attribute_builder("builtin.TensorMapOOBKind")
+def _tensormapoobkind(x, context):
+    return IntegerAttr.get(IntegerType.get_signless(32, context=context), int(x))
+
+
+ at register_attribute_builder("builtin.TensorMapSwizzleKind")
+def _tensormapswizzlekind(x, context):
+    return IntegerAttr.get(IntegerType.get_signless(32, context=context), int(x))
diff --git a/mlir/python/mlir/dialects/nvvm.py b/mlir/python/mlir/dialects/nvvm.py
index 9477de39c9ead7..21bf24cb73fdab 100644
--- a/mlir/python/mlir/dialects/nvvm.py
+++ b/mlir/python/mlir/dialects/nvvm.py
@@ -2,5 +2,81 @@
 #  See https://llvm.org/LICENSE.txt for license information.
 #  SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 
+from ..ir import IntegerAttr, IntegerType, register_attribute_builder
 from ._nvvm_ops_gen import *
 from ._nvvm_enum_gen import *
+
+
+ at register_attribute_builder("builtin.LoadCacheModifierKind")
+def _loadcachemodifierkind(x, context):
+    return IntegerAttr.get(IntegerType.get_signless(32, context=context), int(x))
+
+
+ at register_attribute_builder("builtin.MMAB1Op")
+def _mmab1op(x, context):
+    return IntegerAttr.get(IntegerType.get_signless(32, context=context), int(x))
+
+
+ at register_attribute_builder("builtin.MMAFrag")
+def _mmafrag(x, context):
+    return IntegerAttr.get(IntegerType.get_signless(32, context=context), int(x))
+
+
+ at register_attribute_builder("builtin.MMAIntOverflow")
+def _mmaintoverflow(x, context):
+    return IntegerAttr.get(IntegerType.get_signless(32, context=context), int(x))
+
+
+ at register_attribute_builder("builtin.MMALayout")
+def _mmalayout(x, context):
+    return IntegerAttr.get(IntegerType.get_signless(32, context=context), int(x))
+
+
+ at register_attribute_builder("builtin.MMATypes")
+def _mmatypes(x, context):
+    return IntegerAttr.get(IntegerType.get_signless(32, context=context), int(x))
+
+
+ at register_attribute_builder("builtin.MemScopeKind")
+def _memscopekind(x, context):
+    return IntegerAttr.get(IntegerType.get_signless(32, context=context), int(x))
+
+
+ at register_attribute_builder("builtin.ProxyKind")
+def _proxykind(x, context):
+    return IntegerAttr.get(IntegerType.get_signless(32, context=context), int(x))
+
+
+ at register_attribute_builder("builtin.ReduxKind")
+def _reduxkind(x, context):
+    return IntegerAttr.get(IntegerType.get_signless(32, context=context), int(x))
+
+
+ at register_attribute_builder("builtin.SetMaxRegisterAction")
+def _setmaxregisteraction(x, context):
+    return IntegerAttr.get(IntegerType.get_signless(32, context=context), int(x))
+
+
+ at register_attribute_builder("builtin.SharedSpace")
+def _sharedspace(x, context):
+    return IntegerAttr.get(IntegerType.get_signless(32, context=context), int(x))
+
+
+ at register_attribute_builder("builtin.ShflKind")
+def _shflkind(x, context):
+    return IntegerAttr.get(IntegerType.get_signless(32, context=context), int(x))
+
+
+ at register_attribute_builder("builtin.WGMMAScaleIn")
+def _wgmmascalein(x, context):
+    return IntegerAttr.get(IntegerType.get_signless(32, context=context), int(x))
+
+
+ at register_attribute_builder("builtin.WGMMAScaleOut")
+def _wgmmascaleout(x, context):
+    return IntegerAttr.get(IntegerType.get_signless(32, context=context), int(x))
+
+
+ at register_attribute_builder("builtin.WGMMATypes")
+def _wgmmatypes(x, context):
+    return IntegerAttr.get(IntegerType.get_signless(32, context=context), int(x))
diff --git a/mlir/python/mlir/dialects/sparse_tensor.py b/mlir/python/mlir/dialects/sparse_tensor.py
index 209ecc95fa8fc8..8f1b83f9d514fd 100644
--- a/mlir/python/mlir/dialects/sparse_tensor.py
+++ b/mlir/python/mlir/dialects/sparse_tensor.py
@@ -2,7 +2,23 @@
 #  See https://llvm.org/LICENSE.txt for license information.
 #  SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 
+from ..ir import IntegerAttr, IntegerType, register_attribute_builder
 from ._sparse_tensor_ops_gen import *
 from ._sparse_tensor_enum_gen import *
 from .._mlir_libs._mlirDialectsSparseTensor import *
 from .._mlir_libs import _mlirSparseTensorPasses as _cextSparseTensorPasses
+
+
+ at register_attribute_builder("builtin.SparseTensorCrdTransDirectionEnum")
+def _sparsetensorcrdtransdirectionenum(x, context):
+    return IntegerAttr.get(IntegerType.get_signless(32, context=context), int(x))
+
+
+ at register_attribute_...
[truncated]

``````````

</details>


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

More information about the Mlir-commits mailing list