code: def conv2d_winograd_comp(
      cfg,
      data,
      kernel,
      strides,
      padding,
      dilation,
      out_dtype,
      pre_computed,
      layout,
  ):
      """Compute declaration for winograd
  
      Parameters
      ----------
      cfg: ConfigEntity
          The config for this template
  
      data: tvm.te.Tensor
          4-D or 5-D Data tensor with shape NCHW or NCHW4c
  
      kernel: tvm.te.Tensor
          4-D or 5-D tensor with shape OIHW or OIHW4o
  
      strides: int or a list/tuple of two ints
          stride size, or [stride_height, stride_width]
  
      padding: int or a list/tuple of 2 or 4 ints
          padding size, or
          [pad_height, pad_width] for 2 ints, or
          [pad_top, pad_left, pad_bottom, pad_right] for 4 ints
  
      dilation: int or a list/tuple of two ints
          dilation size, or [dilation_height, dilation_width]
  
      out_dtype: str
          The output type. This is used for mixed precision.
  
      pre_computed: bool
          Flag if weights were pre computed if true or the weights should be
          computed in runtime
  
      layout: str
          NHWC or NCHW values are accepted
  
      Returns
      -------
      output: tvm.te.Tensor
          4-D or 5-D with shape NCHW or NCHW4c
      """
      assert layout in ("NCHW", "NHWC")
      tile_size = infer_tile_size(data, layout)
  
      if isinstance(dilation, int):
          dilation_h = dilation_w = dilation
      else:
          dilation_h, dilation_w = dilation
      HSTR, WSTR = (
          (strides, strides) if isinstance(strides, int) else strides
      )
  
      convert_from4d = False
      if len(data.shape) == 4:
          convert_from4d = True
          if layout == "NCHW":
              N, DCI, H, W = get_const_tuple(data.shape)
          else:
              N, H, W, DCI = get_const_tuple(data.shape)
          if not pre_computed:
              if layout == "NCHW":
                  out_channels, CI, KH, KW = get_const_tuple(
                      kernel.shape
                  )
              else:
                  KH, KW, CI, out_channels = get_const_tuple(
                      kernel.shape
                  )
          else:
              alpha, _, CI, out_channels = get_const_tuple(kernel.shape)
              KH = KW = alpha + 1 - tile_size
  
          (
              in_channel_chunks,
              in_channel_block,
              in_channel_tail,
          ) = split_to_chunks(CI, 4)
          (
              out_channel_chunks,
              out_channel_block,
              out_channel_tail,
          ) = split_to_chunks(out_channels, 4)
          if autotvm.GLOBAL_SCOPE.in_tuning is True:
              if layout == "NCHW":
                  dshape = (
                      N,
                      in_channel_chunks,
                      H,
                      W,
                      in_channel_block,
                  )
              else:
                  dshape = (
                      N,
                      H,
                      W,
                      in_channel_chunks,
                      in_channel_block,
                  )
              if (
                  not pre_computed
              ):  # kernel tensor is raw tensor, do strict check
                  if layout == "NCHW":
                      kshape = (
                          out_channel_chunks,
                          CI,
                          KH,
                          KW,
                          out_channel_block,
                      )
                  else:
                      kshape = (
                          KH,
                          KW,
                          CI,
                          out_channel_chunks,
                          out_channel_block,
                      )
              else:
                  kshape = (
                      alpha,
                      alpha,
                      CI,
                      out_channel_chunks,
                      out_channel_block,
                  )
              data = tvm.te.placeholder(
                  dshape, data.dtype, name="data_placeholder"
              )
              kernel = tvm.te.placeholder(
                  kshape, kernel.dtype, name="kernel_placeholder"
              )
          else:
              data = pack_input(
                  data,
                  layout,
                  N,
                  in_channel_chunks,
                  in_channel_block,
                  in_channel_tail,
                  H,
                  W,
              )
              kernel_layout = "OIHW" if layout == "NCHW" else "HWIO"
              if (
                  not pre_computed
              ):  # kernel tensor is raw tensor, do strict check
                  kernel = pack_filter(
                      kernel,
                      kernel_layout,
                      out_channel_chunks,
                      out_channel_block,
                      out_channel_tail,
                      CI,
                      in_channel_chunks,
                      in_channel_block,
                      in_channel_tail,
                      KH,
                      KW,
                  )
              else:
                  kernel = pack_filter(
                      kernel,
                      "HWIO",
                      out_channel_chunks,
                      out_channel_block,
                      out_channel_tail,
                      CI,
                      in_channel_chunks,
                      in_channel_block,
                      in_channel_tail,
                      alpha,
                      alpha,
                  )
      if layout == "NCHW":
          N, DCI, H, W, CB = get_const_tuple(data.shape)
      else:
          N, H, W, DCI, CB = get_const_tuple(data.shape)
      if (
          not pre_computed
      ):  # kernel tensor is raw tensor, do strict check
          if layout == "NCHW":
              CO, CI, KH, KW, COB = get_const_tuple(kernel.shape)
          else:
              KH, KW, CI, CO, COB = get_const_tuple(kernel.shape)
          alpha = KW + tile_size - 1
          assert HSTR == 1 and WSTR == 1 and KH == KW
      else:
          alpha, _, CI, CO, COB = get_const_tuple(kernel.shape)
          KH = KW = alpha + 1 - tile_size
          assert (
              HSTR == 1
              and WSTR == 1
              and dilation_h == 1
              and dilation_w == 1
          )
  
      if isinstance(N, tvm.tir.Any):
          N = tvm.te.size_var("n")
  
      if not isinstance(H, int) or not isinstance(W, int):
          raise RuntimeError(
              "adreno winograd conv2d doesn't support dynamic input\
                             height or width."
          )
  
      pt, pl, pb, pr = nn.get_pad_tuple(padding, (KH, KW))
      if layout == "NCHW":
          data_pad = nn.pad(
              data,
              (0, 0, pt, pl, 0),
              (0, 0, pb, pr, 0),
              name="data_pad",
          )
      else:
          data_pad = nn.pad(
              data,
              (0, pt, pl, 0, 0),
              (0, pb, pr, 0, 0),
              name="data_pad",
          )
  
      r = KW
      m = tile_size
      A, B, G = winograd_transform_matrices(m, r, data.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")
          if layout == "NCHW":
              kernel_pack = te.compute(
                  (alpha, alpha, CI, CO, COB),
                  lambda eps, nu, ci, co, cob: te.sum(
                      kernel[co][ci][r_kh][r_kw][cob]
                      * G[eps][r_kh]
                      * G[nu][r_kw],
                      axis=[r_kh, r_kw],
                  ),
                  name="kernel_pack",
              )
          else:
              kernel_pack = te.compute(
                  (alpha, alpha, CI, CO, COB),
                  lambda eps, nu, ci, co, cob: te.sum(
                      kernel[r_kh][r_kw][ci][co][cob]
                      * G[eps][r_kh]
                      * G[nu][r_kw],
                      axis=[r_kh, r_kw],
                  ),
                  name="kernel_pack",
              )
      else:
          kernel_pack = kernel
  
      idxdiv = tvm.tir.indexdiv
      idxmod = tvm.tir.indexmod
      if layout == "NCHW":
          N, CI, _, _, CB = get_const_tuple(data.shape)
      else:
          N, _, _, CI, CB = get_const_tuple(data.shape)
  
      # pack input tile
      if layout == "NCHW":
          input_tile = te.compute(
              (alpha, alpha, CI, P, CB),
              lambda eps, nu, c, p, cb: data_pad[idxdiv(p, (nH * nW))][
                  c
              ][idxmod(idxdiv(p, nW), nH) * m + eps][
                  idxmod(p, nW) * m + nu
              ][
                  cb
              ],
              name="d",
          )
      else:
          input_tile = te.compute(
              (alpha, alpha, CI, P, CB),
              lambda eps, nu, c, p, cb: data_pad[idxdiv(p, (nH * nW))][
                  idxmod(idxdiv(p, nW), nH) * m + eps
              ][idxmod(p, nW) * m + nu][c][cb],
              name="d",
          )
  
      # transform data
      r_a = te.reduce_axis((0, alpha), "r_a")
      r_b = te.reduce_axis((0, alpha), "r_a")
      data_pack = te.compute(
          (P, CI, alpha, alpha, CB),
          lambda p, ci, eps, nu, cb: te.sum(
              input_tile[r_a][r_b][ci][p][cb]
              * B[r_a][eps]
              * B[r_b][nu],
              axis=[r_a, r_b],
          ),
          name="data_pack",
      )
  
      # repack transformed data
      data_pack_trans = te.compute(
          (alpha, alpha, CI, P, CB),
          lambda eps, nu, c, p, cb: data_pack[p][c][eps][nu][cb],
          name="data_pack_trans",
      )
  
      # do batch gemm
      ci = te.reduce_axis((0, CI), name="ci")
      cb = te.reduce_axis((0, CB), name="cb")
      bgemm = te.compute(
          (alpha, alpha, CO, P, COB),
          lambda eps, nu, co, p, cob: te.sum(
              (
                  kernel_pack[eps][nu][ci * CB + cb][co][cob]
                  * data_pack_trans[eps][nu][ci][p][cb]
              ).astype(out_dtype),
              axis=[ci, cb],
          ),
          name="bgemm",
      )
  
      # inverse transform
      r_a = te.reduce_axis((0, alpha), "r_a")
      r_b = te.reduce_axis((0, alpha), "r_a")
      inverse = te.compute(
          (CO, P, m, m, COB),
          lambda co, p, vh, vw, cob: te.sum(
              bgemm[r_a][r_b][co][p][cob]
              * (A[r_a][vh] * A[r_b][vw]).astype(out_dtype),
              axis=[r_a, r_b],
          ),
          name="inverse",
      )
  
      # output
      if layout == "NCHW":
          if convert_from4d and autotvm.GLOBAL_SCOPE.in_tuning is False:
              output = te.compute(
                  (N, out_channels, H, W),
                  lambda n, c, h, w: inverse[c // CB][
                      n * nH * nW + idxdiv(h, m) * nW + idxdiv(w, m)
                  ][idxmod(h, m)][idxmod(w, m)][c % CB].astype(
                      out_dtype
                  ),
                  name="output",
                  tag="dummy_compute_at",
              )
          else:
              output = te.compute(
                  (N, CO, H, W, COB),
                  lambda n, co, h, w, cob: inverse[co][
                      n * nH * nW + idxdiv(h, m) * nW + idxdiv(w, m)
                  ][idxmod(h, m)][idxmod(w, m)][cob].astype(out_dtype),
                  name="output",
                  tag="dummy_compute_at",
              )
      else:
          if convert_from4d and autotvm.GLOBAL_SCOPE.in_tuning is False:
              output = te.compute(
                  (N, H, W, out_channels),
                  lambda n, h, w, c: inverse[c // CB][
                      n * nH * nW + idxdiv(h, m) * nW + idxdiv(w, m)
                  ][idxmod(h, m)][idxmod(w, m)][c % CB].astype(
                      out_dtype
                  ),
                  name="output",
                  tag="dummy_compute_at",
              )
          else:
              output = te.compute(
                  (N, H, W, CO, COB),
                  lambda n, h, w, co, cob: inverse[co][
                      n * nH * nW + idxdiv(h, m) * nW + idxdiv(w, m)
                  ][idxmod(h, m)][idxmod(w, m)][cob].astype(out_dtype),
                  name="output",
                  tag="dummy_compute_at",
              )
  
      if isinstance(N, int):
          cfg.add_flop(2 * N * CO * COB * H * W * CI * CB * KH * KW)
  
      return output
  
file_path: tvm/python/tvm/topi/adreno/conv2d_winograd_common.py
imports: import tvm
  from tvm import te
  from tvm import autotvm
  from tvm.topi import nn
  from tvm.topi.utils import get_const_int, get_const_tuple, traverse_inline
  from ..nn.winograd_util import winograd_transform_matrices
  from .utils import split_to_chunks, pack_input, pack_filter, bind_data_copy, get_texture_storage, infer_tile_size
label_desc: Use image node_id 1 for calling a global function with example usage: conv2d_winograd_comp(cfg, data, kernel, strides, padding, dilation, out_dtype, pre_computed, layout) and returns: output