def METHOD_NAME(
    shape_nhwc,
    window_shape,
    stride,
    padding,
    dtype,
    storage_scope="global",
):
    """
    Maxpool2d TE wherein the input activation is defined by its
    logical NHWC shape. The packed physical layout for the
    activation is nhwc8h8w32c.
    """
    block_h, block_w, block_c = get_block_shape()
    shape = get_packed_shape(shape_nhwc)
    logical_output_shape = (
        shape_nhwc[0],
        (shape_nhwc[1] - window_shape[0] + padding[0] + padding[1]) // stride[0] + 1,
        (shape_nhwc[2] - window_shape[1] + padding[2] + padding[3]) // stride[0] + 1,
        shape_nhwc[3],
    )
    output_shape = get_packed_shape(logical_output_shape)
    _, height, width, _ = shape_nhwc
    placeholder_x = te.placeholder(shape_nhwc, dtype=dtype)
    # Combination of padding required by maxpool operator and padding to evenly divisible
    # number of blocks. Note that this padding should be inlined in the schedule so
    # as to avoid input copying.
    pad_h = (block_h - ((height + padding[1]) % block_h)) % block_h
    pad_w = (block_w - ((width + padding[3]) % block_w)) % block_w
    x_pad = topi.nn.pad(
        placeholder_x, [0, padding[0], padding[2], 0], [0, pad_h, pad_w, 0], pad_value=0
    )
    # Calculate packed layout
    x_packed = te.compute(
        shape,
        lambda n, ho, wo, co, hi, wi, ci: x_pad[
            n, ho * block_h + hi, wo * block_w + wi, co * block_c + ci
        ],
    )
    reduce_h = te.reduce_axis((0, window_shape[0]), name="rh")
    reduce_w = te.reduce_axis((0, window_shape[1]), name="rw")
    def compute(batch, h_outer, w_outer, c_outer, h_inner, w_inner, c_inner):
        # Construct blockized strided maxpool height indices
        h = h_outer * block_h + h_inner
        h_contig = h * stride[0] + reduce_h
        h_block_id = h_contig // block_h
        h_block_offset = h_contig % block_h
        # Construct blockized strided maxpool width indices
        w_idx = w_outer * block_w + w_inner
        w_contig = w_idx * stride[1] + reduce_w
        w_block_id = w_contig // block_w
        w_block_offset = w_contig % block_w
        return te.max(
            x_packed[
                batch, h_block_id, w_block_id, c_outer, h_block_offset, w_block_offset, c_inner
            ],
            axis=[reduce_h, reduce_w],
        )
    compute_y = te.compute(output_shape, compute)
    schedule = te.create_schedule(compute_y.op)
    # Ensure the padding and array packing is performed inline
    schedule[x_pad].compute_inline()
    schedule[x_packed].compute_inline()
    binds = {}
    if storage_scope and storage_scope != "global":
        with tvm.transform.PassContext():
            x_buffer = tvm.tir.decl_buffer(shape, name="Xb", dtype=dtype, scope=storage_scope)
            y_buffer = tvm.tir.decl_buffer(
                output_shape, name="Yb", dtype=dtype, scope=storage_scope
            )
            binds = {placeholder_x: x_buffer, compute_y: y_buffer}
    return (schedule, [placeholder_x, compute_y], binds)