wantErr string
	}
	execTests := []execTest{
		// Okay:
		{[]any{"Brad", 31}, ""},
		{[]any{"Brad", int64(31)}, ""},
		{[]any{"Bob", "32"}, ""},
		{[]any{7, 9}, ""},

		// Invalid conversions:
		{[]any{"Brad", int64(0xFFFFFFFF)}, "sql: converting argument $2 type: sql/driver: value 4294967295 overflows int32"},

    // CHECK: [[GRAD_SQUARE:%.*]] = "tf.Mul"([[GRAD]], [[GRAD]]) : (tensor<f32>, tensor<f32>) -> tensor<f32>
    // CHECK: [[NEW_ACC:%.*]] = "tf.AddV2"([[OLD_ACC]], [[GRAD_SQUARE]]) : (tensor<*xf32>, tensor<f32>) -> tensor<*xf32>
    // CHECK: [[LR_MULTIPLY:%.*]] = "tf.Mul"([[LR]], [[GRAD]]) : (tensor<f32>, tensor<f32>) -> tensor<f32>

     (CreateTFReadVariableOp $src_op, $grad, $ms_resource),
     (TF_AddV2Op:$ms_new
       (TF_MulOp
         (TF_MulOp $grad, $grad),
         (TF_SubOp $one, $rho)
       ),
       (TF_MulOp
          (CreateTFReadVariableOp $src_op, $grad, $ms_resource),
          $rho
       )
     ),
     (TF_AssignVariableOp $ms_resource, $ms_new, (CreateConstBoolAttrFalse)),
     // mg = grad * (one - rho) + mg * rho;

  if (TF_GetCode(status) != TF_OK) return;
  if (!grad) return;

  status->status = g->graph.AddFunctionDef(grad->record->fdef(),
                                           grad->record->stack_traces());
  if (TF_GetCode(status) != TF_OK) return;

  tensorflow::GradientDef gdef;
  gdef.set_function_name(func->record->fdef().signature().name());
  gdef.set_gradient_func(grad->record->fdef().signature().name());

(Rsh8Ux(64|32|16|8)  x y) && shiftIsBounded(v) => (SRD (MOVBZreg x) y)
(Rsh64x(64|32|16|8)  x y) && shiftIsBounded(v) => (SRAD x y)
(Rsh32x(64|32|16|8)  x y) && shiftIsBounded(v) => (SRAW x y)
(Rsh16x(64|32|16|8)  x y) && shiftIsBounded(v) => (SRAD (MOVHreg x) y)
(Rsh8x(64|32|16|8)   x y) && shiftIsBounded(v) => (SRAD (MOVBreg x) y)

// Unbounded shifts. Go shifts saturate to 0 or -1 when shifting beyond the number of

func.func @max_pool_grad_valid(%orig_input: tensor<10x24x24x64xf32>, %orig_output: tensor<10x12x12x64xf32>, %grad: tensor<10x12x12x64xf32>) -> tensor<10x24x24x64xf32> {
  // CHECK: %[[ZERO:.*]] = mhlo.constant dense<0.000000e+00> : tensor<f32>

	// arm64:"TST\t[$]63",-"UDIV",-"ASR",-"AND"
	// ppc64x:"ANDCC",-"RLDICL",-"SRAD",-"CMP"
	a := n1%64 == 0 // signed divisible

	// 386:"TESTL\t[$]63",-"DIVL",-"SHRL"
	// amd64:"TESTQ\t[$]63",-"DIVQ",-"SHRQ"
	// arm:"AND\t[$]63",-".*udiv",-"SRA"
	// arm64:"TST\t[$]63",-"UDIV",-"ASR",-"AND"
	// ppc64x:"ANDCC",-"RLDICL",-"SRAD",-"CMP"
	b := n2%64 != 0 // signed indivisible

	return a, b
}

                                          grad_func, flib_def,
                                          visited_functions));
    GradientDef grad;
    grad.set_function_name(function_name);
    grad.set_gradient_func(grad_func.getName().str());
    TF_RETURN_IF_ERROR(flib_def->AddGradientDef(grad));
  }

  auto stateful_string = mlir::TF::TensorFlowDialect::GetStatefulAttrName();

                                          grad_func, flib_def,
                                          visited_functions));
    GradientDef grad;
    grad.set_function_name(function_name);
    grad.set_gradient_func(grad_func.getName().str());
    TF_RETURN_IF_ERROR(flib_def->AddGradientDef(grad));
  }

  auto stateful_string = mlir::TF::TensorFlowDialect::GetStatefulAttrName();

	return uint64(v) >> 16
}

func rshConst64Ux64Overflow8(v uint8) uint64 {
	// riscv64:"MOV\t\\$0,",-"SRL"
	return uint64(v) >> 8
}

func rshConst64x64(v int64) int64 {
	// ppc64x:"SRAD"
	// riscv64:"SRAI\t",-"OR",-"SLTIU"
	return v >> uint64(33)
}

func rshConst64x64Overflow32(v int32) int64 {
	// riscv64:"SRAIW",-"SLLI",-"SRAI\t"
	return int64(v) >> 32
}