// CHECK:  return %4 : tensor<i32>
}

// CHECK-LABEL:   func @convert_argmin
func.func @convert_argmin(%arg0: tensor<4x32x256xf32>) -> (tensor<4x32xf32>, tensor<4x32xi32>) {
  %0 = mhlo.constant dense<0x7F800000> : tensor<f32>
  %1 = mhlo.constant dense<0> : tensor<i32>
  %2 = "mhlo.iota"() <{iota_dimension = 0 : i64}> : () -> tensor<256xi32>

			t.Errorf("got %#x, want %#x", x, nan)
		}
	}

	// infinities
	{
		const inf = uint32(0x7f800000) // +∞
		if x := math.Float32bits(math.Float32frombits(inf)); x != inf {
			t.Errorf("got %#x, want %#x", x, inf)
		}
	}
	{
		const negInf = uint32(0xff800000) // -∞
		if x := math.Float32bits(math.Float32frombits(negInf)); x != negInf {
			t.Errorf("got %#x, want %#x", x, negInf)
		}
	}

// CHECK:           return %[[VAL_3]] : tensor<1xf32>
// CHECK:         }
func.func @convert_reduce_to_min(%arg0: tensor<1x256xf32>) -> tensor<1xf32> {
  // "0x7F800000" represents INF for f32.
  %0 = mhlo.constant dense<0x7F800000> : tensor<f32>
  %1 = "mhlo.reduce"(%arg0, %0) ({
  ^bb0(%arg1: tensor<f32>, %arg2: tensor<f32>):
    %2 = mhlo.minimum %arg1, %arg2 : tensor<f32>

		return false
	}
	v := p.From.Val.(float64)
	if float64(float32(v)) != v {
		return false
	}
	// Secondly, is this value a normal value?
	ival := int64(math.Float32bits(float32(v)))
	isDenorm := ival&0x7F800000 == 0 && ival&0x007FFFFF != 0
	if !isDenorm {
		p.As = AXXSPLTIDP
		p.From.Type = obj.TYPE_CONST
		p.From.Offset = ival
		// Convert REG_Fx into equivalent REG_VSx
		p.To.Reg = REG_VS0 + (p.To.Reg & 31)
	}

	b := math.Float64bits(f)
	m := b & ((1 << 52) - 1) // mantissa (a.k.a. significand)
	//          | sign                  | exponent   | mantissa       |
	r := uint32(((b >> 32) & (1 << 31)) | 0x7f800000 | (m >> (52 - 23)))
	return math.Float32frombits(r)
}

// extend32Fto64F converts a float32 value to a float64 value preserving the bit
// pattern of the mantissa.
func extend32Fto64F(f float32) float64 {

  // CHECK: return %[[RESULT]]
  func.return %0 : tensor<2x3xf32>
}

// CHECK-LABEL: @is_inf
func.func @is_inf(%arg0: tensor<3x4xf32>) -> tensor<3x4xi1> {
  // CHECK: %[[INF:.*]] = "tf.Const"() <{value = dense<0x7F800000> : tensor<f32>}> : () -> tensor<f32>
  // CHECK: %[[ABS:.*]] = "tf.Abs"(%arg0) : (tensor<3x4xf32>) -> tensor<3x4xf32>

	MFD_HUGE_16MB                               = 0x60000000
	MFD_HUGE_1GB                                = 0x78000000
	MFD_HUGE_1MB                                = 0x50000000
	MFD_HUGE_256MB                              = 0x70000000
	MFD_HUGE_2GB                                = 0x7c000000
	MFD_HUGE_2MB                                = 0x54000000
	MFD_HUGE_32MB                               = 0x64000000

//go:build linux && (ppc64 || ppc64le)

package cpu

// HWCAP/HWCAP2 bits. These are exposed by the kernel.
const (
	// ISA Level
	_PPC_FEATURE2_ARCH_2_07 = 0x80000000
	_PPC_FEATURE2_ARCH_3_00 = 0x00800000

	// CPU features
	_PPC_FEATURE2_DARN = 0x00200000
	_PPC_FEATURE2_SCV  = 0x00100000
)

func doinit() {
	// HWCAP2 feature bits
	PPC64.IsPOWER8 = isSet(hwCap2, _PPC_FEATURE2_ARCH_2_07)

	MFD_HUGETLB                    = 0x4
	MFD_HUGE_16GB                  = -0x78000000
	MFD_HUGE_16MB                  = 0x60000000
	MFD_HUGE_1GB                   = 0x78000000
	MFD_HUGE_1MB                   = 0x50000000
	MFD_HUGE_256MB                 = 0x70000000
	MFD_HUGE_2GB                   = 0x7c000000
	MFD_HUGE_2MB                   = 0x54000000
	MFD_HUGE_32MB                  = 0x64000000

        0x00802080, 0x00800081, 0x00800001, 0x00002001,
        0x00000000, 0x00802000, 0x00802000, 0x00802081,
        0x00000081, 0x00000000, 0x00800080, 0x00800001,
        0x00000001, 0x00002000, 0x00800000, 0x00802001,
        0x00000080, 0x00800000, 0x00002001, 0x00002080,
        0x00800081, 0x00000001, 0x00002080, 0x00800080,
        0x00002000, 0x00802080, 0x00802081, 0x00000081,
        0x00800080, 0x00800001, 0x00802000, 0x00802081,