return false;
    }
  }
  return true;
}

// It returns "true" when Value $iota is obtained from the following mlir code:
//
// $iota = "mhlo.iota"(){iota_dimension = $dimensions[0]},
//
// where $dimensions must have size 1 and iota can have rank>=1.
// It usually used for matching rank 1 iota since the iotaOp will be folded to
// IotaOp + BroadCastInDimOp except for the case when result shape is rank 1.

namespace mlir {
namespace odml {
// The following 5 different forms of mhlo::iota will be matched:
// 1. IotaOp.
// 2. IotaOp + BroadCastInDim.
// 3. IotaOp + Reshape.
// 4. Constant (folded Iota) + BroadCastInDim.
// 5. Constant (folded result).
// Moreover, the dimensions has to match the iota_dimension.
bool MatchIota(DenseIntElementsAttr dimensions, Value iota);
}  // namespace odml
}  // namespace mlir

// so callers should just use these enums to indicate what they want to do.
type IptablesCmd int

const (
	IPTables        IptablesCmd = iota
	IPTablesSave    IptablesCmd = iota
	IPTablesRestore IptablesCmd = iota

    // Check that the iota init is zero.
    DenseElementsAttr iota_init;
    if (!matchPattern(reduce_op.getInitValues().back(), m_Constant(&iota_init)))
      return failure();
    if (iota_init.getValues<APInt>()[0] != 0) return failure();

    // Verify that the second argument is an Iota op along the same dimension
    // as the reduction.
    Value iota = reduce_op.getInputs().back();

type traceExperiment uint8

const (
	// traceNoExperiment indicates no experiment.
	traceNoExperiment traceExperiment = iota

	// traceExperimentAllocFree is an experiment to add alloc/free events to the trace.
	traceExperimentAllocFree
)

// Experimental events.
const (
	_ traceEv = 127 + iota

	// Experimental events for ExperimentAllocFree.

	// Experimental heap span events. IDs map reversibly to base addresses.

// Use of this source code is governed by a BSD-style
// license that can be found in the LICENSE file.

package types

// BasicKind describes the kind of basic type.
type BasicKind int

const (
	Invalid BasicKind = iota // type is invalid

	// predeclared types
	Bool
	Int
	Int8
	Int16
	Int32
	Int64
	Uint
	Uint8
	Uint16
	Uint32
	Uint64
	Uintptr
	Float32
	Float64
	Complex64
	Complex128

const (
	NSNAME = 8
	NSYM   = 50
	NREG   = 32 // number of general registers
	NFREG  = 32 // number of floating point registers
)

const (
	REG_R0 = obj.RBaseLOONG64 + iota // must be a multiple of 32
	REG_R1
	REG_R2
	REG_R3
	REG_R4
	REG_R5
	REG_R6
	REG_R7
	REG_R8
	REG_R9
	REG_R10
	REG_R11
	REG_R12
	REG_R13
	REG_R14
	REG_R15
	REG_R16
	REG_R17

func (check *Checker) constDecl(obj *Const, typ, init syntax.Expr, inherited bool) {
	assert(obj.typ == nil)

	// use the correct value of iota and errpos
	defer func(iota constant.Value, errpos syntax.Pos) {
		check.iota = iota
		check.errpos = errpos
	}(check.iota, check.errpos)
	check.iota = obj.val
	check.errpos = nopos

	// provide valid constant value under all circumstances
	obj.val = constant.MakeUnknown()

type SchedReqs struct {
	Thread    Constraint
	Proc      Constraint
	Goroutine Constraint
}

// Constraint represents a various presence requirements.
type Constraint uint8

const (
	MustNotHave Constraint = iota
	MayHave
	MustHave
)

// UserGoReqs is a common requirement among events that are running
// or are close to running user code.

func.func @convert_argmax(%arg0: tensor<4x32x256xf32>) -> (tensor<4x32xf32>, tensor<4x32xi32>) {
  %0 = mhlo.constant dense<0xFF800000> : tensor<f32>
  %1 = mhlo.constant dense<0> : tensor<i32>
  %2 = "mhlo.iota"() <{iota_dimension = 0 : i64}> : () -> tensor<256xi32>
  %3 = "mhlo.broadcast_in_dim"(%2) <{broadcast_dimensions = dense<2> : tensor<1xi64>}> : (tensor<256xi32>) -> tensor<4x32x256xi32>
  %4:2 = "mhlo.reduce"(%arg0, %3, %0, %1) ({