code: def test_gemm_mma_m8n8k4_row_col_fp64pf64fp64():
      sch = tvm.tir.Schedule(gemm_mma_m8n8k4_row_col_fp64pf64fp64)
      cuda_mod = tvm.build(sch.mod, target="cuda")
  
      A_np = np.random.uniform(-1, 1, [8, 4]).astype("float64")
      B_np = np.random.uniform(-1, 1, [8, 4]).astype("float64")
      C_np = np.zeros([8, 8]).astype("float64")
  
      ctx = tvm.cuda()
      A_tvm = tvm.nd.array(A_np, ctx)
      B_tvm = tvm.nd.array(B_np, ctx)
      C_tvm = tvm.nd.array(C_np, ctx)
  
      cuda_mod(A_tvm, B_tvm, C_tvm)
  
      golden = np.matmul(
          A_np.astype("float64"), B_np.astype("float64").T
      )
  
      C_numpy = C_tvm.numpy()
  
      tvm.testing.assert_allclose(golden, C_numpy, atol=1e-3, rtol=1e-3)
  
imports: import sys
  import pytest
  import tvm
  from tvm.script import tir
  import numpy
  import tvm.testing
label_desc: Use image node_id 2 for calling a global function with example usage: test_gemm_mma_m8n8k4_row_col_fp64pf64fp64() without return types