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

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

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


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

@_ods_cext.register_operation(_Dialect)
class DeviceAsyncCopyOp(_ods_ir.OpView):
  OPERATION_NAME = "nvgpu.device_async_copy"

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

  _ODS_REGIONS = (0, True)

  def __init__(self, asyncToken, dst, dstIndices, src, srcIndices, dstElements, *, srcElements=None, bypassL1=None, loc=None, ip=None):
    operands = []
    results = []
    attributes = {}
    regions = None
    operands.append(_get_op_result_or_value(dst))
    operands.append(_get_op_results_or_values(dstIndices))
    operands.append(_get_op_result_or_value(src))
    operands.append(_get_op_results_or_values(srcIndices))
    operands.append(_get_op_result_or_value(srcElements) if srcElements is not None else None)
    _ods_context = _ods_get_default_loc_context(loc)
    attributes["dstElements"] = (dstElements if (
    isinstance(dstElements, _ods_ir.Attribute) or
    not _ods_ir.AttrBuilder.contains('IndexAttr')) else
      _ods_ir.AttrBuilder.get('IndexAttr')(dstElements, context=_ods_context))
    if bool(bypassL1): attributes["bypassL1"] = _ods_ir.UnitAttr.get(
      _ods_get_default_loc_context(loc))
    results.append(asyncToken)
    _ods_successors = None
    super().__init__(self.build_generic(attributes=attributes, results=results, operands=operands, successors=_ods_successors, regions=regions, loc=loc, ip=ip))

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

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

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

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

  @builtins.property
  def srcElements(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 dstElements(self):
    return self.operation.attributes["dstElements"]

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

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

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

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

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

def device_async_copy(async_token, dst, dst_indices, src, src_indices, dst_elements, *, src_elements=None, bypass_l1=None, loc=None, ip=None) -> _ods_ir.Value:
  return _get_op_result_or_op_results(DeviceAsyncCopyOp(asyncToken=async_token, dst=dst, dstIndices=dst_indices, src=src, srcIndices=src_indices, dstElements=dst_elements, srcElements=src_elements, bypassL1=bypass_l1, loc=loc, ip=ip))

@_ods_cext.register_operation(_Dialect)
class DeviceAsyncCreateGroupOp(_ods_ir.OpView):
  OPERATION_NAME = "nvgpu.device_async_create_group"

  _ODS_REGIONS = (0, True)

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

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

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

def device_async_create_group(async_token, input_tokens, *, loc=None, ip=None) -> _ods_ir.Value:
  return _get_op_result_or_op_results(DeviceAsyncCreateGroupOp(asyncToken=async_token, inputTokens=input_tokens, loc=loc, ip=ip))

@_ods_cext.register_operation(_Dialect)
class DeviceAsyncWaitOp(_ods_ir.OpView):
  OPERATION_NAME = "nvgpu.device_async_wait"

  _ODS_REGIONS = (0, True)

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

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

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

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

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

def device_async_wait(async_dependencies, *, num_groups=None, loc=None, ip=None) -> _ods_ir.Operation:
  return _get_op_result_or_op_results(DeviceAsyncWaitOp(asyncDependencies=async_dependencies, numGroups=num_groups, loc=loc, ip=ip))

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

  _ODS_REGIONS = (0, True)

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

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

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

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

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

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

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

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

def ldmatrix(res, src_memref, indices, transpose, num_tiles, *, loc=None, ip=None) -> _ods_ir.Value:
  return _get_op_result_or_op_results(LdMatrixOp(res=res, srcMemref=src_memref, indices=indices, transpose=transpose, numTiles=num_tiles, loc=loc, ip=ip))

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

  _ODS_REGIONS = (0, True)

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

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

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

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

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

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

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

  _ODS_REGIONS = (0, True)

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

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

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

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

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

def mbarrier_arrive_nocomplete(token, barriers, mbar_id, count, *, loc=None, ip=None) -> _ods_ir.Value:
  return _get_op_result_or_op_results(MBarrierArriveNoCompleteOp(token=token, barriers=barriers, mbarId=mbar_id, count=count, loc=loc, ip=ip))

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

  _ODS_REGIONS = (0, True)

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

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

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

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

def mbarrier_arrive(token, barriers, mbar_id, *, loc=None, ip=None) -> _ods_ir.Value:
  return _get_op_result_or_op_results(MBarrierArriveOp(token=token, barriers=barriers, mbarId=mbar_id, loc=loc, ip=ip))

@_ods_cext.register_operation(_Dialect)
class MBarrierCreateOp(_ods_ir.OpView):
  OPERATION_NAME = "nvgpu.mbarrier.create"

  _ODS_REGIONS = (0, True)

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

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

def mbarrier_create(barriers, *, loc=None, ip=None) -> _ods_ir.Value:
  return _get_op_result_or_op_results(MBarrierCreateOp(barriers=barriers, loc=loc, ip=ip))

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

  _ODS_REGIONS = (0, True)

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

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

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

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

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

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

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

  _ODS_REGIONS = (0, True)

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

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

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

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

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

def mbarrier_test_wait(wait_complete, barriers, token, mbar_id, *, loc=None, ip=None) -> _ods_ir.Value:
  return _get_op_result_or_op_results(MBarrierTestWaitOp(waitComplete=wait_complete, barriers=barriers, token=token, mbarId=mbar_id, loc=loc, ip=ip))

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

  _ODS_REGIONS = (0, True)

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

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

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

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

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

def mbarrier_try_wait_parity(barriers, phase_parity, ticks, mbar_id, *, loc=None, ip=None) -> _ods_ir.Operation:
  return _get_op_result_or_op_results(MBarrierTryWaitParityOp(barriers=barriers, phaseParity=phase_parity, ticks=ticks, mbarId=mbar_id, loc=loc, ip=ip))

@_ods_cext.register_operation(_Dialect)
class MmaSparseSyncOp(_ods_ir.OpView):
  OPERATION_NAME = "nvgpu.mma.sp.sync"

  _ODS_REGIONS = (0, True)

  def __init__(self, res, matrixA, matrixB, matrixC, sparseMetadata, mmaShape, *, sparsitySelector=None, tf32Enabled=None, loc=None, ip=None):
    operands = []
    results = []
    attributes = {}
    regions = None
    operands.append(_get_op_result_or_value(matrixA))
    operands.append(_get_op_result_or_value(matrixB))
    operands.append(_get_op_result_or_value(matrixC))
    operands.append(_get_op_result_or_value(sparseMetadata))
    _ods_context = _ods_get_default_loc_context(loc)
    attributes["mmaShape"] = (mmaShape if (
    isinstance(mmaShape, _ods_ir.Attribute) or
    not _ods_ir.AttrBuilder.contains('I64ArrayAttr')) else
      _ods_ir.AttrBuilder.get('I64ArrayAttr')(mmaShape, context=_ods_context))
    if sparsitySelector is not None: attributes["sparsitySelector"] = (sparsitySelector if (
        isinstance(sparsitySelector, _ods_ir.Attribute) or
        not _ods_ir.AttrBuilder.contains('I32Attr')) else
          _ods_ir.AttrBuilder.get('I32Attr')(sparsitySelector, context=_ods_context))
    if bool(tf32Enabled): attributes["tf32Enabled"] = _ods_ir.UnitAttr.get(
      _ods_get_default_loc_context(loc))
    results.append(res)
    _ods_successors = None
    super().__init__(self.build_generic(attributes=attributes, results=results, operands=operands, successors=_ods_successors, regions=regions, loc=loc, ip=ip))

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

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

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

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

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

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

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

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

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

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

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

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

def mma_sp_sync(res, matrix_a, matrix_b, matrix_c, sparse_metadata, mma_shape, *, sparsity_selector=None, tf32_enabled=None, loc=None, ip=None) -> _ods_ir.Value:
  return _get_op_result_or_op_results(MmaSparseSyncOp(res=res, matrixA=matrix_a, matrixB=matrix_b, matrixC=matrix_c, sparseMetadata=sparse_metadata, mmaShape=mma_shape, sparsitySelector=sparsity_selector, tf32Enabled=tf32_enabled, loc=loc, ip=ip))

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

  _ODS_REGIONS = (0, True)

  def __init__(self, res, matrixA, matrixB, matrixC, mmaShape, *, tf32Enabled=None, loc=None, ip=None):
    operands = []
    results = []
    attributes = {}
    regions = None
    operands.append(_get_op_result_or_value(matrixA))
    operands.append(_get_op_result_or_value(matrixB))
    operands.append(_get_op_result_or_value(matrixC))
    _ods_context = _ods_get_default_loc_context(loc)
    attributes["mmaShape"] = (mmaShape if (
    isinstance(mmaShape, _ods_ir.Attribute) or
    not _ods_ir.AttrBuilder.contains('I64ArrayAttr')) else
      _ods_ir.AttrBuilder.get('I64ArrayAttr')(mmaShape, context=_ods_context))
    if bool(tf32Enabled): attributes["tf32Enabled"] = _ods_ir.UnitAttr.get(
      _ods_get_default_loc_context(loc))
    results.append(res)
    _ods_successors = None
    super().__init__(self.build_generic(attributes=attributes, results=results, operands=operands, successors=_ods_successors, regions=regions, loc=loc, ip=ip))

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

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

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

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

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

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

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

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

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

def mma_sync(res, matrix_a, matrix_b, matrix_c, mma_shape, *, tf32_enabled=None, loc=None, ip=None) -> _ods_ir.Value:
  return _get_op_result_or_op_results(MmaSyncOp(res=res, matrixA=matrix_a, matrixB=matrix_b, matrixC=matrix_c, mmaShape=mma_shape, tf32Enabled=tf32_enabled, loc=loc, ip=ip))

@_ods_cext.register_operation(_Dialect)
class TmaAsyncLoadOp(_ods_ir.OpView):
  OPERATION_NAME = "nvgpu.tma.async.load"

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

  _ODS_REGIONS = (0, True)

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

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

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

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

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

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

  @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 predicate(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

def tma_async_load(dst, barriers, tensor_map_descriptor, coordinates, mbar_id, *, multicast_mask=None, predicate=None, loc=None, ip=None) -> _ods_ir.Operation:
  return _get_op_result_or_op_results(TmaAsyncLoadOp(dst=dst, barriers=barriers, tensorMapDescriptor=tensor_map_descriptor, coordinates=coordinates, mbarId=mbar_id, multicastMask=multicast_mask, predicate=predicate, loc=loc, ip=ip))

@_ods_cext.register_operation(_Dialect)
class TmaAsyncStoreOp(_ods_ir.OpView):
  OPERATION_NAME = "nvgpu.tma.async.store"

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

  _ODS_REGIONS = (0, True)

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

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

  @builtins.property
  def tensorMapDescriptor(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 tma_async_store(src, tensor_map_descriptor, coordinates, *, predicate=None, loc=None, ip=None) -> _ods_ir.Operation:
  return _get_op_result_or_op_results(TmaAsyncStoreOp(src=src, tensorMapDescriptor=tensor_map_descriptor, coordinates=coordinates, predicate=predicate, loc=loc, ip=ip))

@_ods_cext.register_operation(_Dialect)
class TmaCreateDescriptorOp(_ods_ir.OpView):
  OPERATION_NAME = "nvgpu.tma.create.descriptor"

  _ODS_REGIONS = (0, True)

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

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

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

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

def tma_create_descriptor(tensor_map, tensor, box_dimensions, *, loc=None, ip=None) -> _ods_ir.Value:
  return _get_op_result_or_op_results(TmaCreateDescriptorOp(tensorMap=tensor_map, tensor=tensor, boxDimensions=box_dimensions, loc=loc, ip=ip))

@_ods_cext.register_operation(_Dialect)
class TmaPrefetchOp(_ods_ir.OpView):
  OPERATION_NAME = "nvgpu.tma.prefetch.descriptor"

  _ODS_REGIONS = (0, True)

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

  @builtins.property
  def tensorMapDescriptor(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 tma_prefetch_descriptor(tensor_map_descriptor, *, predicate=None, loc=None, ip=None) -> _ods_ir.Operation:
  return _get_op_result_or_op_results(TmaPrefetchOp(tensorMapDescriptor=tensor_map_descriptor, predicate=predicate, loc=loc, ip=ip))

@_ods_cext.register_operation(_Dialect)
class WarpgroupGenerateDescriptorOp(_ods_ir.OpView):
  OPERATION_NAME = "nvgpu.warpgroup.generate.descriptor"

  _ODS_REGIONS = (0, True)

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

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

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

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

def warpgroup_generate_descriptor(descriptor, tensor, tensor_map, *, loc=None, ip=None) -> _ods_ir.Value:
  return _get_op_result_or_op_results(WarpgroupGenerateDescriptorOp(descriptor=descriptor, tensor=tensor, tensorMap=tensor_map, loc=loc, ip=ip))

@_ods_cext.register_operation(_Dialect)
class WarpgroupMmaInitAccumulatorOp(_ods_ir.OpView):
  OPERATION_NAME = "nvgpu.warpgroup.mma.init.accumulator"

  _ODS_REGIONS = (0, True)

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

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

def warpgroup_mma_init_accumulator(matrix_c, *, loc=None, ip=None) -> _ods_ir.Value:
  return _get_op_result_or_op_results(WarpgroupMmaInitAccumulatorOp(matrixC=matrix_c, loc=loc, ip=ip))

@_ods_cext.register_operation(_Dialect)
class WarpgroupMmaOp(_ods_ir.OpView):
  OPERATION_NAME = "nvgpu.warpgroup.mma"

  _ODS_REGIONS = (0, True)

  def __init__(self, matrixD, descriptorA, descriptorB, matrixC, *, waitGroup=None, transposeA=None, transposeB=None, loc=None, ip=None):
    operands = []
    results = []
    attributes = {}
    regions = None
    operands.append(_get_op_result_or_value(descriptorA))
    operands.append(_get_op_result_or_value(descriptorB))
    operands.append(_get_op_result_or_value(matrixC))
    _ods_context = _ods_get_default_loc_context(loc)
    if waitGroup is not None: attributes["waitGroup"] = (waitGroup if (
        isinstance(waitGroup, _ods_ir.Attribute) or
        not _ods_ir.AttrBuilder.contains('I32Attr')) else
          _ods_ir.AttrBuilder.get('I32Attr')(waitGroup, context=_ods_context))
    if bool(transposeA): attributes["transposeA"] = _ods_ir.UnitAttr.get(
      _ods_get_default_loc_context(loc))
    if bool(transposeB): attributes["transposeB"] = _ods_ir.UnitAttr.get(
      _ods_get_default_loc_context(loc))
    results.append(matrixD)
    _ods_successors = None
    super().__init__(self.build_generic(attributes=attributes, results=results, operands=operands, successors=_ods_successors, regions=regions, loc=loc, ip=ip))

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

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

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

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

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

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

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

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

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

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

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

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

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

def warpgroup_mma(matrix_d, descriptor_a, descriptor_b, matrix_c, *, wait_group=None, transpose_a=None, transpose_b=None, loc=None, ip=None) -> _ods_ir.Value:
  return _get_op_result_or_op_results(WarpgroupMmaOp(matrixD=matrix_d, descriptorA=descriptor_a, descriptorB=descriptor_b, matrixC=matrix_c, waitGroup=wait_group, transposeA=transpose_a, transposeB=transpose_b, loc=loc, ip=ip))

@_ods_cext.register_operation(_Dialect)
class WarpgroupMmaStoreOp(_ods_ir.OpView):
  OPERATION_NAME = "nvgpu.warpgroup.mma.store"

  _ODS_REGIONS = (0, True)

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

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

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

def warpgroup_mma_store(matrix_d, dst_memref, *, loc=None, ip=None) -> _ods_ir.Operation:
  return _get_op_result_or_op_results(WarpgroupMmaStoreOp(matrixD=matrix_d, dstMemref=dst_memref, loc=loc, ip=ip))