func.func @skip_nan_inf_constant(%arg0: tensor<?x112x112x64xf32>) -> tensor<?x56x56x64xf32> {
  // CHECK-DAG: %[[cst0:.*]] = stablehlo.constant dense<0xFF800000> : tensor<f32
  // CHECK-DAG: %[[cst1:.*]] = stablehlo.constant dense<0x7FC00000> : tensor<f32>
  // CHECK-DAG: %[[cst2:.*]] = stablehlo.constant dense<6.000000e+00> : tensor<f32>
  // CHECK-DAG: %[[cst3:.*]] = stablehlo.constant dense<0.000000e+00> : tensor<f32>

	{
		const nan = uint32(0xffbfffff) // sNaN
		if x := math.Float32bits(math.Float32frombits(nan)); x != nan {
			t.Errorf("got %#x, want %#x", x, nan)
		}
	}

	// quiet NaNs
	{
		const nan = uint32(0x7fc00000) // qNaN
		if x := math.Float32bits(math.Float32frombits(nan)); x != nan {
			t.Errorf("got %#x, want %#x", x, nan)
		}
	}
	{
		const nan = uint32(0x7fffffff) // qNaN

	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

	want        string // File name of data with the expected output with input available.
	wantNoInput string // File name of the expected output when no input is available.
}

const ml = 0x7fc00000 // Maximum length token. Used to reduce the size of writeBlockTests

var writeBlockTests = []huffTest{
	{
		input:       "testdata/huffman-null-max.in",
		want:        "testdata/huffman-null-max.%s.expect",

  // 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>

		return -1
	}

	ei := math.Float64bits(e)
	l := uint32(ei)
	h := uint32(ei >> 32)

	if l != 0 || h&0xffff != 0 {
		return -1
	}
	h1 := h & 0x7fc00000
	if h1 != 0x40000000 && h1 != 0x3fc00000 {
		return -1
	}
	n := 0

	// sign bit (a)
	if h&0x80000000 != 0 {
		n |= 1 << 7
	}

	// exp sign bit (b)
	if h1 == 0x3fc00000 {

	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

	{0xffe00c00, 0x78800c00, LDRSH, instArgs{arg_Xt, arg_Xns_mem_wb_imm9_1_signed}, nil},
	// LDRSH <Wt>, [<Xn|SP>{, #<pimm>}]
	{0xffc00000, 0x79c00000, LDRSH, instArgs{arg_Wt, arg_Xns_mem_optional_imm12_2_unsigned}, nil},
	// LDRSH <Xt>, [<Xn|SP>{, #<pimm>}]
	{0xffc00000, 0x79800000, LDRSH, instArgs{arg_Xt, arg_Xns_mem_optional_imm12_2_unsigned}, nil},
	// LDRSH <Wt>, [<Xn|SP>, (<Wm>|<Xm>) {, <extend> {<amount>}}]

	arch := sys.ArchARM64

	theArch := ld.Arch{
		Funcalign:  funcAlign,
		Maxalign:   maxAlign,
		Minalign:   minAlign,
		Dwarfregsp: dwarfRegSP,
		Dwarfreglr: dwarfRegLR,
		TrampLimit: 0x7c00000, // 26-bit signed offset * 4, leave room for PLT etc.

		Adddynrel:        adddynrel,
		Archinit:         archinit,
		Archreloc:        archreloc,
		Archrelocvariant: archrelocvariant,
		Extreloc:         extreloc,

	PT_HIOS ProgType = 0x6fffffff /* Last OS-specific. */

	PT_LOPROC ProgType = 0x70000000 /* First processor-specific type. */

	PT_ARM_ARCHEXT ProgType = 0x70000000 /* Architecture compatibility */
	PT_ARM_EXIDX   ProgType = 0x70000001 /* Exception unwind tables */

	PT_AARCH64_ARCHEXT ProgType = 0x70000000 /* Architecture compatibility */
	PT_AARCH64_UNWIND  ProgType = 0x70000001 /* Exception unwind tables */