code: def hwnc_tensorcore_cuda(
      cfg, Input, Filter, stride, padding, dilation, out_dtype="int32"
  ):
      """Compute declaration for tensorcore"""
      assert isinstance(stride, int) or len(stride) == 2
      assert isinstance(dilation, int) or len(dilation) == 2
  
      if isinstance(stride, int):
          stride_h = stride_w = stride
      else:
          stride_h, stride_w = stride
  
      if isinstance(dilation, int):
          dilation_h = dilation_w = dilation
      else:
          dilation_h, dilation_w = dilation
  
      in_dtype = Input.dtype
  
      if in_dtype in ["int4", "uint4"]:
          wmma_n = wmma_m = 8
          wmma_k = 32
      else:
          wmma_m = 8
          wmma_n = 32
          wmma_k = 16
  
      pre_computed = len(Filter.shape) == 6
      in_height, in_width, batch, in_channels = get_const_tuple(
          Input.shape
      )
      if pre_computed:
          (
              kernel_h,
              kernel_w,
              oc_chunk,
              _,
              oc_block_factor,
              _,
          ) = get_const_tuple(Filter.shape)
          num_filter = oc_block_factor * oc_chunk
      else:
          kernel_h, kernel_w, num_filter, _ = get_const_tuple(
              Filter.shape
          )
  
      if in_dtype in ["int4", "uint4"]:
          assert (
              batch % 8 == 0
              and in_channels % 32 == 0
              and num_filter % 8 == 0
          )
      else:
          assert (
              batch % 8 == 0
              and in_channels % 16 == 0
              and num_filter % 32 == 0
          ), (
              "The shape of (batch, in_channels, num_filter) "
              "must be multiple of (8, 16, 32) for int8, "
              "and (8, 32, 8) for int4"
          )
  
      # compute the output shape
      dilated_kernel_h = (kernel_h - 1) * dilation_h + 1
      dilated_kernel_w = (kernel_w - 1) * dilation_w + 1
  
      pad_top, pad_left, pad_down, pad_right = get_pad_tuple(
          padding, (dilated_kernel_h, dilated_kernel_w)
      )
  
      out_channels = num_filter
      out_height = simplify(
          (in_height - dilated_kernel_h + pad_top + pad_down)
          // stride_h
          + 1
      )
      out_width = simplify(
          (in_width - dilated_kernel_w + pad_left + pad_right)
          // stride_w
          + 1
      )
  
      cfg.add_flop(
          2
          * batch
          * out_height
          * out_width
          * out_channels
          * in_channels
          * kernel_h
          * kernel_w
      )
  
      # Input feature map: (H, W, N, IC, n, ic)
      data_shape = (
          in_height,
          in_width,
          batch // wmma_m,
          in_channels // wmma_k,
          wmma_m,
          wmma_k,
      )
  
      # Kernel: (H, W, OC, IC, oc, ic)
      kernel_shape = (
          kernel_h,
          kernel_w,
          out_channels // wmma_n,
          in_channels // wmma_k,
          wmma_n,
          wmma_k,
      )
  
      # Reduction axes
      kh = te.reduce_axis((0, kernel_h), name="kh")
      kw = te.reduce_axis((0, kernel_w), name="kw")
      ic = te.reduce_axis((0, in_channels // wmma_k), name="ic")
      ii = te.reduce_axis((0, wmma_k), name="ii")
  
      if pre_computed:
          packed_kernel = Filter
      else:
          packed_kernel = te.compute(
              kernel_shape,
              lambda kh, kw, o, i, oo, ii: Filter[
                  kh, kw, o * wmma_n + oo, i * wmma_k + ii
              ],
              name="packed_kernel",
          )
  
      packed_data = te.compute(
          data_shape,
          lambda h, w, n, i, nn, ii: Input[
              h, w, n * wmma_m + nn, i * wmma_k + ii
          ],
      )
  
      pad_before = [pad_top, pad_left, 0, 0, 0, 0]
      pad_after = [pad_down, pad_right, 0, 0, 0, 0]
      pad_data = pad(
          packed_data, pad_before, pad_after, name="pad_data"
      )
  
      Conv = te.compute(
          (
              out_height,
              out_width,
              batch // wmma_m,
              out_channels // wmma_n,
              wmma_m,
              wmma_n,
          ),
          lambda h, w, n, o, nn, oo: te.sum(
              (
                  pad_data[
                      h * stride_h + kh,
                      w * stride_w + kw,
                      n,
                      ic,
                      nn,
                      ii,
                  ].astype("int32")
                  * packed_kernel[kh, kw, o, ic, oo, ii].astype("int32")
              ),
              axis=[ic, kh, kw, ii],
          ),
          name="Conv",
          tag="conv2d_HWNCnc_tensorcore",
      )
      return Conv
  
file_path: tvm/python/tvm/topi/cuda/conv2d_hwnc_tensorcore.py
imports: import tvm
  from tvm import te
  from tvm import autotvm
  from tvm.target import Target
  from tvm.topi.cuda.injective import schedule_injective_from_existing
  from ..utils import get_const_tuple, traverse_inline, simplify, tag
  from ..nn.pad import pad
  from ..nn.utils import get_pad_tuple
  from .tensor_intrin import intrin_wmma_load_matrix_A
  from .tensor_intrin import intrin_wmma_load_matrix_W
  from .tensor_intrin import intrin_wmma_store_matrix
  from .tensor_intrin import intrin_wmma_gemm
label_desc: Use image node_id 3 for calling a global function with example usage: hwnc_tensorcore_cuda(cfg, Input, Filter, stride, padding, dilation, out_dtype) and returns: Conv