for i := ppc64.REG_R0; i <= ppc64.REG_R31; i++ {
		register[obj.Rconv(i)] = int16(i)
	}
	for i := ppc64.REG_F0; i <= ppc64.REG_F31; i++ {
		register[obj.Rconv(i)] = int16(i)
	}
	for i := ppc64.REG_V0; i <= ppc64.REG_V31; i++ {
		register[obj.Rconv(i)] = int16(i)
	}
	for i := ppc64.REG_VS0; i <= ppc64.REG_VS63; i++ {
		register[obj.Rconv(i)] = int16(i)
	}
	for i := ppc64.REG_A0; i <= ppc64.REG_A7; i++ {

				ir.NewIndexExpr(base.Pos, globals, ir.NewInt(base.Pos, int64(i))), lname(f)), val)
			init.Append(typecheck.Stmt(r))
		}
		// globals[i].beg = uintptr(unsafe.Pointer(&n))
		c = tconv(typecheck.NodAddr(n), types.Types[types.TUNSAFEPTR])
		c = tconv(c, types.Types[types.TUINTPTR])
		setField("beg", c, i)
		// Assign globals[i].size.
		g := n.(*ir.Name)
		size := g.Type().Size()

// support non-constant weights.
def ConvertTFConv2DToXLAConvOp : Pat<
  (TF_Conv2DOp:$conv
    (TF_SubOp (TF_CastOp $input, $truncate), $input_zp),
    (TF_CastOp (TF_IdentityOp $filter), $truncate1),
    $strides, $use_cudnn, $padding, $explicit_padding,
    IsDataFormatNHWC:$data_format, $dilations),
  (CreateXLAConvOpFromTFConv2DOp
    $input, $filter, $input_zp, $conv, $strides,
    $dilations, $padding, $explicit_padding),

    %zp_offset: tensor<?x2x2x1xi32>, %bias: tensor<1xi32>
  ) -> tensor<?x2x2x1xi32> {
  // CHECK-DAG: %[[conv:.*]] = mhlo.convolution
  // CHECK-DAG: %[[combined:.*]] = chlo.broadcast_add %[[zp_offset:.*]], %[[bias:.*]]
  // CHECK-DAG: %[[result:.*]] = chlo.broadcast_add %[[conv]], %[[combined]]
  // CHECK: return %[[result]]
  %0 = mhlo.convolution(%lhs, %rhs)
      dim_numbers = [b, 0, 1, f]x[0, 1, i, o]->[b, 0, 1, f],

def IsNull : Constraint<CPred<"!$0">>;

// This pattern optimizes:
//   conv/dot_general + a + b -> conv/dot_general + (a + b)
//   conv/dot_general - a - b -> conv/dot_general - (a + b)
//   conv/dot_general + a - b -> conv/dot_general + (a - b)
//   conv/dot_general - a + b -> conv/dot_general - (a - b)
foreach OpsTuple = [
    [CHLO_BroadcastAddOp, CHLO_BroadcastAddOp, CHLO_BroadcastAddOp],

Michael Levesque-Dion <******@****.***> 1706075999 -0800

// CHECK: %[[CONV:.+]] = stablehlo.convolution(%[[TRANSPOSE_0]], %[[CONST]]) dim_numbers = [b, 0, 1, f]x[0, 1, i, o]->[b, 0, 1, f], window = {pad = {{\[\[}}1, 1], [1, 1]]} {batch_group_count = 1 : i64, feature_group_count = 1 : i64} : (tensor<1x4x4x8xf32>, tensor<3x3x8x8xf32>) -> tensor<1x4x4x8xf32>
// CHECK: %[[TRANSPOSE_1:.+]] = stablehlo.transpose %[[CONV]], dims = [0, 3, 1, 2] : (tensor<1x4x4x8xf32>) -> tensor<1x8x4x4xf32>

	// if uint(newLen) <= uint(oldCap)
	nif := ir.NewIfStmt(base.Pos, nil, nil, nil)
	nuint := typecheck.Conv(newLen, types.Types[types.TUINT])
	scapuint := typecheck.Conv(oldCap, types.Types[types.TUINT])
	nif.Cond = ir.NewBinaryExpr(base.Pos, ir.OLE, nuint, scapuint)
	nif.Likely = true

	// then { s = s[:newLen] }

set unit=%value:~-1%
rem assume the unit is specified
set conv=%value:~0,-1%

if "%unit%" == "k" goto kilo
if "%unit%" == "K" goto kilo
if "%unit%" == "m" goto mega
if "%unit%" == "M" goto mega
if "%unit%" == "g" goto giga
if "%unit%" == "G" goto giga

rem no unit found, must be bytes; consider the whole value
set conv=%value%
rem convert to KB
set /a conv=%conv% / 1024
:kilo
rem convert to MB

	"MAXREG", // [MAXREG]
}

func init() {
	obj.RegisterRegister(REG_AL, REG_AL+len(Register), rconv)
	obj.RegisterOpcode(obj.ABaseAMD64, Anames)
	obj.RegisterRegisterList(obj.RegListX86Lo, obj.RegListX86Hi, rlconv)
	obj.RegisterOpSuffix("386", opSuffixString)
	obj.RegisterOpSuffix("amd64", opSuffixString)
}

func rconv(r int) string {
	if REG_AL <= r && r-REG_AL < len(Register) {
		return Register[r-REG_AL]
	}