data.h_in = op_type.getShape()[2];
  data.w_in = op_type.getShape()[3];

  std::vector<int32_t> kernel_size;
  GetI32VectorFromDenseI64CompositeAttr(composite_attrs, "kernel_size",
                                        &kernel_size);
  data.kh = kernel_size[0];
  data.kw = kernel_size[1];

  std::vector<int32_t> padding;
  GetI32VectorFromDenseI64CompositeAttr(composite_attrs, "padding", &padding);

                                    ConstantAttr<RankedI32ElementsAttr<[4]>,"{0, 2, 3, 1}">)),
                            /*filter_height*/(GetI32At<0> (GetAsVectorAttr<"kernel_size"> $attrs)),
                            /*filter_width*/(GetI32At<1> (GetAsVectorAttr<"kernel_size"> $attrs)),
                            /*padding*/(GetAvgPoolOpPadAttr $old_val),
                            /*stride_h*/(GetI32At<0> (GetAsVectorAttr<"stride"> $attrs)),

func.func @avg_pool2d_1(%arg0: tensor<1x3x6x6xf32>) -> (tensor<*xf32>) {
  %0 = mhlo.composite "aten.avg_pool2d.default" %arg0 {composite_attributes = {ceil_mode = false, count_include_pad = true, divisor_override = "py_None", kernel_size = dense<3> : tensor<2xi64>, padding = dense<0> : tensor<2xi64>, stride = dense<1> : tensor<2xi64>}, decomposition = @XlaCallModule_aten.avg_pool2d.default.impl_0} : (tensor<1x3x6x6xf32>) -> tensor<1x3x4x4xf32>

    kernel of shape [7, 7, 3, 64] to [4, 4, 12, 84].

    ```python
    conv0 = tf.compat.v1.layers.Conv2D(
     filters=filters,
     kernel_size=kernel_size,
     strides=2,
     padding=('SAME' if strides == 1 else 'VALID'),
     use_bias=False,
     kernel_initializer=tf.variance_scaling_initializer(),
     data_format=data_format)

    }
  }
  return success();
}

// Handles padding before convolution for space to depth transform.
LogicalResult HandlePad(TF::PadOp op, int32_t kernel_size, int32_t block_size) {
  auto ranked_type = mlir::dyn_cast<RankedTensorType>(op.getInput().getType());
  if (!ranked_type) return failure();
  auto pad_input_shape = ranked_type.getShape();
  Location loc = op.getLoc();

    TF_Buffer* buf = TF_GetRegisteredKernelsForOp(op_name, status);
    EXPECT_EQ(TF_OK, TF_GetCode(status));
    KernelList list;
    list.ParseFromArray(buf->data, buf->length);
    ASSERT_EQ(1, list.kernel_size());
    ASSERT_EQ(device_name, list.kernel(0).device_type());
    TF_DeleteBuffer(buf);
    TF_DeleteStatus(status);
  }

  {
    Status status;
    std::unique_ptr<OpKernel> kernel =

  EXPECT_EQ(TF_OK, TF_GetCode(status)) << TF_Message(status);
  KernelList kernel_list;
  kernel_list.ParseFromArray(kernel_list_buf->data, kernel_list_buf->length);
  ASSERT_GT(kernel_list.kernel_size(), 0);
  TF_DeleteBuffer(kernel_list_buf);
  TF_DeleteStatus(status);
}

TEST(TestKernel, TestGetRegisteredKernelsForOp) {
  TF_Status* status = TF_NewStatus();
  TF_Buffer* kernel_list_buf =

      // 2. Output size respects the `VALID` padding scenario
      if ((padding[2 * i] == padding[2 * i + 1]) &&
          (((kernel_size - 1) != padding[2 * i]) ||
           (output_size != (stride * (input_size - 1)) + kernel_size))) {
        // padding[2 * i] == padding[2 * i + 1] means equal padding is applied
        // on both sides of a spatial dimension.