def METHOD_NAME(cfg, s, pad_data, kernel, conv, layout):
    """schedule depthwise_conv2d"""
    assert layout in ("NCHW", "NHWC")
    max_unroll = 16
    vec_size = [1, 2, 4, 8, 16]
    ##### space definition begin #####
    if layout == "NCHW":
        n, c, h, w = s[conv].op.axis
    else:
        n, h, w, c = s[conv].op.axis
    bc, tc, ci = cfg.define_split("tile_c", c, num_outputs=3)
    bh, th, hi = cfg.define_split("tile_y", h, num_outputs=3)
    bw, tw, wi = cfg.define_split("tile_x", w, num_outputs=3)
    cfg.define_annotate("ann_spatial", [ci, hi, wi], policy="try_unroll_vec")
    # fallback support
    if cfg.is_fallback:
        if layout == "NCHW":
            ref_log = autotvm.tophub.load_reference_log(
                "mali", "rk3399", "depthwise_conv2d_nchw.mali"
            )
            cfg.fallback_with_reference_log(ref_log)
        else:
            cfg.fallback_split("tile_c", [-1, 4, 2])
            cfg.fallback_split("tile_y", [-1, 4, 2])
            cfg.fallback_split("tile_x", [-1, 4, 2])
    ###### space definition end ######
    # schedule padding
    if layout == "NCHW":
        n, c, h, w = s[pad_data].op.axis
        z, y, x = c, h, w
        z_factor, y_factor, x_factor = cfg["tile_c"].size[1], 1, 1
    else:
        n, h, w, c = s[pad_data].op.axis
        z, y, x = h, w, c
        z_factor, y_factor, x_factor = 1, 1, cfg["tile_c"].size[1]
    tile_and_bind3d(s, pad_data, z, y, x, z_factor, y_factor, x_factor)
    # schedule dilation
    if isinstance(kernel.op, tvm.te.ComputeOp) and "dilate" in kernel.op.tag:
        s[kernel].compute_inline()
    # schedule conv
    if conv.op not in s.outputs:
        s[conv].set_scope("local")
        OL = conv
        output = s.outputs[0].output(0)
    else:
        OL = s.cache_write(conv, "local")
        output = conv
    if layout == "NCHW":
        n, c, h, w = s[output].op.axis
    else:
        n, h, w, c = s[output].op.axis
    bc, tc, ci = cfg["tile_c"].apply(s, output, c)
    bh, th, hi = cfg["tile_y"].apply(s, output, h)
    bw, tw, wi = cfg["tile_x"].apply(s, output, w)
    if layout == "NCHW":
        bz, tz, by, ty, bx, tx = bc, tc, bh, th, bw, tw
    else:
        bz, tz, by, ty, bx, tx = bh, th, bw, tw, bc, tc
    bz = s[output].fuse(n, bz)
    s[output].bind(bz, te.thread_axis("blockIdx.z"))
    s[output].bind(tz, te.thread_axis("threadIdx.z"))
    s[output].bind(by, te.thread_axis("blockIdx.y"))
    s[output].bind(ty, te.thread_axis("threadIdx.y"))
    s[output].bind(bx, te.thread_axis("blockIdx.x"))
    s[output].bind(tx, te.thread_axis("threadIdx.x"))
    di, dj = s[OL].op.reduce_axis
    s[OL].unroll(di)
    s[OL].unroll(dj)
    s[OL].compute_at(s[output], tx)
    if layout == "NCHW":
        n, ci, hi, wi = s[OL].op.axis
    else:
        n, hi, wi, ci = s[OL].op.axis
    cfg["ann_spatial"].apply(
        s,
        OL,
        [ci, hi, wi],
        axis_lens=[cfg["tile_c"].size[2], cfg["tile_y"].size[2], cfg["tile_x"].size[2]],
        max_unroll=max_unroll,
        vec_size=vec_size,
        cfg=cfg,
    )