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(N, tvm.tir.Any):
          N = tvm.te.size_var("n")
      if not isinstance(IH, int) or not isinstance(IW, int):
          raise RuntimeError(
              "ARM winograd conv2d doesn't support dynamic input height or width."
          )
  
      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 = 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"
      )
  
      idxd = tvm.tir.indexdiv
      idxm = tvm.tir.indexmod
  
      r = KW
      m = tile_size
      alpha = m + r - 1
      A, B, G = winograd_transform_matrices(m, r, out_dtype)
  
      K = CO
      C = CI
  
      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
  
      # TODO(@kevinthesun): Support tuning/optimization for dynamic shape.
      tile_p = P if isinstance(N, int) else nH * nW
      cfg.define_split(
          "tile_p",
          cfg.axis(tile_p),
          num_outputs=2,
          filter=lambda x: x.size[-1] <= 16,
      )
      cfg.define_split(
          "tile_k",
          cfg.axis(K),
          num_outputs=2,
          filter=lambda x: x.size[-1] <= 16,
      )
      VP = cfg["tile_p"].size[-1]
      VK = cfg["tile_k"].size[-1]
  
      # pack input tile
      input_tile = te.compute(
          (C, idxd(P, VP), alpha, alpha, VP),
          lambda c, b, eps, nu, bb: data_pad[
              idxd(b * VP + bb, nH * nW),
              c,
              idxm(idxd(b * VP + bb, nW), nH) * m + eps,
              idxm(b * VP + bb, nW) * m + nu,
          ],
          name="d",
      )
  
      if autotvm.GLOBAL_SCOPE.in_tuning:
          VC = cfg["tile_k"].size[-1]
          kvshape = (
              KH + tile_size - 1,
              KW + tile_size - 1,
              idxd(CO, VC),
              CI,
              VC,
          )
          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, idxd(K, VK), C, VK),
                  lambda eps, nu, k, c, kk: te.sum(
                      kernel[k * VK + kk][c][r_kh][r_kw].astype(
                          out_dtype
                      )
                      * G[eps][r_kh]
                      * G[nu][r_kw],
                      axis=[r_kh, r_kw],
                  ),
                  name="U",
              )
  
      # transform image
      r_eps = te.reduce_axis((0, alpha), "r_eps")
      r_nu = te.reduce_axis((0, alpha), "r_nu")
      V = te.compute(
          (alpha, alpha, idxd(P, VP), C, VP),
          lambda eps, nu, b, c, bb: te.sum(
              input_tile[c][b][r_eps][r_nu][bb].astype(out_dtype)
              * B[r_eps][eps]
              * B[r_nu][nu],
              axis=[r_eps, r_nu],
          ),
          name="V",
      )
  
      # batch gemm
      c = te.reduce_axis((0, C), name="c")
      M = te.compute(
          (alpha, alpha, K, P),
          lambda eps, nu, k, b: te.sum(
              U[eps][nu][idxd(k, VK)][c][idxm(k, VK)]
              * V[eps][nu][idxd(b, VP)][c][idxm(b, VP)],
              axis=c,
          ),
          name="M",
      )
  
      # inverse transform
      r_eps = te.reduce_axis((0, alpha), "r_eps")
      r_nu = te.reduce_axis((0, alpha), "r_nu")
      Y = te.compute(
          (K, P, m, m),
          lambda k, b, vh, vw: te.sum(
              M[r_eps][r_nu][k][b] * A[r_eps][vh] * A[r_nu][vw],
              axis=[r_eps, r_nu],
          ),
          name="Y",
      )
  
      # unpack output
      output = te.compute(
          (N, K, H, W),
          lambda n, k, h, w: Y[k][
              n * nH * nW + idxd(h, m) * nW + idxd(w, m),
              idxm(h, m),
              idxm(w, m),
          ],
          name="output",
          tag="winograd_conv2d_output",
      )
  
      # we have to manually assign effective GFLOP for winograd
      if isinstance(N, int):
          cfg.add_flop(2 * N * K * H * W * KH * KW * C)
      return output
  
imports: from __future__ import absolute_import
  import tvm
  from tvm import te
  from tvm import autotvm
  import tvm.contrib.nnpack
  from ..utils import traverse_inline, get_const_tuple
  from ..None import nn
  from ..nn.utils import get_const_int, get_pad_tuple
  from ..nn.winograd_util import winograd_transform_matrices
  from .arm_utils import get_tiling_B_transformed
  from .conv2d_spatial_pack import conv2d_spatial_pack_nchw, conv2d_spatial_pack_nhwc, schedule_conv2d_spatial_pack_nchw, schedule_conv2d_spatial_pack_nhwc
  from .conv2d_gemm import compute_conv2d_gemm_without_weight_transform, schedule_conv2d_gemm_interleaved, schedule_conv2d_gemm_native
  from .mprofile.dsp.conv2d import conv2d_nhwc_dsp_compute, conv2d_nhwc_dsp_schedule
label_desc: Use image node_id 7 for calling a global function with example usage: _decl_winograd(cfg, data, kernel, strides, padding, dilation, out_dtype, tile_size) and returns: output