TF_ASSERT_OK(scope.status());

    // Construct grad inputs.
    auto dx = Const(scope, 4, {4, 2});
    auto dy = Const(scope, 5, {4, 2});
    auto dz = Const(scope, 6, {4, 2});

    if (expected) {
      // Construct backward graph.
      // NOTE: We should only expect the grad function for unpack in the
      // gradients graph, based on the requested grad outputs.
      auto unpack_grad = Stack(scope, {dx, dy, dz});

TEXT    ·atanhAsm(SB), NOSPLIT, $0-16
	FMOVD   x+0(FP), F0
	MOVD    $·atanhrodataL10<>+0(SB), R5
	LGDR    F0, R1
	WORD    $0xC0393FEF //iilf  %r3,1072693247
	BYTE    $0xFF
	BYTE    $0xFF
	SRAD    $32, R1
	WORD    $0xB9170021 //llgtr %r2,%r1
	MOVW    R2, R6
	MOVW    R3, R7
	CMPBGT  R6, R7, L2
	WORD    $0xC0392FFF //iilf  %r3,805306367
	BYTE    $0xFF
	BYTE    $0xFF
	MOVW    R2, R6
	MOVW    R3, R7

TEXT	·logAsm(SB), NOSPLIT, $0-16
	FMOVD	x+0(FP), F0
	MOVD	$·logrodataL21<>+0(SB), R9
	MOVH	$0x8006, R4
	LGDR	F0, R1
	MOVD	$0x3FF0000000000000, R6
	SRAD	$48, R1, R1
	MOVD	$0x40F03E8000000000, R8
	SUBW	R1, R4
	RISBGZ	$32, $59, $0, R4, R2
	RISBGN	$0, $15, $48, R2, R6
	RISBGN	$16, $31, $32, R2, R8
	MOVW	R1, R7
	CMPBGT	R7, $22, L17
	LTDBR	F0, F0

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>

TEXT	·cbrtAsm(SB), NOSPLIT, $0-16
	FMOVD	x+0(FP), F0
	MOVD	$·cbrtrodataL9<>+0(SB), R9
	LGDR	F0, R2
	WORD	$0xC039000F	//iilf	%r3,1048575
	BYTE	$0xFF
	BYTE	$0xFF
	SRAD	$32, R2
	WORD	$0xB9170012	//llgtr	%r1,%r2
	MOVW	R1, R6
	MOVW	R3, R7
	CMPBLE	R6, R7, L2
	WORD	$0xC0397FEF	//iilf	%r3,2146435071
	BYTE	$0xFF
	BYTE	$0xFF
	MOVW	R3, R7
	CMPBLE	R6, R7, L8
L1:

	// 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
}

		return true
	case OR, ORCC, ORC, ORCCC, AND, ANDCC, ANDC, ANDCCC, XOR, XORCC, NAND, NANDCC, EQV, EQVCC, NOR, NORCC:
		return true
	case SLW, SLWCC, SLD, SLDCC, SRW, SRAW, SRWCC, SRAWCC, SRD, SRDCC, SRAD, SRADCC:
		return true
	}
	return false
}

// revCondMap maps a conditional register bit to its inverse, if possible.
var revCondMap = map[string]string{
	"LT": "GE", "GT": "LE", "EQ": "NE",
}

// Adds a copy of function `func` and optionally its gradient function `grad`
// to `g`. Once `func`/`grad` is added to `g`, it can be called by creating
// an operation using the function's name.
// Any changes to `func`/`grad` (including deleting it) done after this method
// returns, won't affect the copy of `func`/`grad` in `g`.
// If `func` or `grad` are already in `g`, TF_GraphCopyFunction has no

std::vector<std::pair<string, string>> GetGradDefs(
    const tensorflow::GraphDef& graph_def) {
  std::vector<std::pair<string, string>> grads;
  for (const tensorflow::GradientDef& grad : graph_def.library().gradient()) {
    grads.emplace_back(grad.function_name(), grad.gradient_func());
  }
  std::sort(grads.begin(), grads.end());
  return grads;
}

std::vector<string> GetFuncNames(const tensorflow::GraphDef& graph_def) {

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