# 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_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.OPERATION_NAME, self._ODS_REGIONS, self._ODS_OPERAND_SEGMENTS, self._ODS_RESULT_SEGMENTS, 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 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(barrierId)
    operands.append(numberOfThreads)
    _ods_context = _ods_get_default_loc_context(loc)
    _ods_successors = None
    super().__init__(self.OPERATION_NAME, self._ODS_REGIONS, self._ODS_OPERAND_SEGMENTS, self._ODS_RESULT_SEGMENTS, 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 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(barrierId)
    operands.append(numberOfThreads)
    _ods_context = _ods_get_default_loc_context(loc)
    _ods_successors = None
    super().__init__(self.OPERATION_NAME, self._ODS_REGIONS, self._ODS_OPERAND_SEGMENTS, self._ODS_RESULT_SEGMENTS, 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 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, *, range=None, loc=None, ip=None):
    operands = []
    results = []
    attributes = {}
    regions = None
    _ods_context = _ods_get_default_loc_context(loc)
    if range is not None: attributes["range"] = (range if (
        isinstance(range, _ods_ir.Attribute) or
        not _ods_ir.AttrBuilder.contains('LLVM_ConstantRangeAttr')) else
          _ods_ir.AttrBuilder.get('LLVM_ConstantRangeAttr')(range, context=_ods_context))
    results.append(res)
    _ods_successors = None
    super().__init__(self.OPERATION_NAME, self._ODS_REGIONS, self._ODS_OPERAND_SEGMENTS, self._ODS_RESULT_SEGMENTS, attributes=attributes, results=results, operands=operands, successors=_ods_successors, regions=regions, loc=loc, ip=ip)

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

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

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

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

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

@_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, *, range=None, loc=None, ip=None):
    operands = []
    results = []
    attributes = {}
    regions = None
    _ods_context = _ods_get_default_loc_context(loc)
    if range is not None: attributes["range"] = (range if (
        isinstance(range, _ods_ir.Attribute) or
        not _ods_ir.AttrBuilder.contains('LLVM_ConstantRangeAttr')) else
          _ods_ir.AttrBuilder.get('LLVM_ConstantRangeAttr')(range, context=_ods_context))
    results.append(res)
    _ods_successors = None
    super().__init__(self.OPERATION_NAME, self._ODS_REGIONS, self._ODS_OPERAND_SEGMENTS, self._ODS_RESULT_SEGMENTS, attributes=attributes, results=results, operands=operands, successors=_ods_successors, regions=regions, loc=loc, ip=ip)

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

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

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

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

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

@_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, *, range=None, loc=None, ip=None):
    operands = []
    results = []
    attributes = {}
    regions = None
    _ods_context = _ods_get_default_loc_context(loc)
    if range is not None: attributes["range"] = (range if (
        isinstance(range, _ods_ir.Attribute) or
        not _ods_ir.AttrBuilder.contains('LLVM_ConstantRangeAttr')) else
          _ods_ir.AttrBuilder.get('LLVM_ConstantRangeAttr')(range, context=_ods_context))
    results.append(res)
    _ods_successors = None
    super().__init__(self.OPERATION_NAME, self._ODS_REGIONS, self._ODS_OPERAND_SEGMENTS, self._ODS_RESULT_SEGMENTS, attributes=attributes, results=results, operands=operands, successors=_ods_successors, regions=regions, loc=loc, ip=ip)

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

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

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

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

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

@_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, *, range=None, loc=None, ip=None):
    operands = []
    results = []
    attributes = {}
    regions = None
    _ods_context = _ods_get_default_loc_context(loc)
    if range is not None: attributes["range"] = (range if (
        isinstance(range, _ods_ir.Attribute) or
        not _ods_ir.AttrBuilder.contains('LLVM_ConstantRangeAttr')) else
          _ods_ir.AttrBuilder.get('LLVM_ConstantRangeAttr')(range, context=_ods_context))
    results.append(res)
    _ods_successors = None
    super().__init__(self.OPERATION_NAME, self._ODS_REGIONS, self._ODS_OPERAND_SEGMENTS, self._ODS_RESULT_SEGMENTS, attributes=attributes, results=results, operands=operands, successors=_ods_successors, regions=regions, loc=loc, ip=ip)

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

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

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

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

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

@_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, *, range=None, loc=None, ip=None):
    operands = []
    results = []
    attributes = {}
    regions = None
    _ods_context = _ods_get_default_loc_context(loc)
    if range is not None: attributes["range"] = (range if (
        isinstance(range, _ods_ir.Attribute) or
        not _ods_ir.AttrBuilder.contains('LLVM_ConstantRangeAttr')) else
          _ods_ir.AttrBuilder.get('LLVM_ConstantRangeAttr')(range, context=_ods_context))
    results.append(res)
    _ods_successors = None
    super().__init__(self.OPERATION_NAME, self._ODS_REGIONS, self._ODS_OPERAND_SEGMENTS, self._ODS_RESULT_SEGMENTS, attributes=attributes, results=results, operands=operands, successors=_ods_successors, regions=regions, loc=loc, ip=ip)

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

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

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

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

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

@_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, *, range=None, loc=None, ip=None):
    operands = []
    results = []
    attributes = {}
    regions = None
    _ods_context = _ods_get_default_loc_context(loc)
    if range is not None: attributes["range"] = (range if (
        isinstance(range, _ods_ir.Attribute) or
        not _ods_ir.AttrBuilder.contains('LLVM_ConstantRangeAttr')) else
          _ods_ir.AttrBuilder.get('LLVM_ConstantRangeAttr')(range, context=_ods_context))
    results.append(res)
    _ods_successors = None
    super().__init__(self.OPERATION_NAME, self._ODS_REGIONS, self._ODS_OPERAND_SEGMENTS, self._ODS_RESULT_SEGMENTS, attributes=attributes, results=results, operands=operands, successors=_ods_successors, regions=regions, loc=loc, ip=ip)

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

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

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

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

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

@_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, *, range=None, loc=None, ip=None):
    operands = []
    results = []
    attributes = {}
    regions = None
    _ods_context = _ods_get_default_loc_context(loc)
    if range is not None: attributes["range"] = (range if (
        isinstance(range, _ods_ir.Attribute) or
        not _ods_ir.AttrBuilder.contains('LLVM_ConstantRangeAttr')) else
          _ods_ir.AttrBuilder.get('LLVM_ConstantRangeAttr')(range, context=_ods_context))
    results.append(res)
    _ods_successors = None
    super().__init__(self.OPERATION_NAME, self._ODS_REGIONS, self._ODS_OPERAND_SEGMENTS, self._ODS_RESULT_SEGMENTS, attributes=attributes, results=results, operands=operands, successors=_ods_successors, regions=regions, loc=loc, ip=ip)

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

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

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

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

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

@_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, *, range=None, loc=None, ip=None):
    operands = []
    results = []
    attributes = {}
    regions = None
    _ods_context = _ods_get_default_loc_context(loc)
    if range is not None: attributes["range"] = (range if (
        isinstance(range, _ods_ir.Attribute) or
        not _ods_ir.AttrBuilder.contains('LLVM_ConstantRangeAttr')) else
          _ods_ir.AttrBuilder.get('LLVM_ConstantRangeAttr')(range, context=_ods_context))
    results.append(res)
    _ods_successors = None
    super().__init__(self.OPERATION_NAME, self._ODS_REGIONS, self._ODS_OPERAND_SEGMENTS, self._ODS_RESULT_SEGMENTS, attributes=attributes, results=results, operands=operands, successors=_ods_successors, regions=regions, loc=loc, ip=ip)

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

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

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

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

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

@_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, *, range=None, loc=None, ip=None):
    operands = []
    results = []
    attributes = {}
    regions = None
    _ods_context = _ods_get_default_loc_context(loc)
    if range is not None: attributes["range"] = (range if (
        isinstance(range, _ods_ir.Attribute) or
        not _ods_ir.AttrBuilder.contains('LLVM_ConstantRangeAttr')) else
          _ods_ir.AttrBuilder.get('LLVM_ConstantRangeAttr')(range, context=_ods_context))
    results.append(res)
    _ods_successors = None
    super().__init__(self.OPERATION_NAME, self._ODS_REGIONS, self._ODS_OPERAND_SEGMENTS, self._ODS_RESULT_SEGMENTS, attributes=attributes, results=results, operands=operands, successors=_ods_successors, regions=regions, loc=loc, ip=ip)

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

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

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

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

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

@_ods_cext.register_operation(_Dialect)
class Breakpoint(_ods_ir.OpView):
  OPERATION_NAME = "nvvm.breakpoint"

  _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.OPERATION_NAME, self._ODS_REGIONS, self._ODS_OPERAND_SEGMENTS, self._ODS_RESULT_SEGMENTS, attributes=attributes, results=results, operands=operands, successors=_ods_successors, regions=regions, loc=loc, ip=ip)

def breakpoint(*, loc=None, ip=None) -> _ods_ir.Operation:
  return Breakpoint(loc=loc, ip=ip)

@_ods_cext.register_operation(_Dialect)
class BulkStoreOp(_ods_ir.OpView):
  OPERATION_NAME = "nvvm.st.bulk"

  _ODS_REGIONS = (0, True)

  def __init__(self, addr, size, *, initVal=None, loc=None, ip=None):
    operands = []
    results = []
    attributes = {}
    regions = None
    operands.append(addr)
    operands.append(size)
    _ods_context = _ods_get_default_loc_context(loc)
    if initVal is not None: attributes["initVal"] = (initVal if (
        isinstance(initVal, _ods_ir.Attribute) or
        not _ods_ir.AttrBuilder.contains('I64Attr')) else
          _ods_ir.AttrBuilder.get('I64Attr')(initVal, context=_ods_context))
    _ods_successors = None
    super().__init__(self.OPERATION_NAME, self._ODS_REGIONS, self._ODS_OPERAND_SEGMENTS, self._ODS_RESULT_SEGMENTS, 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 size(self):
    return self.operation.operands[1]

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

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

def st_bulk(addr, size, *, init_val=None, loc=None, ip=None) -> _ods_ir.Operation:
  return BulkStoreOp(addr=addr, size=size, initVal=init_val, 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.OPERATION_NAME, self._ODS_REGIONS, self._ODS_OPERAND_SEGMENTS, self._ODS_RESULT_SEGMENTS, 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 Clock64Op(res=res, loc=loc, ip=ip).result

@_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.OPERATION_NAME, self._ODS_REGIONS, self._ODS_OPERAND_SEGMENTS, self._ODS_RESULT_SEGMENTS, 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 ClockOp(res=res, loc=loc, ip=ip).result

@_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.OPERATION_NAME, self._ODS_REGIONS, self._ODS_OPERAND_SEGMENTS, self._ODS_RESULT_SEGMENTS, 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 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.OPERATION_NAME, self._ODS_REGIONS, self._ODS_OPERAND_SEGMENTS, self._ODS_RESULT_SEGMENTS, 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 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, *, range=None, loc=None, ip=None):
    operands = []
    results = []
    attributes = {}
    regions = None
    _ods_context = _ods_get_default_loc_context(loc)
    if range is not None: attributes["range"] = (range if (
        isinstance(range, _ods_ir.Attribute) or
        not _ods_ir.AttrBuilder.contains('LLVM_ConstantRangeAttr')) else
          _ods_ir.AttrBuilder.get('LLVM_ConstantRangeAttr')(range, context=_ods_context))
    results.append(res)
    _ods_successors = None
    super().__init__(self.OPERATION_NAME, self._ODS_REGIONS, self._ODS_OPERAND_SEGMENTS, self._ODS_RESULT_SEGMENTS, attributes=attributes, results=results, operands=operands, successors=_ods_successors, regions=regions, loc=loc, ip=ip)

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

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

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

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

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

@_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, *, range=None, loc=None, ip=None):
    operands = []
    results = []
    attributes = {}
    regions = None
    _ods_context = _ods_get_default_loc_context(loc)
    if range is not None: attributes["range"] = (range if (
        isinstance(range, _ods_ir.Attribute) or
        not _ods_ir.AttrBuilder.contains('LLVM_ConstantRangeAttr')) else
          _ods_ir.AttrBuilder.get('LLVM_ConstantRangeAttr')(range, context=_ods_context))
    results.append(res)
    _ods_successors = None
    super().__init__(self.OPERATION_NAME, self._ODS_REGIONS, self._ODS_OPERAND_SEGMENTS, self._ODS_RESULT_SEGMENTS, attributes=attributes, results=results, operands=operands, successors=_ods_successors, regions=regions, loc=loc, ip=ip)

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

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

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

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

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

@_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, *, range=None, loc=None, ip=None):
    operands = []
    results = []
    attributes = {}
    regions = None
    _ods_context = _ods_get_default_loc_context(loc)
    if range is not None: attributes["range"] = (range if (
        isinstance(range, _ods_ir.Attribute) or
        not _ods_ir.AttrBuilder.contains('LLVM_ConstantRangeAttr')) else
          _ods_ir.AttrBuilder.get('LLVM_ConstantRangeAttr')(range, context=_ods_context))
    results.append(res)
    _ods_successors = None
    super().__init__(self.OPERATION_NAME, self._ODS_REGIONS, self._ODS_OPERAND_SEGMENTS, self._ODS_RESULT_SEGMENTS, attributes=attributes, results=results, operands=operands, successors=_ods_successors, regions=regions, loc=loc, ip=ip)

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

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

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

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

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

@_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, *, range=None, loc=None, ip=None):
    operands = []
    results = []
    attributes = {}
    regions = None
    _ods_context = _ods_get_default_loc_context(loc)
    if range is not None: attributes["range"] = (range if (
        isinstance(range, _ods_ir.Attribute) or
        not _ods_ir.AttrBuilder.contains('LLVM_ConstantRangeAttr')) else
          _ods_ir.AttrBuilder.get('LLVM_ConstantRangeAttr')(range, context=_ods_context))
    results.append(res)
    _ods_successors = None
    super().__init__(self.OPERATION_NAME, self._ODS_REGIONS, self._ODS_OPERAND_SEGMENTS, self._ODS_RESULT_SEGMENTS, attributes=attributes, results=results, operands=operands, successors=_ods_successors, regions=regions, loc=loc, ip=ip)

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

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

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

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

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

@_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, *, range=None, loc=None, ip=None):
    operands = []
    results = []
    attributes = {}
    regions = None
    _ods_context = _ods_get_default_loc_context(loc)
    if range is not None: attributes["range"] = (range if (
        isinstance(range, _ods_ir.Attribute) or
        not _ods_ir.AttrBuilder.contains('LLVM_ConstantRangeAttr')) else
          _ods_ir.AttrBuilder.get('LLVM_ConstantRangeAttr')(range, context=_ods_context))
    results.append(res)
    _ods_successors = None
    super().__init__(self.OPERATION_NAME, self._ODS_REGIONS, self._ODS_OPERAND_SEGMENTS, self._ODS_RESULT_SEGMENTS, attributes=attributes, results=results, operands=operands, successors=_ods_successors, regions=regions, loc=loc, ip=ip)

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

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

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

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

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

@_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, *, range=None, loc=None, ip=None):
    operands = []
    results = []
    attributes = {}
    regions = None
    _ods_context = _ods_get_default_loc_context(loc)
    if range is not None: attributes["range"] = (range if (
        isinstance(range, _ods_ir.Attribute) or
        not _ods_ir.AttrBuilder.contains('LLVM_ConstantRangeAttr')) else
          _ods_ir.AttrBuilder.get('LLVM_ConstantRangeAttr')(range, context=_ods_context))
    results.append(res)
    _ods_successors = None
    super().__init__(self.OPERATION_NAME, self._ODS_REGIONS, self._ODS_OPERAND_SEGMENTS, self._ODS_RESULT_SEGMENTS, attributes=attributes, results=results, operands=operands, successors=_ods_successors, regions=regions, loc=loc, ip=ip)

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

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

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

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

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

@_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, *, range=None, loc=None, ip=None):
    operands = []
    results = []
    attributes = {}
    regions = None
    _ods_context = _ods_get_default_loc_context(loc)
    if range is not None: attributes["range"] = (range if (
        isinstance(range, _ods_ir.Attribute) or
        not _ods_ir.AttrBuilder.contains('LLVM_ConstantRangeAttr')) else
          _ods_ir.AttrBuilder.get('LLVM_ConstantRangeAttr')(range, context=_ods_context))
    results.append(res)
    _ods_successors = None
    super().__init__(self.OPERATION_NAME, self._ODS_REGIONS, self._ODS_OPERAND_SEGMENTS, self._ODS_RESULT_SEGMENTS, attributes=attributes, results=results, operands=operands, successors=_ods_successors, regions=regions, loc=loc, ip=ip)

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

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

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

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

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

@_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, *, range=None, loc=None, ip=None):
    operands = []
    results = []
    attributes = {}
    regions = None
    _ods_context = _ods_get_default_loc_context(loc)
    if range is not None: attributes["range"] = (range if (
        isinstance(range, _ods_ir.Attribute) or
        not _ods_ir.AttrBuilder.contains('LLVM_ConstantRangeAttr')) else
          _ods_ir.AttrBuilder.get('LLVM_ConstantRangeAttr')(range, context=_ods_context))
    results.append(res)
    _ods_successors = None
    super().__init__(self.OPERATION_NAME, self._ODS_REGIONS, self._ODS_OPERAND_SEGMENTS, self._ODS_RESULT_SEGMENTS, attributes=attributes, results=results, operands=operands, successors=_ods_successors, regions=regions, loc=loc, ip=ip)

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

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

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

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

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

@_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, *, range=None, loc=None, ip=None):
    operands = []
    results = []
    attributes = {}
    regions = None
    _ods_context = _ods_get_default_loc_context(loc)
    if range is not None: attributes["range"] = (range if (
        isinstance(range, _ods_ir.Attribute) or
        not _ods_ir.AttrBuilder.contains('LLVM_ConstantRangeAttr')) else
          _ods_ir.AttrBuilder.get('LLVM_ConstantRangeAttr')(range, context=_ods_context))
    results.append(res)
    _ods_successors = None
    super().__init__(self.OPERATION_NAME, self._ODS_REGIONS, self._ODS_OPERAND_SEGMENTS, self._ODS_RESULT_SEGMENTS, attributes=attributes, results=results, operands=operands, successors=_ods_successors, regions=regions, loc=loc, ip=ip)

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

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

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

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

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

@_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, *, range=None, loc=None, ip=None):
    operands = []
    results = []
    attributes = {}
    regions = None
    _ods_context = _ods_get_default_loc_context(loc)
    if range is not None: attributes["range"] = (range if (
        isinstance(range, _ods_ir.Attribute) or
        not _ods_ir.AttrBuilder.contains('LLVM_ConstantRangeAttr')) else
          _ods_ir.AttrBuilder.get('LLVM_ConstantRangeAttr')(range, context=_ods_context))
    results.append(res)
    _ods_successors = None
    super().__init__(self.OPERATION_NAME, self._ODS_REGIONS, self._ODS_OPERAND_SEGMENTS, self._ODS_RESULT_SEGMENTS, attributes=attributes, results=results, operands=operands, successors=_ods_successors, regions=regions, loc=loc, ip=ip)

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

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

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

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

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

@_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, *, range=None, loc=None, ip=None):
    operands = []
    results = []
    attributes = {}
    regions = None
    _ods_context = _ods_get_default_loc_context(loc)
    if range is not None: attributes["range"] = (range if (
        isinstance(range, _ods_ir.Attribute) or
        not _ods_ir.AttrBuilder.contains('LLVM_ConstantRangeAttr')) else
          _ods_ir.AttrBuilder.get('LLVM_ConstantRangeAttr')(range, context=_ods_context))
    results.append(res)
    _ods_successors = None
    super().__init__(self.OPERATION_NAME, self._ODS_REGIONS, self._ODS_OPERAND_SEGMENTS, self._ODS_RESULT_SEGMENTS, attributes=attributes, results=results, operands=operands, successors=_ods_successors, regions=regions, loc=loc, ip=ip)

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

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

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

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

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

@_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.OPERATION_NAME, self._ODS_REGIONS, self._ODS_OPERAND_SEGMENTS, self._ODS_RESULT_SEGMENTS, 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 ClusterWaitOp(aligned=aligned, loc=loc, ip=ip)

@_ods_cext.register_operation(_Dialect)
class ConvertBF16x2ToF8x2Op(_ods_ir.OpView):
  OPERATION_NAME = "nvvm.convert.bf16x2.to.f8x2"

  _ODS_REGIONS = (0, True)

  def __init__(self, dst, type_, a, *, rnd=None, sat=None, loc=None, ip=None):
    operands = []
    results = []
    attributes = {}
    regions = None
    operands.append(a)
    _ods_context = _ods_get_default_loc_context(loc)
    attributes["type"] = (type_ if (
    isinstance(type_, _ods_ir.Attribute) or
    not _ods_ir.AttrBuilder.contains('ConvertFP8TypeAttr')) else
      _ods_ir.AttrBuilder.get('ConvertFP8TypeAttr')(type_, context=_ods_context))
    if rnd is not None: attributes["rnd"] = (rnd if (
        isinstance(rnd, _ods_ir.Attribute) or
        not _ods_ir.AttrBuilder.contains('FPRoundingModeAttr')) else
          _ods_ir.AttrBuilder.get('FPRoundingModeAttr')(rnd, context=_ods_context))
    if sat is not None: attributes["sat"] = (sat if (
        isinstance(sat, _ods_ir.Attribute) or
        not _ods_ir.AttrBuilder.contains('SaturationModeAttr')) else
          _ods_ir.AttrBuilder.get('SaturationModeAttr')(sat, context=_ods_context))
    results.append(dst)
    _ods_successors = None
    super().__init__(self.OPERATION_NAME, self._ODS_REGIONS, self._ODS_OPERAND_SEGMENTS, self._ODS_RESULT_SEGMENTS, attributes=attributes, results=results, operands=operands, successors=_ods_successors, regions=regions, loc=loc, ip=ip)

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

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

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

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

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

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

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

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

def convert_bf16x2_to_f8x2(dst, type_, a, *, rnd=None, sat=None, loc=None, ip=None) -> _ods_ir.Value:
  return ConvertBF16x2ToF8x2Op(dst=dst, type_=type_, a=a, rnd=rnd, sat=sat, loc=loc, ip=ip).result

@_ods_cext.register_operation(_Dialect)
class ConvertF16x2ToF8x2Op(_ods_ir.OpView):
  OPERATION_NAME = "nvvm.convert.f16x2.to.f8x2"

  _ODS_REGIONS = (0, True)

  def __init__(self, dst, type_, a, *, relu=None, loc=None, ip=None):
    operands = []
    results = []
    attributes = {}
    regions = None
    operands.append(a)
    _ods_context = _ods_get_default_loc_context(loc)
    attributes["type"] = (type_ if (
    isinstance(type_, _ods_ir.Attribute) or
    not _ods_ir.AttrBuilder.contains('ConvertFP8TypeAttr')) else
      _ods_ir.AttrBuilder.get('ConvertFP8TypeAttr')(type_, context=_ods_context))
    if relu is not None: attributes["relu"] = (relu if (
        isinstance(relu, _ods_ir.Attribute) or
        not _ods_ir.AttrBuilder.contains('BoolAttr')) else
          _ods_ir.AttrBuilder.get('BoolAttr')(relu, context=_ods_context))
    results.append(dst)
    _ods_successors = None
    super().__init__(self.OPERATION_NAME, self._ODS_REGIONS, self._ODS_OPERAND_SEGMENTS, self._ODS_RESULT_SEGMENTS, attributes=attributes, results=results, operands=operands, successors=_ods_successors, regions=regions, loc=loc, ip=ip)

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

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

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

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

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

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

def convert_f16x2_to_f8x2(dst, type_, a, *, relu=None, loc=None, ip=None) -> _ods_ir.Value:
  return ConvertF16x2ToF8x2Op(dst=dst, type_=type_, a=a, relu=relu, loc=loc, ip=ip).result

@_ods_cext.register_operation(_Dialect)
class ConvertF32x2ToF6x2Op(_ods_ir.OpView):
  OPERATION_NAME = "nvvm.convert.f32x2.to.f6x2"

  _ODS_REGIONS = (0, True)

  def __init__(self, dst, type_, a, b, *, relu=None, loc=None, ip=None):
    operands = []
    results = []
    attributes = {}
    regions = None
    operands.append(a)
    operands.append(b)
    _ods_context = _ods_get_default_loc_context(loc)
    attributes["type"] = (type_ if (
    isinstance(type_, _ods_ir.Attribute) or
    not _ods_ir.AttrBuilder.contains('ConvertFP6TypeAttr')) else
      _ods_ir.AttrBuilder.get('ConvertFP6TypeAttr')(type_, context=_ods_context))
    if relu is not None: attributes["relu"] = (relu if (
        isinstance(relu, _ods_ir.Attribute) or
        not _ods_ir.AttrBuilder.contains('BoolAttr')) else
          _ods_ir.AttrBuilder.get('BoolAttr')(relu, context=_ods_context))
    results.append(dst)
    _ods_successors = None
    super().__init__(self.OPERATION_NAME, self._ODS_REGIONS, self._ODS_OPERAND_SEGMENTS, self._ODS_RESULT_SEGMENTS, attributes=attributes, results=results, operands=operands, successors=_ods_successors, regions=regions, loc=loc, ip=ip)

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

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

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

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

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

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

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

def convert_f32x2_to_f6x2(dst, type_, a, b, *, relu=None, loc=None, ip=None) -> _ods_ir.Value:
  return ConvertF32x2ToF6x2Op(dst=dst, type_=type_, a=a, b=b, relu=relu, loc=loc, ip=ip).result

@_ods_cext.register_operation(_Dialect)
class ConvertF32x2ToF8x2Op(_ods_ir.OpView):
  OPERATION_NAME = "nvvm.convert.f32x2.to.f8x2"

  _ODS_REGIONS = (0, True)

  def __init__(self, dst, type_, a, b, *, rnd=None, sat=None, relu=None, loc=None, ip=None):
    operands = []
    results = []
    attributes = {}
    regions = None
    operands.append(a)
    operands.append(b)
    _ods_context = _ods_get_default_loc_context(loc)
    attributes["type"] = (type_ if (
    isinstance(type_, _ods_ir.Attribute) or
    not _ods_ir.AttrBuilder.contains('ConvertFP8TypeAttr')) else
      _ods_ir.AttrBuilder.get('ConvertFP8TypeAttr')(type_, context=_ods_context))
    if rnd is not None: attributes["rnd"] = (rnd if (
        isinstance(rnd, _ods_ir.Attribute) or
        not _ods_ir.AttrBuilder.contains('FPRoundingModeAttr')) else
          _ods_ir.AttrBuilder.get('FPRoundingModeAttr')(rnd, context=_ods_context))
    if sat is not None: attributes["sat"] = (sat if (
        isinstance(sat, _ods_ir.Attribute) or
        not _ods_ir.AttrBuilder.contains('SaturationModeAttr')) else
          _ods_ir.AttrBuilder.get('SaturationModeAttr')(sat, context=_ods_context))
    if relu is not None: attributes["relu"] = (relu if (
        isinstance(relu, _ods_ir.Attribute) or
        not _ods_ir.AttrBuilder.contains('BoolAttr')) else
          _ods_ir.AttrBuilder.get('BoolAttr')(relu, context=_ods_context))
    results.append(dst)
    _ods_successors = None
    super().__init__(self.OPERATION_NAME, self._ODS_REGIONS, self._ODS_OPERAND_SEGMENTS, self._ODS_RESULT_SEGMENTS, attributes=attributes, results=results, operands=operands, successors=_ods_successors, regions=regions, loc=loc, ip=ip)

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

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

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

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

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

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

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

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

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

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

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

def convert_f32x2_to_f8x2(dst, type_, a, b, *, rnd=None, sat=None, relu=None, loc=None, ip=None) -> _ods_ir.Value:
  return ConvertF32x2ToF8x2Op(dst=dst, type_=type_, a=a, b=b, rnd=rnd, sat=sat, relu=relu, loc=loc, ip=ip).result

@_ods_cext.register_operation(_Dialect)
class ConvertFloatToTF32Op(_ods_ir.OpView):
  OPERATION_NAME = "nvvm.convert.float.to.tf32"

  _ODS_REGIONS = (0, True)

  def __init__(self, res, src, *, rnd=None, sat=None, relu=None, loc=None, ip=None):
    operands = []
    results = []
    attributes = {}
    regions = None
    operands.append(src)
    _ods_context = _ods_get_default_loc_context(loc)
    if rnd is not None: attributes["rnd"] = (rnd if (
        isinstance(rnd, _ods_ir.Attribute) or
        not _ods_ir.AttrBuilder.contains('FPRoundingModeAttr')) else
          _ods_ir.AttrBuilder.get('FPRoundingModeAttr')(rnd, context=_ods_context))
    if sat is not None: attributes["sat"] = (sat if (
        isinstance(sat, _ods_ir.Attribute) or
        not _ods_ir.AttrBuilder.contains('SaturationModeAttr')) else
          _ods_ir.AttrBuilder.get('SaturationModeAttr')(sat, context=_ods_context))
    if relu is not None: attributes["relu"] = (relu if (
        isinstance(relu, _ods_ir.Attribute) or
        not _ods_ir.AttrBuilder.contains('BoolAttr')) else
          _ods_ir.AttrBuilder.get('BoolAttr')(relu, context=_ods_context))
    results.append(res)
    _ods_successors = None
    super().__init__(self.OPERATION_NAME, self._ODS_REGIONS, self._ODS_OPERAND_SEGMENTS, self._ODS_RESULT_SEGMENTS, attributes=attributes, results=results, operands=operands, successors=_ods_successors, regions=regions, loc=loc, ip=ip)

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

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

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

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

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

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

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

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

def convert_float_to_tf32(res, src, *, rnd=None, sat=None, relu=None, loc=None, ip=None) -> _ods_ir.Value:
  return ConvertFloatToTF32Op(res=res, src=src, rnd=rnd, sat=sat, relu=relu, loc=loc, ip=ip).result

@_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.OPERATION_NAME, self._ODS_REGIONS, self._ODS_OPERAND_SEGMENTS, self._ODS_RESULT_SEGMENTS, 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 CpAsyncBulkCommitGroupOp(loc=loc, ip=ip)

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

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

  _ODS_REGIONS = (0, True)

  def __init__(self, dstMem, srcMem, mbar, size, *, multicastMask=None, l2CacheHint=None, loc=None, ip=None):
    operands = []
    results = []
    attributes = {}
    regions = None
    operands.append(dstMem)
    operands.append(srcMem)
    operands.append(mbar)
    operands.append(size)
    operands.append(multicastMask)
    operands.append(l2CacheHint)
    _ods_context = _ods_get_default_loc_context(loc)
    _ods_successors = None
    super().__init__(self.OPERATION_NAME, self._ODS_REGIONS, self._ODS_OPERAND_SEGMENTS, self._ODS_RESULT_SEGMENTS, 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 srcMem(self):
    operand_range = _ods_segmented_accessor(
         self.operation.operands,
         self.operation.attributes["operandSegmentSizes"], 1)
    return operand_range[0]

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

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

  @builtins.property
  def multicastMask(self):
    operand_range = _ods_segmented_accessor(
         self.operation.operands,
         self.operation.attributes["operandSegmentSizes"], 4)
    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"], 5)
    return operand_range[0] if len(operand_range) > 0 else None

def cp_async_bulk_shared_cluster_global(dst_mem, src_mem, mbar, size, *, multicast_mask=None, l2_cache_hint=None, loc=None, ip=None) -> _ods_ir.Operation:
  return CpAsyncBulkGlobalToSharedClusterOp(dstMem=dst_mem, srcMem=src_mem, mbar=mbar, size=size, multicastMask=multicast_mask, l2CacheHint=l2_cache_hint, loc=loc, ip=ip)

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

  _ODS_REGIONS = (0, True)

  def __init__(self, srcMem, size, *, l2CacheHint=None, loc=None, ip=None):
    operands = []
    results = []
    attributes = {}
    regions = None
    operands.append(srcMem)
    operands.append(size)
    if l2CacheHint is not None: operands.append(l2CacheHint)
    _ods_context = _ods_get_default_loc_context(loc)
    _ods_successors = None
    super().__init__(self.OPERATION_NAME, self._ODS_REGIONS, self._ODS_OPERAND_SEGMENTS, self._ODS_RESULT_SEGMENTS, attributes=attributes, results=results, operands=operands, successors=_ods_successors, regions=regions, loc=loc, ip=ip)

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

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

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

def cp_async_bulk_prefetch(src_mem, size, *, l2_cache_hint=None, loc=None, ip=None) -> _ods_ir.Operation:
  return CpAsyncBulkPrefetchOp(srcMem=src_mem, size=size, l2CacheHint=l2_cache_hint, loc=loc, ip=ip)

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

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

  _ODS_REGIONS = (0, True)

  def __init__(self, dstMem, srcMem, size, *, l2CacheHint=None, byteMask=None, loc=None, ip=None):
    operands = []
    results = []
    attributes = {}
    regions = None
    operands.append(dstMem)
    operands.append(srcMem)
    operands.append(size)
    operands.append(l2CacheHint)
    operands.append(byteMask)
    _ods_context = _ods_get_default_loc_context(loc)
    _ods_successors = None
    super().__init__(self.OPERATION_NAME, self._ODS_REGIONS, self._ODS_OPERAND_SEGMENTS, self._ODS_RESULT_SEGMENTS, 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 srcMem(self):
    operand_range = _ods_segmented_accessor(
         self.operation.operands,
         self.operation.attributes["operandSegmentSizes"], 1)
    return operand_range[0]

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

  @builtins.property
  def l2CacheHint(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

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

def cp_async_bulk_global_shared_cta(dst_mem, src_mem, size, *, l2_cache_hint=None, byte_mask=None, loc=None, ip=None) -> _ods_ir.Operation:
  return CpAsyncBulkSharedCTAToGlobalOp(dstMem=dst_mem, srcMem=src_mem, size=size, l2CacheHint=l2_cache_hint, byteMask=byte_mask, loc=loc, ip=ip)

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

  _ODS_REGIONS = (0, True)

  def __init__(self, dstMem, srcMem, mbar, size, *, loc=None, ip=None):
    operands = []
    results = []
    attributes = {}
    regions = None
    operands.append(dstMem)
    operands.append(srcMem)
    operands.append(mbar)
    operands.append(size)
    _ods_context = _ods_get_default_loc_context(loc)
    _ods_successors = None
    super().__init__(self.OPERATION_NAME, self._ODS_REGIONS, self._ODS_OPERAND_SEGMENTS, self._ODS_RESULT_SEGMENTS, attributes=attributes, results=results, operands=operands, successors=_ods_successors, regions=regions, loc=loc, ip=ip)

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

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

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

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

def cp_async_bulk_shared_cluster_shared_cta(dst_mem, src_mem, mbar, size, *, loc=None, ip=None) -> _ods_ir.Operation:
  return CpAsyncBulkSharedCTAToSharedClusterOp(dstMem=dst_mem, srcMem=src_mem, mbar=mbar, size=size, 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(dstMem)
    operands.append(tmaDescriptor)
    operands.append(_get_op_results_or_values(coordinates))
    operands.append(mbar)
    operands.append(_get_op_results_or_values(im2colOffsets))
    operands.append(multicastMask)
    operands.append(l2CacheHint)
    operands.append(predicate)
    _ods_context = _ods_get_default_loc_context(loc)
    _ods_successors = None
    super().__init__(self.OPERATION_NAME, self._ODS_REGIONS, self._ODS_OPERAND_SEGMENTS, self._ODS_RESULT_SEGMENTS, 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 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 CpAsyncBulkTensorPrefetchOp(_ods_ir.OpView):
  OPERATION_NAME = "nvvm.cp.async.bulk.tensor.prefetch"

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

  _ODS_REGIONS = (0, True)

  def __init__(self, tmaDescriptor, coordinates, im2colOffsets, *, l2CacheHint=None, loc=None, ip=None):
    operands = []
    results = []
    attributes = {}
    regions = None
    operands.append(tmaDescriptor)
    operands.append(_get_op_results_or_values(coordinates))
    operands.append(_get_op_results_or_values(im2colOffsets))
    operands.append(l2CacheHint)
    _ods_context = _ods_get_default_loc_context(loc)
    _ods_successors = None
    super().__init__(self.OPERATION_NAME, self._ODS_REGIONS, self._ODS_OPERAND_SEGMENTS, self._ODS_RESULT_SEGMENTS, 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 coordinates(self):
    operand_range = _ods_segmented_accessor(
         self.operation.operands,
         self.operation.attributes["operandSegmentSizes"], 1)
    return operand_range

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

  @builtins.property
  def l2CacheHint(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_prefetch(tma_descriptor, coordinates, im2col_offsets, *, l2_cache_hint=None, loc=None, ip=None) -> _ods_ir.Operation:
  return CpAsyncBulkTensorPrefetchOp(tmaDescriptor=tma_descriptor, coordinates=coordinates, im2colOffsets=im2col_offsets, l2CacheHint=l2_cache_hint, loc=loc, ip=ip)

@_ods_cext.register_operation(_Dialect)
class CpAsyncBulkTensorReduceOp(_ods_ir.OpView):
  OPERATION_NAME = "nvvm.cp.async.bulk.tensor.reduce"

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

  _ODS_REGIONS = (0, True)

  def __init__(self, tmaDescriptor, srcMem, redKind, coordinates, *, mode=None, l2CacheHint=None, loc=None, ip=None):
    operands = []
    results = []
    attributes = {}
    regions = None
    operands.append(tmaDescriptor)
    operands.append(srcMem)
    operands.append(_get_op_results_or_values(coordinates))
    operands.append(l2CacheHint)
    _ods_context = _ods_get_default_loc_context(loc)
    attributes["redKind"] = (redKind if (
    isinstance(redKind, _ods_ir.Attribute) or
    not _ods_ir.AttrBuilder.contains('TMAReduxKindAttr')) else
      _ods_ir.AttrBuilder.get('TMAReduxKindAttr')(redKind, context=_ods_context))
    if mode is not None: attributes["mode"] = (mode if (
        isinstance(mode, _ods_ir.Attribute) or
        not _ods_ir.AttrBuilder.contains('TMAStoreModeAttr')) else
          _ods_ir.AttrBuilder.get('TMAStoreModeAttr')(mode, context=_ods_context))
    _ods_successors = None
    super().__init__(self.OPERATION_NAME, self._ODS_REGIONS, self._ODS_OPERAND_SEGMENTS, self._ODS_RESULT_SEGMENTS, 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 l2CacheHint(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

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

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

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

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

def cp_async_bulk_tensor_reduce(tma_descriptor, src_mem, red_kind, coordinates, *, mode=None, l2_cache_hint=None, loc=None, ip=None) -> _ods_ir.Operation:
  return CpAsyncBulkTensorReduceOp(tmaDescriptor=tma_descriptor, srcMem=src_mem, redKind=red_kind, coordinates=coordinates, mode=mode, l2CacheHint=l2_cache_hint, 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(tmaDescriptor)
    operands.append(srcMem)
    operands.append(_get_op_results_or_values(coordinates))
    operands.append(predicate)
    _ods_context = _ods_get_default_loc_context(loc)
    _ods_successors = None
    super().__init__(self.OPERATION_NAME, self._ODS_REGIONS, self._ODS_OPERAND_SEGMENTS, self._ODS_RESULT_SEGMENTS, 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 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.OPERATION_NAME, self._ODS_REGIONS, self._ODS_OPERAND_SEGMENTS, self._ODS_RESULT_SEGMENTS, 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 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.OPERATION_NAME, self._ODS_REGIONS, self._ODS_OPERAND_SEGMENTS, self._ODS_RESULT_SEGMENTS, 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 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(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.OPERATION_NAME, self._ODS_REGIONS, self._ODS_OPERAND_SEGMENTS, self._ODS_RESULT_SEGMENTS, 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 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(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.OPERATION_NAME, self._ODS_REGIONS, self._ODS_OPERAND_SEGMENTS, self._ODS_RESULT_SEGMENTS, 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 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(dst)
    operands.append(src)
    if cpSize is not None: operands.append(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.OPERATION_NAME, self._ODS_REGIONS, self._ODS_OPERAND_SEGMENTS, self._ODS_RESULT_SEGMENTS, 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 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.OPERATION_NAME, self._ODS_REGIONS, self._ODS_OPERAND_SEGMENTS, self._ODS_RESULT_SEGMENTS, 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 CpAsyncWaitGroupOp(n=n, loc=loc, ip=ip)

@_ods_cext.register_operation(_Dialect)
class DotAccumulate2WayOp(_ods_ir.OpView):
  OPERATION_NAME = "nvvm.dot.accumulate.2way"

  _ODS_REGIONS = (0, True)

  def __init__(self, res, a, a_type, b, b_type, c, b_hi, *, loc=None, ip=None):
    operands = []
    results = []
    attributes = {}
    regions = None
    operands.append(a)
    operands.append(b)
    operands.append(c)
    _ods_context = _ods_get_default_loc_context(loc)
    attributes["a_type"] = (a_type if (
    isinstance(a_type, _ods_ir.Attribute) or
    not _ods_ir.AttrBuilder.contains('DotAccumulateTypeAttr')) else
      _ods_ir.AttrBuilder.get('DotAccumulateTypeAttr')(a_type, context=_ods_context))
    attributes["b_type"] = (b_type if (
    isinstance(b_type, _ods_ir.Attribute) or
    not _ods_ir.AttrBuilder.contains('DotAccumulateTypeAttr')) else
      _ods_ir.AttrBuilder.get('DotAccumulateTypeAttr')(b_type, context=_ods_context))
    attributes["b_hi"] = (b_hi if (
    isinstance(b_hi, _ods_ir.Attribute) or
    not _ods_ir.AttrBuilder.contains('BoolAttr')) else
      _ods_ir.AttrBuilder.get('BoolAttr')(b_hi, context=_ods_context))
    results.append(res)
    _ods_successors = None
    super().__init__(self.OPERATION_NAME, self._ODS_REGIONS, self._ODS_OPERAND_SEGMENTS, self._ODS_RESULT_SEGMENTS, attributes=attributes, results=results, operands=operands, successors=_ods_successors, regions=regions, loc=loc, ip=ip)

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

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

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

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

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

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

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

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

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

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

def dot_accumulate_2way(res, a, a_type, b, b_type, c, b_hi, *, loc=None, ip=None) -> _ods_ir.Value:
  return DotAccumulate2WayOp(res=res, a=a, a_type=a_type, b=b, b_type=b_type, c=c, b_hi=b_hi, loc=loc, ip=ip).result

@_ods_cext.register_operation(_Dialect)
class DotAccumulate4WayOp(_ods_ir.OpView):
  OPERATION_NAME = "nvvm.dot.accumulate.4way"

  _ODS_REGIONS = (0, True)

  def __init__(self, res, a, a_type, b, b_type, c, *, loc=None, ip=None):
    operands = []
    results = []
    attributes = {}
    regions = None
    operands.append(a)
    operands.append(b)
    operands.append(c)
    _ods_context = _ods_get_default_loc_context(loc)
    attributes["a_type"] = (a_type if (
    isinstance(a_type, _ods_ir.Attribute) or
    not _ods_ir.AttrBuilder.contains('DotAccumulateTypeAttr')) else
      _ods_ir.AttrBuilder.get('DotAccumulateTypeAttr')(a_type, context=_ods_context))
    attributes["b_type"] = (b_type if (
    isinstance(b_type, _ods_ir.Attribute) or
    not _ods_ir.AttrBuilder.contains('DotAccumulateTypeAttr')) else
      _ods_ir.AttrBuilder.get('DotAccumulateTypeAttr')(b_type, context=_ods_context))
    results.append(res)
    _ods_successors = None
    super().__init__(self.OPERATION_NAME, self._ODS_REGIONS, self._ODS_OPERAND_SEGMENTS, self._ODS_RESULT_SEGMENTS, attributes=attributes, results=results, operands=operands, successors=_ods_successors, regions=regions, loc=loc, ip=ip)

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

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

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

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

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

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

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

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

def dot_accumulate_4way(res, a, a_type, b, b_type, c, *, loc=None, ip=None) -> _ods_ir.Value:
  return DotAccumulate4WayOp(res=res, a=a, a_type=a_type, b=b, b_type=b_type, c=c, loc=loc, ip=ip).result

@_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.OPERATION_NAME, self._ODS_REGIONS, self._ODS_OPERAND_SEGMENTS, self._ODS_RESULT_SEGMENTS, 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 ElectSyncOp(pred=pred, loc=loc, ip=ip).result

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

  _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.OPERATION_NAME, self._ODS_REGIONS, self._ODS_OPERAND_SEGMENTS, self._ODS_RESULT_SEGMENTS, 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_envreg0(res, *, loc=None, ip=None) -> _ods_ir.Value:
  return EnvReg0Op(res=res, loc=loc, ip=ip).result

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

  _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.OPERATION_NAME, self._ODS_REGIONS, self._ODS_OPERAND_SEGMENTS, self._ODS_RESULT_SEGMENTS, 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_envreg1(res, *, loc=None, ip=None) -> _ods_ir.Value:
  return EnvReg1Op(res=res, loc=loc, ip=ip).result

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

  _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.OPERATION_NAME, self._ODS_REGIONS, self._ODS_OPERAND_SEGMENTS, self._ODS_RESULT_SEGMENTS, 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_envreg2(res, *, loc=None, ip=None) -> _ods_ir.Value:
  return EnvReg2Op(res=res, loc=loc, ip=ip).result

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

  _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.OPERATION_NAME, self._ODS_REGIONS, self._ODS_OPERAND_SEGMENTS, self._ODS_RESULT_SEGMENTS, 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_envreg3(res, *, loc=None, ip=None) -> _ods_ir.Value:
  return EnvReg3Op(res=res, loc=loc, ip=ip).result

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

  _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.OPERATION_NAME, self._ODS_REGIONS, self._ODS_OPERAND_SEGMENTS, self._ODS_RESULT_SEGMENTS, 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_envreg4(res, *, loc=None, ip=None) -> _ods_ir.Value:
  return EnvReg4Op(res=res, loc=loc, ip=ip).result

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

  _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.OPERATION_NAME, self._ODS_REGIONS, self._ODS_OPERAND_SEGMENTS, self._ODS_RESULT_SEGMENTS, 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_envreg5(res, *, loc=None, ip=None) -> _ods_ir.Value:
  return EnvReg5Op(res=res, loc=loc, ip=ip).result

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

  _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.OPERATION_NAME, self._ODS_REGIONS, self._ODS_OPERAND_SEGMENTS, self._ODS_RESULT_SEGMENTS, 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_envreg6(res, *, loc=None, ip=None) -> _ods_ir.Value:
  return EnvReg6Op(res=res, loc=loc, ip=ip).result

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

  _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.OPERATION_NAME, self._ODS_REGIONS, self._ODS_OPERAND_SEGMENTS, self._ODS_RESULT_SEGMENTS, 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_envreg7(res, *, loc=None, ip=None) -> _ods_ir.Value:
  return EnvReg7Op(res=res, loc=loc, ip=ip).result

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

  _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.OPERATION_NAME, self._ODS_REGIONS, self._ODS_OPERAND_SEGMENTS, self._ODS_RESULT_SEGMENTS, 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_envreg8(res, *, loc=None, ip=None) -> _ods_ir.Value:
  return EnvReg8Op(res=res, loc=loc, ip=ip).result

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

  _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.OPERATION_NAME, self._ODS_REGIONS, self._ODS_OPERAND_SEGMENTS, self._ODS_RESULT_SEGMENTS, 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_envreg9(res, *, loc=None, ip=None) -> _ods_ir.Value:
  return EnvReg9Op(res=res, loc=loc, ip=ip).result

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

  _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.OPERATION_NAME, self._ODS_REGIONS, self._ODS_OPERAND_SEGMENTS, self._ODS_RESULT_SEGMENTS, 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_envreg10(res, *, loc=None, ip=None) -> _ods_ir.Value:
  return EnvReg10Op(res=res, loc=loc, ip=ip).result

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

  _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.OPERATION_NAME, self._ODS_REGIONS, self._ODS_OPERAND_SEGMENTS, self._ODS_RESULT_SEGMENTS, 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_envreg11(res, *, loc=None, ip=None) -> _ods_ir.Value:
  return EnvReg11Op(res=res, loc=loc, ip=ip).result

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

  _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.OPERATION_NAME, self._ODS_REGIONS, self._ODS_OPERAND_SEGMENTS, self._ODS_RESULT_SEGMENTS, 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_envreg12(res, *, loc=None, ip=None) -> _ods_ir.Value:
  return EnvReg12Op(res=res, loc=loc, ip=ip).result

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

  _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.OPERATION_NAME, self._ODS_REGIONS, self._ODS_OPERAND_SEGMENTS, self._ODS_RESULT_SEGMENTS, 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_envreg13(res, *, loc=None, ip=None) -> _ods_ir.Value:
  return EnvReg13Op(res=res, loc=loc, ip=ip).result

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

  _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.OPERATION_NAME, self._ODS_REGIONS, self._ODS_OPERAND_SEGMENTS, self._ODS_RESULT_SEGMENTS, 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_envreg14(res, *, loc=None, ip=None) -> _ods_ir.Value:
  return EnvReg14Op(res=res, loc=loc, ip=ip).result

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

  _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.OPERATION_NAME, self._ODS_REGIONS, self._ODS_OPERAND_SEGMENTS, self._ODS_RESULT_SEGMENTS, 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_envreg15(res, *, loc=None, ip=None) -> _ods_ir.Value:
  return EnvReg15Op(res=res, loc=loc, ip=ip).result

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

  _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.OPERATION_NAME, self._ODS_REGIONS, self._ODS_OPERAND_SEGMENTS, self._ODS_RESULT_SEGMENTS, 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_envreg16(res, *, loc=None, ip=None) -> _ods_ir.Value:
  return EnvReg16Op(res=res, loc=loc, ip=ip).result

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

  _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.OPERATION_NAME, self._ODS_REGIONS, self._ODS_OPERAND_SEGMENTS, self._ODS_RESULT_SEGMENTS, 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_envreg17(res, *, loc=None, ip=None) -> _ods_ir.Value:
  return EnvReg17Op(res=res, loc=loc, ip=ip).result

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

  _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.OPERATION_NAME, self._ODS_REGIONS, self._ODS_OPERAND_SEGMENTS, self._ODS_RESULT_SEGMENTS, 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_envreg18(res, *, loc=None, ip=None) -> _ods_ir.Value:
  return EnvReg18Op(res=res, loc=loc, ip=ip).result

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

  _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.OPERATION_NAME, self._ODS_REGIONS, self._ODS_OPERAND_SEGMENTS, self._ODS_RESULT_SEGMENTS, 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_envreg19(res, *, loc=None, ip=None) -> _ods_ir.Value:
  return EnvReg19Op(res=res, loc=loc, ip=ip).result

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

  _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.OPERATION_NAME, self._ODS_REGIONS, self._ODS_OPERAND_SEGMENTS, self._ODS_RESULT_SEGMENTS, 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_envreg20(res, *, loc=None, ip=None) -> _ods_ir.Value:
  return EnvReg20Op(res=res, loc=loc, ip=ip).result

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

  _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.OPERATION_NAME, self._ODS_REGIONS, self._ODS_OPERAND_SEGMENTS, self._ODS_RESULT_SEGMENTS, 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_envreg21(res, *, loc=None, ip=None) -> _ods_ir.Value:
  return EnvReg21Op(res=res, loc=loc, ip=ip).result

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

  _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.OPERATION_NAME, self._ODS_REGIONS, self._ODS_OPERAND_SEGMENTS, self._ODS_RESULT_SEGMENTS, 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_envreg22(res, *, loc=None, ip=None) -> _ods_ir.Value:
  return EnvReg22Op(res=res, loc=loc, ip=ip).result

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

  _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.OPERATION_NAME, self._ODS_REGIONS, self._ODS_OPERAND_SEGMENTS, self._ODS_RESULT_SEGMENTS, 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_envreg23(res, *, loc=None, ip=None) -> _ods_ir.Value:
  return EnvReg23Op(res=res, loc=loc, ip=ip).result

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

  _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.OPERATION_NAME, self._ODS_REGIONS, self._ODS_OPERAND_SEGMENTS, self._ODS_RESULT_SEGMENTS, 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_envreg24(res, *, loc=None, ip=None) -> _ods_ir.Value:
  return EnvReg24Op(res=res, loc=loc, ip=ip).result

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

  _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.OPERATION_NAME, self._ODS_REGIONS, self._ODS_OPERAND_SEGMENTS, self._ODS_RESULT_SEGMENTS, 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_envreg25(res, *, loc=None, ip=None) -> _ods_ir.Value:
  return EnvReg25Op(res=res, loc=loc, ip=ip).result

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

  _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.OPERATION_NAME, self._ODS_REGIONS, self._ODS_OPERAND_SEGMENTS, self._ODS_RESULT_SEGMENTS, 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_envreg26(res, *, loc=None, ip=None) -> _ods_ir.Value:
  return EnvReg26Op(res=res, loc=loc, ip=ip).result

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

  _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.OPERATION_NAME, self._ODS_REGIONS, self._ODS_OPERAND_SEGMENTS, self._ODS_RESULT_SEGMENTS, 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_envreg27(res, *, loc=None, ip=None) -> _ods_ir.Value:
  return EnvReg27Op(res=res, loc=loc, ip=ip).result

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

  _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.OPERATION_NAME, self._ODS_REGIONS, self._ODS_OPERAND_SEGMENTS, self._ODS_RESULT_SEGMENTS, 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_envreg28(res, *, loc=None, ip=None) -> _ods_ir.Value:
  return EnvReg28Op(res=res, loc=loc, ip=ip).result

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

  _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.OPERATION_NAME, self._ODS_REGIONS, self._ODS_OPERAND_SEGMENTS, self._ODS_RESULT_SEGMENTS, 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_envreg29(res, *, loc=None, ip=None) -> _ods_ir.Value:
  return EnvReg29Op(res=res, loc=loc, ip=ip).result

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

  _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.OPERATION_NAME, self._ODS_REGIONS, self._ODS_OPERAND_SEGMENTS, self._ODS_RESULT_SEGMENTS, 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_envreg30(res, *, loc=None, ip=None) -> _ods_ir.Value:
  return EnvReg30Op(res=res, loc=loc, ip=ip).result

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

  _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.OPERATION_NAME, self._ODS_REGIONS, self._ODS_OPERAND_SEGMENTS, self._ODS_RESULT_SEGMENTS, 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_envreg31(res, *, loc=None, ip=None) -> _ods_ir.Value:
  return EnvReg31Op(res=res, loc=loc, ip=ip).result

@_ods_cext.register_operation(_Dialect)
class Exit(_ods_ir.OpView):
  OPERATION_NAME = "nvvm.exit"

  _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.OPERATION_NAME, self._ODS_REGIONS, self._ODS_OPERAND_SEGMENTS, self._ODS_RESULT_SEGMENTS, attributes=attributes, results=results, operands=operands, successors=_ods_successors, regions=regions, loc=loc, ip=ip)

def exit(*, loc=None, ip=None) -> _ods_ir.Operation:
  return Exit(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.OPERATION_NAME, self._ODS_REGIONS, self._ODS_OPERAND_SEGMENTS, self._ODS_RESULT_SEGMENTS, 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 FenceMbarrierInitOp(loc=loc, ip=ip)

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

  _ODS_REGIONS = (0, True)

  def __init__(self, scope, addr, size, *, fromProxy=None, toProxy=None, loc=None, ip=None):
    operands = []
    results = []
    attributes = {}
    regions = None
    operands.append(addr)
    operands.append(size)
    _ods_context = _ods_get_default_loc_context(loc)
    attributes["scope"] = (scope if (
    isinstance(scope, _ods_ir.Attribute) or
    not _ods_ir.AttrBuilder.contains('MemScopeKindAttr')) else
      _ods_ir.AttrBuilder.get('MemScopeKindAttr')(scope, context=_ods_context))
    if fromProxy is not None: attributes["fromProxy"] = (fromProxy if (
        isinstance(fromProxy, _ods_ir.Attribute) or
        not _ods_ir.AttrBuilder.contains('ProxyKindAttr')) else
          _ods_ir.AttrBuilder.get('ProxyKindAttr')(fromProxy, context=_ods_context))
    if toProxy is not None: attributes["toProxy"] = (toProxy if (
        isinstance(toProxy, _ods_ir.Attribute) or
        not _ods_ir.AttrBuilder.contains('ProxyKindAttr')) else
          _ods_ir.AttrBuilder.get('ProxyKindAttr')(toProxy, context=_ods_context))
    _ods_successors = None
    super().__init__(self.OPERATION_NAME, self._ODS_REGIONS, self._ODS_OPERAND_SEGMENTS, self._ODS_RESULT_SEGMENTS, 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 size(self):
    return self.operation.operands[1]

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

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

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

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

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

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

def fence_proxy_acquire(scope, addr, size, *, from_proxy=None, to_proxy=None, loc=None, ip=None) -> _ods_ir.Operation:
  return FenceProxyAcquireOp(scope=scope, addr=addr, size=size, fromProxy=from_proxy, toProxy=to_proxy, 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.OPERATION_NAME, self._ODS_REGIONS, self._ODS_OPERAND_SEGMENTS, self._ODS_RESULT_SEGMENTS, 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 FenceProxyOp(kind=kind, space=space, loc=loc, ip=ip)

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

  _ODS_REGIONS = (0, True)

  def __init__(self, scope, *, fromProxy=None, toProxy=None, loc=None, ip=None):
    operands = []
    results = []
    attributes = {}
    regions = None
    _ods_context = _ods_get_default_loc_context(loc)
    attributes["scope"] = (scope if (
    isinstance(scope, _ods_ir.Attribute) or
    not _ods_ir.AttrBuilder.contains('MemScopeKindAttr')) else
      _ods_ir.AttrBuilder.get('MemScopeKindAttr')(scope, context=_ods_context))
    if fromProxy is not None: attributes["fromProxy"] = (fromProxy if (
        isinstance(fromProxy, _ods_ir.Attribute) or
        not _ods_ir.AttrBuilder.contains('ProxyKindAttr')) else
          _ods_ir.AttrBuilder.get('ProxyKindAttr')(fromProxy, context=_ods_context))
    if toProxy is not None: attributes["toProxy"] = (toProxy if (
        isinstance(toProxy, _ods_ir.Attribute) or
        not _ods_ir.AttrBuilder.contains('ProxyKindAttr')) else
          _ods_ir.AttrBuilder.get('ProxyKindAttr')(toProxy, context=_ods_context))
    _ods_successors = None
    super().__init__(self.OPERATION_NAME, self._ODS_REGIONS, self._ODS_OPERAND_SEGMENTS, self._ODS_RESULT_SEGMENTS, attributes=attributes, results=results, operands=operands, successors=_ods_successors, regions=regions, loc=loc, ip=ip)

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

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

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

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

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

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

def fence_proxy_release(scope, *, from_proxy=None, to_proxy=None, loc=None, ip=None) -> _ods_ir.Operation:
  return FenceProxyReleaseOp(scope=scope, fromProxy=from_proxy, toProxy=to_proxy, 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.OPERATION_NAME, self._ODS_REGIONS, self._ODS_OPERAND_SEGMENTS, self._ODS_RESULT_SEGMENTS, 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 FenceScClusterOp(loc=loc, ip=ip)

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

  _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.OPERATION_NAME, self._ODS_REGIONS, self._ODS_OPERAND_SEGMENTS, self._ODS_RESULT_SEGMENTS, 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_globaltimer(res, *, loc=None, ip=None) -> _ods_ir.Value:
  return GlobalTimerOp(res=res, loc=loc, ip=ip).result

@_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, *, range=None, loc=None, ip=None):
    operands = []
    results = []
    attributes = {}
    regions = None
    _ods_context = _ods_get_default_loc_context(loc)
    if range is not None: attributes["range"] = (range if (
        isinstance(range, _ods_ir.Attribute) or
        not _ods_ir.AttrBuilder.contains('LLVM_ConstantRangeAttr')) else
          _ods_ir.AttrBuilder.get('LLVM_ConstantRangeAttr')(range, context=_ods_context))
    results.append(res)
    _ods_successors = None
    super().__init__(self.OPERATION_NAME, self._ODS_REGIONS, self._ODS_OPERAND_SEGMENTS, self._ODS_RESULT_SEGMENTS, attributes=attributes, results=results, operands=operands, successors=_ods_successors, regions=regions, loc=loc, ip=ip)

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

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

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

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

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

@_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, *, range=None, loc=None, ip=None):
    operands = []
    results = []
    attributes = {}
    regions = None
    _ods_context = _ods_get_default_loc_context(loc)
    if range is not None: attributes["range"] = (range if (
        isinstance(range, _ods_ir.Attribute) or
        not _ods_ir.AttrBuilder.contains('LLVM_ConstantRangeAttr')) else
          _ods_ir.AttrBuilder.get('LLVM_ConstantRangeAttr')(range, context=_ods_context))
    results.append(res)
    _ods_successors = None
    super().__init__(self.OPERATION_NAME, self._ODS_REGIONS, self._ODS_OPERAND_SEGMENTS, self._ODS_RESULT_SEGMENTS, attributes=attributes, results=results, operands=operands, successors=_ods_successors, regions=regions, loc=loc, ip=ip)

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

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

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

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

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

@_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, *, range=None, loc=None, ip=None):
    operands = []
    results = []
    attributes = {}
    regions = None
    _ods_context = _ods_get_default_loc_context(loc)
    if range is not None: attributes["range"] = (range if (
        isinstance(range, _ods_ir.Attribute) or
        not _ods_ir.AttrBuilder.contains('LLVM_ConstantRangeAttr')) else
          _ods_ir.AttrBuilder.get('LLVM_ConstantRangeAttr')(range, context=_ods_context))
    results.append(res)
    _ods_successors = None
    super().__init__(self.OPERATION_NAME, self._ODS_REGIONS, self._ODS_OPERAND_SEGMENTS, self._ODS_RESULT_SEGMENTS, attributes=attributes, results=results, operands=operands, successors=_ods_successors, regions=regions, loc=loc, ip=ip)

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

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

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

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

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

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

  _ODS_REGIONS = (0, True)

  def __init__(self, res, *, range=None, loc=None, ip=None):
    operands = []
    results = []
    attributes = {}
    regions = None
    _ods_context = _ods_get_default_loc_context(loc)
    if range is not None: attributes["range"] = (range if (
        isinstance(range, _ods_ir.Attribute) or
        not _ods_ir.AttrBuilder.contains('LLVM_ConstantRangeAttr')) else
          _ods_ir.AttrBuilder.get('LLVM_ConstantRangeAttr')(range, context=_ods_context))
    results.append(res)
    _ods_successors = None
    super().__init__(self.OPERATION_NAME, self._ODS_REGIONS, self._ODS_OPERAND_SEGMENTS, self._ODS_RESULT_SEGMENTS, attributes=attributes, results=results, operands=operands, successors=_ods_successors, regions=regions, loc=loc, ip=ip)

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

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

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

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

def read_ptx_sreg_gridid(res, *, range=None, loc=None, ip=None) -> _ods_ir.Value:
  return GridIdOp(res=res, range=range, loc=loc, ip=ip).result

@_ods_cext.register_operation(_Dialect)
class GriddepcontrolLaunchDependentsOp(_ods_ir.OpView):
  OPERATION_NAME = "nvvm.griddepcontrol.launch.dependents"

  _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.OPERATION_NAME, self._ODS_REGIONS, self._ODS_OPERAND_SEGMENTS, self._ODS_RESULT_SEGMENTS, attributes=attributes, results=results, operands=operands, successors=_ods_successors, regions=regions, loc=loc, ip=ip)

def griddepcontrol_launch_dependents(*, loc=None, ip=None) -> _ods_ir.Operation:
  return GriddepcontrolLaunchDependentsOp(loc=loc, ip=ip)

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

  _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.OPERATION_NAME, self._ODS_REGIONS, self._ODS_OPERAND_SEGMENTS, self._ODS_RESULT_SEGMENTS, attributes=attributes, results=results, operands=operands, successors=_ods_successors, regions=regions, loc=loc, ip=ip)

def griddepcontrol_wait(*, loc=None, ip=None) -> _ods_ir.Operation:
  return GriddepcontrolWaitOp(loc=loc, ip=ip)

@_ods_cext.register_operation(_Dialect)
class InlinePtxOp(_ods_ir.OpView):
  OPERATION_NAME = "nvvm.inline_ptx"

  _ODS_OPERAND_SEGMENTS = [-1,0,]

  _ODS_REGIONS = (0, True)

  def __init__(self, writeOnlyArgs, readOnlyArgs, ptxCode, *, predicate=None, loc=None, ip=None):
    operands = []
    results = []
    attributes = {}
    regions = None
    operands.append(_get_op_results_or_values(readOnlyArgs))
    operands.append(predicate)
    _ods_context = _ods_get_default_loc_context(loc)
    attributes["ptxCode"] = (ptxCode if (
    isinstance(ptxCode, _ods_ir.Attribute) or
    not _ods_ir.AttrBuilder.contains('StrAttr')) else
      _ods_ir.AttrBuilder.get('StrAttr')(ptxCode, context=_ods_context))
    results.extend(writeOnlyArgs)
    _ods_successors = None
    super().__init__(self.OPERATION_NAME, self._ODS_REGIONS, self._ODS_OPERAND_SEGMENTS, self._ODS_RESULT_SEGMENTS, attributes=attributes, results=results, operands=operands, successors=_ods_successors, regions=regions, loc=loc, ip=ip)

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

  @builtins.property
  def predicate(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

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

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

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

def inline_ptx(write_only_args, read_only_args, ptx_code, *, predicate=None, loc=None, ip=None) -> _ods_ir.Value:
  return _get_op_result_or_op_results(InlinePtxOp(writeOnlyArgs=write_only_args, readOnlyArgs=read_only_args, ptxCode=ptx_code, predicate=predicate, 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, *, range=None, loc=None, ip=None):
    operands = []
    results = []
    attributes = {}
    regions = None
    _ods_context = _ods_get_default_loc_context(loc)
    if range is not None: attributes["range"] = (range if (
        isinstance(range, _ods_ir.Attribute) or
        not _ods_ir.AttrBuilder.contains('LLVM_ConstantRangeAttr')) else
          _ods_ir.AttrBuilder.get('LLVM_ConstantRangeAttr')(range, context=_ods_context))
    results.append(res)
    _ods_successors = None
    super().__init__(self.OPERATION_NAME, self._ODS_REGIONS, self._ODS_OPERAND_SEGMENTS, self._ODS_RESULT_SEGMENTS, attributes=attributes, results=results, operands=operands, successors=_ods_successors, regions=regions, loc=loc, ip=ip)

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

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

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

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

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

@_ods_cext.register_operation(_Dialect)
class LaneMaskEqOp(_ods_ir.OpView):
  OPERATION_NAME = "nvvm.read.ptx.sreg.lanemask.eq"

  _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.OPERATION_NAME, self._ODS_REGIONS, self._ODS_OPERAND_SEGMENTS, self._ODS_RESULT_SEGMENTS, 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_lanemask_eq(res, *, loc=None, ip=None) -> _ods_ir.Value:
  return LaneMaskEqOp(res=res, loc=loc, ip=ip).result

@_ods_cext.register_operation(_Dialect)
class LaneMaskGeOp(_ods_ir.OpView):
  OPERATION_NAME = "nvvm.read.ptx.sreg.lanemask.ge"

  _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.OPERATION_NAME, self._ODS_REGIONS, self._ODS_OPERAND_SEGMENTS, self._ODS_RESULT_SEGMENTS, 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_lanemask_ge(res, *, loc=None, ip=None) -> _ods_ir.Value:
  return LaneMaskGeOp(res=res, loc=loc, ip=ip).result

@_ods_cext.register_operation(_Dialect)
class LaneMaskGtOp(_ods_ir.OpView):
  OPERATION_NAME = "nvvm.read.ptx.sreg.lanemask.gt"

  _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.OPERATION_NAME, self._ODS_REGIONS, self._ODS_OPERAND_SEGMENTS, self._ODS_RESULT_SEGMENTS, 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_lanemask_gt(res, *, loc=None, ip=None) -> _ods_ir.Value:
  return LaneMaskGtOp(res=res, loc=loc, ip=ip).result

@_ods_cext.register_operation(_Dialect)
class LaneMaskLeOp(_ods_ir.OpView):
  OPERATION_NAME = "nvvm.read.ptx.sreg.lanemask.le"

  _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.OPERATION_NAME, self._ODS_REGIONS, self._ODS_OPERAND_SEGMENTS, self._ODS_RESULT_SEGMENTS, 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_lanemask_le(res, *, loc=None, ip=None) -> _ods_ir.Value:
  return LaneMaskLeOp(res=res, loc=loc, ip=ip).result

@_ods_cext.register_operation(_Dialect)
class LaneMaskLtOp(_ods_ir.OpView):
  OPERATION_NAME = "nvvm.read.ptx.sreg.lanemask.lt"

  _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.OPERATION_NAME, self._ODS_REGIONS, self._ODS_OPERAND_SEGMENTS, self._ODS_RESULT_SEGMENTS, 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_lanemask_lt(res, *, loc=None, ip=None) -> _ods_ir.Value:
  return LaneMaskLtOp(res=res, loc=loc, ip=ip).result

@_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(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.OPERATION_NAME, self._ODS_REGIONS, self._ODS_OPERAND_SEGMENTS, self._ODS_RESULT_SEGMENTS, 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 LdMatrixOp(res=res, ptr=ptr, num=num, layout=layout, loc=loc, ip=ip).result

@_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(addr)
    operands.append(txcount)
    if predicate is not None: operands.append(predicate)
    _ods_context = _ods_get_default_loc_context(loc)
    _ods_successors = None
    super().__init__(self.OPERATION_NAME, self._ODS_REGIONS, self._ODS_OPERAND_SEGMENTS, self._ODS_RESULT_SEGMENTS, 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 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(addr)
    operands.append(txcount)
    if predicate is not None: operands.append(predicate)
    _ods_context = _ods_get_default_loc_context(loc)
    _ods_successors = None
    super().__init__(self.OPERATION_NAME, self._ODS_REGIONS, self._ODS_OPERAND_SEGMENTS, self._ODS_RESULT_SEGMENTS, 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 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(addr)
    operands.append(count)
    _ods_context = _ods_get_default_loc_context(loc)
    results.append(res)
    _ods_successors = None
    super().__init__(self.OPERATION_NAME, self._ODS_REGIONS, self._ODS_OPERAND_SEGMENTS, self._ODS_RESULT_SEGMENTS, 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 MBarrierArriveNocompleteOp(res=res, addr=addr, count=count, loc=loc, ip=ip).result

@_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(addr)
    operands.append(count)
    _ods_context = _ods_get_default_loc_context(loc)
    results.append(res)
    _ods_successors = None
    super().__init__(self.OPERATION_NAME, self._ODS_REGIONS, self._ODS_OPERAND_SEGMENTS, self._ODS_RESULT_SEGMENTS, 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 MBarrierArriveNocompleteSharedOp(res=res, addr=addr, count=count, loc=loc, ip=ip).result

@_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(addr)
    _ods_context = _ods_get_default_loc_context(loc)
    results.append(res)
    _ods_successors = None
    super().__init__(self.OPERATION_NAME, self._ODS_REGIONS, self._ODS_OPERAND_SEGMENTS, self._ODS_RESULT_SEGMENTS, 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 MBarrierArriveOp(res=res, addr=addr, loc=loc, ip=ip).result

@_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(addr)
    _ods_context = _ods_get_default_loc_context(loc)
    results.append(res)
    _ods_successors = None
    super().__init__(self.OPERATION_NAME, self._ODS_REGIONS, self._ODS_OPERAND_SEGMENTS, self._ODS_RESULT_SEGMENTS, 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 MBarrierArriveSharedOp(res=res, addr=addr, loc=loc, ip=ip).result

@_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(addr)
    operands.append(count)
    if predicate is not None: operands.append(predicate)
    _ods_context = _ods_get_default_loc_context(loc)
    _ods_successors = None
    super().__init__(self.OPERATION_NAME, self._ODS_REGIONS, self._ODS_OPERAND_SEGMENTS, self._ODS_RESULT_SEGMENTS, 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 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(addr)
    operands.append(count)
    if predicate is not None: operands.append(predicate)
    _ods_context = _ods_get_default_loc_context(loc)
    _ods_successors = None
    super().__init__(self.OPERATION_NAME, self._ODS_REGIONS, self._ODS_OPERAND_SEGMENTS, self._ODS_RESULT_SEGMENTS, 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 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(addr)
    _ods_context = _ods_get_default_loc_context(loc)
    _ods_successors = None
    super().__init__(self.OPERATION_NAME, self._ODS_REGIONS, self._ODS_OPERAND_SEGMENTS, self._ODS_RESULT_SEGMENTS, 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 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(addr)
    _ods_context = _ods_get_default_loc_context(loc)
    _ods_successors = None
    super().__init__(self.OPERATION_NAME, self._ODS_REGIONS, self._ODS_OPERAND_SEGMENTS, self._ODS_RESULT_SEGMENTS, 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 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(addr)
    operands.append(state)
    _ods_context = _ods_get_default_loc_context(loc)
    results.append(res)
    _ods_successors = None
    super().__init__(self.OPERATION_NAME, self._ODS_REGIONS, self._ODS_OPERAND_SEGMENTS, self._ODS_RESULT_SEGMENTS, 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 MBarrierTestWaitOp(res=res, addr=addr, state=state, loc=loc, ip=ip).result

@_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(addr)
    operands.append(state)
    _ods_context = _ods_get_default_loc_context(loc)
    results.append(res)
    _ods_successors = None
    super().__init__(self.OPERATION_NAME, self._ODS_REGIONS, self._ODS_OPERAND_SEGMENTS, self._ODS_RESULT_SEGMENTS, 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 MBarrierTestWaitSharedOp(res=res, addr=addr, state=state, loc=loc, ip=ip).result

@_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(addr)
    operands.append(phase)
    operands.append(ticks)
    _ods_context = _ods_get_default_loc_context(loc)
    _ods_successors = None
    super().__init__(self.OPERATION_NAME, self._ODS_REGIONS, self._ODS_OPERAND_SEGMENTS, self._ODS_RESULT_SEGMENTS, 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 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(addr)
    operands.append(phase)
    operands.append(ticks)
    _ods_context = _ods_get_default_loc_context(loc)
    _ods_successors = None
    super().__init__(self.OPERATION_NAME, self._ODS_REGIONS, self._ODS_OPERAND_SEGMENTS, self._ODS_RESULT_SEGMENTS, 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 MBarrierTryWaitParitySharedOp(addr=addr, phase=phase, ticks=ticks, loc=loc, ip=ip)

@_ods_cext.register_operation(_Dialect)
class MapaOp(_ods_ir.OpView):
  OPERATION_NAME = "nvvm.mapa"

  _ODS_REGIONS = (0, True)

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

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

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

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

def mapa(res, a, b, *, loc=None, ip=None) -> _ods_ir.Value:
  return MapaOp(res=res, a=a, b=b, loc=loc, ip=ip).result

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

  _ODS_REGIONS = (0, True)

  def __init__(self, res, thread_mask, val, kind, *, loc=None, ip=None):
    operands = []
    results = []
    attributes = {}
    regions = None
    operands.append(thread_mask)
    operands.append(val)
    _ods_context = _ods_get_default_loc_context(loc)
    attributes["kind"] = (kind if (
    isinstance(kind, _ods_ir.Attribute) or
    not _ods_ir.AttrBuilder.contains('MatchSyncKindAttr')) else
      _ods_ir.AttrBuilder.get('MatchSyncKindAttr')(kind, context=_ods_context))
    results.append(res)
    _ods_successors = None
    super().__init__(self.OPERATION_NAME, self._ODS_REGIONS, self._ODS_OPERAND_SEGMENTS, self._ODS_RESULT_SEGMENTS, 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 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 match_sync(res, thread_mask, val, kind, *, loc=None, ip=None) -> _ods_ir.Value:
  return MatchSyncOp(res=res, thread_mask=thread_mask, val=val, kind=kind, loc=loc, ip=ip).result

@_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.OPERATION_NAME, self._ODS_REGIONS, self._ODS_OPERAND_SEGMENTS, self._ODS_RESULT_SEGMENTS, 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 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).result

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

  _ODS_REGIONS = (0, True)

  def __init__(self, cacheLevel, addr, *, uniform=None, evictPriority=None, loc=None, ip=None):
    operands = []
    results = []
    attributes = {}
    regions = None
    operands.append(addr)
    _ods_context = _ods_get_default_loc_context(loc)
    attributes["cacheLevel"] = (cacheLevel if (
    isinstance(cacheLevel, _ods_ir.Attribute) or
    not _ods_ir.AttrBuilder.contains('PrefetchCacheLevelAttr')) else
      _ods_ir.AttrBuilder.get('PrefetchCacheLevelAttr')(cacheLevel, context=_ods_context))
    if bool(uniform): attributes["uniform"] = _ods_ir.UnitAttr.get(
      _ods_get_default_loc_context(loc))
    if evictPriority is not None: attributes["evictPriority"] = (evictPriority if (
        isinstance(evictPriority, _ods_ir.Attribute) or
        not _ods_ir.AttrBuilder.contains('CacheEvictionPriorityAttr')) else
          _ods_ir.AttrBuilder.get('CacheEvictionPriorityAttr')(evictPriority, context=_ods_context))
    _ods_successors = None
    super().__init__(self.OPERATION_NAME, self._ODS_REGIONS, self._ODS_OPERAND_SEGMENTS, self._ODS_RESULT_SEGMENTS, 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 cacheLevel(self):
    return self.operation.attributes["cacheLevel"]

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

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

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

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

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

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

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

def prefetch(cache_level, addr, *, uniform=None, evict_priority=None, loc=None, ip=None) -> _ods_ir.Operation:
  return PrefetchOp(cacheLevel=cache_level, addr=addr, uniform=uniform, evictPriority=evict_priority, 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(tmaDescriptor)
    if predicate is not None: operands.append(predicate)
    _ods_context = _ods_get_default_loc_context(loc)
    _ods_successors = None
    super().__init__(self.OPERATION_NAME, self._ODS_REGIONS, self._ODS_OPERAND_SEGMENTS, self._ODS_RESULT_SEGMENTS, 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 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(arg)
    _ods_context = _ods_get_default_loc_context(loc)
    results.append(res)
    _ods_successors = None
    super().__init__(self.OPERATION_NAME, self._ODS_REGIONS, self._ODS_OPERAND_SEGMENTS, self._ODS_RESULT_SEGMENTS, 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 RcpApproxFtzF32Op(res=res, arg=arg, loc=loc, ip=ip).result

@_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, *, abs=None, nan=None, loc=None, ip=None):
    operands = []
    results = []
    attributes = {}
    regions = None
    operands.append(val)
    operands.append(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))
    if abs is not None: attributes["abs"] = (abs if (
        isinstance(abs, _ods_ir.Attribute) or
        not _ods_ir.AttrBuilder.contains('BoolAttr')) else
          _ods_ir.AttrBuilder.get('BoolAttr')(abs, context=_ods_context))
    if nan is not None: attributes["nan"] = (nan if (
        isinstance(nan, _ods_ir.Attribute) or
        not _ods_ir.AttrBuilder.contains('BoolAttr')) else
          _ods_ir.AttrBuilder.get('BoolAttr')(nan, context=_ods_context))
    results.append(res)
    _ods_successors = None
    super().__init__(self.OPERATION_NAME, self._ODS_REGIONS, self._ODS_OPERAND_SEGMENTS, self._ODS_RESULT_SEGMENTS, 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 abs(self):
    return self.operation.attributes["abs"]

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

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

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

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

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

@_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.OPERATION_NAME, self._ODS_REGIONS, self._ODS_OPERAND_SEGMENTS, self._ODS_RESULT_SEGMENTS, 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 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(thread_mask)
    operands.append(val)
    operands.append(offset)
    operands.append(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.OPERATION_NAME, self._ODS_REGIONS, self._ODS_OPERAND_SEGMENTS, self._ODS_RESULT_SEGMENTS, 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 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).result

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

  _ODS_REGIONS = (0, True)

  def __init__(self, res, *, range=None, loc=None, ip=None):
    operands = []
    results = []
    attributes = {}
    regions = None
    _ods_context = _ods_get_default_loc_context(loc)
    if range is not None: attributes["range"] = (range if (
        isinstance(range, _ods_ir.Attribute) or
        not _ods_ir.AttrBuilder.contains('LLVM_ConstantRangeAttr')) else
          _ods_ir.AttrBuilder.get('LLVM_ConstantRangeAttr')(range, context=_ods_context))
    results.append(res)
    _ods_successors = None
    super().__init__(self.OPERATION_NAME, self._ODS_REGIONS, self._ODS_OPERAND_SEGMENTS, self._ODS_RESULT_SEGMENTS, attributes=attributes, results=results, operands=operands, successors=_ods_successors, regions=regions, loc=loc, ip=ip)

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

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

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

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

def read_ptx_sreg_nsmid(res, *, range=None, loc=None, ip=None) -> _ods_ir.Value:
  return SmDimOp(res=res, range=range, loc=loc, ip=ip).result

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

  _ODS_REGIONS = (0, True)

  def __init__(self, res, *, range=None, loc=None, ip=None):
    operands = []
    results = []
    attributes = {}
    regions = None
    _ods_context = _ods_get_default_loc_context(loc)
    if range is not None: attributes["range"] = (range if (
        isinstance(range, _ods_ir.Attribute) or
        not _ods_ir.AttrBuilder.contains('LLVM_ConstantRangeAttr')) else
          _ods_ir.AttrBuilder.get('LLVM_ConstantRangeAttr')(range, context=_ods_context))
    results.append(res)
    _ods_successors = None
    super().__init__(self.OPERATION_NAME, self._ODS_REGIONS, self._ODS_OPERAND_SEGMENTS, self._ODS_RESULT_SEGMENTS, attributes=attributes, results=results, operands=operands, successors=_ods_successors, regions=regions, loc=loc, ip=ip)

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

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

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

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

def read_ptx_sreg_smid(res, *, range=None, loc=None, ip=None) -> _ods_ir.Value:
  return SmIdOp(res=res, range=range, loc=loc, ip=ip).result

@_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(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.OPERATION_NAME, self._ODS_REGIONS, self._ODS_OPERAND_SEGMENTS, self._ODS_RESULT_SEGMENTS, 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 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(mask)
    _ods_context = _ods_get_default_loc_context(loc)
    _ods_successors = None
    super().__init__(self.OPERATION_NAME, self._ODS_REGIONS, self._ODS_OPERAND_SEGMENTS, self._ODS_RESULT_SEGMENTS, 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 SyncWarpOp(mask=mask, loc=loc, ip=ip)

@_ods_cext.register_operation(_Dialect)
class Tcgen05AllocOp(_ods_ir.OpView):
  OPERATION_NAME = "nvvm.tcgen05.alloc"

  _ODS_REGIONS = (0, True)

  def __init__(self, addr, nCols, *, group=None, loc=None, ip=None):
    operands = []
    results = []
    attributes = {}
    regions = None
    operands.append(addr)
    operands.append(nCols)
    _ods_context = _ods_get_default_loc_context(loc)
    if group is not None: attributes["group"] = (group if (
        isinstance(group, _ods_ir.Attribute) or
        not _ods_ir.AttrBuilder.contains('Tcgen05GroupKindAttr')) else
          _ods_ir.AttrBuilder.get('Tcgen05GroupKindAttr')(group, context=_ods_context))
    _ods_successors = None
    super().__init__(self.OPERATION_NAME, self._ODS_REGIONS, self._ODS_OPERAND_SEGMENTS, self._ODS_RESULT_SEGMENTS, 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 nCols(self):
    return self.operation.operands[1]

  @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 tcgen05_alloc(addr, n_cols, *, group=None, loc=None, ip=None) -> _ods_ir.Operation:
  return Tcgen05AllocOp(addr=addr, nCols=n_cols, group=group, loc=loc, ip=ip)

@_ods_cext.register_operation(_Dialect)
class Tcgen05CommitOp(_ods_ir.OpView):
  OPERATION_NAME = "nvvm.tcgen05.commit"

  _ODS_REGIONS = (0, True)

  def __init__(self, addr, *, multicastMask=None, group=None, loc=None, ip=None):
    operands = []
    results = []
    attributes = {}
    regions = None
    operands.append(addr)
    if multicastMask is not None: operands.append(multicastMask)
    _ods_context = _ods_get_default_loc_context(loc)
    if group is not None: attributes["group"] = (group if (
        isinstance(group, _ods_ir.Attribute) or
        not _ods_ir.AttrBuilder.contains('Tcgen05GroupKindAttr')) else
          _ods_ir.AttrBuilder.get('Tcgen05GroupKindAttr')(group, context=_ods_context))
    _ods_successors = None
    super().__init__(self.OPERATION_NAME, self._ODS_REGIONS, self._ODS_OPERAND_SEGMENTS, self._ODS_RESULT_SEGMENTS, 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 multicastMask(self):
    return None if len(self.operation.operands) < 2 else self.operation.operands[1]

  @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 tcgen05_commit(addr, *, multicast_mask=None, group=None, loc=None, ip=None) -> _ods_ir.Operation:
  return Tcgen05CommitOp(addr=addr, multicastMask=multicast_mask, group=group, loc=loc, ip=ip)

@_ods_cext.register_operation(_Dialect)
class Tcgen05CpOp(_ods_ir.OpView):
  OPERATION_NAME = "nvvm.tcgen05.cp"

  _ODS_REGIONS = (0, True)

  def __init__(self, shape, taddr, smem_desc, *, group=None, multicast=None, srcFormat=None, loc=None, ip=None):
    operands = []
    results = []
    attributes = {}
    regions = None
    operands.append(taddr)
    operands.append(smem_desc)
    _ods_context = _ods_get_default_loc_context(loc)
    attributes["shape"] = (shape if (
    isinstance(shape, _ods_ir.Attribute) or
    not _ods_ir.AttrBuilder.contains('Tcgen05CpShapeAttr')) else
      _ods_ir.AttrBuilder.get('Tcgen05CpShapeAttr')(shape, context=_ods_context))
    if group is not None: attributes["group"] = (group if (
        isinstance(group, _ods_ir.Attribute) or
        not _ods_ir.AttrBuilder.contains('Tcgen05GroupKindAttr')) else
          _ods_ir.AttrBuilder.get('Tcgen05GroupKindAttr')(group, context=_ods_context))
    if multicast is not None: attributes["multicast"] = (multicast if (
        isinstance(multicast, _ods_ir.Attribute) or
        not _ods_ir.AttrBuilder.contains('Tcgen05CpMulticastAttr')) else
          _ods_ir.AttrBuilder.get('Tcgen05CpMulticastAttr')(multicast, context=_ods_context))
    if srcFormat is not None: attributes["srcFormat"] = (srcFormat if (
        isinstance(srcFormat, _ods_ir.Attribute) or
        not _ods_ir.AttrBuilder.contains('Tcgen05CpSrcFormatAttr')) else
          _ods_ir.AttrBuilder.get('Tcgen05CpSrcFormatAttr')(srcFormat, context=_ods_context))
    _ods_successors = None
    super().__init__(self.OPERATION_NAME, self._ODS_REGIONS, self._ODS_OPERAND_SEGMENTS, self._ODS_RESULT_SEGMENTS, attributes=attributes, results=results, operands=operands, successors=_ods_successors, regions=regions, loc=loc, ip=ip)

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

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

  @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 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 multicast(self):
    return self.operation.attributes["multicast"]

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

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

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

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

def tcgen05_cp(shape, taddr, smem_desc, *, group=None, multicast=None, src_format=None, loc=None, ip=None) -> _ods_ir.Operation:
  return Tcgen05CpOp(shape=shape, taddr=taddr, smem_desc=smem_desc, group=group, multicast=multicast, srcFormat=src_format, loc=loc, ip=ip)

@_ods_cext.register_operation(_Dialect)
class Tcgen05DeallocOp(_ods_ir.OpView):
  OPERATION_NAME = "nvvm.tcgen05.dealloc"

  _ODS_REGIONS = (0, True)

  def __init__(self, taddr, nCols, *, group=None, loc=None, ip=None):
    operands = []
    results = []
    attributes = {}
    regions = None
    operands.append(taddr)
    operands.append(nCols)
    _ods_context = _ods_get_default_loc_context(loc)
    if group is not None: attributes["group"] = (group if (
        isinstance(group, _ods_ir.Attribute) or
        not _ods_ir.AttrBuilder.contains('Tcgen05GroupKindAttr')) else
          _ods_ir.AttrBuilder.get('Tcgen05GroupKindAttr')(group, context=_ods_context))
    _ods_successors = None
    super().__init__(self.OPERATION_NAME, self._ODS_REGIONS, self._ODS_OPERAND_SEGMENTS, self._ODS_RESULT_SEGMENTS, attributes=attributes, results=results, operands=operands, successors=_ods_successors, regions=regions, loc=loc, ip=ip)

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

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

  @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 tcgen05_dealloc(taddr, n_cols, *, group=None, loc=None, ip=None) -> _ods_ir.Operation:
  return Tcgen05DeallocOp(taddr=taddr, nCols=n_cols, group=group, loc=loc, ip=ip)

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

  _ODS_REGIONS = (0, True)

  def __init__(self, kind, *, 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('Tcgen05FenceKindAttr')) else
      _ods_ir.AttrBuilder.get('Tcgen05FenceKindAttr')(kind, context=_ods_context))
    _ods_successors = None
    super().__init__(self.OPERATION_NAME, self._ODS_REGIONS, self._ODS_OPERAND_SEGMENTS, self._ODS_RESULT_SEGMENTS, 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

def tcgen05_fence(kind, *, loc=None, ip=None) -> _ods_ir.Operation:
  return Tcgen05FenceOp(kind=kind, loc=loc, ip=ip)

@_ods_cext.register_operation(_Dialect)
class Tcgen05LdOp(_ods_ir.OpView):
  OPERATION_NAME = "nvvm.tcgen05.ld"

  _ODS_REGIONS = (0, True)

  def __init__(self, res, shape, tmemAddr, *, pack=None, offset=None, loc=None, ip=None):
    operands = []
    results = []
    attributes = {}
    regions = None
    operands.append(tmemAddr)
    if offset is not None: operands.append(offset)
    _ods_context = _ods_get_default_loc_context(loc)
    if bool(pack): attributes["pack"] = _ods_ir.UnitAttr.get(
      _ods_get_default_loc_context(loc))
    attributes["shape"] = (shape if (
    isinstance(shape, _ods_ir.Attribute) or
    not _ods_ir.AttrBuilder.contains('Tcgen05LdStShapeAttr')) else
      _ods_ir.AttrBuilder.get('Tcgen05LdStShapeAttr')(shape, context=_ods_context))
    results.append(res)
    _ods_successors = None
    super().__init__(self.OPERATION_NAME, self._ODS_REGIONS, self._ODS_OPERAND_SEGMENTS, self._ODS_RESULT_SEGMENTS, attributes=attributes, results=results, operands=operands, successors=_ods_successors, regions=regions, loc=loc, ip=ip)

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

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

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

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

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

  @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 res(self):
    return self.operation.results[0]

def tcgen05_ld(res, shape, tmem_addr, *, pack=None, offset=None, loc=None, ip=None) -> _ods_ir.Value:
  return Tcgen05LdOp(res=res, shape=shape, tmemAddr=tmem_addr, pack=pack, offset=offset, loc=loc, ip=ip).result

@_ods_cext.register_operation(_Dialect)
class Tcgen05MmaSmemDescOp(_ods_ir.OpView):
  OPERATION_NAME = "nvvm.tcgen05.mma_smem_desc"

  _ODS_REGIONS = (0, True)

  def __init__(self, res, startAddr, leadingDimOffset, strideDimOffset, baseOffset, leadingDimMode, swizzleMode, *, loc=None, ip=None):
    operands = []
    results = []
    attributes = {}
    regions = None
    operands.append(startAddr)
    operands.append(leadingDimOffset)
    operands.append(strideDimOffset)
    operands.append(baseOffset)
    operands.append(leadingDimMode)
    operands.append(swizzleMode)
    _ods_context = _ods_get_default_loc_context(loc)
    results.append(res)
    _ods_successors = None
    super().__init__(self.OPERATION_NAME, self._ODS_REGIONS, self._ODS_OPERAND_SEGMENTS, self._ODS_RESULT_SEGMENTS, attributes=attributes, results=results, operands=operands, successors=_ods_successors, regions=regions, loc=loc, ip=ip)

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

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

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

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

  @builtins.property
  def leadingDimMode(self):
    return self.operation.operands[4]

  @builtins.property
  def swizzleMode(self):
    return self.operation.operands[5]

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

def tcgen05_mma_smem_desc(res, start_addr, leading_dim_offset, stride_dim_offset, base_offset, leading_dim_mode, swizzle_mode, *, loc=None, ip=None) -> _ods_ir.Value:
  return Tcgen05MmaSmemDescOp(res=res, startAddr=start_addr, leadingDimOffset=leading_dim_offset, strideDimOffset=stride_dim_offset, baseOffset=base_offset, leadingDimMode=leading_dim_mode, swizzleMode=swizzle_mode, loc=loc, ip=ip).result

@_ods_cext.register_operation(_Dialect)
class Tcgen05RelinquishAllocPermitOp(_ods_ir.OpView):
  OPERATION_NAME = "nvvm.tcgen05.relinquish_alloc_permit"

  _ODS_REGIONS = (0, True)

  def __init__(self, *, group=None, loc=None, ip=None):
    operands = []
    results = []
    attributes = {}
    regions = None
    _ods_context = _ods_get_default_loc_context(loc)
    if group is not None: attributes["group"] = (group if (
        isinstance(group, _ods_ir.Attribute) or
        not _ods_ir.AttrBuilder.contains('Tcgen05GroupKindAttr')) else
          _ods_ir.AttrBuilder.get('Tcgen05GroupKindAttr')(group, context=_ods_context))
    _ods_successors = None
    super().__init__(self.OPERATION_NAME, self._ODS_REGIONS, self._ODS_OPERAND_SEGMENTS, self._ODS_RESULT_SEGMENTS, 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 tcgen05_relinquish_alloc_permit(*, group=None, loc=None, ip=None) -> _ods_ir.Operation:
  return Tcgen05RelinquishAllocPermitOp(group=group, loc=loc, ip=ip)

@_ods_cext.register_operation(_Dialect)
class Tcgen05ShiftOp(_ods_ir.OpView):
  OPERATION_NAME = "nvvm.tcgen05.shift"

  _ODS_REGIONS = (0, True)

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

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

  @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 tcgen05_shift(taddr, *, group=None, loc=None, ip=None) -> _ods_ir.Operation:
  return Tcgen05ShiftOp(taddr=taddr, group=group, loc=loc, ip=ip)

@_ods_cext.register_operation(_Dialect)
class Tcgen05StOp(_ods_ir.OpView):
  OPERATION_NAME = "nvvm.tcgen05.st"

  _ODS_REGIONS = (0, True)

  def __init__(self, shape, tmemAddr, val, *, unpack=None, offset=None, loc=None, ip=None):
    operands = []
    results = []
    attributes = {}
    regions = None
    operands.append(tmemAddr)
    operands.append(val)
    if offset is not None: operands.append(offset)
    _ods_context = _ods_get_default_loc_context(loc)
    if bool(unpack): attributes["unpack"] = _ods_ir.UnitAttr.get(
      _ods_get_default_loc_context(loc))
    attributes["shape"] = (shape if (
    isinstance(shape, _ods_ir.Attribute) or
    not _ods_ir.AttrBuilder.contains('Tcgen05LdStShapeAttr')) else
      _ods_ir.AttrBuilder.get('Tcgen05LdStShapeAttr')(shape, context=_ods_context))
    _ods_successors = None
    super().__init__(self.OPERATION_NAME, self._ODS_REGIONS, self._ODS_OPERAND_SEGMENTS, self._ODS_RESULT_SEGMENTS, attributes=attributes, results=results, operands=operands, successors=_ods_successors, regions=regions, loc=loc, ip=ip)

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

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

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

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

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

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

  @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

def tcgen05_st(shape, tmem_addr, val, *, unpack=None, offset=None, loc=None, ip=None) -> _ods_ir.Operation:
  return Tcgen05StOp(shape=shape, tmemAddr=tmem_addr, val=val, unpack=unpack, offset=offset, loc=loc, ip=ip)

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

  _ODS_REGIONS = (0, True)

  def __init__(self, kind, *, 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('Tcgen05WaitKindAttr')) else
      _ods_ir.AttrBuilder.get('Tcgen05WaitKindAttr')(kind, context=_ods_context))
    _ods_successors = None
    super().__init__(self.OPERATION_NAME, self._ODS_REGIONS, self._ODS_OPERAND_SEGMENTS, self._ODS_RESULT_SEGMENTS, 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

def tcgen05_wait(kind, *, loc=None, ip=None) -> _ods_ir.Operation:
  return Tcgen05WaitOp(kind=kind, 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, *, range=None, loc=None, ip=None):
    operands = []
    results = []
    attributes = {}
    regions = None
    _ods_context = _ods_get_default_loc_context(loc)
    if range is not None: attributes["range"] = (range if (
        isinstance(range, _ods_ir.Attribute) or
        not _ods_ir.AttrBuilder.contains('LLVM_ConstantRangeAttr')) else
          _ods_ir.AttrBuilder.get('LLVM_ConstantRangeAttr')(range, context=_ods_context))
    results.append(res)
    _ods_successors = None
    super().__init__(self.OPERATION_NAME, self._ODS_REGIONS, self._ODS_OPERAND_SEGMENTS, self._ODS_RESULT_SEGMENTS, attributes=attributes, results=results, operands=operands, successors=_ods_successors, regions=regions, loc=loc, ip=ip)

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

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

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

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

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

@_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, *, range=None, loc=None, ip=None):
    operands = []
    results = []
    attributes = {}
    regions = None
    _ods_context = _ods_get_default_loc_context(loc)
    if range is not None: attributes["range"] = (range if (
        isinstance(range, _ods_ir.Attribute) or
        not _ods_ir.AttrBuilder.contains('LLVM_ConstantRangeAttr')) else
          _ods_ir.AttrBuilder.get('LLVM_ConstantRangeAttr')(range, context=_ods_context))
    results.append(res)
    _ods_successors = None
    super().__init__(self.OPERATION_NAME, self._ODS_REGIONS, self._ODS_OPERAND_SEGMENTS, self._ODS_RESULT_SEGMENTS, attributes=attributes, results=results, operands=operands, successors=_ods_successors, regions=regions, loc=loc, ip=ip)

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

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

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

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

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

@_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, *, range=None, loc=None, ip=None):
    operands = []
    results = []
    attributes = {}
    regions = None
    _ods_context = _ods_get_default_loc_context(loc)
    if range is not None: attributes["range"] = (range if (
        isinstance(range, _ods_ir.Attribute) or
        not _ods_ir.AttrBuilder.contains('LLVM_ConstantRangeAttr')) else
          _ods_ir.AttrBuilder.get('LLVM_ConstantRangeAttr')(range, context=_ods_context))
    results.append(res)
    _ods_successors = None
    super().__init__(self.OPERATION_NAME, self._ODS_REGIONS, self._ODS_OPERAND_SEGMENTS, self._ODS_RESULT_SEGMENTS, attributes=attributes, results=results, operands=operands, successors=_ods_successors, regions=regions, loc=loc, ip=ip)

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

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

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

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

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

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

  _ODS_REGIONS = (0, True)

  def __init__(self, res, mask, pred, kind, *, loc=None, ip=None):
    operands = []
    results = []
    attributes = {}
    regions = None
    operands.append(mask)
    operands.append(pred)
    _ods_context = _ods_get_default_loc_context(loc)
    attributes["kind"] = (kind if (
    isinstance(kind, _ods_ir.Attribute) or
    not _ods_ir.AttrBuilder.contains('VoteSyncKindAttr')) else
      _ods_ir.AttrBuilder.get('VoteSyncKindAttr')(kind, context=_ods_context))
    results.append(res)
    _ods_successors = None
    super().__init__(self.OPERATION_NAME, self._ODS_REGIONS, self._ODS_OPERAND_SEGMENTS, self._ODS_RESULT_SEGMENTS, 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 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 vote_sync(res, mask, pred, kind, *, loc=None, ip=None) -> _ods_ir.Value:
  return VoteSyncOp(res=res, mask=mask, pred=pred, kind=kind, loc=loc, ip=ip).result

@_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(ptr)
    operands.append(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.OPERATION_NAME, self._ODS_REGIONS, self._ODS_OPERAND_SEGMENTS, self._ODS_RESULT_SEGMENTS, 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 WMMALoadOp(res=res, ptr=ptr, stride=stride, m=m, n=n, k=k, layout=layout, eltype=eltype, frag=frag, loc=loc, ip=ip).result

@_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.OPERATION_NAME, self._ODS_REGIONS, self._ODS_OPERAND_SEGMENTS, self._ODS_RESULT_SEGMENTS, 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 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).result

@_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(ptr)
    operands.extend(_get_op_results_or_values(args))
    operands.append(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.OPERATION_NAME, self._ODS_REGIONS, self._ODS_OPERAND_SEGMENTS, self._ODS_RESULT_SEGMENTS, 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 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 WarpDimOp(_ods_ir.OpView):
  OPERATION_NAME = "nvvm.read.ptx.sreg.nwarpid"

  _ODS_REGIONS = (0, True)

  def __init__(self, res, *, range=None, loc=None, ip=None):
    operands = []
    results = []
    attributes = {}
    regions = None
    _ods_context = _ods_get_default_loc_context(loc)
    if range is not None: attributes["range"] = (range if (
        isinstance(range, _ods_ir.Attribute) or
        not _ods_ir.AttrBuilder.contains('LLVM_ConstantRangeAttr')) else
          _ods_ir.AttrBuilder.get('LLVM_ConstantRangeAttr')(range, context=_ods_context))
    results.append(res)
    _ods_successors = None
    super().__init__(self.OPERATION_NAME, self._ODS_REGIONS, self._ODS_OPERAND_SEGMENTS, self._ODS_RESULT_SEGMENTS, attributes=attributes, results=results, operands=operands, successors=_ods_successors, regions=regions, loc=loc, ip=ip)

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

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

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

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

def read_ptx_sreg_nwarpid(res, *, range=None, loc=None, ip=None) -> _ods_ir.Value:
  return WarpDimOp(res=res, range=range, loc=loc, ip=ip).result

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

  _ODS_REGIONS = (0, True)

  def __init__(self, res, *, range=None, loc=None, ip=None):
    operands = []
    results = []
    attributes = {}
    regions = None
    _ods_context = _ods_get_default_loc_context(loc)
    if range is not None: attributes["range"] = (range if (
        isinstance(range, _ods_ir.Attribute) or
        not _ods_ir.AttrBuilder.contains('LLVM_ConstantRangeAttr')) else
          _ods_ir.AttrBuilder.get('LLVM_ConstantRangeAttr')(range, context=_ods_context))
    results.append(res)
    _ods_successors = None
    super().__init__(self.OPERATION_NAME, self._ODS_REGIONS, self._ODS_OPERAND_SEGMENTS, self._ODS_RESULT_SEGMENTS, attributes=attributes, results=results, operands=operands, successors=_ods_successors, regions=regions, loc=loc, ip=ip)

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

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

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

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

def read_ptx_sreg_warpid(res, *, range=None, loc=None, ip=None) -> _ods_ir.Value:
  return WarpIdOp(res=res, range=range, loc=loc, ip=ip).result

@_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, *, range=None, loc=None, ip=None):
    operands = []
    results = []
    attributes = {}
    regions = None
    _ods_context = _ods_get_default_loc_context(loc)
    if range is not None: attributes["range"] = (range if (
        isinstance(range, _ods_ir.Attribute) or
        not _ods_ir.AttrBuilder.contains('LLVM_ConstantRangeAttr')) else
          _ods_ir.AttrBuilder.get('LLVM_ConstantRangeAttr')(range, context=_ods_context))
    results.append(res)
    _ods_successors = None
    super().__init__(self.OPERATION_NAME, self._ODS_REGIONS, self._ODS_OPERAND_SEGMENTS, self._ODS_RESULT_SEGMENTS, attributes=attributes, results=results, operands=operands, successors=_ods_successors, regions=regions, loc=loc, ip=ip)

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

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

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

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

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

@_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.OPERATION_NAME, self._ODS_REGIONS, self._ODS_OPERAND_SEGMENTS, self._ODS_RESULT_SEGMENTS, 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 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.OPERATION_NAME, self._ODS_REGIONS, self._ODS_OPERAND_SEGMENTS, self._ODS_RESULT_SEGMENTS, 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 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(inouts)
    operands.append(descriptorA)
    operands.append(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.OPERATION_NAME, self._ODS_REGIONS, self._ODS_OPERAND_SEGMENTS, self._ODS_RESULT_SEGMENTS, 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 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).result

@_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('I64Attr')) else
      _ods_ir.AttrBuilder.get('I64Attr')(group, context=_ods_context))
    _ods_successors = None
    super().__init__(self.OPERATION_NAME, self._ODS_REGIONS, self._ODS_OPERAND_SEGMENTS, self._ODS_RESULT_SEGMENTS, 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 WgmmaWaitGroupSyncOp(group=group, loc=loc, ip=ip)