code: def gen_test_bin(dir, M, N, K, exe="test", algo_id=0):
      test_src = f"""
  int main(int argc, char **argv) {{
    int M = {M}, N = {N}, K = {K};
  
    // initialize CUDA handles
    CUdevice dev;
    CUcontext ctx;
    CUstream stream;
    CUdeviceptr A, B, C;
    CUresult err = 0;
    cuInit(0);
    cuDeviceGet(&dev, 0);
    cuCtxCreate(&ctx, 0, dev);
    cuMemAlloc(&A, M * K * 2);
    cuMemAlloc(&B, K * N * 2);
    cuMemAlloc(&C, M * N * 4);
    cuStreamCreate(&stream, 0);
    load_matmul_fp16();
  
    // initialize input data
    int16_t hA[M*K];
    int16_t hB[K*N];
    memset(hA, 0, M*K*2);
    memset(hB, 0, K*N*2);
    read_csv_to_buffer(argv[1], hA, M*K);
    read_csv_to_buffer(argv[2], hB, K*N);
    cuMemcpyHtoD(A, hA, M*K*2);
    cuMemcpyHtoD(B, hB, K*N*2);
  
    // launch kernel
    cuStreamSynchronize(stream);
    CUresult ret;
    int algo_id = {algo_id};
    if (algo_id == 0) {{
      ret = matmul_fp16_default(stream, C, A, B, M, N, K, N, 1, K, 1, N, 1);
    }} else {{
      ret = matmul_fp16(stream, C, A, B, M, N, K, N, 1, K, 1, N, 1, {algo_id});
    }}
    if (ret != 0) fprintf(stderr, "kernel launch failed\\n");
    assert(ret == 0);
  
    cuStreamSynchronize(stream);
  
    // read data
    int32_t hC[M*N];
    memset(hC, 0, M*N*4);
    cuMemcpyDtoH(hC, C, M*N*4);
    write_buffer_to_csv(argv[3], hC, M*N);
  
    // free cuda handles
    unload_matmul_fp16();
    cuMemFree(A);
    cuMemFree(B);
    cuMemFree(C);
    cuCtxDestroy(ctx);
  }}
  """
      src = test_utils_src + test_src
      with open(os.path.join(dir, "test.c"), "w") as file:
          file.write(src)
      subprocess.run(
          ["gcc"]
          + [
              "test.c",
              "-I",
              include_dir[0],
              "-L",
              library_dir[0],
              "-l",
              "cuda",
              "-L",
              dir,
              "-l",
              "kernel",
              "-o",
              exe,
          ],
          check=True,
          cwd=dir,
      )
  
imports: import glob
  import os
  import subprocess
  import sys
  import tempfile
  import numpy
  import triton
  from triton.backends.nvidia.driver import include_dir, library_dir
label_desc: Use image node_id 2 for calling a global function with example usage: gen_test_bin(dir, M, N, K, exe, algo_id) without return types