CALL 'public final fun EQEQ (arg0: kotlin.Any?, arg1: kotlin.Any?): kotlin.Boolean declared in kotlin.internal.ir' type=kotlin.Boolean origin=EQEQ
                            arg0: CALL 'public final fun rem (other: kotlin.Int): kotlin.Int [operator] declared in kotlin.Int' type=kotlin.Int origin=PERC
                              $this: GET_VAR 'it: kotlin.Int declared in <root>.CodeFragment.run.<anonymous>' type=kotlin.Int origin=null

	"MULLWCC",
	"MULLWVCC",
	"MULLWV",
	"NAND",
	"NANDCC",
	"NEG",
	"NEGCC",
	"NEGVCC",
	"NEGV",
	"NOR",
	"NORCC",
	"OR",
	"ORCC",
	"ORN",
	"ORNCC",
	"ORIS",
	"REM",
	"REMU",
	"RFI",
	"RLWMI",
	"RLWMICC",
	"RLWNM",
	"RLWNMCC",
	"CLRLSLWI",
	"SLW",
	"SLWCC",
	"SRW",
	"SRAW",
	"SRAWCC",
	"SRWCC",
	"STBCCC",
	"STHCCC",
	"STSW",

  @SuppressWarnings("fallthrough")
  public static long divide(long p, long q, RoundingMode mode) {
    checkNotNull(mode);
    long div = p / q; // throws if q == 0
    long rem = p - q * div; // equals p % q

    if (rem == 0) {
      return div;
    }

    /*
     * Normal Java division rounds towards 0, consistently with RoundingMode.DOWN. We just have to

  @SuppressWarnings("fallthrough")
  public static long divide(long p, long q, RoundingMode mode) {
    checkNotNull(mode);
    long div = p / q; // throws if q == 0
    long rem = p - q * div; // equals p % q

    if (rem == 0) {
      return div;
    }

    /*
     * Normal Java division rounds towards 0, consistently with RoundingMode.DOWN. We just have to

		// large clear
		// auxint = remaining bytes after loop (rem)
		// arg0 = address of dst memory (in R1, changed as a side effect)
		// arg1 = pointer to last address to zero in loop + 256
		// arg2 = mem
		// returns mem
		//
		// clear: CLEAR $256, 0(R1)
		//        MOVD  $256(R1), R1
		//        CMP   R1, Rarg2
		//        BNE   clear
		//	  CLEAR $rem, 0(R1) // if rem > 0
		{
			name:      "LoweredZero",

		// arg2 should be src+size,
		//
		// mvc: MVC  $256, 0(R2), 0(R1)
		//      MOVD $256(R1), R1
		//      MOVD $256(R2), R2
		//      CMP  R2, Rarg2
		//      BNE  mvc
		//      MVC  $rem, 0(R2), 0(R1) // if rem > 0
		// arg2 is the last address to move in the loop + 256
		mvc := s.Prog(s390x.AMVC)
		mvc.From.Type = obj.TYPE_CONST
		mvc.From.Offset = 256

	binop{"*", func(x, y *big.Int) *big.Int { return new(big.Int).Mul(x, y) }},
	binop{"/", func(x, y *big.Int) *big.Int { return new(big.Int).Quo(x, y) }},
	binop{"%", func(x, y *big.Int) *big.Int { return new(big.Int).Rem(x, y) }},
	binop{"&", func(x, y *big.Int) *big.Int { return new(big.Int).And(x, y) }},
	binop{"|", func(x, y *big.Int) *big.Int { return new(big.Int).Or(x, y) }},

		lastMajorHasGoMod bool
	)
	for i, v := range incompatible {
		major := semver.Major(v)

		if major != lastMajor {
			rem := incompatible[i:]
			j := sort.Search(len(rem), func(j int) bool {
				return semver.Major(rem[j]) != major
			})
			latestAtMajor := rem[j-1]

			var err error
			lastMajor = major
			lastMajorHasGoMod, err = versionHasGoMod(latestAtMajor)
			if err != nil {

	{_Operator, "+", Add, precAdd},
	{_Operator, "-", Sub, precAdd},
	{_Operator, "|", Or, precAdd},
	{_Operator, "^", Xor, precAdd},

	{_Star, "*", Mul, precMul},
	{_Operator, "/", Div, precMul},
	{_Operator, "%", Rem, precMul},
	{_Operator, "&", And, precMul},
	{_Operator, "&^", AndNot, precMul},
	{_Operator, "<<", Shl, precMul},
	{_Operator, ">>", Shr, precMul},

	// assignment operations
	{_AssignOp, "+=", Add, precAdd},

func.func @floormod_broadcast_numerator(%arg0: tensor<3xi32>, %arg1: tensor<2x3xi32>) -> tensor<2x3xi32> {
  // CHECK-DAG: [[REM:%.+]] = chlo.broadcast_remainder %arg0, %arg1 {broadcast_dimensions = array<i64: 1>}
  // CHECK-DAG: [[ZL:%.+]] = mhlo.constant dense<0>
  // CHECK-DAG: [[CMP1:%.+]] = chlo.broadcast_compare [[REM]], [[ZL]] {comparison_direction = #chlo<comparison_direction NE>}
  // CHECK-DAG: [[ZR:%.+]] = mhlo.constant dense<0>