code: def get_mma_load_intrin(
      m_dim: int,
      n_dim: int,
      k_dim: int,
      dtype: str,
      shared_scope: str,
      is_b: bool,
      is_col_major: bool,
  ) -> Tuple[PrimFunc, PrimFunc]:
      """Generator of mma ldmatrix intrins"""
      mma_fragment_scope = f"m16n8k8.matrix{'B' if is_b else 'A'}"
      frag_m, frag_n = (k_dim, n_dim) if is_b else (m_dim, k_dim)
      trans = (not is_col_major) if is_b else is_col_major
      if is_col_major:
          frag_m, frag_n = frag_n, frag_m
      get_index = get_index_B if is_b else get_index_A
      get_tx_index = (
          (lambda tx, s0: (tx % 8) * s0 + (tx // 8) * 8)
          if trans
          else (lambda tx, s0: tx * s0)
      )
  
      @T.prim_func
      def mma_load_desc(a: T.handle, c: T.handle) -> None:
          src = T.match_buffer(
              a,
              (frag_m, frag_n),
              dtype,
              align=64,
              offset_factor=1,
              scope=shared_scope,
          )
          dst = T.match_buffer(
              c,
              (frag_m, frag_n),
              dtype,
              align=64,
              offset_factor=1,
              scope=mma_fragment_scope,
          )
  
          with T.block("root"):
              T.reads(src[0:frag_m, 0:frag_n])
              T.writes(dst[0:frag_m, 0:frag_n])
              for i, j in T.grid(frag_m, frag_n):
                  with T.block("root"):
                      vi, vj = T.axis.remap("SS", [i, j])
                      dst[vi, vj] = src[vi, vj]
  
      @T.prim_func
      def mma_load_impl(a: T.handle, c: T.handle) -> None:
          s0 = T.int32()
          s1 = T.int32()
          src = T.match_buffer(
              a,
              (frag_m, frag_n),
              dtype,
              align=64,
              offset_factor=1,
              scope=shared_scope,
              strides=[s0, s1],
          )
          d0 = T.int32()
          d1 = T.int32()
          dst = T.match_buffer(
              c,
              (frag_m, frag_n),
              dtype,
              align=64,
              offset_factor=1,
              scope=mma_fragment_scope,
              strides=[d0, d1],
          )
  
          with T.block("root"):
              T.reads(src[0:frag_m, 0:frag_n])
              T.writes(dst[0:frag_m, 0:frag_n])
  
              for tx in T.thread_binding(0, WARP_SIZE, "threadIdx.x"):
                  T.evaluate(
                      T.ptx_ldmatrix(
                          trans,
                          4,  # Always load 4 matrices
                          ".b16",
                          dst.data,
                          get_index(dst.elem_offset, d0),
                          src.access_ptr("r"),
                          get_tx_index(tx, s0),
                          dtype=dtype,
                      )
                  )
  
      return mma_load_desc, mma_load_impl
  
imports: from typing import Dict, Optional, Tuple, Literal
  from tvm._ffi import register_func
  from tvm.runtime import convert
  from tvm.script import tir
  from tvm.tir.function import PrimFunc
  from tvm.tir import Cast, IntImm, TensorIntrin
label_desc: Use image node_id 21 for calling a global function with example usage: get_mma_load_intrin(m_dim, n_dim, k_dim, dtype, shared_scope, is_b, is_col_major) and returns: mma_load_desc,  mma_load_impl