n.To.Reg = x86.REG_DX
			}

			j.To.SetTarget(n)
			j2.To.SetTarget(s.Pc())
		}

	case ssa.Op386HMULL, ssa.Op386HMULLU:
		// the frontend rewrites constant division by 8/16/32 bit integers into
		// HMUL by a constant
		// SSA rewrites generate the 64 bit versions

		// Arg[0] is already in AX as it's the only register we allow
		// and DX is the only output we care about (the high bits)
		p := s.Prog(v.Op.Asm())

		{name: "UMULH", argLength: 2, reg: gp21, asm: "UMULH", commutative: true},                                     // (arg0 * arg1) >> 64, unsigned
		{name: "MULL", argLength: 2, reg: gp21, asm: "SMULL", commutative: true},                                      // arg0 * arg1, signed, 32-bit mult results in 64-bit

  // CHECK-NEXT: %[[mul:.*]] = mhlo.multiply %[[grad]], %[[sub]] : tensor<8x8x8x8xf32>
  // CHECK-NEXT: mhlo.constant dense<[0, 1, 2]> : tensor<3xi64>
  // CHECK-NEXT: %[[cmul:.*]] = mhlo.convert %[[mul]] : tensor<8x8x8x8xf32>
  // CHECK-NEXT: %[[init:.*]] = mhlo.constant dense<-0.000000e+00> : tensor<f32>

		ir.OPTRLIT,
		ir.OCONV,
		ir.OCONVIFACE,
		ir.OCONVNOP,
		ir.ODOT,
		ir.OEQ,
		ir.ONE,
		ir.OLT,
		ir.OLE,
		ir.OGT,
		ir.OGE,
		ir.OKEY,
		ir.OSTRUCTKEY,
		ir.OLEN,
		ir.OMUL,
		ir.OLSH,
		ir.ORSH,
		ir.OAND,
		ir.OANDNOT,
		ir.ONEW,
		ir.ONOT,
		ir.OBITNOT,
		ir.OPLUS,
		ir.ONEG,
		ir.OOROR,
		ir.OPAREN,
		ir.ORUNESTR,
		ir.OREAL,
		ir.OIMAG,

		y = z
		ch <- 1
	}()
	<-ch
	<-ch
}

func TestRaceDivConst(t *testing.T) {
	var x, y, z uint32
	_ = x
	ch := make(chan int, 2)

	go func() {
		x = y / 3 // involves only a HMUL node
		ch <- 1
	}()
	go func() {
		y = z
		ch <- 1
	}()
	<-ch
	<-ch
}

func TestRaceMod(t *testing.T) {
	var x, y, z int
	_ = x
	ch := make(chan int, 2)

	go func() {

			rel := obj.Addrel(c.cursym)

			rel.Off = int32(c.pc)
			rel.Siz = 4
			rel.Sym = p.To.Sym
			rel.Add = p.To.Offset
			rel.Type = objabi.R_ADDR
			o1 = 0
		}

	case 15: /* mul/mneg/umulh/umull r,[r,]r; madd/msub/fmadd/fmsub/fnmadd/fnmsub Rm,Ra,Rn,Rd */
		o1 = c.oprrr(p, p.As)

		rf := int(p.From.Reg)
		rt := int(p.To.Reg)
		var r int
		var ra int
		if p.From3Type() == obj.TYPE_REG {

	UMADDL R0, R20, R17, R17                   // 3152a09b
	UMSUBL R22, R4, R3, R7                     // 6790b69b
	UMNEGL R3, R19, R1                         // 61fea39b
	UMULH R24, R20, R24                        // 987ed89b
	UMULL R19, R22, R19                        // d37eb39b
	UXTBW R2, R6                               // 461c0053
	UXTHW R7, R20                              // f43c0053

		n := n.(*ir.AddrExpr)
		n.X = walkExpr(n.X, init)
		return n

	case ir.ODEREF:
		n := n.(*ir.StarExpr)
		n.X = walkExpr(n.X, init)
		return n

	case ir.OMAKEFACE, ir.OAND, ir.OANDNOT, ir.OSUB, ir.OMUL, ir.OADD, ir.OOR, ir.OXOR, ir.OLSH, ir.ORSH,
		ir.OUNSAFEADD:
		n := n.(*ir.BinaryExpr)
		n.X = walkExpr(n.X, init)
		n.Y = walkExpr(n.Y, init)
		return n

	case ir.OUNSAFESLICE:
		n := n.(*ir.BinaryExpr)

(Select1 (Sub32carry <t> x y)) => (SGTU <typ.Bool> (SUB <t.FieldType(0)> x y) x)
(Sub32withcarry <t> x y c) => (SUB (SUB <t> x y) c)

(Mul(32|16|8) ...) => (MUL ...)
(Mul(32|64)F ...) => (MUL(F|D) ...)

(Hmul(32|32u) x y) => (Select0 (MUL(T|TU) x y))
(Mul32uhilo ...) => (MULTU ...)

(Div32 x y) => (Select1 (DIV x y))
(Div32u x y) => (Select1 (DIVU x y))

	case AMOVF,
		AMOVW,
		AMOVWL,
		AMOVWR:
		sz = 4
		ld = 1

	case AMOVD,
		AMOVV,
		AMOVVL,
		AMOVVR:
		sz = 8
		ld = 1

	case ADIV,
		ADIVU,
		AMUL,
		AMULU,
		AREM,
		AREMU,
		ADIVV,
		ADIVVU,
		AMULV,
		AMULVU,
		AREMV,
		AREMVU:
		s.set.cc = E_HILO
		fallthrough
	case AADD,
		AADDU,
		AADDV,