code: def intrin_wmma_store_matrix(
      strides_dst, strides_from, shape, out_dtype, A_shape, C_shape
  ):
      """Intrin function for storing the results from wmma.accumulator to shared"""
      wmma_m, wmma_n, wmma_k = shape
      A = te.placeholder(A_shape, name="A", dtype=out_dtype)
      BA = tvm.tir.decl_buffer(
          A.shape,
          A.dtype,
          scope="wmma.accumulator",
          strides=strides_from,
          data_alignment=32,
          offset_factor=8,
      )
      C = te.compute(C_shape, lambda *i: A(*i), name="C")
      BC = tvm.tir.decl_buffer(
          C.shape,
          C.dtype,
          scope="shared",
          strides=strides_dst,
          data_alignment=32,
          offset_factor=8,
      )
  
      def intrin_func(ins, outs):
          ib = tvm.tir.ir_builder.create()
  
          BA = ins[0]
          BC = outs[0]
          row = wmma_m * wmma_n
          warp_index = (
              BA.elem_offset // row + BA.elem_offset % row // wmma_n
          )
          ib.emit(
              tvm.tir.call_intrin(
                  "handle",
                  "tir.tvm_store_matrix_sync",
                  BA.data,
                  wmma_m,
                  wmma_n,
                  wmma_k,
                  warp_index,
                  BC.access_ptr("w"),
                  strides_dst[0],
                  "row_major",
              )
          )
          return ib.get()
  
      return te.decl_tensor_intrin(
          C.op, intrin_func, binds={A: BA, C: BC}
      )
  
file_path: tvm/python/tvm/topi/cuda/tensor_intrin.py
imports: import tvm
  from tvm import te
  from ..utils import is_target
label_desc: Use image node_id 4 for calling a global function with example usage: intrin_wmma_store_matrix(strides_dst, strides_from, shape, out_dtype, A_shape, C_shape) and returns: te, ib