code: def test_unrolled_vectorization():
      dtype = "float32"
      target = "cuda"
  
      # Compute declaration
      N = 128
      A = te.placeholder((N, N), name="A")
      B = te.placeholder((N, N), name="B")
      k = te.reduce_axis((0, N), name="k")
      C = te.compute(
          (N, N),
          lambda i, j: te.sum(A[i][k] * B[k][j], axis=[k]),
          name="C",
      )
  
      # Schedule
      s = te.create_schedule([C.op])
      CC = s.cache_write(C, "local")
      i, j = s[C].op.axis
      bx, tx, ii, ji = s[C].tile(i, j, 1, 2)
      s[C].bind(bx, te.thread_axis("blockIdx.x"))
      s[C].bind(tx, te.thread_axis("threadIdx.x"))
      s[C].vectorize(ji)
      s[CC].compute_at(s[C], tx)
      i, j = s[CC].op.axis
      k = s[CC].op.reduce_axis[0]
      ko, ki = s[CC].split(k, 2)
      s[CC].unroll(ki)
      s[CC].vectorize(j)
  
      # Check correctness
      dev = tvm.device(target)
      a_tvm = tvm.nd.array(np.ones((N, N)).astype(dtype), device=dev)
      b_tvm = tvm.nd.array(np.ones((N, N)).astype(dtype), device=dev)
      c_tvm = tvm.nd.empty((N, N), device=dev)
      func_tvm = tvm.build(s, [A, B, C], target=target)
      func_tvm(a_tvm, b_tvm, c_tvm)
      c_np = c_tvm.numpy()
      tvm.testing.assert_allclose(c_np, N * np.ones((N, N)))
  
imports: import os
  import re
  import tvm
  from tvm import te
  import numpy
  from tvm import topi
  from tvm.contrib.nvcc import have_fp16, have_int8, have_bf16
  from tvm.contrib import utils
  from tvm.script import tir
  import tvm.testing
  import pytest
label_desc: Use image node_id 27 for calling a global function with example usage: test_unrolled_vectorization() without return types