code: def intrin_wmma_gemm(shape):
      n, m, l = shape
      A = te.placeholder((n, l), name="A", dtype="float16")
      B = te.placeholder((l, m), name="B", dtype="float16")
      k = te.reduce_axis((0, l), name="k")
      C = te.compute(
          (n, m),
          lambda ii, jj: te.sum(
              A[ii, k].astype("float") * B[k, jj].astype("float"),
              axis=k,
          ),
          name="C",
      )
      BA = tvm.tir.decl_buffer(
          A.shape,
          A.dtype,
          name="BA",
          scope="wmma.matrix_a",
          data_alignment=32,
          offset_factor=n * l,
      )
      BB = tvm.tir.decl_buffer(
          B.shape,
          B.dtype,
          name="BB",
          scope="wmma.matrix_b",
          data_alignment=32,
          offset_factor=l * m,
      )
      BC = tvm.tir.decl_buffer(
          C.shape,
          C.dtype,
          name="BC",
          scope="wmma.accumulator",
          data_alignment=32,
          offset_factor=n * m,
      )
  
      def intrin_func(ins, outs):
          BA, BB = ins
          (BC,) = outs
  
          def init():
              ib = tvm.tir.ir_builder.create()
              ib.emit(
                  tvm.tir.call_intrin(
                      "handle",
                      "tir.tvm_fill_fragment",
                      BC.data,
                      n,
                      m,
                      l,
                      BC.elem_offset // (n * m),
                      0.0,
                  )
              )
              return ib.get()
  
          def update():
              ib = tvm.tir.ir_builder.create()
              ib.emit(
                  tvm.tir.call_intrin(
                      "handle",
                      "tir.tvm_mma_sync",
                      BC.data,
                      BC.elem_offset // (n * m),
                      BA.data,
                      BA.elem_offset // (n * l),
                      BB.data,
                      BB.elem_offset // (l * m),
                      BC.data,
                      BC.elem_offset // (n * m),
                  )
              )
              return ib.get()
  
          return update(), init(), update()
  
      return te.decl_tensor_intrin(
          C.op, intrin_func, binds={A: BA, B: BB, C: BC}
      )
  
file_path: tvm/tests/python/te/test_te_schedule_tensor_core.py
imports: import tvm
  from tvm import te
  import numpy
  from tvm.topi.testing import conv2d_nhwc_python
  import tvm.testing
label_desc: Use image node_id 2 for calling a global function with example usage: intrin_wmma_gemm(shape) and returns: te, update,  init,  update, ib, ib