code: def _decl_winograd(
      cfg,
      data,
      kernel,
      strides,
      padding,
      dilation,
      out_dtype,
      tile_size,
  ):
      N, CI, IH, IW = get_const_tuple(data.shape)
      if isinstance(dilation, int):
          dilation_h = dilation_w = dilation
      else:
          dilation_h, dilation_w = dilation
  
      if len(kernel.shape) == 4:
          if dilation_h != 1 or dilation_w != 1:
              kernel = nn.dilate(kernel, (1, 1, dilation_h, dilation_w))
          pre_computed = False
          CO, _, KH, KW = get_const_tuple(kernel.shape)
      else:
          assert (dilation_h, dilation_w) == (
              1,
              1,
          ), "Does not support dilation"
          pre_computed = True
          H_CAT, W_CAT, CO, CI, VC = get_const_tuple(kernel.shape)
          CO *= VC
          KH, KW = H_CAT - tile_size + 1, W_CAT - tile_size + 1
      HSTR, WSTR = (
          strides
          if isinstance(strides, (tuple, list))
          else (strides, strides)
      )
      pt, pl, pb, pr = nn.get_pad_tuple(padding, (KH, KW))
  
      assert KH == 3 and KW == 3 and HSTR == 1 and WSTR == 1
      data_pad = nn.pad(
          data, (0, 0, pt, pl), (0, 0, pb, pr), name="data_pad"
      )
  
      r = KW
      m = tile_size
      alpha = m + r - 1
      A, B, G = winograd_transform_matrices(m, r, out_dtype)
  
      H = (IH + pt + pb - 3) // HSTR + 1
      W = (IW + pl + pr - 3) // WSTR + 1
      nH, nW = (H + m - 1) // m, (W + m - 1) // m
      P = N * nH * nW
  
      ##### space definition begin #####
      tile_bna_candidates = [1, 2, 4, 8, 16]
      factors = get_factors(CO)
      cfg.define_knob(
          "tile_bna", [x for x in tile_bna_candidates if x in factors]
      )
      cfg.define_knob("tile_bnb", [1, 2, 4, 8, 16])
      cfg.define_split("tile_t1", CI, num_outputs=2, max_factor=128)
      cfg.define_split("tile_t2", CO, num_outputs=2, max_factor=128)
      cfg.define_split("c_unroll", CI, num_outputs=2, max_factor=8)
      cfg.define_knob("yt", [1, 2, 4, 8, 16, 32])
      ##### space definition end #####
  
      if cfg.is_fallback:
          cfg["tile_bnb"].val = 4
          cfg["tile_bna"].val = 4
          while CO % cfg["tile_bna"].val != 0:
              cfg["tile_bna"].val //= 2
          cfg["yt"].val = 8
          cfg.fallback_split("tile_t1", [-1, 128])
          cfg.fallback_split("tile_t2", [-1, 128])
          cfg.fallback_split("c_unroll", [-1, 8])
  
      bna = cfg["tile_bna"].val
      bnb = cfg["tile_bnb"].val
  
      P_round = (P + bnb - 1) // bnb * bnb
      assert CO % bna == 0 and P_round % bnb == 0
  
      # pack input tile
      input_tile = te.compute(
          (CI, P_round // bnb, alpha, alpha, bnb),
          lambda ci, b, eps, nu, bb: tvm.tir.if_then_else(
              b * bnb + bb < P,
              data_pad[(b * bnb + bb) // (nH * nW)][ci][
                  (b * bnb + bb) // nW % nH * m + eps
              ][(b * bnb + bb) % nW * m + nu],
              tvm.tir.const(0, data_pad.dtype),
          ),
          name="d",
      )
  
      if autotvm.GLOBAL_SCOPE.in_tuning:
          kvshape = (alpha, alpha, CO // bna, CI, bna)
          U = tvm.te.placeholder(kvshape, kernel.dtype, name="U")
      else:
          # transform kernel
          if pre_computed:
              U = kernel
          else:
              r_kh = te.reduce_axis((0, KH), "r_kh")
              r_kw = te.reduce_axis((0, KW), "r_kw")
              U = te.compute(
                  (alpha, alpha, CO // bna, CI, bna),
                  lambda eps, nu, co, ci, vco: te.sum(
                      kernel[co * bna + vco][ci][r_kh][r_kw]
                      * G[eps][r_kh]
                      * G[nu][r_kw],
                      axis=[r_kh, r_kw],
                  ),
                  name="U",
              )
  
      # transform image
      r_a = te.reduce_axis((0, alpha), "r_a")
      r_b = te.reduce_axis((0, alpha), "r_b")
      V = te.compute(
          (alpha, alpha, P_round // bnb, CI, bnb),
          lambda eps, nu, p, ci, vp: te.sum(
              input_tile[ci][p][r_a][r_b][vp]
              * B[r_a][eps]
              * B[r_b][nu],
              axis=[r_a, r_b],
          ),
          name="V",
      )
  
      idxdiv = tvm.tir.indexdiv
      idxmod = tvm.tir.indexmod
  
      # batch gemm
      ci = te.reduce_axis((0, CI), name="c")
      M = te.compute(
          (alpha, alpha, CO, P_round),
          lambda eps, nu, co, p: te.sum(
              U[eps][nu][idxdiv(co, bna)][ci][idxmod(co, bna)]
              * V[eps][nu][idxdiv(p, bnb)][ci][idxmod(p, bnb)],
              axis=ci,
          ),
          name="M",
      )
  
      r_a = te.reduce_axis((0, alpha), "r_a")
      r_b = te.reduce_axis((0, alpha), "r_b")
      Y = te.compute(
          (CO, P, m, m),
          lambda co, p, vh, vw: te.sum(
              M[r_a][r_b][co][p] * A[r_a][vh] * A[r_b][vw],
              axis=[r_a, r_b],
          ),
          name="Y",
      )
  
      # unpack output
      output = te.compute(
          (N, CO, H, W),
          lambda n, co, h, w: Y[
              co,
              n * nH * nW + idxdiv(h, m) * nW + idxdiv(w, m),
              idxmod(h, m),
              idxmod(w, m),
          ]
          # The following hack term is used to make the padding in batch gemm ("M")
          # effective, otherwise the padding will be eliminated by bound inference.
          # Use `tvm.tir.Mul` instead of `*` to avoid issues in const folding.
          + tvm.tir.Mul(
              tvm.tir.const(0, out_dtype),
              M[alpha - 1][alpha - 1][CO - 1][P_round - 1],
          ),
          name="output",
          tag="winograd_conv2d_output",
      )
  
      # we have to manually assign effective GFLOP for winograd
      cfg.add_flop(2 * N * CO * H * W * KH * KW * CI)
      return output
  
imports: import logging
  import tvm
  from tvm import te
  from tvm import relay
  from tvm import autotvm
  from tvm.autotvm.task.space import get_factors
  from ..utils import traverse_inline, get_const_int, get_const_tuple
  from ..None import nn
  from ..nn.winograd_util import winograd_transform_matrices
  from ..nn.conv2d import conv2d_winograd_nhwc, _conv2d_winograd_nhwc_impl
  from ..arm_cpu.conv2d_spatial_pack import conv2d_spatial_pack_nchw
  from ..arm_cpu.conv2d_spatial_pack import conv2d_spatial_pack_nhwc
label_desc: Use image node_id 9 for calling a global function with example usage: _decl_winograd(cfg, data, kernel, strides, padding, dilation, out_dtype, tile_size) and returns: output