code: def dense_tensorcore_cuda(data, weight, bias=None, out_dtype=None):
      """Dense tensorcore operator on CUDA"""
      assert (
          len(data.shape) == 2 and len(weight.shape) == 2
      ), "only support 2-dim dense"
      if bias is not None:
          assert len(bias.shape) == 1
      if out_dtype is None:
          out_dtype = data.dtype
      batch, in_dim = get_const_tuple(data.shape)
      out_dim, _ = get_const_tuple(weight.shape)
  
      assert data.dtype == weight.dtype
      assert data.dtype in ["float16", "int8", "uint8", "int4", "uint4"]
      if data.dtype in ["float16", "int8", "uint8"]:
          assert (
              (
                  batch % 8 == 0
                  and in_dim % 16 == 0
                  and out_dim % 32 == 0
              )
              or (
                  batch % 16 == 0
                  and in_dim % 16 == 0
                  and out_dim % 16 == 0
              )
              or (
                  batch % 32 == 0
                  and in_dim % 16 == 0
                  and out_dim % 8 == 0
              )
          ), (
              "The shape of (batch, in_dim, out_dim) "
              "must be multiple of (16, 16, 16) or (32, 16, 8) or (8, 16, 32) for now"
          )
      else:
          assert (
              batch % 8 == 0 and in_dim % 32 == 0 and out_dim % 8 == 0
          ), "The shape of (batch, in_dim, out_dim) must be multiple of (8, 32, 8)"
  
      k = te.reduce_axis((0, in_dim), name="k")
      matmul = te.compute(
          (batch, out_dim),
          lambda i, j: te.sum(
              data[i, k].astype(out_dtype)
              * weight[j, k].astype(out_dtype),
              axis=k,
          ),
          name="T_dense",
          tag="dense_tensorcore",
      )
      if bias is not None:
          matmul = te.compute(
              (batch, out_dim),
              lambda i, j: matmul[i, j] + bias[j].astype(out_dtype),
              tag=tag.BROADCAST,
          )
      return matmul
  
file_path: tvm/python/tvm/topi/cuda/dense_tensorcore.py
imports: from __future__ import absolute_import
  import tvm
  from tvm import te, autotvm
  from ..None import tag
  from ..utils import traverse_inline, get_const_tuple
  from .tensor_intrin import intrin_wmma_load_matrix_A, intrin_wmma_load_matrix_W, intrin_wmma_store_matrix, intrin_wmma_gemm
label_desc: Use image node_id 3 for calling a global function with example usage: dense_tensorcore_cuda(data, weight, bias, out_dtype) and returns: matmul