return 0x0c1<<20 | 0x3<<15 // FCMP.SLT.S

	case ASQRTF:
		return 0x4511 << 10
	case ASQRTD:
		return 0x4512 << 10
	}

	if a < 0 {
		c.ctxt.Diag("bad rrr opcode -%v", -a)
	} else {
		c.ctxt.Diag("bad rrr opcode %v", a)
	}
	return 0
}

func (c *ctxt0) oprr(a obj.As) uint32 {
	switch a {
	case ACLO:
		return 0x4 << 10
	case ACLZ:
		return 0x5 << 10

64-bit variant ldr, str, stur => MOVD;
32-bit variant str, stur, ldrsw => MOVW;
32-bit variant ldr => MOVWU;
ldrb => MOVBU; ldrh => MOVHU;
ldrsb, sturb, strb => MOVB;
ldrsh, sturh, strh =>  MOVH.

4. Go moves conditions into opcode suffix, like BLT.

5. Go adds a V prefix for most floating-point and SIMD instructions, except cryptographic extension
instructions and floating-point(scalar) instructions.

Examples:

  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 ") {

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

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

//	}
	NOP	F2

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

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

//
// RET
//
//	LRETRN	comma // asm doesn't support the trailing comma.
//	{
//		outcode(int($1), &nullgen, 0, &nullgen);
//	}
	SYSCALL
	BEQ	R1, 2(PC)
	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,