EXPECT_TRUE(compilation_result.ok());

  // Ensure that the compilation result contains a constant.
  EXPECT_THAT(compilation_result,
              ComputationProtoContains("opcode:.*constant"));
}

TEST(LegalizeTFTest, SkipsTensorListSetItemIfDimensionsTooLarge) {
  static constexpr char kTensorListSetItemDimensionTooLarge[] = R"(

		case isU32 && p.From.Offset&0xFFFF == 0 && p.From.Offset != 0:
			p.As = AORIS
			p.From.Offset >>= 16
			p.Reg = REG_R0

		case isS32 || isU32 || isS34:
			// The assembler can generate this opcode in 1 (on Power10) or 2 opcodes.

		// Otherwise, see if the large constant can be generated with 2 instructions. If not, load it from memory.
		default:

		if p.Link != nil {
			size = p.Link.Pc - p.Pc
		} else if p.Isize != 0 {
			size = int64(p.Isize)
		}
		var code []byte
		if p.Pc < int64(len(text.P)) {
			code = text.P[p.Pc:]
			if size < int64(len(code)) {
				code = code[:size]
			}
		}
		codeHex := fmt.Sprintf("%x", code)
		if codeHex == "" {
			codeHex = "empty"
		}
		ok := false
		for _, hex := range strings.Split(hexes, " or ") {

	// stored to ABI-defined locations.  Each non-memory input or output fits in a register.
	//
	// Subsequent architecture-specific lowering only changes the opcode.

	{name: "ClosureCall", argLength: -1, aux: "CallOff", call: true}, // arg0=code pointer, arg1=context ptr, arg2..argN-1 are register inputs, argN=memory.  auxint=arg size.  Returns Result of register results, plus memory.

					args[1] = fmt.Sprintf("$%f", math.Float32frombits(binary.LittleEndian.Uint32(buf)))
				default:
					panic(fmt.Sprintf("wrong FP register: %v", inst))
				}
			}
		}
	}

	// Move addressing mode into opcode suffix.
	suffix := ""
	switch inst.Op &^ 15 {
	case PLD, PLI, PLD_W:
		if mem, ok := inst.Args[0].(Mem); ok {
			args[0], suffix = memOpTrans(mem)
		} else {
			panic(fmt.Sprintf("illegal instruction: %v", inst))

//	{
//		outcode(int($1), &nullgen, 0, &$3);
//	}
	NOP	R2

//	LNOP ',' freg // asm doesn't support the leading comma.
//	{
//		outcode(int($1), &nullgen, 0, &$3);
//	}
	NOP	F2

//	LNOP imm
//	{
//		outcode(int($1), &$2, 0, &nullgen);
//	}
	NOP	$4

//
// special
//
	SYSCALL
	BREAK
	// overloaded cache opcode:
	BREAK	R1, (R1)

//
// RET
//

	AXVCVSXDDP
	AXVCVSXWDP
	AXVCVUXDDP
	AXVCVUXWDP
	AXVCVSXDSP
	AXVCVSXWSP
	AXVCVUXDSP
	AXVCVUXWSP
	AXSMAXJDP
	AXSMINJDP
	ALASTAOUT // The last instruction in this list. Also the first opcode generated by ppc64map.

	// aliases
	ABR   = obj.AJMP
	ABL   = obj.ACALL
	ALAST = ALASTGEN // The final enumerated instruction value + 1. This is used to size the oprange table.

}

bool RootInstructionIsTuple(const xla::HloModule& hlo_module) {
  xla::HloInstruction* root_instruction =
      hlo_module.entry_computation()->root_instruction();

  return root_instruction->opcode() == xla::HloOpcode::kTuple;
}

};  // namespace

LogicalResult Tf2XlaRewriter::RewriteOp(Operation* op,
                                        PatternRewriter& rewriter,

// Code generated by ./parse.py -go rv64_a rv64_d rv64_f rv64_i rv64_m rv64_q rv64_zba rv64_zbb rv64_zbs rv_a rv_d rv_f rv_i rv_m rv_q rv_zba rv_zbb rv_zbs rv_s rv_system rv_zicsr; DO NOT EDIT.
package riscv

import "cmd/internal/obj"

type inst struct {
	opcode uint32
	funct3 uint32
	rs2    uint32
	csr    int64
	funct7 uint32
}

func encode(a obj.As) *inst {
	switch a {
	case AADD:
		return &inst{0x33, 0x0, 0x0, 0, 0x0}

	// Recursively decompose elem phi.
	decomposeUserPhi(elem)
}

// MaxStruct is the maximum number of fields a struct
// can have and still be SSAable.
const MaxStruct = 4

// StructMakeOp returns the opcode to construct a struct with the
// given number of fields.
func StructMakeOp(nf int) Op {
	switch nf {
	case 0:
		return OpStructMake0
	case 1:
		return OpStructMake1
	case 2:
		return OpStructMake2