//				0: R_LARCH_PCALA_HI20	local.moduledata
	o(0x1a000004)
	rel, _ := initfunc.AddRel(objabi.R_LOONG64_ADDR_HI)
	rel.SetOff(0)
	rel.SetSiz(4)
	rel.SetSym(ctxt.Moduledata)

	//	4:	02c00084	addi.d	$a0, $a0, 0
	//				4: R_LARCH_PCALA_LO12	local.moduledata
	o(0x02c00084)
	rel2, _ := initfunc.AddRel(objabi.R_LOONG64_ADDR_LO)
	rel2.SetOff(4)
	rel2.SetSiz(4)
	rel2.SetSym(ctxt.Moduledata)

	//	8:	50000000	b	0

 * <<#rel4.8:switch_to_publishing_plugins,Maven Publish and Ivy Publish Plugins>> that now support digital signatures with the <<signing_plugin#signing_plugin,Signing Plugin>>.
 * Use native <<#rel5.0:bom_import,BOM import>> in your builds.
 * The <<worker_api.adoc#worker_api,Worker API>> for enabling units of work to run in parallel.

				continue
			}
			val64 := sym.Value + uint64(rela.Addend)
			f.ByteOrder.PutUint64(dst[rela.Off:rela.Off+8], val64)
		case R_X86_64_32:
			if rela.Off+4 >= uint64(len(dst)) || rela.Addend < 0 {
				continue
			}
			val32 := uint32(sym.Value) + uint32(rela.Addend)
			f.ByteOrder.PutUint32(dst[rela.Off:rela.Off+4], val32)
		}
	}

	return nil
}

    "tf.Yield"(%t0, %t1, %t2) : (tensor<2xf32>, tensor<2xf32>, tensor<2xf32>) -> ()
    }, {
     %e0 = "tf.Neg"(%arg1) : (tensor<2xf32>) -> tensor<2xf32>
     %e1 = "tf.Relu"(%arg1) : (tensor<2xf32>) -> tensor<2xf32>
     %e2 = "tf.Sin"(%arg1) : (tensor<2xf32>) -> tensor<2xf32>
     "tf.Yield"(%e0, %e1, %e2) : (tensor<2xf32>, tensor<2xf32>, tensor<2xf32>) -> ()

names specified as strings (e.g. "Relu"). They are performed in order, where the
(first) input to each op is the output of the preceding op. The first input and
the output of each fused_op must be of type T.

Currently supported fused_op combinations are: [X] and [X,A], where X is one of
{"BiasAdd","FusedBatchNorm"} and A is one of {"Elu","Relu","Relu6"}.

// CHECK-PER-TENSOR: return %[[UNIFORM_QUANTIZE_0]] : tensor<?x3x4x2x!quant.uniform<i8:f32, {{.*}}>>

// -----

// Tests that fused pattern for convolution + bias + relu with
// dynamic batch dimension is properly quantized.

// Note that this checks for identical condition as
// quantize_conv_with_bias_dynamic_fn, omitting stablehlo.maximum.

	o(0x90000000)
	o(0x91000000)
	rel, _ := initfunc.AddRel(objabi.R_ADDRARM64)
	rel.SetOff(0)
	rel.SetSiz(8)
	rel.SetSym(ctxt.Moduledata)

	// 8:	14000000 	b	0 <runtime.addmoduledata>
	// 	8: R_AARCH64_CALL26	runtime.addmoduledata
	o(0x14000000)
	rel2, _ := initfunc.AddRel(objabi.R_CALLARM64)
	rel2.SetOff(8)
	rel2.SetSiz(4)
	rel2.SetSym(addmoduledata)
}

// 1) `is_stateless` is true turns off automatic ordering and we purely rely on
//    instance_key to distinguish collective groups. In this case, ordering
//    tokens are irrelevant. Each collective group should have a unique
//    instance_key at runtime.
// 2) If at least one ordering token is present, then we purely rely on ordering
//    tokens for side effect modeling and ignore the op-based effect

			// these fields in the 'reloc' phase.
			rela := ldr.MakeSymbolUpdater(syms.Rela)
			rela.AddAddrPlus(target.Arch, s, int64(r.Off()))
			if r.Siz() == 8 {
				rela.AddUint64(target.Arch, elf.R_INFO(0, uint32(elf.R_X86_64_RELATIVE)))
			} else {
				ldr.Errorf(s, "unexpected relocation for dynamic symbol %s", ldr.SymName(targ))
			}
			rela.AddAddrPlus(target.Arch, targ, int64(r.Add()))

// CHECK:         }
func.func @const() -> tensor<2xi32> {
  %0 = mhlo.constant dense<0> : tensor<2xi32>
  func.return %0 : tensor<2xi32>
}

// CHECK-LABEL:   func @relu(
// CHECK-SAME:               %[[VAL_0:.*]]: tensor<1xi32>) -> tensor<1xi32> {
// CHECK:           %[[VAL_1:.*]] = "tf.Const"() <{value = dense<0> : tensor<i32>}> : () -> tensor<i32>