# Autogenerated by mlir-tblgen; don't manually edit.

from ._ods_common import _cext as _ods_cext
from ._ods_common import (
    equally_sized_accessor as _ods_equally_sized_accessor,
    get_default_loc_context as _ods_get_default_loc_context,
    get_op_result_or_op_results as _get_op_result_or_op_results,
    get_op_result_or_value as _get_op_result_or_value,
    get_op_results_or_values as _get_op_results_or_values,
    segmented_accessor as _ods_segmented_accessor,
)
_ods_ir = _ods_cext.ir

import builtins
from typing import Sequence as _Sequence, Union as _Union


@_ods_cext.register_dialect
class _Dialect(_ods_ir.Dialect):
  DIALECT_NAMESPACE = "nvvm"

@_ods_cext.register_operation(_Dialect)
class Barrier0Op(_ods_ir.OpView):
  OPERATION_NAME = "nvvm.barrier0"

  _ODS_REGIONS = (0, True)

  def __init__(self, *, loc=None, ip=None):
    operands = []
    results = []
    attributes = {}
    regions = None
    _ods_context = _ods_get_default_loc_context(loc)
    _ods_successors = None
    super().__init__(self.build_generic(attributes=attributes, results=results, operands=operands, successors=_ods_successors, regions=regions, loc=loc, ip=ip))

def barrier0(*, loc=None, ip=None) -> _ods_ir.Operation:
  return _get_op_result_or_op_results(Barrier0Op(loc=loc, ip=ip))

@_ods_cext.register_operation(_Dialect)
class BarrierArriveOp(_ods_ir.OpView):
  OPERATION_NAME = "nvvm.barrier.arrive"

  _ODS_REGIONS = (0, True)

  def __init__(self, numberOfThreads, *, barrierId=None, loc=None, ip=None):
    operands = []
    results = []
    attributes = {}
    regions = None
    if barrierId is not None: operands.append(_get_op_result_or_value(barrierId))
    operands.append(_get_op_result_or_value(numberOfThreads))
    _ods_context = _ods_get_default_loc_context(loc)
    _ods_successors = None
    super().__init__(self.build_generic(attributes=attributes, results=results, operands=operands, successors=_ods_successors, regions=regions, loc=loc, ip=ip))

  @builtins.property
  def barrierId(self):
    return None if len(self.operation.operands) < 2 else self.operation.operands[0]

  @builtins.property
  def numberOfThreads(self):
    _ods_variadic_group_length = len(self.operation.operands) - 2 + 1
    return self.operation.operands[1 + _ods_variadic_group_length - 1]

def barrier_arrive(number_of_threads, *, barrier_id=None, loc=None, ip=None) -> _ods_ir.Operation:
  return _get_op_result_or_op_results(BarrierArriveOp(numberOfThreads=number_of_threads, barrierId=barrier_id, loc=loc, ip=ip))

@_ods_cext.register_operation(_Dialect)
class BarrierOp(_ods_ir.OpView):
  OPERATION_NAME = "nvvm.barrier"

  _ODS_OPERAND_SEGMENTS = [0,0,]

  _ODS_REGIONS = (0, True)

  def __init__(self, *, barrierId=None, numberOfThreads=None, loc=None, ip=None):
    operands = []
    results = []
    attributes = {}
    regions = None
    operands.append(_get_op_result_or_value(barrierId) if barrierId is not None else None)
    operands.append(_get_op_result_or_value(numberOfThreads) if numberOfThreads is not None else None)
    _ods_context = _ods_get_default_loc_context(loc)
    _ods_successors = None
    super().__init__(self.build_generic(attributes=attributes, results=results, operands=operands, successors=_ods_successors, regions=regions, loc=loc, ip=ip))

  @builtins.property
  def barrierId(self):
    operand_range = _ods_segmented_accessor(
         self.operation.operands,
         self.operation.attributes["operandSegmentSizes"], 0)
    return operand_range[0] if len(operand_range) > 0 else None

  @builtins.property
  def numberOfThreads(self):
    operand_range = _ods_segmented_accessor(
         self.operation.operands,
         self.operation.attributes["operandSegmentSizes"], 1)
    return operand_range[0] if len(operand_range) > 0 else None

def barrier(*, barrier_id=None, number_of_threads=None, loc=None, ip=None) -> _ods_ir.Operation:
  return _get_op_result_or_op_results(BarrierOp(barrierId=barrier_id, numberOfThreads=number_of_threads, loc=loc, ip=ip))

@_ods_cext.register_operation(_Dialect)
class BlockDimXOp(_ods_ir.OpView):
  OPERATION_NAME = "nvvm.read.ptx.sreg.ntid.x"

  _ODS_REGIONS = (0, True)

  def __init__(self, res, *, loc=None, ip=None):
    operands = []
    results = []
    attributes = {}
    regions = None
    _ods_context = _ods_get_default_loc_context(loc)
    results.append(res)
    _ods_successors = None
    super().__init__(self.build_generic(attributes=attributes, results=results, operands=operands, successors=_ods_successors, regions=regions, loc=loc, ip=ip))

  @builtins.property
  def res(self):
    return self.operation.results[0]

def read_ptx_sreg_ntid_x(res, *, loc=None, ip=None) -> _ods_ir.Value:
  return _get_op_result_or_op_results(BlockDimXOp(res=res, loc=loc, ip=ip))

@_ods_cext.register_operation(_Dialect)
class BlockDimYOp(_ods_ir.OpView):
  OPERATION_NAME = "nvvm.read.ptx.sreg.ntid.y"

  _ODS_REGIONS = (0, True)

  def __init__(self, res, *, loc=None, ip=None):
    operands = []
    results = []
    attributes = {}
    regions = None
    _ods_context = _ods_get_default_loc_context(loc)
    results.append(res)
    _ods_successors = None
    super().__init__(self.build_generic(attributes=attributes, results=results, operands=operands, successors=_ods_successors, regions=regions, loc=loc, ip=ip))

  @builtins.property
  def res(self):
    return self.operation.results[0]

def read_ptx_sreg_ntid_y(res, *, loc=None, ip=None) -> _ods_ir.Value:
  return _get_op_result_or_op_results(BlockDimYOp(res=res, loc=loc, ip=ip))

@_ods_cext.register_operation(_Dialect)
class BlockDimZOp(_ods_ir.OpView):
  OPERATION_NAME = "nvvm.read.ptx.sreg.ntid.z"

  _ODS_REGIONS = (0, True)

  def __init__(self, res, *, loc=None, ip=None):
    operands = []
    results = []
    attributes = {}
    regions = None
    _ods_context = _ods_get_default_loc_context(loc)
    results.append(res)
    _ods_successors = None
    super().__init__(self.build_generic(attributes=attributes, results=results, operands=operands, successors=_ods_successors, regions=regions, loc=loc, ip=ip))

  @builtins.property
  def res(self):
    return self.operation.results[0]

def read_ptx_sreg_ntid_z(res, *, loc=None, ip=None) -> _ods_ir.Value:
  return _get_op_result_or_op_results(BlockDimZOp(res=res, loc=loc, ip=ip))

@_ods_cext.register_operation(_Dialect)
class BlockIdXOp(_ods_ir.OpView):
  OPERATION_NAME = "nvvm.read.ptx.sreg.ctaid.x"

  _ODS_REGIONS = (0, True)

  def __init__(self, res, *, loc=None, ip=None):
    operands = []
    results = []
    attributes = {}
    regions = None
    _ods_context = _ods_get_default_loc_context(loc)
    results.append(res)
    _ods_successors = None
    super().__init__(self.build_generic(attributes=attributes, results=results, operands=operands, successors=_ods_successors, regions=regions, loc=loc, ip=ip))

  @builtins.property
  def res(self):
    return self.operation.results[0]

def read_ptx_sreg_ctaid_x(res, *, loc=None, ip=None) -> _ods_ir.Value:
  return _get_op_result_or_op_results(BlockIdXOp(res=res, loc=loc, ip=ip))

@_ods_cext.register_operation(_Dialect)
class BlockIdYOp(_ods_ir.OpView):
  OPERATION_NAME = "nvvm.read.ptx.sreg.ctaid.y"

  _ODS_REGIONS = (0, True)

  def __init__(self, res, *, loc=None, ip=None):
    operands = []
    results = []
    attributes = {}
    regions = None
    _ods_context = _ods_get_default_loc_context(loc)
    results.append(res)
    _ods_successors = None
    super().__init__(self.build_generic(attributes=attributes, results=results, operands=operands, successors=_ods_successors, regions=regions, loc=loc, ip=ip))

  @builtins.property
  def res(self):
    return self.operation.results[0]

def read_ptx_sreg_ctaid_y(res, *, loc=None, ip=None) -> _ods_ir.Value:
  return _get_op_result_or_op_results(BlockIdYOp(res=res, loc=loc, ip=ip))

@_ods_cext.register_operation(_Dialect)
class BlockIdZOp(_ods_ir.OpView):
  OPERATION_NAME = "nvvm.read.ptx.sreg.ctaid.z"

  _ODS_REGIONS = (0, True)

  def __init__(self, res, *, loc=None, ip=None):
    operands = []
    results = []
    attributes = {}
    regions = None
    _ods_context = _ods_get_default_loc_context(loc)
    results.append(res)
    _ods_successors = None
    super().__init__(self.build_generic(attributes=attributes, results=results, operands=operands, successors=_ods_successors, regions=regions, loc=loc, ip=ip))

  @builtins.property
  def res(self):
    return self.operation.results[0]

def read_ptx_sreg_ctaid_z(res, *, loc=None, ip=None) -> _ods_ir.Value:
  return _get_op_result_or_op_results(BlockIdZOp(res=res, loc=loc, ip=ip))

@_ods_cext.register_operation(_Dialect)
class BlockInClusterIdXOp(_ods_ir.OpView):
  OPERATION_NAME = "nvvm.read.ptx.sreg.cluster.ctaid.x"

  _ODS_REGIONS = (0, True)

  def __init__(self, res, *, loc=None, ip=None):
    operands = []
    results = []
    attributes = {}
    regions = None
    _ods_context = _ods_get_default_loc_context(loc)
    results.append(res)
    _ods_successors = None
    super().__init__(self.build_generic(attributes=attributes, results=results, operands=operands, successors=_ods_successors, regions=regions, loc=loc, ip=ip))

  @builtins.property
  def res(self):
    return self.operation.results[0]

def read_ptx_sreg_cluster_ctaid_x(res, *, loc=None, ip=None) -> _ods_ir.Value:
  return _get_op_result_or_op_results(BlockInClusterIdXOp(res=res, loc=loc, ip=ip))

@_ods_cext.register_operation(_Dialect)
class BlockInClusterIdYOp(_ods_ir.OpView):
  OPERATION_NAME = "nvvm.read.ptx.sreg.cluster.ctaid.y"

  _ODS_REGIONS = (0, True)

  def __init__(self, res, *, loc=None, ip=None):
    operands = []
    results = []
    attributes = {}
    regions = None
    _ods_context = _ods_get_default_loc_context(loc)
    results.append(res)
    _ods_successors = None
    super().__init__(self.build_generic(attributes=attributes, results=results, operands=operands, successors=_ods_successors, regions=regions, loc=loc, ip=ip))

  @builtins.property
  def res(self):
    return self.operation.results[0]

def read_ptx_sreg_cluster_ctaid_y(res, *, loc=None, ip=None) -> _ods_ir.Value:
  return _get_op_result_or_op_results(BlockInClusterIdYOp(res=res, loc=loc, ip=ip))

@_ods_cext.register_operation(_Dialect)
class BlockInClusterIdZOp(_ods_ir.OpView):
  OPERATION_NAME = "nvvm.read.ptx.sreg.cluster.ctaid.z"

  _ODS_REGIONS = (0, True)

  def __init__(self, res, *, loc=None, ip=None):
    operands = []
    results = []
    attributes = {}
    regions = None
    _ods_context = _ods_get_default_loc_context(loc)
    results.append(res)
    _ods_successors = None
    super().__init__(self.build_generic(attributes=attributes, results=results, operands=operands, successors=_ods_successors, regions=regions, loc=loc, ip=ip))

  @builtins.property
  def res(self):
    return self.operation.results[0]

def read_ptx_sreg_cluster_ctaid_z(res, *, loc=None, ip=None) -> _ods_ir.Value:
  return _get_op_result_or_op_results(BlockInClusterIdZOp(res=res, loc=loc, ip=ip))

@_ods_cext.register_operation(_Dialect)
class Clock64Op(_ods_ir.OpView):
  OPERATION_NAME = "nvvm.read.ptx.sreg.clock64"

  _ODS_REGIONS = (0, True)

  def __init__(self, res, *, loc=None, ip=None):
    operands = []
    results = []
    attributes = {}
    regions = None
    _ods_context = _ods_get_default_loc_context(loc)
    results.append(res)
    _ods_successors = None
    super().__init__(self.build_generic(attributes=attributes, results=results, operands=operands, successors=_ods_successors, regions=regions, loc=loc, ip=ip))

  @builtins.property
  def res(self):
    return self.operation.results[0]

def read_ptx_sreg_clock64(res, *, loc=None, ip=None) -> _ods_ir.Value:
  return _get_op_result_or_op_results(Clock64Op(res=res, loc=loc, ip=ip))

@_ods_cext.register_operation(_Dialect)
class ClockOp(_ods_ir.OpView):
  OPERATION_NAME = "nvvm.read.ptx.sreg.clock"

  _ODS_REGIONS = (0, True)

  def __init__(self, res, *, loc=None, ip=None):
    operands = []
    results = []
    attributes = {}
    regions = None
    _ods_context = _ods_get_default_loc_context(loc)
    results.append(res)
    _ods_successors = None
    super().__init__(self.build_generic(attributes=attributes, results=results, operands=operands, successors=_ods_successors, regions=regions, loc=loc, ip=ip))

  @builtins.property
  def res(self):
    return self.operation.results[0]

def read_ptx_sreg_clock(res, *, loc=None, ip=None) -> _ods_ir.Value:
  return _get_op_result_or_op_results(ClockOp(res=res, loc=loc, ip=ip))

@_ods_cext.register_operation(_Dialect)
class ClusterArriveOp(_ods_ir.OpView):
  OPERATION_NAME = "nvvm.cluster.arrive"

  _ODS_REGIONS = (0, True)

  def __init__(self, *, aligned=None, loc=None, ip=None):
    operands = []
    results = []
    attributes = {}
    regions = None
    _ods_context = _ods_get_default_loc_context(loc)
    if bool(aligned): attributes["aligned"] = _ods_ir.UnitAttr.get(
      _ods_get_default_loc_context(loc))
    _ods_successors = None
    super().__init__(self.build_generic(attributes=attributes, results=results, operands=operands, successors=_ods_successors, regions=regions, loc=loc, ip=ip))

  @builtins.property
  def aligned(self):
    return "aligned" in self.operation.attributes

  @aligned.setter
  def aligned(self, value):
    if bool(value):
      self.operation.attributes["aligned"] = _ods_ir.UnitAttr.get()
    elif "aligned" in self.operation.attributes:
      del self.operation.attributes["aligned"]

  @aligned.deleter
  def aligned(self):
    del self.operation.attributes["aligned"]

def cluster_arrive(*, aligned=None, loc=None, ip=None) -> _ods_ir.Operation:
  return _get_op_result_or_op_results(ClusterArriveOp(aligned=aligned, loc=loc, ip=ip))

@_ods_cext.register_operation(_Dialect)
class ClusterArriveRelaxedOp(_ods_ir.OpView):
  OPERATION_NAME = "nvvm.cluster.arrive.relaxed"

  _ODS_REGIONS = (0, True)

  def __init__(self, *, aligned=None, loc=None, ip=None):
    operands = []
    results = []
    attributes = {}
    regions = None
    _ods_context = _ods_get_default_loc_context(loc)
    if bool(aligned): attributes["aligned"] = _ods_ir.UnitAttr.get(
      _ods_get_default_loc_context(loc))
    _ods_successors = None
    super().__init__(self.build_generic(attributes=attributes, results=results, operands=operands, successors=_ods_successors, regions=regions, loc=loc, ip=ip))

  @builtins.property
  def aligned(self):
    return "aligned" in self.operation.attributes

  @aligned.setter
  def aligned(self, value):
    if bool(value):
      self.operation.attributes["aligned"] = _ods_ir.UnitAttr.get()
    elif "aligned" in self.operation.attributes:
      del self.operation.attributes["aligned"]

  @aligned.deleter
  def aligned(self):
    del self.operation.attributes["aligned"]

def cluster_arrive_relaxed(*, aligned=None, loc=None, ip=None) -> _ods_ir.Operation:
  return _get_op_result_or_op_results(ClusterArriveRelaxedOp(aligned=aligned, loc=loc, ip=ip))

@_ods_cext.register_operation(_Dialect)
class ClusterDim(_ods_ir.OpView):
  OPERATION_NAME = "nvvm.read.ptx.sreg.cluster.nctarank"

  _ODS_REGIONS = (0, True)

  def __init__(self, res, *, loc=None, ip=None):
    operands = []
    results = []
    attributes = {}
    regions = None
    _ods_context = _ods_get_default_loc_context(loc)
    results.append(res)
    _ods_successors = None
    super().__init__(self.build_generic(attributes=attributes, results=results, operands=operands, successors=_ods_successors, regions=regions, loc=loc, ip=ip))

  @builtins.property
  def res(self):
    return self.operation.results[0]

def read_ptx_sreg_cluster_nctarank(res, *, loc=None, ip=None) -> _ods_ir.Value:
  return _get_op_result_or_op_results(ClusterDim(res=res, loc=loc, ip=ip))

@_ods_cext.register_operation(_Dialect)
class ClusterDimBlocksXOp(_ods_ir.OpView):
  OPERATION_NAME = "nvvm.read.ptx.sreg.cluster.nctaid.x"

  _ODS_REGIONS = (0, True)

  def __init__(self, res, *, loc=None, ip=None):
    operands = []
    results = []
    attributes = {}
    regions = None
    _ods_context = _ods_get_default_loc_context(loc)
    results.append(res)
    _ods_successors = None
    super().__init__(self.build_generic(attributes=attributes, results=results, operands=operands, successors=_ods_successors, regions=regions, loc=loc, ip=ip))

  @builtins.property
  def res(self):
    return self.operation.results[0]

def read_ptx_sreg_cluster_nctaid_x(res, *, loc=None, ip=None) -> _ods_ir.Value:
  return _get_op_result_or_op_results(ClusterDimBlocksXOp(res=res, loc=loc, ip=ip))

@_ods_cext.register_operation(_Dialect)
class ClusterDimBlocksYOp(_ods_ir.OpView):
  OPERATION_NAME = "nvvm.read.ptx.sreg.cluster.nctaid.y"

  _ODS_REGIONS = (0, True)

  def __init__(self, res, *, loc=None, ip=None):
    operands = []
    results = []
    attributes = {}
    regions = None
    _ods_context = _ods_get_default_loc_context(loc)
    results.append(res)
    _ods_successors = None
    super().__init__(self.build_generic(attributes=attributes, results=results, operands=operands, successors=_ods_successors, regions=regions, loc=loc, ip=ip))

  @builtins.property
  def res(self):
    return self.operation.results[0]

def read_ptx_sreg_cluster_nctaid_y(res, *, loc=None, ip=None) -> _ods_ir.Value:
  return _get_op_result_or_op_results(ClusterDimBlocksYOp(res=res, loc=loc, ip=ip))

@_ods_cext.register_operation(_Dialect)
class ClusterDimBlocksZOp(_ods_ir.OpView):
  OPERATION_NAME = "nvvm.read.ptx.sreg.cluster.nctaid.z"

  _ODS_REGIONS = (0, True)

  def __init__(self, res, *, loc=None, ip=None):
    operands = []
    results = []
    attributes = {}
    regions = None
    _ods_context = _ods_get_default_loc_context(loc)
    results.append(res)
    _ods_successors = None
    super().__init__(self.build_generic(attributes=attributes, results=results, operands=operands, successors=_ods_successors, regions=regions, loc=loc, ip=ip))

  @builtins.property
  def res(self):
    return self.operation.results[0]

def read_ptx_sreg_cluster_nctaid_z(res, *, loc=None, ip=None) -> _ods_ir.Value:
  return _get_op_result_or_op_results(ClusterDimBlocksZOp(res=res, loc=loc, ip=ip))

@_ods_cext.register_operation(_Dialect)
class ClusterDimXOp(_ods_ir.OpView):
  OPERATION_NAME = "nvvm.read.ptx.sreg.nclusterid.x"

  _ODS_REGIONS = (0, True)

  def __init__(self, res, *, loc=None, ip=None):
    operands = []
    results = []
    attributes = {}
    regions = None
    _ods_context = _ods_get_default_loc_context(loc)
    results.append(res)
    _ods_successors = None
    super().__init__(self.build_generic(attributes=attributes, results=results, operands=operands, successors=_ods_successors, regions=regions, loc=loc, ip=ip))

  @builtins.property
  def res(self):
    return self.operation.results[0]

def read_ptx_sreg_nclusterid_x(res, *, loc=None, ip=None) -> _ods_ir.Value:
  return _get_op_result_or_op_results(ClusterDimXOp(res=res, loc=loc, ip=ip))

@_ods_cext.register_operation(_Dialect)
class ClusterDimYOp(_ods_ir.OpView):
  OPERATION_NAME = "nvvm.read.ptx.sreg.nclusterid.y"

  _ODS_REGIONS = (0, True)

  def __init__(self, res, *, loc=None, ip=None):
    operands = []
    results = []
    attributes = {}
    regions = None
    _ods_context = _ods_get_default_loc_context(loc)
    results.append(res)
    _ods_successors = None
    super().__init__(self.build_generic(attributes=attributes, results=results, operands=operands, successors=_ods_successors, regions=regions, loc=loc, ip=ip))

  @builtins.property
  def res(self):
    return self.operation.results[0]

def read_ptx_sreg_nclusterid_y(res, *, loc=None, ip=None) -> _ods_ir.Value:
  return _get_op_result_or_op_results(ClusterDimYOp(res=res, loc=loc, ip=ip))

@_ods_cext.register_operation(_Dialect)
class ClusterDimZOp(_ods_ir.OpView):
  OPERATION_NAME = "nvvm.read.ptx.sreg.nclusterid.z"

  _ODS_REGIONS = (0, True)

  def __init__(self, res, *, loc=None, ip=None):
    operands = []
    results = []
    attributes = {}
    regions = None
    _ods_context = _ods_get_default_loc_context(loc)
    results.append(res)
    _ods_successors = None
    super().__init__(self.build_generic(attributes=attributes, results=results, operands=operands, successors=_ods_successors, regions=regions, loc=loc, ip=ip))

  @builtins.property
  def res(self):
    return self.operation.results[0]

def read_ptx_sreg_nclusterid_z(res, *, loc=None, ip=None) -> _ods_ir.Value:
  return _get_op_result_or_op_results(ClusterDimZOp(res=res, loc=loc, ip=ip))

@_ods_cext.register_operation(_Dialect)
class ClusterId(_ods_ir.OpView):
  OPERATION_NAME = "nvvm.read.ptx.sreg.cluster.ctarank"

  _ODS_REGIONS = (0, True)

  def __init__(self, res, *, loc=None, ip=None):
    operands = []
    results = []
    attributes = {}
    regions = None
    _ods_context = _ods_get_default_loc_context(loc)
    results.append(res)
    _ods_successors = None
    super().__init__(self.build_generic(attributes=attributes, results=results, operands=operands, successors=_ods_successors, regions=regions, loc=loc, ip=ip))

  @builtins.property
  def res(self):
    return self.operation.results[0]

def read_ptx_sreg_cluster_ctarank(res, *, loc=None, ip=None) -> _ods_ir.Value:
  return _get_op_result_or_op_results(ClusterId(res=res, loc=loc, ip=ip))

@_ods_cext.register_operation(_Dialect)
class ClusterIdXOp(_ods_ir.OpView):
  OPERATION_NAME = "nvvm.read.ptx.sreg.clusterid.x"

  _ODS_REGIONS = (0, True)

  def __init__(self, res, *, loc=None, ip=None):
    operands = []
    results = []
    attributes = {}
    regions = None
    _ods_context = _ods_get_default_loc_context(loc)
    results.append(res)
    _ods_successors = None
    super().__init__(self.build_generic(attributes=attributes, results=results, operands=operands, successors=_ods_successors, regions=regions, loc=loc, ip=ip))

  @builtins.property
  def res(self):
    return self.operation.results[0]

def read_ptx_sreg_clusterid_x(res, *, loc=None, ip=None) -> _ods_ir.Value:
  return _get_op_result_or_op_results(ClusterIdXOp(res=res, loc=loc, ip=ip))

@_ods_cext.register_operation(_Dialect)
class ClusterIdYOp(_ods_ir.OpView):
  OPERATION_NAME = "nvvm.read.ptx.sreg.clusterid.y"

  _ODS_REGIONS = (0, True)

  def __init__(self, res, *, loc=None, ip=None):
    operands = []
    results = []
    attributes = {}
    regions = None
    _ods_context = _ods_get_default_loc_context(loc)
    results.append(res)
    _ods_successors = None
    super().__init__(self.build_generic(attributes=attributes, results=results, operands=operands, successors=_ods_successors, regions=regions, loc=loc, ip=ip))

  @builtins.property
  def res(self):
    return self.operation.results[0]

def read_ptx_sreg_clusterid_y(res, *, loc=None, ip=None) -> _ods_ir.Value:
  return _get_op_result_or_op_results(ClusterIdYOp(res=res, loc=loc, ip=ip))

@_ods_cext.register_operation(_Dialect)
class ClusterIdZOp(_ods_ir.OpView):
  OPERATION_NAME = "nvvm.read.ptx.sreg.clusterid.z"

  _ODS_REGIONS = (0, True)

  def __init__(self, res, *, loc=None, ip=None):
    operands = []
    results = []
    attributes = {}
    regions = None
    _ods_context = _ods_get_default_loc_context(loc)
    results.append(res)
    _ods_successors = None
    super().__init__(self.build_generic(attributes=attributes, results=results, operands=operands, successors=_ods_successors, regions=regions, loc=loc, ip=ip))

  @builtins.property
  def res(self):
    return self.operation.results[0]

def read_ptx_sreg_clusterid_z(res, *, loc=None, ip=None) -> _ods_ir.Value:
  return _get_op_result_or_op_results(ClusterIdZOp(res=res, loc=loc, ip=ip))

@_ods_cext.register_operation(_Dialect)
class ClusterWaitOp(_ods_ir.OpView):
  OPERATION_NAME = "nvvm.cluster.wait"

  _ODS_REGIONS = (0, True)

  def __init__(self, *, aligned=None, loc=None, ip=None):
    operands = []
    results = []
    attributes = {}
    regions = None
    _ods_context = _ods_get_default_loc_context(loc)
    if bool(aligned): attributes["aligned"] = _ods_ir.UnitAttr.get(
      _ods_get_default_loc_context(loc))
    _ods_successors = None
    super().__init__(self.build_generic(attributes=attributes, results=results, operands=operands, successors=_ods_successors, regions=regions, loc=loc, ip=ip))

  @builtins.property
  def aligned(self):
    return "aligned" in self.operation.attributes

  @aligned.setter
  def aligned(self, value):
    if bool(value):
      self.operation.attributes["aligned"] = _ods_ir.UnitAttr.get()
    elif "aligned" in self.operation.attributes:
      del self.operation.attributes["aligned"]

  @aligned.deleter
  def aligned(self):
    del self.operation.attributes["aligned"]

def cluster_wait(*, aligned=None, loc=None, ip=None) -> _ods_ir.Operation:
  return _get_op_result_or_op_results(ClusterWaitOp(aligned=aligned, loc=loc, ip=ip))

@_ods_cext.register_operation(_Dialect)
class CpAsyncBulkCommitGroupOp(_ods_ir.OpView):
  OPERATION_NAME = "nvvm.cp.async.bulk.commit.group"

  _ODS_REGIONS = (0, True)

  def __init__(self, *, loc=None, ip=None):
    operands = []
    results = []
    attributes = {}
    regions = None
    _ods_context = _ods_get_default_loc_context(loc)
    _ods_successors = None
    super().__init__(self.build_generic(attributes=attributes, results=results, operands=operands, successors=_ods_successors, regions=regions, loc=loc, ip=ip))

def cp_async_bulk_commit_group(*, loc=None, ip=None) -> _ods_ir.Operation:
  return _get_op_result_or_op_results(CpAsyncBulkCommitGroupOp(loc=loc, ip=ip))

@_ods_cext.register_operation(_Dialect)
class CpAsyncBulkTensorGlobalToSharedClusterOp(_ods_ir.OpView):
  OPERATION_NAME = "nvvm.cp.async.bulk.tensor.shared.cluster.global"

  _ODS_OPERAND_SEGMENTS = [1,1,-1,1,-1,0,0,0,]

  _ODS_REGIONS = (0, True)

  def __init__(self, dstMem, tmaDescriptor, coordinates, mbar, im2colOffsets, *, multicastMask=None, l2CacheHint=None, predicate=None, loc=None, ip=None):
    operands = []
    results = []
    attributes = {}
    regions = None
    operands.append(_get_op_result_or_value(dstMem))
    operands.append(_get_op_result_or_value(tmaDescriptor))
    operands.append(_get_op_results_or_values(coordinates))
    operands.append(_get_op_result_or_value(mbar))
    operands.append(_get_op_results_or_values(im2colOffsets))
    operands.append(_get_op_result_or_value(multicastMask) if multicastMask is not None else None)
    operands.append(_get_op_result_or_value(l2CacheHint) if l2CacheHint is not None else None)
    operands.append(_get_op_result_or_value(predicate) if predicate is not None else None)
    _ods_context = _ods_get_default_loc_context(loc)
    _ods_successors = None
    super().__init__(self.build_generic(attributes=attributes, results=results, operands=operands, successors=_ods_successors, regions=regions, loc=loc, ip=ip))

  @builtins.property
  def dstMem(self):
    operand_range = _ods_segmented_accessor(
         self.operation.operands,
         self.operation.attributes["operandSegmentSizes"], 0)
    return operand_range[0]

  @builtins.property
  def tmaDescriptor(self):
    operand_range = _ods_segmented_accessor(
         self.operation.operands,
         self.operation.attributes["operandSegmentSizes"], 1)
    return operand_range[0]

  @builtins.property
  def coordinates(self):
    operand_range = _ods_segmented_accessor(
         self.operation.operands,
         self.operation.attributes["operandSegmentSizes"], 2)
    return operand_range

  @builtins.property
  def mbar(self):
    operand_range = _ods_segmented_accessor(
         self.operation.operands,
         self.operation.attributes["operandSegmentSizes"], 3)
    return operand_range[0]

  @builtins.property
  def im2colOffsets(self):
    operand_range = _ods_segmented_accessor(
         self.operation.operands,
         self.operation.attributes["operandSegmentSizes"], 4)
    return operand_range

  @builtins.property
  def multicastMask(self):
    operand_range = _ods_segmented_accessor(
         self.operation.operands,
         self.operation.attributes["operandSegmentSizes"], 5)
    return operand_range[0] if len(operand_range) > 0 else None

  @builtins.property
  def l2CacheHint(self):
    operand_range = _ods_segmented_accessor(
         self.operation.operands,
         self.operation.attributes["operandSegmentSizes"], 6)
    return operand_range[0] if len(operand_range) > 0 else None

  @builtins.property
  def predicate(self):
    operand_range = _ods_segmented_accessor(
         self.operation.operands,
         self.operation.attributes["operandSegmentSizes"], 7)
    return operand_range[0] if len(operand_range) > 0 else None

def cp_async_bulk_tensor_shared_cluster_global(dst_mem, tma_descriptor, coordinates, mbar, im2col_offsets, *, multicast_mask=None, l2_cache_hint=None, predicate=None, loc=None, ip=None) -> _ods_ir.Operation:
  return _get_op_result_or_op_results(CpAsyncBulkTensorGlobalToSharedClusterOp(dstMem=dst_mem, tmaDescriptor=tma_descriptor, coordinates=coordinates, mbar=mbar, im2colOffsets=im2col_offsets, multicastMask=multicast_mask, l2CacheHint=l2_cache_hint, predicate=predicate, loc=loc, ip=ip))

@_ods_cext.register_operation(_Dialect)
class CpAsyncBulkTensorSharedCTAToGlobalOp(_ods_ir.OpView):
  OPERATION_NAME = "nvvm.cp.async.bulk.tensor.global.shared.cta"

  _ODS_OPERAND_SEGMENTS = [1,1,-1,0,]

  _ODS_REGIONS = (0, True)

  def __init__(self, tmaDescriptor, srcMem, coordinates, *, predicate=None, loc=None, ip=None):
    operands = []
    results = []
    attributes = {}
    regions = None
    operands.append(_get_op_result_or_value(tmaDescriptor))
    operands.append(_get_op_result_or_value(srcMem))
    operands.append(_get_op_results_or_values(coordinates))
    operands.append(_get_op_result_or_value(predicate) if predicate is not None else None)
    _ods_context = _ods_get_default_loc_context(loc)
    _ods_successors = None
    super().__init__(self.build_generic(attributes=attributes, results=results, operands=operands, successors=_ods_successors, regions=regions, loc=loc, ip=ip))

  @builtins.property
  def tmaDescriptor(self):
    operand_range = _ods_segmented_accessor(
         self.operation.operands,
         self.operation.attributes["operandSegmentSizes"], 0)
    return operand_range[0]

  @builtins.property
  def srcMem(self):
    operand_range = _ods_segmented_accessor(
         self.operation.operands,
         self.operation.attributes["operandSegmentSizes"], 1)
    return operand_range[0]

  @builtins.property
  def coordinates(self):
    operand_range = _ods_segmented_accessor(
         self.operation.operands,
         self.operation.attributes["operandSegmentSizes"], 2)
    return operand_range

  @builtins.property
  def predicate(self):
    operand_range = _ods_segmented_accessor(
         self.operation.operands,
         self.operation.attributes["operandSegmentSizes"], 3)
    return operand_range[0] if len(operand_range) > 0 else None

def cp_async_bulk_tensor_global_shared_cta(tma_descriptor, src_mem, coordinates, *, predicate=None, loc=None, ip=None) -> _ods_ir.Operation:
  return _get_op_result_or_op_results(CpAsyncBulkTensorSharedCTAToGlobalOp(tmaDescriptor=tma_descriptor, srcMem=src_mem, coordinates=coordinates, predicate=predicate, loc=loc, ip=ip))

@_ods_cext.register_operation(_Dialect)
class CpAsyncBulkWaitGroupOp(_ods_ir.OpView):
  OPERATION_NAME = "nvvm.cp.async.bulk.wait_group"

  _ODS_REGIONS = (0, True)

  def __init__(self, group, *, read=None, loc=None, ip=None):
    operands = []
    results = []
    attributes = {}
    regions = None
    _ods_context = _ods_get_default_loc_context(loc)
    attributes["group"] = (group if (
    isinstance(group, _ods_ir.Attribute) or
    not _ods_ir.AttrBuilder.contains('I32Attr')) else
      _ods_ir.AttrBuilder.get('I32Attr')(group, context=_ods_context))
    if bool(read): attributes["read"] = _ods_ir.UnitAttr.get(
      _ods_get_default_loc_context(loc))
    _ods_successors = None
    super().__init__(self.build_generic(attributes=attributes, results=results, operands=operands, successors=_ods_successors, regions=regions, loc=loc, ip=ip))

  @builtins.property
  def group(self):
    return self.operation.attributes["group"]

  @group.setter
  def group(self, value):
    if value is None:
      raise ValueError("'None' not allowed as value for mandatory attributes")
    self.operation.attributes["group"] = value

  @builtins.property
  def read(self):
    return "read" in self.operation.attributes

  @read.setter
  def read(self, value):
    if bool(value):
      self.operation.attributes["read"] = _ods_ir.UnitAttr.get()
    elif "read" in self.operation.attributes:
      del self.operation.attributes["read"]

  @read.deleter
  def read(self):
    del self.operation.attributes["read"]

def cp_async_bulk_wait_group(group, *, read=None, loc=None, ip=None) -> _ods_ir.Operation:
  return _get_op_result_or_op_results(CpAsyncBulkWaitGroupOp(group=group, read=read, loc=loc, ip=ip))

@_ods_cext.register_operation(_Dialect)
class CpAsyncCommitGroupOp(_ods_ir.OpView):
  OPERATION_NAME = "nvvm.cp.async.commit.group"

  _ODS_REGIONS = (0, True)

  def __init__(self, *, loc=None, ip=None):
    operands = []
    results = []
    attributes = {}
    regions = None
    _ods_context = _ods_get_default_loc_context(loc)
    _ods_successors = None
    super().__init__(self.build_generic(attributes=attributes, results=results, operands=operands, successors=_ods_successors, regions=regions, loc=loc, ip=ip))

def cp_async_commit_group(*, loc=None, ip=None) -> _ods_ir.Operation:
  return _get_op_result_or_op_results(CpAsyncCommitGroupOp(loc=loc, ip=ip))

@_ods_cext.register_operation(_Dialect)
class CpAsyncMBarrierArriveOp(_ods_ir.OpView):
  OPERATION_NAME = "nvvm.cp.async.mbarrier.arrive"

  _ODS_REGIONS = (0, True)

  def __init__(self, addr, *, noinc=None, loc=None, ip=None):
    operands = []
    results = []
    attributes = {}
    regions = None
    operands.append(_get_op_result_or_value(addr))
    _ods_context = _ods_get_default_loc_context(loc)
    if noinc is not None: attributes["noinc"] = (noinc if (
        isinstance(noinc, _ods_ir.Attribute) or
        not _ods_ir.AttrBuilder.contains('I1Attr')) else
          _ods_ir.AttrBuilder.get('I1Attr')(noinc, context=_ods_context))
    _ods_successors = None
    super().__init__(self.build_generic(attributes=attributes, results=results, operands=operands, successors=_ods_successors, regions=regions, loc=loc, ip=ip))

  @builtins.property
  def addr(self):
    return self.operation.operands[0]

  @builtins.property
  def noinc(self):
    return self.operation.attributes["noinc"]

  @noinc.setter
  def noinc(self, value):
    if value is None:
      raise ValueError("'None' not allowed as value for mandatory attributes")
    self.operation.attributes["noinc"] = value

def cp_async_mbarrier_arrive(addr, *, noinc=None, loc=None, ip=None) -> _ods_ir.Operation:
  return _get_op_result_or_op_results(CpAsyncMBarrierArriveOp(addr=addr, noinc=noinc, loc=loc, ip=ip))

@_ods_cext.register_operation(_Dialect)
class CpAsyncMBarrierArriveSharedOp(_ods_ir.OpView):
  OPERATION_NAME = "nvvm.cp.async.mbarrier.arrive.shared"

  _ODS_REGIONS = (0, True)

  def __init__(self, addr, *, noinc=None, loc=None, ip=None):
    operands = []
    results = []
    attributes = {}
    regions = None
    operands.append(_get_op_result_or_value(addr))
    _ods_context = _ods_get_default_loc_context(loc)
    if noinc is not None: attributes["noinc"] = (noinc if (
        isinstance(noinc, _ods_ir.Attribute) or
        not _ods_ir.AttrBuilder.contains('I1Attr')) else
          _ods_ir.AttrBuilder.get('I1Attr')(noinc, context=_ods_context))
    _ods_successors = None
    super().__init__(self.build_generic(attributes=attributes, results=results, operands=operands, successors=_ods_successors, regions=regions, loc=loc, ip=ip))

  @builtins.property
  def addr(self):
    return self.operation.operands[0]

  @builtins.property
  def noinc(self):
    return self.operation.attributes["noinc"]

  @noinc.setter
  def noinc(self, value):
    if value is None:
      raise ValueError("'None' not allowed as value for mandatory attributes")
    self.operation.attributes["noinc"] = value

def cp_async_mbarrier_arrive_shared(addr, *, noinc=None, loc=None, ip=None) -> _ods_ir.Operation:
  return _get_op_result_or_op_results(CpAsyncMBarrierArriveSharedOp(addr=addr, noinc=noinc, loc=loc, ip=ip))

@_ods_cext.register_operation(_Dialect)
class CpAsyncOp(_ods_ir.OpView):
  OPERATION_NAME = "nvvm.cp.async.shared.global"

  _ODS_REGIONS = (0, True)

  def __init__(self, dst, src, size, modifier, *, cpSize=None, loc=None, ip=None):
    operands = []
    results = []
    attributes = {}
    regions = None
    operands.append(_get_op_result_or_value(dst))
    operands.append(_get_op_result_or_value(src))
    if cpSize is not None: operands.append(_get_op_result_or_value(cpSize))
    _ods_context = _ods_get_default_loc_context(loc)
    attributes["size"] = (size if (
    isinstance(size, _ods_ir.Attribute) or
    not _ods_ir.AttrBuilder.contains('I32Attr')) else
      _ods_ir.AttrBuilder.get('I32Attr')(size, context=_ods_context))
    attributes["modifier"] = (modifier if (
    isinstance(modifier, _ods_ir.Attribute) or
    not _ods_ir.AttrBuilder.contains('LoadCacheModifierAttr')) else
      _ods_ir.AttrBuilder.get('LoadCacheModifierAttr')(modifier, context=_ods_context))
    _ods_successors = None
    super().__init__(self.build_generic(attributes=attributes, results=results, operands=operands, successors=_ods_successors, regions=regions, loc=loc, ip=ip))

  @builtins.property
  def dst(self):
    return self.operation.operands[0]

  @builtins.property
  def src(self):
    return self.operation.operands[1]

  @builtins.property
  def cpSize(self):
    return None if len(self.operation.operands) < 3 else self.operation.operands[2]

  @builtins.property
  def size(self):
    return self.operation.attributes["size"]

  @size.setter
  def size(self, value):
    if value is None:
      raise ValueError("'None' not allowed as value for mandatory attributes")
    self.operation.attributes["size"] = value

  @builtins.property
  def modifier(self):
    return self.operation.attributes["modifier"]

  @modifier.setter
  def modifier(self, value):
    if value is None:
      raise ValueError("'None' not allowed as value for mandatory attributes")
    self.operation.attributes["modifier"] = value

def cp_async_shared_global(dst, src, size, modifier, *, cp_size=None, loc=None, ip=None) -> _ods_ir.Operation:
  return _get_op_result_or_op_results(CpAsyncOp(dst=dst, src=src, size=size, modifier=modifier, cpSize=cp_size, loc=loc, ip=ip))

@_ods_cext.register_operation(_Dialect)
class CpAsyncWaitGroupOp(_ods_ir.OpView):
  OPERATION_NAME = "nvvm.cp.async.wait.group"

  _ODS_REGIONS = (0, True)

  def __init__(self, n, *, loc=None, ip=None):
    operands = []
    results = []
    attributes = {}
    regions = None
    _ods_context = _ods_get_default_loc_context(loc)
    attributes["n"] = (n if (
    isinstance(n, _ods_ir.Attribute) or
    not _ods_ir.AttrBuilder.contains('I32Attr')) else
      _ods_ir.AttrBuilder.get('I32Attr')(n, context=_ods_context))
    _ods_successors = None
    super().__init__(self.build_generic(attributes=attributes, results=results, operands=operands, successors=_ods_successors, regions=regions, loc=loc, ip=ip))

  @builtins.property
  def n(self):
    return self.operation.attributes["n"]

  @n.setter
  def n(self, value):
    if value is None:
      raise ValueError("'None' not allowed as value for mandatory attributes")
    self.operation.attributes["n"] = value

def cp_async_wait_group(n, *, loc=None, ip=None) -> _ods_ir.Operation:
  return _get_op_result_or_op_results(CpAsyncWaitGroupOp(n=n, loc=loc, ip=ip))

@_ods_cext.register_operation(_Dialect)
class ElectSyncOp(_ods_ir.OpView):
  OPERATION_NAME = "nvvm.elect.sync"

  _ODS_REGIONS = (0, True)

  def __init__(self, pred, *, loc=None, ip=None):
    operands = []
    results = []
    attributes = {}
    regions = None
    _ods_context = _ods_get_default_loc_context(loc)
    results.append(pred)
    _ods_successors = None
    super().__init__(self.build_generic(attributes=attributes, results=results, operands=operands, successors=_ods_successors, regions=regions, loc=loc, ip=ip))

  @builtins.property
  def pred(self):
    return self.operation.results[0]

def elect_sync(pred, *, loc=None, ip=None) -> _ods_ir.Value:
  return _get_op_result_or_op_results(ElectSyncOp(pred=pred, loc=loc, ip=ip))

@_ods_cext.register_operation(_Dialect)
class FenceMbarrierInitOp(_ods_ir.OpView):
  OPERATION_NAME = "nvvm.fence.mbarrier.init"

  _ODS_REGIONS = (0, True)

  def __init__(self, *, loc=None, ip=None):
    operands = []
    results = []
    attributes = {}
    regions = None
    _ods_context = _ods_get_default_loc_context(loc)
    _ods_successors = None
    super().__init__(self.build_generic(attributes=attributes, results=results, operands=operands, successors=_ods_successors, regions=regions, loc=loc, ip=ip))

def fence_mbarrier_init(*, loc=None, ip=None) -> _ods_ir.Operation:
  return _get_op_result_or_op_results(FenceMbarrierInitOp(loc=loc, ip=ip))

@_ods_cext.register_operation(_Dialect)
class FenceProxyOp(_ods_ir.OpView):
  OPERATION_NAME = "nvvm.fence.proxy"

  _ODS_REGIONS = (0, True)

  def __init__(self, kind, *, space=None, loc=None, ip=None):
    operands = []
    results = []
    attributes = {}
    regions = None
    _ods_context = _ods_get_default_loc_context(loc)
    attributes["kind"] = (kind if (
    isinstance(kind, _ods_ir.Attribute) or
    not _ods_ir.AttrBuilder.contains('ProxyKindAttr')) else
      _ods_ir.AttrBuilder.get('ProxyKindAttr')(kind, context=_ods_context))
    if space is not None: attributes["space"] = (space if (
        isinstance(space, _ods_ir.Attribute) or
        not _ods_ir.AttrBuilder.contains('SharedSpaceAttr')) else
          _ods_ir.AttrBuilder.get('SharedSpaceAttr')(space, context=_ods_context))
    _ods_successors = None
    super().__init__(self.build_generic(attributes=attributes, results=results, operands=operands, successors=_ods_successors, regions=regions, loc=loc, ip=ip))

  @builtins.property
  def kind(self):
    return self.operation.attributes["kind"]

  @kind.setter
  def kind(self, value):
    if value is None:
      raise ValueError("'None' not allowed as value for mandatory attributes")
    self.operation.attributes["kind"] = value

  @builtins.property
  def space(self):
    if "space" not in self.operation.attributes:
      return None
    return self.operation.attributes["space"]

  @space.setter
  def space(self, value):
    if value is not None:
      self.operation.attributes["space"] = value
    elif "space" in self.operation.attributes:
      del self.operation.attributes["space"]

  @space.deleter
  def space(self):
    del self.operation.attributes["space"]

def fence_proxy(kind, *, space=None, loc=None, ip=None) -> _ods_ir.Operation:
  return _get_op_result_or_op_results(FenceProxyOp(kind=kind, space=space, loc=loc, ip=ip))

@_ods_cext.register_operation(_Dialect)
class FenceScClusterOp(_ods_ir.OpView):
  OPERATION_NAME = "nvvm.fence.sc.cluster"

  _ODS_REGIONS = (0, True)

  def __init__(self, *, loc=None, ip=None):
    operands = []
    results = []
    attributes = {}
    regions = None
    _ods_context = _ods_get_default_loc_context(loc)
    _ods_successors = None
    super().__init__(self.build_generic(attributes=attributes, results=results, operands=operands, successors=_ods_successors, regions=regions, loc=loc, ip=ip))

def fence_sc_cluster(*, loc=None, ip=None) -> _ods_ir.Operation:
  return _get_op_result_or_op_results(FenceScClusterOp(loc=loc, ip=ip))

@_ods_cext.register_operation(_Dialect)
class GridDimXOp(_ods_ir.OpView):
  OPERATION_NAME = "nvvm.read.ptx.sreg.nctaid.x"

  _ODS_REGIONS = (0, True)

  def __init__(self, res, *, loc=None, ip=None):
    operands = []
    results = []
    attributes = {}
    regions = None
    _ods_context = _ods_get_default_loc_context(loc)
    results.append(res)
    _ods_successors = None
    super().__init__(self.build_generic(attributes=attributes, results=results, operands=operands, successors=_ods_successors, regions=regions, loc=loc, ip=ip))

  @builtins.property
  def res(self):
    return self.operation.results[0]

def read_ptx_sreg_nctaid_x(res, *, loc=None, ip=None) -> _ods_ir.Value:
  return _get_op_result_or_op_results(GridDimXOp(res=res, loc=loc, ip=ip))

@_ods_cext.register_operation(_Dialect)
class GridDimYOp(_ods_ir.OpView):
  OPERATION_NAME = "nvvm.read.ptx.sreg.nctaid.y"

  _ODS_REGIONS = (0, True)

  def __init__(self, res, *, loc=None, ip=None):
    operands = []
    results = []
    attributes = {}
    regions = None
    _ods_context = _ods_get_default_loc_context(loc)
    results.append(res)
    _ods_successors = None
    super().__init__(self.build_generic(attributes=attributes, results=results, operands=operands, successors=_ods_successors, regions=regions, loc=loc, ip=ip))

  @builtins.property
  def res(self):
    return self.operation.results[0]

def read_ptx_sreg_nctaid_y(res, *, loc=None, ip=None) -> _ods_ir.Value:
  return _get_op_result_or_op_results(GridDimYOp(res=res, loc=loc, ip=ip))

@_ods_cext.register_operation(_Dialect)
class GridDimZOp(_ods_ir.OpView):
  OPERATION_NAME = "nvvm.read.ptx.sreg.nctaid.z"

  _ODS_REGIONS = (0, True)

  def __init__(self, res, *, loc=None, ip=None):
    operands = []
    results = []
    attributes = {}
    regions = None
    _ods_context = _ods_get_default_loc_context(loc)
    results.append(res)
    _ods_successors = None
    super().__init__(self.build_generic(attributes=attributes, results=results, operands=operands, successors=_ods_successors, regions=regions, loc=loc, ip=ip))

  @builtins.property
  def res(self):
    return self.operation.results[0]

def read_ptx_sreg_nctaid_z(res, *, loc=None, ip=None) -> _ods_ir.Value:
  return _get_op_result_or_op_results(GridDimZOp(res=res, loc=loc, ip=ip))

@_ods_cext.register_operation(_Dialect)
class LaneIdOp(_ods_ir.OpView):
  OPERATION_NAME = "nvvm.read.ptx.sreg.laneid"

  _ODS_REGIONS = (0, True)

  def __init__(self, res, *, loc=None, ip=None):
    operands = []
    results = []
    attributes = {}
    regions = None
    _ods_context = _ods_get_default_loc_context(loc)
    results.append(res)
    _ods_successors = None
    super().__init__(self.build_generic(attributes=attributes, results=results, operands=operands, successors=_ods_successors, regions=regions, loc=loc, ip=ip))

  @builtins.property
  def res(self):
    return self.operation.results[0]

def read_ptx_sreg_laneid(res, *, loc=None, ip=None) -> _ods_ir.Value:
  return _get_op_result_or_op_results(LaneIdOp(res=res, loc=loc, ip=ip))

@_ods_cext.register_operation(_Dialect)
class LdMatrixOp(_ods_ir.OpView):
  OPERATION_NAME = "nvvm.ldmatrix"

  _ODS_REGIONS = (0, True)

  def __init__(self, res, ptr, num, layout, *, loc=None, ip=None):
    operands = []
    results = []
    attributes = {}
    regions = None
    operands.append(_get_op_result_or_value(ptr))
    _ods_context = _ods_get_default_loc_context(loc)
    attributes["num"] = (num if (
    isinstance(num, _ods_ir.Attribute) or
    not _ods_ir.AttrBuilder.contains('I32Attr')) else
      _ods_ir.AttrBuilder.get('I32Attr')(num, context=_ods_context))
    attributes["layout"] = (layout if (
    isinstance(layout, _ods_ir.Attribute) or
    not _ods_ir.AttrBuilder.contains('MMALayoutAttr')) else
      _ods_ir.AttrBuilder.get('MMALayoutAttr')(layout, context=_ods_context))
    results.append(res)
    _ods_successors = None
    super().__init__(self.build_generic(attributes=attributes, results=results, operands=operands, successors=_ods_successors, regions=regions, loc=loc, ip=ip))

  @builtins.property
  def ptr(self):
    return self.operation.operands[0]

  @builtins.property
  def num(self):
    return self.operation.attributes["num"]

  @num.setter
  def num(self, value):
    if value is None:
      raise ValueError("'None' not allowed as value for mandatory attributes")
    self.operation.attributes["num"] = value

  @builtins.property
  def layout(self):
    return self.operation.attributes["layout"]

  @layout.setter
  def layout(self, value):
    if value is None:
      raise ValueError("'None' not allowed as value for mandatory attributes")
    self.operation.attributes["layout"] = value

  @builtins.property
  def res(self):
    return self.operation.results[0]

def ldmatrix(res, ptr, num, layout, *, loc=None, ip=None) -> _ods_ir.Value:
  return _get_op_result_or_op_results(LdMatrixOp(res=res, ptr=ptr, num=num, layout=layout, loc=loc, ip=ip))

@_ods_cext.register_operation(_Dialect)
class MBarrierArriveExpectTxOp(_ods_ir.OpView):
  OPERATION_NAME = "nvvm.mbarrier.arrive.expect_tx"

  _ODS_REGIONS = (0, True)

  def __init__(self, addr, txcount, *, predicate=None, loc=None, ip=None):
    operands = []
    results = []
    attributes = {}
    regions = None
    operands.append(_get_op_result_or_value(addr))
    operands.append(_get_op_result_or_value(txcount))
    if predicate is not None: operands.append(_get_op_result_or_value(predicate))
    _ods_context = _ods_get_default_loc_context(loc)
    _ods_successors = None
    super().__init__(self.build_generic(attributes=attributes, results=results, operands=operands, successors=_ods_successors, regions=regions, loc=loc, ip=ip))

  @builtins.property
  def addr(self):
    return self.operation.operands[0]

  @builtins.property
  def txcount(self):
    return self.operation.operands[1]

  @builtins.property
  def predicate(self):
    return None if len(self.operation.operands) < 3 else self.operation.operands[2]

def mbarrier_arrive_expect_tx(addr, txcount, *, predicate=None, loc=None, ip=None) -> _ods_ir.Operation:
  return _get_op_result_or_op_results(MBarrierArriveExpectTxOp(addr=addr, txcount=txcount, predicate=predicate, loc=loc, ip=ip))

@_ods_cext.register_operation(_Dialect)
class MBarrierArriveExpectTxSharedOp(_ods_ir.OpView):
  OPERATION_NAME = "nvvm.mbarrier.arrive.expect_tx.shared"

  _ODS_REGIONS = (0, True)

  def __init__(self, addr, txcount, *, predicate=None, loc=None, ip=None):
    operands = []
    results = []
    attributes = {}
    regions = None
    operands.append(_get_op_result_or_value(addr))
    operands.append(_get_op_result_or_value(txcount))
    if predicate is not None: operands.append(_get_op_result_or_value(predicate))
    _ods_context = _ods_get_default_loc_context(loc)
    _ods_successors = None
    super().__init__(self.build_generic(attributes=attributes, results=results, operands=operands, successors=_ods_successors, regions=regions, loc=loc, ip=ip))

  @builtins.property
  def addr(self):
    return self.operation.operands[0]

  @builtins.property
  def txcount(self):
    return self.operation.operands[1]

  @builtins.property
  def predicate(self):
    return None if len(self.operation.operands) < 3 else self.operation.operands[2]

def mbarrier_arrive_expect_tx_shared(addr, txcount, *, predicate=None, loc=None, ip=None) -> _ods_ir.Operation:
  return _get_op_result_or_op_results(MBarrierArriveExpectTxSharedOp(addr=addr, txcount=txcount, predicate=predicate, loc=loc, ip=ip))

@_ods_cext.register_operation(_Dialect)
class MBarrierArriveNocompleteOp(_ods_ir.OpView):
  OPERATION_NAME = "nvvm.mbarrier.arrive.nocomplete"

  _ODS_REGIONS = (0, True)

  def __init__(self, res, addr, count, *, loc=None, ip=None):
    operands = []
    results = []
    attributes = {}
    regions = None
    operands.append(_get_op_result_or_value(addr))
    operands.append(_get_op_result_or_value(count))
    _ods_context = _ods_get_default_loc_context(loc)
    results.append(res)
    _ods_successors = None
    super().__init__(self.build_generic(attributes=attributes, results=results, operands=operands, successors=_ods_successors, regions=regions, loc=loc, ip=ip))

  @builtins.property
  def addr(self):
    return self.operation.operands[0]

  @builtins.property
  def count(self):
    return self.operation.operands[1]

  @builtins.property
  def res(self):
    return self.operation.results[0]

def mbarrier_arrive_nocomplete(res, addr, count, *, loc=None, ip=None) -> _ods_ir.Value:
  return _get_op_result_or_op_results(MBarrierArriveNocompleteOp(res=res, addr=addr, count=count, loc=loc, ip=ip))

@_ods_cext.register_operation(_Dialect)
class MBarrierArriveNocompleteSharedOp(_ods_ir.OpView):
  OPERATION_NAME = "nvvm.mbarrier.arrive.nocomplete.shared"

  _ODS_REGIONS = (0, True)

  def __init__(self, res, addr, count, *, loc=None, ip=None):
    operands = []
    results = []
    attributes = {}
    regions = None
    operands.append(_get_op_result_or_value(addr))
    operands.append(_get_op_result_or_value(count))
    _ods_context = _ods_get_default_loc_context(loc)
    results.append(res)
    _ods_successors = None
    super().__init__(self.build_generic(attributes=attributes, results=results, operands=operands, successors=_ods_successors, regions=regions, loc=loc, ip=ip))

  @builtins.property
  def addr(self):
    return self.operation.operands[0]

  @builtins.property
  def count(self):
    return self.operation.operands[1]

  @builtins.property
  def res(self):
    return self.operation.results[0]

def mbarrier_arrive_nocomplete_shared(res, addr, count, *, loc=None, ip=None) -> _ods_ir.Value:
  return _get_op_result_or_op_results(MBarrierArriveNocompleteSharedOp(res=res, addr=addr, count=count, loc=loc, ip=ip))

@_ods_cext.register_operation(_Dialect)
class MBarrierArriveOp(_ods_ir.OpView):
  OPERATION_NAME = "nvvm.mbarrier.arrive"

  _ODS_REGIONS = (0, True)

  def __init__(self, res, addr, *, loc=None, ip=None):
    operands = []
    results = []
    attributes = {}
    regions = None
    operands.append(_get_op_result_or_value(addr))
    _ods_context = _ods_get_default_loc_context(loc)
    results.append(res)
    _ods_successors = None
    super().__init__(self.build_generic(attributes=attributes, results=results, operands=operands, successors=_ods_successors, regions=regions, loc=loc, ip=ip))

  @builtins.property
  def addr(self):
    return self.operation.operands[0]

  @builtins.property
  def res(self):
    return self.operation.results[0]

def mbarrier_arrive(res, addr, *, loc=None, ip=None) -> _ods_ir.Value:
  return _get_op_result_or_op_results(MBarrierArriveOp(res=res, addr=addr, loc=loc, ip=ip))

@_ods_cext.register_operation(_Dialect)
class MBarrierArriveSharedOp(_ods_ir.OpView):
  OPERATION_NAME = "nvvm.mbarrier.arrive.shared"

  _ODS_REGIONS = (0, True)

  def __init__(self, res, addr, *, loc=None, ip=None):
    operands = []
    results = []
    attributes = {}
    regions = None
    operands.append(_get_op_result_or_value(addr))
    _ods_context = _ods_get_default_loc_context(loc)
    results.append(res)
    _ods_successors = None
    super().__init__(self.build_generic(attributes=attributes, results=results, operands=operands, successors=_ods_successors, regions=regions, loc=loc, ip=ip))

  @builtins.property
  def addr(self):
    return self.operation.operands[0]

  @builtins.property
  def res(self):
    return self.operation.results[0]

def mbarrier_arrive_shared(res, addr, *, loc=None, ip=None) -> _ods_ir.Value:
  return _get_op_result_or_op_results(MBarrierArriveSharedOp(res=res, addr=addr, loc=loc, ip=ip))

@_ods_cext.register_operation(_Dialect)
class MBarrierInitOp(_ods_ir.OpView):
  OPERATION_NAME = "nvvm.mbarrier.init"

  _ODS_REGIONS = (0, True)

  def __init__(self, addr, count, *, predicate=None, loc=None, ip=None):
    operands = []
    results = []
    attributes = {}
    regions = None
    operands.append(_get_op_result_or_value(addr))
    operands.append(_get_op_result_or_value(count))
    if predicate is not None: operands.append(_get_op_result_or_value(predicate))
    _ods_context = _ods_get_default_loc_context(loc)
    _ods_successors = None
    super().__init__(self.build_generic(attributes=attributes, results=results, operands=operands, successors=_ods_successors, regions=regions, loc=loc, ip=ip))

  @builtins.property
  def addr(self):
    return self.operation.operands[0]

  @builtins.property
  def count(self):
    return self.operation.operands[1]

  @builtins.property
  def predicate(self):
    return None if len(self.operation.operands) < 3 else self.operation.operands[2]

def mbarrier_init(addr, count, *, predicate=None, loc=None, ip=None) -> _ods_ir.Operation:
  return _get_op_result_or_op_results(MBarrierInitOp(addr=addr, count=count, predicate=predicate, loc=loc, ip=ip))

@_ods_cext.register_operation(_Dialect)
class MBarrierInitSharedOp(_ods_ir.OpView):
  OPERATION_NAME = "nvvm.mbarrier.init.shared"

  _ODS_REGIONS = (0, True)

  def __init__(self, addr, count, *, predicate=None, loc=None, ip=None):
    operands = []
    results = []
    attributes = {}
    regions = None
    operands.append(_get_op_result_or_value(addr))
    operands.append(_get_op_result_or_value(count))
    if predicate is not None: operands.append(_get_op_result_or_value(predicate))
    _ods_context = _ods_get_default_loc_context(loc)
    _ods_successors = None
    super().__init__(self.build_generic(attributes=attributes, results=results, operands=operands, successors=_ods_successors, regions=regions, loc=loc, ip=ip))

  @builtins.property
  def addr(self):
    return self.operation.operands[0]

  @builtins.property
  def count(self):
    return self.operation.operands[1]

  @builtins.property
  def predicate(self):
    return None if len(self.operation.operands) < 3 else self.operation.operands[2]

def mbarrier_init_shared(addr, count, *, predicate=None, loc=None, ip=None) -> _ods_ir.Operation:
  return _get_op_result_or_op_results(MBarrierInitSharedOp(addr=addr, count=count, predicate=predicate, loc=loc, ip=ip))

@_ods_cext.register_operation(_Dialect)
class MBarrierInvalOp(_ods_ir.OpView):
  OPERATION_NAME = "nvvm.mbarrier.inval"

  _ODS_REGIONS = (0, True)

  def __init__(self, addr, *, loc=None, ip=None):
    operands = []
    results = []
    attributes = {}
    regions = None
    operands.append(_get_op_result_or_value(addr))
    _ods_context = _ods_get_default_loc_context(loc)
    _ods_successors = None
    super().__init__(self.build_generic(attributes=attributes, results=results, operands=operands, successors=_ods_successors, regions=regions, loc=loc, ip=ip))

  @builtins.property
  def addr(self):
    return self.operation.operands[0]

def mbarrier_inval(addr, *, loc=None, ip=None) -> _ods_ir.Operation:
  return _get_op_result_or_op_results(MBarrierInvalOp(addr=addr, loc=loc, ip=ip))

@_ods_cext.register_operation(_Dialect)
class MBarrierInvalSharedOp(_ods_ir.OpView):
  OPERATION_NAME = "nvvm.mbarrier.inval.shared"

  _ODS_REGIONS = (0, True)

  def __init__(self, addr, *, loc=None, ip=None):
    operands = []
    results = []
    attributes = {}
    regions = None
    operands.append(_get_op_result_or_value(addr))
    _ods_context = _ods_get_default_loc_context(loc)
    _ods_successors = None
    super().__init__(self.build_generic(attributes=attributes, results=results, operands=operands, successors=_ods_successors, regions=regions, loc=loc, ip=ip))

  @builtins.property
  def addr(self):
    return self.operation.operands[0]

def mbarrier_inval_shared(addr, *, loc=None, ip=None) -> _ods_ir.Operation:
  return _get_op_result_or_op_results(MBarrierInvalSharedOp(addr=addr, loc=loc, ip=ip))

@_ods_cext.register_operation(_Dialect)
class MBarrierTestWaitOp(_ods_ir.OpView):
  OPERATION_NAME = "nvvm.mbarrier.test.wait"

  _ODS_REGIONS = (0, True)

  def __init__(self, res, addr, state, *, loc=None, ip=None):
    operands = []
    results = []
    attributes = {}
    regions = None
    operands.append(_get_op_result_or_value(addr))
    operands.append(_get_op_result_or_value(state))
    _ods_context = _ods_get_default_loc_context(loc)
    results.append(res)
    _ods_successors = None
    super().__init__(self.build_generic(attributes=attributes, results=results, operands=operands, successors=_ods_successors, regions=regions, loc=loc, ip=ip))

  @builtins.property
  def addr(self):
    return self.operation.operands[0]

  @builtins.property
  def state(self):
    return self.operation.operands[1]

  @builtins.property
  def res(self):
    return self.operation.results[0]

def mbarrier_test_wait(res, addr, state, *, loc=None, ip=None) -> _ods_ir.Value:
  return _get_op_result_or_op_results(MBarrierTestWaitOp(res=res, addr=addr, state=state, loc=loc, ip=ip))

@_ods_cext.register_operation(_Dialect)
class MBarrierTestWaitSharedOp(_ods_ir.OpView):
  OPERATION_NAME = "nvvm.mbarrier.test.wait.shared"

  _ODS_REGIONS = (0, True)

  def __init__(self, res, addr, state, *, loc=None, ip=None):
    operands = []
    results = []
    attributes = {}
    regions = None
    operands.append(_get_op_result_or_value(addr))
    operands.append(_get_op_result_or_value(state))
    _ods_context = _ods_get_default_loc_context(loc)
    results.append(res)
    _ods_successors = None
    super().__init__(self.build_generic(attributes=attributes, results=results, operands=operands, successors=_ods_successors, regions=regions, loc=loc, ip=ip))

  @builtins.property
  def addr(self):
    return self.operation.operands[0]

  @builtins.property
  def state(self):
    return self.operation.operands[1]

  @builtins.property
  def res(self):
    return self.operation.results[0]

def mbarrier_test_wait_shared(res, addr, state, *, loc=None, ip=None) -> _ods_ir.Value:
  return _get_op_result_or_op_results(MBarrierTestWaitSharedOp(res=res, addr=addr, state=state, loc=loc, ip=ip))

@_ods_cext.register_operation(_Dialect)
class MBarrierTryWaitParityOp(_ods_ir.OpView):
  OPERATION_NAME = "nvvm.mbarrier.try_wait.parity"

  _ODS_REGIONS = (0, True)

  def __init__(self, addr, phase, ticks, *, loc=None, ip=None):
    operands = []
    results = []
    attributes = {}
    regions = None
    operands.append(_get_op_result_or_value(addr))
    operands.append(_get_op_result_or_value(phase))
    operands.append(_get_op_result_or_value(ticks))
    _ods_context = _ods_get_default_loc_context(loc)
    _ods_successors = None
    super().__init__(self.build_generic(attributes=attributes, results=results, operands=operands, successors=_ods_successors, regions=regions, loc=loc, ip=ip))

  @builtins.property
  def addr(self):
    return self.operation.operands[0]

  @builtins.property
  def phase(self):
    return self.operation.operands[1]

  @builtins.property
  def ticks(self):
    return self.operation.operands[2]

def mbarrier_try_wait_parity(addr, phase, ticks, *, loc=None, ip=None) -> _ods_ir.Operation:
  return _get_op_result_or_op_results(MBarrierTryWaitParityOp(addr=addr, phase=phase, ticks=ticks, loc=loc, ip=ip))

@_ods_cext.register_operation(_Dialect)
class MBarrierTryWaitParitySharedOp(_ods_ir.OpView):
  OPERATION_NAME = "nvvm.mbarrier.try_wait.parity.shared"

  _ODS_REGIONS = (0, True)

  def __init__(self, addr, phase, ticks, *, loc=None, ip=None):
    operands = []
    results = []
    attributes = {}
    regions = None
    operands.append(_get_op_result_or_value(addr))
    operands.append(_get_op_result_or_value(phase))
    operands.append(_get_op_result_or_value(ticks))
    _ods_context = _ods_get_default_loc_context(loc)
    _ods_successors = None
    super().__init__(self.build_generic(attributes=attributes, results=results, operands=operands, successors=_ods_successors, regions=regions, loc=loc, ip=ip))

  @builtins.property
  def addr(self):
    return self.operation.operands[0]

  @builtins.property
  def phase(self):
    return self.operation.operands[1]

  @builtins.property
  def ticks(self):
    return self.operation.operands[2]

def mbarrier_try_wait_parity_shared(addr, phase, ticks, *, loc=None, ip=None) -> _ods_ir.Operation:
  return _get_op_result_or_op_results(MBarrierTryWaitParitySharedOp(addr=addr, phase=phase, ticks=ticks, loc=loc, ip=ip))

@_ods_cext.register_operation(_Dialect)
class MmaOp(_ods_ir.OpView):
  OPERATION_NAME = "nvvm.mma.sync"

  _ODS_OPERAND_SEGMENTS = [-1,-1,-1,]

  _ODS_REGIONS = (0, True)

  def __init__(self, res, shape, layoutA, layoutB, operandA, operandB, operandC, *, b1Op=None, intOverflowBehavior=None, multiplicandAPtxType=None, multiplicandBPtxType=None, loc=None, ip=None):
    operands = []
    results = []
    attributes = {}
    regions = None
    operands.append(_get_op_results_or_values(operandA))
    operands.append(_get_op_results_or_values(operandB))
    operands.append(_get_op_results_or_values(operandC))
    _ods_context = _ods_get_default_loc_context(loc)
    attributes["shape"] = (shape if (
    isinstance(shape, _ods_ir.Attribute) or
    not _ods_ir.AttrBuilder.contains('NVVM_MMAShapeAttr')) else
      _ods_ir.AttrBuilder.get('NVVM_MMAShapeAttr')(shape, context=_ods_context))
    if b1Op is not None: attributes["b1Op"] = (b1Op if (
        isinstance(b1Op, _ods_ir.Attribute) or
        not _ods_ir.AttrBuilder.contains('MMAB1OpAttr')) else
          _ods_ir.AttrBuilder.get('MMAB1OpAttr')(b1Op, context=_ods_context))
    if intOverflowBehavior is not None: attributes["intOverflowBehavior"] = (intOverflowBehavior if (
        isinstance(intOverflowBehavior, _ods_ir.Attribute) or
        not _ods_ir.AttrBuilder.contains('MMAIntOverflowAttr')) else
          _ods_ir.AttrBuilder.get('MMAIntOverflowAttr')(intOverflowBehavior, context=_ods_context))
    attributes["layoutA"] = (layoutA if (
    isinstance(layoutA, _ods_ir.Attribute) or
    not _ods_ir.AttrBuilder.contains('MMALayoutAttr')) else
      _ods_ir.AttrBuilder.get('MMALayoutAttr')(layoutA, context=_ods_context))
    attributes["layoutB"] = (layoutB if (
    isinstance(layoutB, _ods_ir.Attribute) or
    not _ods_ir.AttrBuilder.contains('MMALayoutAttr')) else
      _ods_ir.AttrBuilder.get('MMALayoutAttr')(layoutB, context=_ods_context))
    if multiplicandAPtxType is not None: attributes["multiplicandAPtxType"] = (multiplicandAPtxType if (
        isinstance(multiplicandAPtxType, _ods_ir.Attribute) or
        not _ods_ir.AttrBuilder.contains('MMATypesAttr')) else
          _ods_ir.AttrBuilder.get('MMATypesAttr')(multiplicandAPtxType, context=_ods_context))
    if multiplicandBPtxType is not None: attributes["multiplicandBPtxType"] = (multiplicandBPtxType if (
        isinstance(multiplicandBPtxType, _ods_ir.Attribute) or
        not _ods_ir.AttrBuilder.contains('MMATypesAttr')) else
          _ods_ir.AttrBuilder.get('MMATypesAttr')(multiplicandBPtxType, context=_ods_context))
    results.append(res)
    _ods_successors = None
    super().__init__(self.build_generic(attributes=attributes, results=results, operands=operands, successors=_ods_successors, regions=regions, loc=loc, ip=ip))

  @builtins.property
  def operandA(self):
    operand_range = _ods_segmented_accessor(
         self.operation.operands,
         self.operation.attributes["operandSegmentSizes"], 0)
    return operand_range

  @builtins.property
  def operandB(self):
    operand_range = _ods_segmented_accessor(
         self.operation.operands,
         self.operation.attributes["operandSegmentSizes"], 1)
    return operand_range

  @builtins.property
  def operandC(self):
    operand_range = _ods_segmented_accessor(
         self.operation.operands,
         self.operation.attributes["operandSegmentSizes"], 2)
    return operand_range

  @builtins.property
  def shape(self):
    return self.operation.attributes["shape"]

  @shape.setter
  def shape(self, value):
    if value is None:
      raise ValueError("'None' not allowed as value for mandatory attributes")
    self.operation.attributes["shape"] = value

  @builtins.property
  def b1Op(self):
    if "b1Op" not in self.operation.attributes:
      return None
    return self.operation.attributes["b1Op"]

  @b1Op.setter
  def b1Op(self, value):
    if value is not None:
      self.operation.attributes["b1Op"] = value
    elif "b1Op" in self.operation.attributes:
      del self.operation.attributes["b1Op"]

  @b1Op.deleter
  def b1Op(self):
    del self.operation.attributes["b1Op"]

  @builtins.property
  def intOverflowBehavior(self):
    if "intOverflowBehavior" not in self.operation.attributes:
      return None
    return self.operation.attributes["intOverflowBehavior"]

  @intOverflowBehavior.setter
  def intOverflowBehavior(self, value):
    if value is not None:
      self.operation.attributes["intOverflowBehavior"] = value
    elif "intOverflowBehavior" in self.operation.attributes:
      del self.operation.attributes["intOverflowBehavior"]

  @intOverflowBehavior.deleter
  def intOverflowBehavior(self):
    del self.operation.attributes["intOverflowBehavior"]

  @builtins.property
  def layoutA(self):
    return self.operation.attributes["layoutA"]

  @layoutA.setter
  def layoutA(self, value):
    if value is None:
      raise ValueError("'None' not allowed as value for mandatory attributes")
    self.operation.attributes["layoutA"] = value

  @builtins.property
  def layoutB(self):
    return self.operation.attributes["layoutB"]

  @layoutB.setter
  def layoutB(self, value):
    if value is None:
      raise ValueError("'None' not allowed as value for mandatory attributes")
    self.operation.attributes["layoutB"] = value

  @builtins.property
  def multiplicandAPtxType(self):
    if "multiplicandAPtxType" not in self.operation.attributes:
      return None
    return self.operation.attributes["multiplicandAPtxType"]

  @multiplicandAPtxType.setter
  def multiplicandAPtxType(self, value):
    if value is not None:
      self.operation.attributes["multiplicandAPtxType"] = value
    elif "multiplicandAPtxType" in self.operation.attributes:
      del self.operation.attributes["multiplicandAPtxType"]

  @multiplicandAPtxType.deleter
  def multiplicandAPtxType(self):
    del self.operation.attributes["multiplicandAPtxType"]

  @builtins.property
  def multiplicandBPtxType(self):
    if "multiplicandBPtxType" not in self.operation.attributes:
      return None
    return self.operation.attributes["multiplicandBPtxType"]

  @multiplicandBPtxType.setter
  def multiplicandBPtxType(self, value):
    if value is not None:
      self.operation.attributes["multiplicandBPtxType"] = value
    elif "multiplicandBPtxType" in self.operation.attributes:
      del self.operation.attributes["multiplicandBPtxType"]

  @multiplicandBPtxType.deleter
  def multiplicandBPtxType(self):
    del self.operation.attributes["multiplicandBPtxType"]

  @builtins.property
  def res(self):
    return self.operation.results[0]

def mma_sync(res, shape, layout_a, layout_b, operand_a, operand_b, operand_c, *, b1_op=None, int_overflow_behavior=None, multiplicand_a_ptx_type=None, multiplicand_b_ptx_type=None, loc=None, ip=None) -> _ods_ir.Value:
  return _get_op_result_or_op_results(MmaOp(res=res, shape=shape, layoutA=layout_a, layoutB=layout_b, operandA=operand_a, operandB=operand_b, operandC=operand_c, b1Op=b1_op, intOverflowBehavior=int_overflow_behavior, multiplicandAPtxType=multiplicand_a_ptx_type, multiplicandBPtxType=multiplicand_b_ptx_type, loc=loc, ip=ip))

@_ods_cext.register_operation(_Dialect)
class PrefetchTensorMapOp(_ods_ir.OpView):
  OPERATION_NAME = "nvvm.prefetch.tensormap"

  _ODS_REGIONS = (0, True)

  def __init__(self, tmaDescriptor, *, predicate=None, loc=None, ip=None):
    operands = []
    results = []
    attributes = {}
    regions = None
    operands.append(_get_op_result_or_value(tmaDescriptor))
    if predicate is not None: operands.append(_get_op_result_or_value(predicate))
    _ods_context = _ods_get_default_loc_context(loc)
    _ods_successors = None
    super().__init__(self.build_generic(attributes=attributes, results=results, operands=operands, successors=_ods_successors, regions=regions, loc=loc, ip=ip))

  @builtins.property
  def tmaDescriptor(self):
    return self.operation.operands[0]

  @builtins.property
  def predicate(self):
    return None if len(self.operation.operands) < 2 else self.operation.operands[1]

def prefetch_tensormap(tma_descriptor, *, predicate=None, loc=None, ip=None) -> _ods_ir.Operation:
  return _get_op_result_or_op_results(PrefetchTensorMapOp(tmaDescriptor=tma_descriptor, predicate=predicate, loc=loc, ip=ip))

@_ods_cext.register_operation(_Dialect)
class RcpApproxFtzF32Op(_ods_ir.OpView):
  OPERATION_NAME = "nvvm.rcp.approx.ftz.f"

  _ODS_REGIONS = (0, True)

  def __init__(self, res, arg, *, loc=None, ip=None):
    operands = []
    results = []
    attributes = {}
    regions = None
    operands.append(_get_op_result_or_value(arg))
    _ods_context = _ods_get_default_loc_context(loc)
    results.append(res)
    _ods_successors = None
    super().__init__(self.build_generic(attributes=attributes, results=results, operands=operands, successors=_ods_successors, regions=regions, loc=loc, ip=ip))

  @builtins.property
  def arg(self):
    return self.operation.operands[0]

  @builtins.property
  def res(self):
    return self.operation.results[0]

def rcp_approx_ftz_f(res, arg, *, loc=None, ip=None) -> _ods_ir.Value:
  return _get_op_result_or_op_results(RcpApproxFtzF32Op(res=res, arg=arg, loc=loc, ip=ip))

@_ods_cext.register_operation(_Dialect)
class ReduxOp(_ods_ir.OpView):
  OPERATION_NAME = "nvvm.redux.sync"

  _ODS_REGIONS = (0, True)

  def __init__(self, res, val, kind, mask_and_clamp, *, loc=None, ip=None):
    operands = []
    results = []
    attributes = {}
    regions = None
    operands.append(_get_op_result_or_value(val))
    operands.append(_get_op_result_or_value(mask_and_clamp))
    _ods_context = _ods_get_default_loc_context(loc)
    attributes["kind"] = (kind if (
    isinstance(kind, _ods_ir.Attribute) or
    not _ods_ir.AttrBuilder.contains('ReduxKindAttr')) else
      _ods_ir.AttrBuilder.get('ReduxKindAttr')(kind, context=_ods_context))
    results.append(res)
    _ods_successors = None
    super().__init__(self.build_generic(attributes=attributes, results=results, operands=operands, successors=_ods_successors, regions=regions, loc=loc, ip=ip))

  @builtins.property
  def val(self):
    return self.operation.operands[0]

  @builtins.property
  def mask_and_clamp(self):
    return self.operation.operands[1]

  @builtins.property
  def kind(self):
    return self.operation.attributes["kind"]

  @kind.setter
  def kind(self, value):
    if value is None:
      raise ValueError("'None' not allowed as value for mandatory attributes")
    self.operation.attributes["kind"] = value

  @builtins.property
  def res(self):
    return self.operation.results[0]

def redux_sync(res, val, kind, mask_and_clamp, *, loc=None, ip=None) -> _ods_ir.Value:
  return _get_op_result_or_op_results(ReduxOp(res=res, val=val, kind=kind, mask_and_clamp=mask_and_clamp, loc=loc, ip=ip))

@_ods_cext.register_operation(_Dialect)
class SetMaxRegisterOp(_ods_ir.OpView):
  OPERATION_NAME = "nvvm.setmaxregister"

  _ODS_REGIONS = (0, True)

  def __init__(self, regCount, action, *, loc=None, ip=None):
    operands = []
    results = []
    attributes = {}
    regions = None
    _ods_context = _ods_get_default_loc_context(loc)
    attributes["regCount"] = (regCount if (
    isinstance(regCount, _ods_ir.Attribute) or
    not _ods_ir.AttrBuilder.contains('I32Attr')) else
      _ods_ir.AttrBuilder.get('I32Attr')(regCount, context=_ods_context))
    attributes["action"] = (action if (
    isinstance(action, _ods_ir.Attribute) or
    not _ods_ir.AttrBuilder.contains('SetMaxRegisterActionAttr')) else
      _ods_ir.AttrBuilder.get('SetMaxRegisterActionAttr')(action, context=_ods_context))
    _ods_successors = None
    super().__init__(self.build_generic(attributes=attributes, results=results, operands=operands, successors=_ods_successors, regions=regions, loc=loc, ip=ip))

  @builtins.property
  def regCount(self):
    return self.operation.attributes["regCount"]

  @regCount.setter
  def regCount(self, value):
    if value is None:
      raise ValueError("'None' not allowed as value for mandatory attributes")
    self.operation.attributes["regCount"] = value

  @builtins.property
  def action(self):
    return self.operation.attributes["action"]

  @action.setter
  def action(self, value):
    if value is None:
      raise ValueError("'None' not allowed as value for mandatory attributes")
    self.operation.attributes["action"] = value

def setmaxregister(reg_count, action, *, loc=None, ip=None) -> _ods_ir.Operation:
  return _get_op_result_or_op_results(SetMaxRegisterOp(regCount=reg_count, action=action, loc=loc, ip=ip))

@_ods_cext.register_operation(_Dialect)
class ShflOp(_ods_ir.OpView):
  OPERATION_NAME = "nvvm.shfl.sync"

  _ODS_REGIONS = (0, True)

  def __init__(self, res, thread_mask, val, offset, mask_and_clamp, kind, *, return_value_and_is_valid=None, loc=None, ip=None):
    operands = []
    results = []
    attributes = {}
    regions = None
    operands.append(_get_op_result_or_value(thread_mask))
    operands.append(_get_op_result_or_value(val))
    operands.append(_get_op_result_or_value(offset))
    operands.append(_get_op_result_or_value(mask_and_clamp))
    _ods_context = _ods_get_default_loc_context(loc)
    attributes["kind"] = (kind if (
    isinstance(kind, _ods_ir.Attribute) or
    not _ods_ir.AttrBuilder.contains('ShflKindAttr')) else
      _ods_ir.AttrBuilder.get('ShflKindAttr')(kind, context=_ods_context))
    if bool(return_value_and_is_valid): attributes["return_value_and_is_valid"] = _ods_ir.UnitAttr.get(
      _ods_get_default_loc_context(loc))
    results.append(res)
    _ods_successors = None
    super().__init__(self.build_generic(attributes=attributes, results=results, operands=operands, successors=_ods_successors, regions=regions, loc=loc, ip=ip))

  @builtins.property
  def thread_mask(self):
    return self.operation.operands[0]

  @builtins.property
  def val(self):
    return self.operation.operands[1]

  @builtins.property
  def offset(self):
    return self.operation.operands[2]

  @builtins.property
  def mask_and_clamp(self):
    return self.operation.operands[3]

  @builtins.property
  def kind(self):
    return self.operation.attributes["kind"]

  @kind.setter
  def kind(self, value):
    if value is None:
      raise ValueError("'None' not allowed as value for mandatory attributes")
    self.operation.attributes["kind"] = value

  @builtins.property
  def return_value_and_is_valid(self):
    return "return_value_and_is_valid" in self.operation.attributes

  @return_value_and_is_valid.setter
  def return_value_and_is_valid(self, value):
    if bool(value):
      self.operation.attributes["return_value_and_is_valid"] = _ods_ir.UnitAttr.get()
    elif "return_value_and_is_valid" in self.operation.attributes:
      del self.operation.attributes["return_value_and_is_valid"]

  @return_value_and_is_valid.deleter
  def return_value_and_is_valid(self):
    del self.operation.attributes["return_value_and_is_valid"]

  @builtins.property
  def res(self):
    return self.operation.results[0]

def shfl_sync(res, thread_mask, val, offset, mask_and_clamp, kind, *, return_value_and_is_valid=None, loc=None, ip=None) -> _ods_ir.Value:
  return _get_op_result_or_op_results(ShflOp(res=res, thread_mask=thread_mask, val=val, offset=offset, mask_and_clamp=mask_and_clamp, kind=kind, return_value_and_is_valid=return_value_and_is_valid, loc=loc, ip=ip))

@_ods_cext.register_operation(_Dialect)
class StMatrixOp(_ods_ir.OpView):
  OPERATION_NAME = "nvvm.stmatrix"

  _ODS_REGIONS = (0, True)

  def __init__(self, ptr, sources, layout, *, loc=None, ip=None):
    operands = []
    results = []
    attributes = {}
    regions = None
    operands.append(_get_op_result_or_value(ptr))
    operands.extend(_get_op_results_or_values(sources))
    _ods_context = _ods_get_default_loc_context(loc)
    attributes["layout"] = (layout if (
    isinstance(layout, _ods_ir.Attribute) or
    not _ods_ir.AttrBuilder.contains('MMALayoutAttr')) else
      _ods_ir.AttrBuilder.get('MMALayoutAttr')(layout, context=_ods_context))
    _ods_successors = None
    super().__init__(self.build_generic(attributes=attributes, results=results, operands=operands, successors=_ods_successors, regions=regions, loc=loc, ip=ip))

  @builtins.property
  def ptr(self):
    return self.operation.operands[0]

  @builtins.property
  def sources(self):
    _ods_variadic_group_length = len(self.operation.operands) - 2 + 1
    return self.operation.operands[1:1 + _ods_variadic_group_length]

  @builtins.property
  def layout(self):
    return self.operation.attributes["layout"]

  @layout.setter
  def layout(self, value):
    if value is None:
      raise ValueError("'None' not allowed as value for mandatory attributes")
    self.operation.attributes["layout"] = value

def stmatrix(ptr, sources, layout, *, loc=None, ip=None) -> _ods_ir.Operation:
  return _get_op_result_or_op_results(StMatrixOp(ptr=ptr, sources=sources, layout=layout, loc=loc, ip=ip))

@_ods_cext.register_operation(_Dialect)
class SyncWarpOp(_ods_ir.OpView):
  OPERATION_NAME = "nvvm.bar.warp.sync"

  _ODS_REGIONS = (0, True)

  def __init__(self, mask, *, loc=None, ip=None):
    operands = []
    results = []
    attributes = {}
    regions = None
    operands.append(_get_op_result_or_value(mask))
    _ods_context = _ods_get_default_loc_context(loc)
    _ods_successors = None
    super().__init__(self.build_generic(attributes=attributes, results=results, operands=operands, successors=_ods_successors, regions=regions, loc=loc, ip=ip))

  @builtins.property
  def mask(self):
    return self.operation.operands[0]

def bar_warp_sync(mask, *, loc=None, ip=None) -> _ods_ir.Operation:
  return _get_op_result_or_op_results(SyncWarpOp(mask=mask, loc=loc, ip=ip))

@_ods_cext.register_operation(_Dialect)
class ThreadIdXOp(_ods_ir.OpView):
  OPERATION_NAME = "nvvm.read.ptx.sreg.tid.x"

  _ODS_REGIONS = (0, True)

  def __init__(self, res, *, loc=None, ip=None):
    operands = []
    results = []
    attributes = {}
    regions = None
    _ods_context = _ods_get_default_loc_context(loc)
    results.append(res)
    _ods_successors = None
    super().__init__(self.build_generic(attributes=attributes, results=results, operands=operands, successors=_ods_successors, regions=regions, loc=loc, ip=ip))

  @builtins.property
  def res(self):
    return self.operation.results[0]

def read_ptx_sreg_tid_x(res, *, loc=None, ip=None) -> _ods_ir.Value:
  return _get_op_result_or_op_results(ThreadIdXOp(res=res, loc=loc, ip=ip))

@_ods_cext.register_operation(_Dialect)
class ThreadIdYOp(_ods_ir.OpView):
  OPERATION_NAME = "nvvm.read.ptx.sreg.tid.y"

  _ODS_REGIONS = (0, True)

  def __init__(self, res, *, loc=None, ip=None):
    operands = []
    results = []
    attributes = {}
    regions = None
    _ods_context = _ods_get_default_loc_context(loc)
    results.append(res)
    _ods_successors = None
    super().__init__(self.build_generic(attributes=attributes, results=results, operands=operands, successors=_ods_successors, regions=regions, loc=loc, ip=ip))

  @builtins.property
  def res(self):
    return self.operation.results[0]

def read_ptx_sreg_tid_y(res, *, loc=None, ip=None) -> _ods_ir.Value:
  return _get_op_result_or_op_results(ThreadIdYOp(res=res, loc=loc, ip=ip))

@_ods_cext.register_operation(_Dialect)
class ThreadIdZOp(_ods_ir.OpView):
  OPERATION_NAME = "nvvm.read.ptx.sreg.tid.z"

  _ODS_REGIONS = (0, True)

  def __init__(self, res, *, loc=None, ip=None):
    operands = []
    results = []
    attributes = {}
    regions = None
    _ods_context = _ods_get_default_loc_context(loc)
    results.append(res)
    _ods_successors = None
    super().__init__(self.build_generic(attributes=attributes, results=results, operands=operands, successors=_ods_successors, regions=regions, loc=loc, ip=ip))

  @builtins.property
  def res(self):
    return self.operation.results[0]

def read_ptx_sreg_tid_z(res, *, loc=None, ip=None) -> _ods_ir.Value:
  return _get_op_result_or_op_results(ThreadIdZOp(res=res, loc=loc, ip=ip))

@_ods_cext.register_operation(_Dialect)
class VoteBallotOp(_ods_ir.OpView):
  OPERATION_NAME = "nvvm.vote.ballot.sync"

  _ODS_REGIONS = (0, True)

  def __init__(self, res, mask, pred, *, loc=None, ip=None):
    operands = []
    results = []
    attributes = {}
    regions = None
    operands.append(_get_op_result_or_value(mask))
    operands.append(_get_op_result_or_value(pred))
    _ods_context = _ods_get_default_loc_context(loc)
    results.append(res)
    _ods_successors = None
    super().__init__(self.build_generic(attributes=attributes, results=results, operands=operands, successors=_ods_successors, regions=regions, loc=loc, ip=ip))

  @builtins.property
  def mask(self):
    return self.operation.operands[0]

  @builtins.property
  def pred(self):
    return self.operation.operands[1]

  @builtins.property
  def res(self):
    return self.operation.results[0]

def vote_ballot_sync(res, mask, pred, *, loc=None, ip=None) -> _ods_ir.Value:
  return _get_op_result_or_op_results(VoteBallotOp(res=res, mask=mask, pred=pred, loc=loc, ip=ip))

@_ods_cext.register_operation(_Dialect)
class WMMALoadOp(_ods_ir.OpView):
  OPERATION_NAME = "nvvm.wmma.load"

  _ODS_REGIONS = (0, True)

  def __init__(self, res, ptr, stride, m, n, k, layout, eltype, frag, *, loc=None, ip=None):
    operands = []
    results = []
    attributes = {}
    regions = None
    operands.append(_get_op_result_or_value(ptr))
    operands.append(_get_op_result_or_value(stride))
    _ods_context = _ods_get_default_loc_context(loc)
    attributes["m"] = (m if (
    isinstance(m, _ods_ir.Attribute) or
    not _ods_ir.AttrBuilder.contains('I32Attr')) else
      _ods_ir.AttrBuilder.get('I32Attr')(m, context=_ods_context))
    attributes["n"] = (n if (
    isinstance(n, _ods_ir.Attribute) or
    not _ods_ir.AttrBuilder.contains('I32Attr')) else
      _ods_ir.AttrBuilder.get('I32Attr')(n, context=_ods_context))
    attributes["k"] = (k if (
    isinstance(k, _ods_ir.Attribute) or
    not _ods_ir.AttrBuilder.contains('I32Attr')) else
      _ods_ir.AttrBuilder.get('I32Attr')(k, context=_ods_context))
    attributes["layout"] = (layout if (
    isinstance(layout, _ods_ir.Attribute) or
    not _ods_ir.AttrBuilder.contains('MMALayoutAttr')) else
      _ods_ir.AttrBuilder.get('MMALayoutAttr')(layout, context=_ods_context))
    attributes["eltype"] = (eltype if (
    isinstance(eltype, _ods_ir.Attribute) or
    not _ods_ir.AttrBuilder.contains('MMATypesAttr')) else
      _ods_ir.AttrBuilder.get('MMATypesAttr')(eltype, context=_ods_context))
    attributes["frag"] = (frag if (
    isinstance(frag, _ods_ir.Attribute) or
    not _ods_ir.AttrBuilder.contains('MMAFragAttr')) else
      _ods_ir.AttrBuilder.get('MMAFragAttr')(frag, context=_ods_context))
    results.append(res)
    _ods_successors = None
    super().__init__(self.build_generic(attributes=attributes, results=results, operands=operands, successors=_ods_successors, regions=regions, loc=loc, ip=ip))

  @builtins.property
  def ptr(self):
    return self.operation.operands[0]

  @builtins.property
  def stride(self):
    return self.operation.operands[1]

  @builtins.property
  def m(self):
    return self.operation.attributes["m"]

  @m.setter
  def m(self, value):
    if value is None:
      raise ValueError("'None' not allowed as value for mandatory attributes")
    self.operation.attributes["m"] = value

  @builtins.property
  def n(self):
    return self.operation.attributes["n"]

  @n.setter
  def n(self, value):
    if value is None:
      raise ValueError("'None' not allowed as value for mandatory attributes")
    self.operation.attributes["n"] = value

  @builtins.property
  def k(self):
    return self.operation.attributes["k"]

  @k.setter
  def k(self, value):
    if value is None:
      raise ValueError("'None' not allowed as value for mandatory attributes")
    self.operation.attributes["k"] = value

  @builtins.property
  def layout(self):
    return self.operation.attributes["layout"]

  @layout.setter
  def layout(self, value):
    if value is None:
      raise ValueError("'None' not allowed as value for mandatory attributes")
    self.operation.attributes["layout"] = value

  @builtins.property
  def eltype(self):
    return self.operation.attributes["eltype"]

  @eltype.setter
  def eltype(self, value):
    if value is None:
      raise ValueError("'None' not allowed as value for mandatory attributes")
    self.operation.attributes["eltype"] = value

  @builtins.property
  def frag(self):
    return self.operation.attributes["frag"]

  @frag.setter
  def frag(self, value):
    if value is None:
      raise ValueError("'None' not allowed as value for mandatory attributes")
    self.operation.attributes["frag"] = value

  @builtins.property
  def res(self):
    return self.operation.results[0]

def wmma_load(res, ptr, stride, m, n, k, layout, eltype, frag, *, loc=None, ip=None) -> _ods_ir.Value:
  return _get_op_result_or_op_results(WMMALoadOp(res=res, ptr=ptr, stride=stride, m=m, n=n, k=k, layout=layout, eltype=eltype, frag=frag, loc=loc, ip=ip))

@_ods_cext.register_operation(_Dialect)
class WMMAMmaOp(_ods_ir.OpView):
  OPERATION_NAME = "nvvm.wmma.mma"

  _ODS_REGIONS = (0, True)

  def __init__(self, res, m, n, k, layoutA, layoutB, eltypeA, eltypeB, args, *, loc=None, ip=None):
    operands = []
    results = []
    attributes = {}
    regions = None
    operands.extend(_get_op_results_or_values(args))
    _ods_context = _ods_get_default_loc_context(loc)
    attributes["m"] = (m if (
    isinstance(m, _ods_ir.Attribute) or
    not _ods_ir.AttrBuilder.contains('I32Attr')) else
      _ods_ir.AttrBuilder.get('I32Attr')(m, context=_ods_context))
    attributes["n"] = (n if (
    isinstance(n, _ods_ir.Attribute) or
    not _ods_ir.AttrBuilder.contains('I32Attr')) else
      _ods_ir.AttrBuilder.get('I32Attr')(n, context=_ods_context))
    attributes["k"] = (k if (
    isinstance(k, _ods_ir.Attribute) or
    not _ods_ir.AttrBuilder.contains('I32Attr')) else
      _ods_ir.AttrBuilder.get('I32Attr')(k, context=_ods_context))
    attributes["layoutA"] = (layoutA if (
    isinstance(layoutA, _ods_ir.Attribute) or
    not _ods_ir.AttrBuilder.contains('MMALayoutAttr')) else
      _ods_ir.AttrBuilder.get('MMALayoutAttr')(layoutA, context=_ods_context))
    attributes["layoutB"] = (layoutB if (
    isinstance(layoutB, _ods_ir.Attribute) or
    not _ods_ir.AttrBuilder.contains('MMALayoutAttr')) else
      _ods_ir.AttrBuilder.get('MMALayoutAttr')(layoutB, context=_ods_context))
    attributes["eltypeA"] = (eltypeA if (
    isinstance(eltypeA, _ods_ir.Attribute) or
    not _ods_ir.AttrBuilder.contains('MMATypesAttr')) else
      _ods_ir.AttrBuilder.get('MMATypesAttr')(eltypeA, context=_ods_context))
    attributes["eltypeB"] = (eltypeB if (
    isinstance(eltypeB, _ods_ir.Attribute) or
    not _ods_ir.AttrBuilder.contains('MMATypesAttr')) else
      _ods_ir.AttrBuilder.get('MMATypesAttr')(eltypeB, context=_ods_context))
    results.append(res)
    _ods_successors = None
    super().__init__(self.build_generic(attributes=attributes, results=results, operands=operands, successors=_ods_successors, regions=regions, loc=loc, ip=ip))

  @builtins.property
  def args(self):
    _ods_variadic_group_length = len(self.operation.operands) - 1 + 1
    return self.operation.operands[0:0 + _ods_variadic_group_length]

  @builtins.property
  def m(self):
    return self.operation.attributes["m"]

  @m.setter
  def m(self, value):
    if value is None:
      raise ValueError("'None' not allowed as value for mandatory attributes")
    self.operation.attributes["m"] = value

  @builtins.property
  def n(self):
    return self.operation.attributes["n"]

  @n.setter
  def n(self, value):
    if value is None:
      raise ValueError("'None' not allowed as value for mandatory attributes")
    self.operation.attributes["n"] = value

  @builtins.property
  def k(self):
    return self.operation.attributes["k"]

  @k.setter
  def k(self, value):
    if value is None:
      raise ValueError("'None' not allowed as value for mandatory attributes")
    self.operation.attributes["k"] = value

  @builtins.property
  def layoutA(self):
    return self.operation.attributes["layoutA"]

  @layoutA.setter
  def layoutA(self, value):
    if value is None:
      raise ValueError("'None' not allowed as value for mandatory attributes")
    self.operation.attributes["layoutA"] = value

  @builtins.property
  def layoutB(self):
    return self.operation.attributes["layoutB"]

  @layoutB.setter
  def layoutB(self, value):
    if value is None:
      raise ValueError("'None' not allowed as value for mandatory attributes")
    self.operation.attributes["layoutB"] = value

  @builtins.property
  def eltypeA(self):
    return self.operation.attributes["eltypeA"]

  @eltypeA.setter
  def eltypeA(self, value):
    if value is None:
      raise ValueError("'None' not allowed as value for mandatory attributes")
    self.operation.attributes["eltypeA"] = value

  @builtins.property
  def eltypeB(self):
    return self.operation.attributes["eltypeB"]

  @eltypeB.setter
  def eltypeB(self, value):
    if value is None:
      raise ValueError("'None' not allowed as value for mandatory attributes")
    self.operation.attributes["eltypeB"] = value

  @builtins.property
  def res(self):
    return self.operation.results[0]

def wmma_mma(res, m, n, k, layout_a, layout_b, eltype_a, eltype_b, args, *, loc=None, ip=None) -> _ods_ir.Value:
  return _get_op_result_or_op_results(WMMAMmaOp(res=res, m=m, n=n, k=k, layoutA=layout_a, layoutB=layout_b, eltypeA=eltype_a, eltypeB=eltype_b, args=args, loc=loc, ip=ip))

@_ods_cext.register_operation(_Dialect)
class WMMAStoreOp(_ods_ir.OpView):
  OPERATION_NAME = "nvvm.wmma.store"

  _ODS_REGIONS = (0, True)

  def __init__(self, ptr, m, n, k, layout, eltype, args, stride, *, loc=None, ip=None):
    operands = []
    results = []
    attributes = {}
    regions = None
    operands.append(_get_op_result_or_value(ptr))
    operands.extend(_get_op_results_or_values(args))
    operands.append(_get_op_result_or_value(stride))
    _ods_context = _ods_get_default_loc_context(loc)
    attributes["m"] = (m if (
    isinstance(m, _ods_ir.Attribute) or
    not _ods_ir.AttrBuilder.contains('I32Attr')) else
      _ods_ir.AttrBuilder.get('I32Attr')(m, context=_ods_context))
    attributes["n"] = (n if (
    isinstance(n, _ods_ir.Attribute) or
    not _ods_ir.AttrBuilder.contains('I32Attr')) else
      _ods_ir.AttrBuilder.get('I32Attr')(n, context=_ods_context))
    attributes["k"] = (k if (
    isinstance(k, _ods_ir.Attribute) or
    not _ods_ir.AttrBuilder.contains('I32Attr')) else
      _ods_ir.AttrBuilder.get('I32Attr')(k, context=_ods_context))
    attributes["layout"] = (layout if (
    isinstance(layout, _ods_ir.Attribute) or
    not _ods_ir.AttrBuilder.contains('MMALayoutAttr')) else
      _ods_ir.AttrBuilder.get('MMALayoutAttr')(layout, context=_ods_context))
    attributes["eltype"] = (eltype if (
    isinstance(eltype, _ods_ir.Attribute) or
    not _ods_ir.AttrBuilder.contains('MMATypesAttr')) else
      _ods_ir.AttrBuilder.get('MMATypesAttr')(eltype, context=_ods_context))
    _ods_successors = None
    super().__init__(self.build_generic(attributes=attributes, results=results, operands=operands, successors=_ods_successors, regions=regions, loc=loc, ip=ip))

  @builtins.property
  def ptr(self):
    return self.operation.operands[0]

  @builtins.property
  def args(self):
    _ods_variadic_group_length = len(self.operation.operands) - 3 + 1
    return self.operation.operands[1:1 + _ods_variadic_group_length]

  @builtins.property
  def stride(self):
    _ods_variadic_group_length = len(self.operation.operands) - 3 + 1
    return self.operation.operands[2 + _ods_variadic_group_length - 1]

  @builtins.property
  def m(self):
    return self.operation.attributes["m"]

  @m.setter
  def m(self, value):
    if value is None:
      raise ValueError("'None' not allowed as value for mandatory attributes")
    self.operation.attributes["m"] = value

  @builtins.property
  def n(self):
    return self.operation.attributes["n"]

  @n.setter
  def n(self, value):
    if value is None:
      raise ValueError("'None' not allowed as value for mandatory attributes")
    self.operation.attributes["n"] = value

  @builtins.property
  def k(self):
    return self.operation.attributes["k"]

  @k.setter
  def k(self, value):
    if value is None:
      raise ValueError("'None' not allowed as value for mandatory attributes")
    self.operation.attributes["k"] = value

  @builtins.property
  def layout(self):
    return self.operation.attributes["layout"]

  @layout.setter
  def layout(self, value):
    if value is None:
      raise ValueError("'None' not allowed as value for mandatory attributes")
    self.operation.attributes["layout"] = value

  @builtins.property
  def eltype(self):
    return self.operation.attributes["eltype"]

  @eltype.setter
  def eltype(self, value):
    if value is None:
      raise ValueError("'None' not allowed as value for mandatory attributes")
    self.operation.attributes["eltype"] = value

def wmma_store(ptr, m, n, k, layout, eltype, args, stride, *, loc=None, ip=None) -> _ods_ir.Operation:
  return _get_op_result_or_op_results(WMMAStoreOp(ptr=ptr, m=m, n=n, k=k, layout=layout, eltype=eltype, args=args, stride=stride, loc=loc, ip=ip))

@_ods_cext.register_operation(_Dialect)
class WarpSizeOp(_ods_ir.OpView):
  OPERATION_NAME = "nvvm.read.ptx.sreg.warpsize"

  _ODS_REGIONS = (0, True)

  def __init__(self, res, *, loc=None, ip=None):
    operands = []
    results = []
    attributes = {}
    regions = None
    _ods_context = _ods_get_default_loc_context(loc)
    results.append(res)
    _ods_successors = None
    super().__init__(self.build_generic(attributes=attributes, results=results, operands=operands, successors=_ods_successors, regions=regions, loc=loc, ip=ip))

  @builtins.property
  def res(self):
    return self.operation.results[0]

def read_ptx_sreg_warpsize(res, *, loc=None, ip=None) -> _ods_ir.Value:
  return _get_op_result_or_op_results(WarpSizeOp(res=res, loc=loc, ip=ip))

@_ods_cext.register_operation(_Dialect)
class WgmmaFenceAlignedOp(_ods_ir.OpView):
  OPERATION_NAME = "nvvm.wgmma.fence.aligned"

  _ODS_REGIONS = (0, True)

  def __init__(self, *, loc=None, ip=None):
    operands = []
    results = []
    attributes = {}
    regions = None
    _ods_context = _ods_get_default_loc_context(loc)
    _ods_successors = None
    super().__init__(self.build_generic(attributes=attributes, results=results, operands=operands, successors=_ods_successors, regions=regions, loc=loc, ip=ip))

def wgmma_fence_aligned(*, loc=None, ip=None) -> _ods_ir.Operation:
  return _get_op_result_or_op_results(WgmmaFenceAlignedOp(loc=loc, ip=ip))

@_ods_cext.register_operation(_Dialect)
class WgmmaGroupSyncAlignedOp(_ods_ir.OpView):
  OPERATION_NAME = "nvvm.wgmma.commit.group.sync.aligned"

  _ODS_REGIONS = (0, True)

  def __init__(self, *, loc=None, ip=None):
    operands = []
    results = []
    attributes = {}
    regions = None
    _ods_context = _ods_get_default_loc_context(loc)
    _ods_successors = None
    super().__init__(self.build_generic(attributes=attributes, results=results, operands=operands, successors=_ods_successors, regions=regions, loc=loc, ip=ip))

def wgmma_commit_group_sync_aligned(*, loc=None, ip=None) -> _ods_ir.Operation:
  return _get_op_result_or_op_results(WgmmaGroupSyncAlignedOp(loc=loc, ip=ip))

@_ods_cext.register_operation(_Dialect)
class WgmmaMmaAsyncOp(_ods_ir.OpView):
  OPERATION_NAME = "nvvm.wgmma.mma_async"

  _ODS_REGIONS = (0, True)

  def __init__(self, results_, inouts, descriptorA, descriptorB, shape, typeA, typeB, typeD, scaleD, scaleA, scaleB, layoutA, layoutB, *, satfinite=None, loc=None, ip=None):
    operands = []
    results = []
    attributes = {}
    regions = None
    operands.append(_get_op_result_or_value(inouts))
    operands.append(_get_op_result_or_value(descriptorA))
    operands.append(_get_op_result_or_value(descriptorB))
    _ods_context = _ods_get_default_loc_context(loc)
    attributes["shape"] = (shape if (
    isinstance(shape, _ods_ir.Attribute) or
    not _ods_ir.AttrBuilder.contains('NVVM_MMAShapeAttr')) else
      _ods_ir.AttrBuilder.get('NVVM_MMAShapeAttr')(shape, context=_ods_context))
    attributes["typeA"] = (typeA if (
    isinstance(typeA, _ods_ir.Attribute) or
    not _ods_ir.AttrBuilder.contains('WGMMATypesAttr')) else
      _ods_ir.AttrBuilder.get('WGMMATypesAttr')(typeA, context=_ods_context))
    attributes["typeB"] = (typeB if (
    isinstance(typeB, _ods_ir.Attribute) or
    not _ods_ir.AttrBuilder.contains('WGMMATypesAttr')) else
      _ods_ir.AttrBuilder.get('WGMMATypesAttr')(typeB, context=_ods_context))
    attributes["typeD"] = (typeD if (
    isinstance(typeD, _ods_ir.Attribute) or
    not _ods_ir.AttrBuilder.contains('WGMMATypesAttr')) else
      _ods_ir.AttrBuilder.get('WGMMATypesAttr')(typeD, context=_ods_context))
    attributes["scaleD"] = (scaleD if (
    isinstance(scaleD, _ods_ir.Attribute) or
    not _ods_ir.AttrBuilder.contains('WGMMAScaleOutAttr')) else
      _ods_ir.AttrBuilder.get('WGMMAScaleOutAttr')(scaleD, context=_ods_context))
    attributes["scaleA"] = (scaleA if (
    isinstance(scaleA, _ods_ir.Attribute) or
    not _ods_ir.AttrBuilder.contains('WGMMAScaleInAttr')) else
      _ods_ir.AttrBuilder.get('WGMMAScaleInAttr')(scaleA, context=_ods_context))
    attributes["scaleB"] = (scaleB if (
    isinstance(scaleB, _ods_ir.Attribute) or
    not _ods_ir.AttrBuilder.contains('WGMMAScaleInAttr')) else
      _ods_ir.AttrBuilder.get('WGMMAScaleInAttr')(scaleB, context=_ods_context))
    attributes["layoutA"] = (layoutA if (
    isinstance(layoutA, _ods_ir.Attribute) or
    not _ods_ir.AttrBuilder.contains('MMALayoutAttr')) else
      _ods_ir.AttrBuilder.get('MMALayoutAttr')(layoutA, context=_ods_context))
    attributes["layoutB"] = (layoutB if (
    isinstance(layoutB, _ods_ir.Attribute) or
    not _ods_ir.AttrBuilder.contains('MMALayoutAttr')) else
      _ods_ir.AttrBuilder.get('MMALayoutAttr')(layoutB, context=_ods_context))
    if satfinite is not None: attributes["satfinite"] = (satfinite if (
        isinstance(satfinite, _ods_ir.Attribute) or
        not _ods_ir.AttrBuilder.contains('MMAIntOverflowAttr')) else
          _ods_ir.AttrBuilder.get('MMAIntOverflowAttr')(satfinite, context=_ods_context))
    results.append(results_)
    _ods_successors = None
    super().__init__(self.build_generic(attributes=attributes, results=results, operands=operands, successors=_ods_successors, regions=regions, loc=loc, ip=ip))

  @builtins.property
  def inouts(self):
    return self.operation.operands[0]

  @builtins.property
  def descriptorA(self):
    return self.operation.operands[1]

  @builtins.property
  def descriptorB(self):
    return self.operation.operands[2]

  @builtins.property
  def shape(self):
    return self.operation.attributes["shape"]

  @shape.setter
  def shape(self, value):
    if value is None:
      raise ValueError("'None' not allowed as value for mandatory attributes")
    self.operation.attributes["shape"] = value

  @builtins.property
  def typeA(self):
    return self.operation.attributes["typeA"]

  @typeA.setter
  def typeA(self, value):
    if value is None:
      raise ValueError("'None' not allowed as value for mandatory attributes")
    self.operation.attributes["typeA"] = value

  @builtins.property
  def typeB(self):
    return self.operation.attributes["typeB"]

  @typeB.setter
  def typeB(self, value):
    if value is None:
      raise ValueError("'None' not allowed as value for mandatory attributes")
    self.operation.attributes["typeB"] = value

  @builtins.property
  def typeD(self):
    return self.operation.attributes["typeD"]

  @typeD.setter
  def typeD(self, value):
    if value is None:
      raise ValueError("'None' not allowed as value for mandatory attributes")
    self.operation.attributes["typeD"] = value

  @builtins.property
  def scaleD(self):
    return self.operation.attributes["scaleD"]

  @scaleD.setter
  def scaleD(self, value):
    if value is None:
      raise ValueError("'None' not allowed as value for mandatory attributes")
    self.operation.attributes["scaleD"] = value

  @builtins.property
  def scaleA(self):
    return self.operation.attributes["scaleA"]

  @scaleA.setter
  def scaleA(self, value):
    if value is None:
      raise ValueError("'None' not allowed as value for mandatory attributes")
    self.operation.attributes["scaleA"] = value

  @builtins.property
  def scaleB(self):
    return self.operation.attributes["scaleB"]

  @scaleB.setter
  def scaleB(self, value):
    if value is None:
      raise ValueError("'None' not allowed as value for mandatory attributes")
    self.operation.attributes["scaleB"] = value

  @builtins.property
  def layoutA(self):
    return self.operation.attributes["layoutA"]

  @layoutA.setter
  def layoutA(self, value):
    if value is None:
      raise ValueError("'None' not allowed as value for mandatory attributes")
    self.operation.attributes["layoutA"] = value

  @builtins.property
  def layoutB(self):
    return self.operation.attributes["layoutB"]

  @layoutB.setter
  def layoutB(self, value):
    if value is None:
      raise ValueError("'None' not allowed as value for mandatory attributes")
    self.operation.attributes["layoutB"] = value

  @builtins.property
  def satfinite(self):
    if "satfinite" not in self.operation.attributes:
      return None
    return self.operation.attributes["satfinite"]

  @satfinite.setter
  def satfinite(self, value):
    if value is not None:
      self.operation.attributes["satfinite"] = value
    elif "satfinite" in self.operation.attributes:
      del self.operation.attributes["satfinite"]

  @satfinite.deleter
  def satfinite(self):
    del self.operation.attributes["satfinite"]

  @builtins.property
  def results_(self):
    return self.operation.results[0]

def wgmma_mma_async(results_, inouts, descriptor_a, descriptor_b, shape, type_a, type_b, type_d, scale_d, scale_a, scale_b, layout_a, layout_b, *, satfinite=None, loc=None, ip=None) -> _ods_ir.Value:
  return _get_op_result_or_op_results(WgmmaMmaAsyncOp(results_=results_, inouts=inouts, descriptorA=descriptor_a, descriptorB=descriptor_b, shape=shape, typeA=type_a, typeB=type_b, typeD=type_d, scaleD=scale_d, scaleA=scale_a, scaleB=scale_b, layoutA=layout_a, layoutB=layout_b, satfinite=satfinite, loc=loc, ip=ip))

@_ods_cext.register_operation(_Dialect)
class WgmmaWaitGroupSyncOp(_ods_ir.OpView):
  OPERATION_NAME = "nvvm.wgmma.wait.group.sync.aligned"

  _ODS_REGIONS = (0, True)

  def __init__(self, group, *, loc=None, ip=None):
    operands = []
    results = []
    attributes = {}
    regions = None
    _ods_context = _ods_get_default_loc_context(loc)
    attributes["group"] = (group if (
    isinstance(group, _ods_ir.Attribute) or
    not _ods_ir.AttrBuilder.contains('I32Attr')) else
      _ods_ir.AttrBuilder.get('I32Attr')(group, context=_ods_context))
    _ods_successors = None
    super().__init__(self.build_generic(attributes=attributes, results=results, operands=operands, successors=_ods_successors, regions=regions, loc=loc, ip=ip))

  @builtins.property
  def group(self):
    return self.operation.attributes["group"]

  @group.setter
  def group(self, value):
    if value is None:
      raise ValueError("'None' not allowed as value for mandatory attributes")
    self.operation.attributes["group"] = value

def wgmma_wait_group_sync_aligned(group, *, loc=None, ip=None) -> _ods_ir.Operation:
  return _get_op_result_or_op_results(WgmmaWaitGroupSyncOp(group=group, loc=loc, ip=ip))