code: def nhwc_winograd_cuda(
      cfg,
      data,
      kernel,
      strides,
      padding,
      dilation,
      out_dtype,
      use_tensorcore,
      pre_computed,
  ):
      """Compute declaration for winograd"""
      tile_size = _infer_tile_size(data, kernel)
      N, H, W, CI = get_const_tuple(data.shape)
  
      if isinstance(N, tvm.tir.Any):
          N = tvm.te.size_var("n")
  
      if not isinstance(H, int) or not isinstance(W, int):
          raise RuntimeError(
              "cuda winograd nhwc conv2d doesn't support dynamic \
                             input height or width."
          )
  
      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
      )
  
      if (
          not pre_computed
      ):  # Kernel tensor is raw tensor, do strict check
          if dilation_h != 1 or dilation_w != 1:
              kernel = nn.dilate(kernel, (dilation_h, dilation_w, 1, 1))
          KH, KW, CI, CO = get_const_tuple(kernel.shape)
          alpha = KW + tile_size - 1
          assert HSTR == 1 and WSTR == 1 and KH == KW
      else:
          # Kernel tensor is pre-transfomred. This op is created by conv2d_alter_op.
          # Dilation is not supported
          alpha, _, CI, CO = 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
          )
  
      pt, pl, pb, pr = nn.get_pad_tuple(padding, (KH, KW))
      data_pad = nn.pad(
          data,
          (0, pt, pl, 0),
          (0, pb, pr, 0),
          name="data_pad",
          attrs={"schedule_rule": "None"},
      )
  
      r = KW
      m = tile_size
      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
  
      # Determine whether the shape is available with tensorcore
      shape_judge = (
          (P % 16 == 0 and CI % 16 == 0 and CO % 16 == 0)
          or (P % 8 == 0 and CI % 16 == 0 and CO % 32 == 0)
          or (P % 32 == 0 and CI % 16 == 0 and CO % 8 == 0)
      )
  
      if shape_judge and use_tensorcore:
          trans_type = "float16"
      else:
          trans_type = data.dtype
  
      # Compute transform matrix
      A, _, _ = winograd_transform_matrices(m, r, out_dtype)
      _, B, G = winograd_transform_matrices(m, r, data.dtype)
  
      # Transform kernel
      if not pre_computed:
          # Check if we are currently tuning, if so we want to avoid counting
          # prepacking in time costs. Just use a placeholder with the packed shape instead.
          if autotvm.GLOBAL_SCOPE.in_tuning:
              kernel_pack = te.placeholder(
                  (alpha, alpha, CI, CO),
                  dtype=kernel.dtype,
                  name="kernel_pack",
              )
          else:
              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(
                      (kernel[r_kh][r_kw][ci][co])
                      * 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
  
      # Pack input tile
      input_tile = te.compute(
          (P, CI, alpha, alpha),
          lambda p, c, eps, nu: data_pad[
              idxdiv(p, (nH * nW)),
              idxmod(idxdiv(p, nW), nH) * m + eps,
              idxmod(p, nW) * m + nu,
              c,
          ],
          name="d",
          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[p][ci][r_a][r_b] * B[r_a][eps] * B[r_b][nu],
              axis=[r_a, r_b],
          ),
          name="data_pack",
      )
  
      # Convert data type of input feature maps and weights for tensorcore
      Transdata = te.compute(
          data_pack.shape,
          lambda eps, nu, p, ci: data_pack[eps, nu, p, ci].astype(
              trans_type
          ),
      )
      TransFilter = te.compute(
          kernel_pack.shape,
          lambda eps, nu, ci, co: kernel_pack[eps, nu, ci, co].astype(
              trans_type
          ),
      )
  
      # 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(
              (Transdata[eps][nu][p][ci]).astype(out_dtype)
              * (TransFilter[eps][nu][ci][co]).astype(out_dtype),
              axis=[ci],
          ),
          name="bgemm",
      )
  
      # Inverse transform
      r_a = te.reduce_axis((0, alpha), "r_a")
      r_b = te.reduce_axis((0, alpha), "r_b")
      inverse = te.compute(
          (P, CO, m, m),
          lambda p, co, vh, vw: te.sum(
              bgemm[r_a][r_b][p][co] * A[r_a][vh] * A[r_b][vw],
              axis=[r_a, r_b],
          ),
          name="inverse",
      )
  
      # Output
      output = te.compute(
          (N, H, W, CO),
          lambda n, h, w, co: inverse[
              n * nH * nW + idxdiv(h, m) * nW + idxdiv(w, m),
              co,
              idxmod(h, m),
              idxmod(w, m),
          ],
          name="output",
          tag="conv2d_nhwc_winograd",
      )
      if isinstance(N, int):
          cfg.add_flop(2 * N * CO * H * W * CI * KH * KW)
      return output
  
file_path: tvm/python/tvm/topi/cuda/conv2d_nhwc_winograd.py
imports: import tvm
  from tvm import autotvm, te
  from ..None import nn
  from ..nn.winograd_util import winograd_transform_matrices
  from ..utils import get_const_int, get_const_tuple, traverse_inline
  from .tensor_intrin import intrin_wmma_gemm, intrin_wmma_load_matrix_A, intrin_wmma_load_matrix_W, intrin_wmma_store_matrix
label_desc: Use image node_id 4 for calling a global function with example usage: nhwc_winograd_cuda(cfg, data, kernel, strides, padding, dilation, out_dtype, use_tensorcore, pre_computed) and returns: output