code: def _conv2d_winograd_nchw_impl(
      data,
      weight,
      strides,
      padding,
      dilation,
      out_dtype,
      tile_size,
      pre_computed=False,
      write_cache_level=None,
      auto_scheduler_rewritten_layout="",
      meta_schedule_original_shape=None,
  ):
      """
      write_cache_level: Optional[int] = None
          The cache level to write to in multi-level tiling rule in MetaSchedule.
      """
      del auto_scheduler_rewritten_layout
  
      N, CI, H, W = get_const_tuple(data.shape)
      if isinstance(dilation, int):
          dilation_h = dilation_w = dilation
      else:
          dilation_h, dilation_w = dilation
      if meta_schedule_original_shape:
          auto_scheduler.rewrite_tensor_shape(
              weight, meta_schedule_original_shape
          )
  
      assert (dilation_h, dilation_w) == (
          1,
          1,
      ), "Does not support dilation"
      HSTR, WSTR = (
          (strides, strides) if isinstance(strides, int) else strides
      )
  
      if (
          not pre_computed
      ):  # kernel tensor is raw tensor, do strict check
          CO, CI, KH, KW = get_const_tuple(weight.shape)
          alpha = KW + tile_size - 1
          assert HSTR == 1 and WSTR == 1 and KH == KW
      else:
          alpha, _, CI, CO = get_const_tuple(weight.shape)
          KH = KW = alpha + 1 - tile_size
          assert (
              HSTR == 1
              and WSTR == 1
              and dilation_h == 1
              and dilation_w == 1
          )
  
      pad_t, pad_l, pad_b, pad_r = get_pad_tuple(padding, (KH, KW))
      assert HSTR == 1 and WSTR == 1 and KH == 3 and KW == 3
  
      pt, pl, pb, pr = get_pad_tuple(padding, (KH, KW))
      data_pad = pad(
          data, (0, 0, pt, pl), (0, 0, pb, pr), name="data_pad"
      )
  
      r = KW
      m = tile_size
      A, B, G = winograd_transform_matrices(m, r, out_dtype)
  
      H = (H + pt + pb - KH) // HSTR + 1
      W = (W + pl + pr - KW) // WSTR + 1
      nH, nW = (H + m - 1) // m, (W + m - 1) // m
  
      P = N * nH * nW if isinstance(N, int) else nH * nW
  
      # transform kernel
      if not pre_computed:
          r_kh = te.reduce_axis((0, KH), name="r_kh")
          r_kw = te.reduce_axis((0, KW), name="r_kw")
          kernel_pack = te.compute(
              (alpha, alpha, CI, CO),
              lambda eps, nu, ci, co: te.sum(
                  weight[co, ci, r_kh, r_kw]
                  * G[eps, r_kh]
                  * G[nu, r_kw],
                  axis=[r_kh, r_kw],
              ),
              name="kernel_pack",
          )
          bgemm_attrs = {}
      else:
          kernel_pack = weight
          bgemm_attrs = {"layout_free_placeholders": [kernel_pack]}
      if write_cache_level is not None:
          if not isinstance(write_cache_level, int):
              bgemm_attrs[
                  "meta_schedule.write_cache_level"
              ] = write_cache_level
          else:
              bgemm_attrs["meta_schedule.write_cache_level"] = [
                  write_cache_level
              ]
  
      # pack data tile
      input_tile = te.compute(
          (CI, P, alpha, alpha),
          lambda ci, p, eps, nu: data_pad[
              p // (nH * nW),
              ci,
              ((p // nW) % nH) * m + eps,
              (p % nW) * m + nu,
          ],
          name="input_tile",
          attrs={"schedule_rule": "None"},
      )
  
      # transform data
      r_a = te.reduce_axis((0, alpha), "r_a")
      r_b = te.reduce_axis((0, alpha), "r_b")
      data_pack = te.compute(
          (alpha, alpha, CI, P),
          lambda eps, nu, ci, p: te.sum(
              input_tile[ci, p, r_a, r_b] * B[r_a, eps] * B[r_b, nu],
              axis=[r_a, r_b],
          ),
          name="data_pack",
          attrs={"schedule_rule": "conv2d_nchw_winograd_data_pack"},
      )
  
      # do batch gemm
      ci = te.reduce_axis((0, CI), name="ci")
      bgemm = te.compute(
          (alpha, alpha, CO, P),
          lambda eps, nu, co, p: te.sum(
              data_pack[eps, nu, ci, p] * kernel_pack[eps, nu, ci, co],
              axis=[ci],
          ),
          name="bgemm",
          attrs=bgemm_attrs,
      )
  
      # inverse transform
      r_a = te.reduce_axis((0, alpha), "r_a")
      r_b = te.reduce_axis((0, alpha), "r_b")
      inverse = te.compute(
          (CO, P, m, m),
          lambda co, p, vh, vw: te.sum(
              bgemm[r_a, r_b, co, p] * A[r_a, vh] * A[r_b, vw],
              axis=[r_a, r_b],
          ),
          name="inverse",
          attrs={"schedule_rule": "conv2d_nchw_winograd_inverse"},
      )
  
      # output
      output = te.compute(
          (N, CO, H, W),
          lambda n, co, h, w: inverse[
              co, n * nH * nW + (h // m) * nW + (w // m), h % m, w % m
          ],
          name="conv2d_winograd",
      )
  
      return output
  
imports: from __future__ import absolute_import
  import re
  from collections import namedtuple
  from typing import Optional, Sequence, Union
  import numpy
  import tvm
  from tvm import auto_scheduler, te
  from ..utils import get_const_int, get_const_tuple, simplify, tag
  from .pad import pad
  from .utils import get_pad_tuple, get_pad_tuple_generic
  from .winograd_util import winograd_transform_matrices
label_desc: Use image node_id 22 for calling a global function with example usage: _conv2d_winograd_nchw_impl(data, weight, strides, padding, dilation, out_dtype, tile_size, pre_computed, write_cache_level, auto_scheduler_rewritten_layout, meta_schedule_original_shape) and returns: output