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;

    /* Given upstream grad U and a Sub op A-B, the gradients are:
     *
     *    dA =  U
     *    dB = -U
     *
     */

    // Grad for A
    DCHECK(grad_outputs[0]);
    grad_inputs[0] = grad_outputs[0];
    grad_inputs[0]->Ref();

    // Grad for B
    // negate the upstream grad
    std::string name = "Neg_Sub_Grad_B";
    TF_RETURN_IF_ERROR(

                 absl::Span<AbstractTensorHandle*> grad_inputs) override {
    // Grad for Softmax Input
    TF_RETURN_IF_ERROR(BroadcastMul(
        ctx, grad_outputs[0], forward_outputs_[1],
        grad_inputs.subspan(0, 1)));  // upstream_grad * local softmax grad

    // Grad for labels is null
    grad_inputs[1] = nullptr;
    return absl::OkStatus();
  }

        return s;
      }
    } else {
      if (!persistent_) {
        trace.backward_function_deleter(trace.backward_function);
      }
      for (Gradient* grad : out_gradients) {
        if (grad != nullptr) {
          vspace.DeleteGradient(grad);
        }
      }
    }
    for (int i = 0, end = in_gradients.size(); i < end; ++i) {
      const int64_t id = trace.input_tensor_id[i];

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

  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>
    // CHECK: %[[RESULT:.*]] = "mhlo.select_and_scatter"(%[[INPUT]], %[[GRAD]], %[[ZERO]]) <{

  if (grads_to_keep.empty()) {
    // Nothing propagated back. Return 'NoGradient'.
    *grad = NoGradient();
  } else if (grads_to_keep.size() == 1) {
    // Just one backprop edge.
    *grad = grads_to_keep[0];
  } else {
    // Otherwise, adds backprop-ed gradients.
    // TODO(andydavis) Use a better accumulator here.
    *grad = ops::AddN(scope_, grads_to_keep);
  }

  return absl::OkStatus();
}

(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