@tf.RegisterGradient('NewConv2D')
def _conv_add_relu_grad(op: ops.Operation, grad):
  act = op.get_attr('act')
  y = op.outputs[0]
  if act == 'RELU':
    grad = gen_nn_ops.relu_grad(grad, y)
  elif act == 'RELU6':
    grad = gen_nn_ops.relu6_grad(grad, y)
  elif act == 'TANH':
    y = math_ops.conj(y)
    grad = gen_math_ops.tanh_grad(y, grad)

  broadcast_shape = tf.shape(y)
  input_value_shape = tf.shape(op.inputs[2])

    AbstractTensorHandle* upstream_grad = grad_outputs[0];
    AbstractTensorHandle* activations = forward_outputs_[0];

    // Calculate Grad
    std::string name = "relu_grad";
    TF_RETURN_IF_ERROR(ReluGrad(ctx, upstream_grad, activations,
                                &grad_inputs[0], name.c_str()));
    return absl::OkStatus();
  }
  ~ReluGradientFunction() override {

 Conv2DBackpropInput 52
 DivNoNan 1
 Equal 1
 FusedBatchNorm 53
 FusedBatchNormGrad 53
 Identity 2
 MatMul 3
 MaxPool 1
 MaxPoolGrad 1
 Mean 1
 Mul 164
 Pad 1
 ReadVariableOp 646
 Relu 49
 ReluGrad 49
 Reshape 2
 ResourceApplyKerasMomentum 161
 ShapeN 50
 Softmax 1
 SparseSoftmaxCrossEntropyWithLogits 1
 Square 55
 Squeeze 1
 Sub 106
 Sum 57
 Tile 1

 Conv2DBackpropFilter 53
 Conv2DBackpropInput 52
 Equal 1
 FusedBatchNormGradV2 53
 FusedBatchNormV2 53
 MatMul 3
 MaxPool 1
 MaxPoolGrad 1
 Mean 1
 Mul 218
 Pad 2
 ReadVariableOp 538
 Relu 49
 ReluGrad 49
 Reshape 2
 ResourceApplyKerasMomentum 161
 Slice 1
 Softmax 1
 SparseSoftmaxCrossEntropyWithLogits 1
 Squeeze 1
 Sum 1
 Tile 1
 Transpose 1
cluster 1 size 815
 AddN 1
 AssignAddVariableOp 1

          (MHLO_ClampOp (MHLO_ConstantOp (GetScalarOfType<0> $input)), $input,
                       (MHLO_ConstantOp (GetScalarOfType<6> $input))),
          [(TF_IntOrFpTensor $input)]>;

// ReluGrad(gradients, features) = gradients * (features > 0)
// The condition that $gradients and $features need to have the same shape is
// implicitly enforced: $zero is created to have the same shape as $features,

            "TanhGrad", "Pow", "SquaredDifference", "ApproximateEqual",
            // Others
            "AddN", "Bitcast", "Cast", "ClipByValue", "Const", "Empty",
            "Identity", "IdentityN", "Relu", "Relu6", "ReluGrad", "Relu6Grad",
            "LeakyReluGrad", "Elu", "EluGrad", "Selu", "SeluGrad", "Select",
            "SelectV2", "Transpose", "ConjugateTranspose",
            "_UnaryOpsComposition", "CollectiveReduceV2",

    %0 = "tf.Selu"(%arg0) : (tensor<1x4x4x3xf32>) -> tensor<1x4x4x3xf32>
    func.return %0 : tensor<1x4x4x3xf32>
}

// CHECK-LABEL: func @selu_grad
// CHECK-SAME: (%[[GRADIENTS:.*]]: tensor<4x8xf32>, %[[FEATURES:.*]]: tensor<4x8xf32>) -> tensor<4x8xf32> {
func.func @selu_grad(%gradients: tensor<4x8xf32>, %features: tensor<4x8xf32>) -> tensor<4x8xf32> {

  %3 = "chlo.broadcast_maximum"(%2, %0) {broadcast_dimensions = array<i64>} : (tensor<?xi32>, tensor<i32>) -> tensor<?xi32>
  func.return %3 : tensor<?xi32>
}

// CHECK-LABEL:   func @relu_grad(
// CHECK-SAME:                    %[[VAL_0:.*]]: tensor<4x8xf32>,
// CHECK-SAME:                    %[[VAL_1:.*]]: tensor<?x?xf32>) -> tensor<4x8xf32> {

  %0 = "tf.Elu"(%arg0) : (tensor<1xf32>) -> tensor<1xf32>
  func.return %0: tensor<1xf32>
}

// -----

// CHECK-LABEL: func @elu_grad
// CHECK-SAME: (%[[GRADIENTS:.*]]: tensor<4x8xf32>, %[[FEATURES:.*]]: tensor<?x?xf32>)
func.func @elu_grad(%gradients: tensor<4x8xf32>, %features: tensor<?x?xf32>) -> tensor<4x8xf32> {
  // CHECK-DAG: %[[ZERO:.*]] = mhlo.constant dense<0.000000e+00> : tensor<f32>

  func.func @testSameOperandsAndResultTypeResolveRefBinary(%lhs: tensor<2x3x?x?xf32>, %rhs: tensor<2x?x5x?xf32>) -> (tensor<?x?x?x?xf32>) {
    // CHECK: (tensor<2x3x?x?xf32>, tensor<2x?x5x?xf32>) -> tensor<2x3x5x?xf32>
    %0 = "tf.ReluGrad"(%lhs, %rhs) : (tensor<2x3x?x?xf32>, tensor<2x?x5x?xf32>) -> tensor<?x?x?x?xf32>
    func.return %0 : tensor<?x?x?x?xf32>
  }

  // CHECK-LABEL: func @test_xla_sharding