def METHOD_NAME(
    kernel, nthreads=1, algorithm=ConvolutionAlgorithm.AUTO, dtype="float32"
):
    """Create an extern op to do inference convolution of 3D tensor data and
    4D tensor kernel and 1D tensor bias with nnpack.
    Parameters
    ----------
    kernel : Tensor
        kernel 4D tensor kernel[output_channels][input_channels][kernel_height]
        [kernel_width] of FP32 elements.
    Returns
    -------
    output : Tensor
        output 4D tensor output[output_channels][input_channels][tile][tile]
        of FP32 elements.
    """
    assert algorithm in (ConvolutionAlgorithm.WT_8x8, ConvolutionAlgorithm.WT_8x8_FP16)
    output_channels, input_channels, _, _ = kernel.shape
    transform_tile_size = 8
    if not isinstance(dtype, str):
        dtype = dtype.dtype
    return te.extern(
        (output_channels, input_channels, transform_tile_size, transform_tile_size),
        [kernel],
        lambda ins, outs: tvm.tir.call_packed(
            "tvm.contrib.nnpack.convolution_inference_weight_transform",
            ins[0],
            outs[0],
            nthreads,
            algorithm,
        ),
        name="transform_kernel",
        dtype=dtype,
    )