code: def _conv2d_winograd_nhwc_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,
  ):
      """Conv2D Winograd implementation in NHWC layout.
      This is a clean version to be used by the auto-scheduler for both CPU and GPU.
  
      Parameters
      ----------
      data : tvm.te.Tensor
          4-D with shape [batch, in_height, in_width, in_channel]
      weight : tvm.te.Tensor
          4-D with shape [filter_height, filter_width, in_channel, num_filter]
      strides : int or a list/tuple of two ints
          stride size, or [stride_height, stride_width]
      padding : int or a list/tuple of two ints
          padding size, or [pad_height, pad_width]
      dilation: int or a list/tuple of two ints
          dilation size, or [dilation_height, dilation_width]
      out_dtype : str, optional
          Specifies the output data type.
      tile_size : int
          The size of the tile to use for the Winograd filter
      pre_computed: bool = False
          Whether the kernel is precomputed
      write_cache_level: Optional[int] = None
          The cache level to write to in multi-level tiling rule in MetaSchedule.
      auto_scheduler_rewritten_layout: str = ""
          The layout after auto-scheduler's layout rewrite pass.
      meta_schedule_original_shape: Optional[List[PrimExpr]] = None
          The original shape of the input tensor.
  
      Returns
      -------
      output : tvm.te.Tensor
          4-D with shape [batch, out_height, out_width, out_channel]
      """
      N, H, W, CI = 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"
      if not pre_computed:
          KH, KW, CI, CO = get_const_tuple(weight.shape)
      else:
          if auto_scheduler_rewritten_layout:
              H_CAT, W_CAT, CO, CI = get_const_tuple(
                  auto_scheduler.get_shape_from_rewritten_layout(
                      auto_scheduler_rewritten_layout,
                      ["eps", "nu", "co", "ci"],
                  )
              )
              auto_scheduler.remove_index_check(weight)
          else:
              H_CAT, W_CAT, CO, CI = get_const_tuple(weight.shape)
  
          KH, KW = H_CAT - tile_size + 1, W_CAT - tile_size + 1
  
      pad_t, pad_l, pad_b, pad_r = get_pad_tuple(padding, (KH, KW))
      HSTR, WSTR = (
          (strides, strides) if isinstance(strides, int) else strides
      )
      assert HSTR == 1 and WSTR == 1 and KH == 3 and KW == 3
  
      r = KW
      m = tile_size
      alpha = m + r - 1
      A, B, G = winograd_transform_matrices(m, r, out_dtype)
  
      H = (H + pad_t + pad_b - KH) // HSTR + 1
      W = (W + pad_l + pad_r - KW) // WSTR + 1
      nH, nW = (H + m - 1) // m, (W + m - 1) // m
      P = N * nH * nW
  
      pad_extra = (nW - 1) * m + alpha - (H + pad_t + pad_b)
      data_pad = pad(
          data,
          (0, pad_t, pad_l, 0),
          (0, pad_b + pad_extra, pad_r + pad_extra, 0),
          name="data_pad",
          attrs={"schedule_rule": "None"},
      )
  
      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, CO, CI),
              lambda eps, nu, co, ci: te.sum(
                  weight[r_kh, r_kw, ci, co]
                  * 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(
          (alpha, alpha, P, CI),
          lambda eps, nu, p, ci: data_pad[
              p // (nH * nW),
              ((p // nW) % nH) * m + eps,
              (p % nW) * m + nu,
              ci,
          ],
          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, P, CI),
          lambda eps, nu, p, ci: te.sum(
              input_tile[r_a, r_b, p, ci] * B[r_a, eps] * B[r_b, nu],
              axis=[r_a, r_b],
          ),
          name="data_pack",
          attrs={
              "auto_scheduler_simplify_const_tensor_indices": [
                  "eps",
                  "nu",
                  "r_a",
                  "r_b",
              ],
              "schedule_rule": "conv2d_nhwc_winograd_data_pack",
          },
      )
  
      # do batch gemm
      ci = te.reduce_axis((0, CI), name="ci")
      bgemm = te.compute(
          (alpha, alpha, P, CO),
          lambda eps, nu, p, co: te.sum(
              data_pack[eps, nu, p, ci] * kernel_pack[eps, nu, co, ci],
              axis=[ci],
          ),
          name="bgemm",
          attrs=bgemm_attrs,
      )
  
      if auto_scheduler_rewritten_layout:
          bgemm = auto_scheduler.rewrite_compute_body(
              bgemm, auto_scheduler_rewritten_layout
          )
  
      # inverse transform
  
      r_a = te.reduce_axis((0, alpha), "r_a")
      r_b = te.reduce_axis((0, alpha), "r_b")
      inverse = te.compute(
          (m, m, P, CO),
          lambda vh, vw, p, co: te.sum(
              bgemm[r_a, r_b, p, co] * A[r_a, vh] * A[r_b, vw],
              axis=[r_a, r_b],
          ),
          name="inverse",
          attrs={
              "auto_scheduler_simplify_const_tensor_indices": [
                  "vh",
                  "vw",
                  "r_a",
                  "r_b",
              ],
              "schedule_rule": "conv2d_nhwc_winograd_inverse",
          },
      )
  
      # output
      output = te.compute(
          (N, H, W, CO),
          lambda n, h, w, co: inverse[
              h % m, w % m, n * nH * nW + (h // m) * nW + (w // m), co
          ],
          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 21 for calling a global function with example usage: _conv2d_winograd_nhwc_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