code: def test_cpu_fusion():
      # pylint: disable=all
      @T.prim_func
      def func(a: T.handle, b: T.handle, c: T.handle) -> None:
          A = T.match_buffer(a, [64, 32], dtype="float32")
          B = T.match_buffer(b, [64, 32], dtype="float32")
          C = T.match_buffer(c, [64, 32], dtype="float32")
          for i, j in T.grid(64, 32):  # type: ignore
              with T.block():
                  T.reads([A[i, j], B[i, j]])  # type: ignore
                  T.writes([B[i, j], C[i, j]])  # type: ignore
                  with T.block("B"):
                      T.reads([A[i, j]])  # type: ignore
                      T.writes([B[i, j]])  # type: ignore
                      B[i, j] = A[i, j]  # type: ignore
                  with T.block("C"):
                      T.reads([B[i, j]])  # type: ignore
                      T.writes([C[i, j]])  # type: ignore
                      C[i, j] = B[i, j]  # type: ignore
  
      # pylint: enable=all
  
      def _create_schedule():
          return tir.Schedule(func, debug_mask="all")
  
      extractor = ms.feature_extractor.PerStoreFeature()
      (feature,) = extractor.extract_from(
          _make_context(tvm.target.Target("llvm")),
          candidates=[_make_candidate(_create_schedule)],
      )
      feature = feature.numpy()
      assert feature.shape == (2, N_FEATURES)
      ## Features for BufferStore(B)
      f = feature[0]
      # Group 1.1: arith
      assert_allclose(
          actual=f[0:16],
          # fmt: off
          desired=[0.0] * 16,
          # fmt: on
          rtol=1e-5,
          atol=1e-5,
      )
      # Group 1.2: vectorize
      assert_allclose(
          actual=f[16:27],
          desired=[
              0.0,
              0.0,
              0.0,
              1.0,
              0.0,
              0.0,
              0.0,
              0.0,
              0.0,
              0.0,
              0.0,
          ],
          rtol=1e-5,
          atol=1e-5,
      )
      # Group 1.3: unroll
      assert_allclose(
          actual=f[27:38],
          desired=[
              0.0,
              0.0,
              0.0,
              1.0,
              0.0,
              0.0,
              0.0,
              0.0,
              0.0,
              0.0,
              0.0,
          ],
          rtol=1e-5,
          atol=1e-5,
      )
      # Group 1.4: parallel
      assert_allclose(
          actual=f[38:49],
          desired=[
              0.0,
              0.0,
              0.0,
              1.0,
              0.0,
              0.0,
              0.0,
              0.0,
              0.0,
              0.0,
              0.0,
          ],
          rtol=1e-5,
          atol=1e-5,
      )
      # Group 1.5: is_gpu, blockIdx.x/y/z, threadIdx.x/y/z, vthread
      assert_allclose(
          actual=f[49:57],
          desired=[0.0, 1.0, 1.0, 1.0, 1.0, 1.0, 1.0, 1.0],
          rtol=1e-5,
          atol=1e-5,
      )
      # Group 2.1: Buffer A
      assert_allclose(
          actual=f[57:75],
          desired=[
              1.0,
              0.0,
              0.0,
              13.000176429748535,
              13.000176429748535,
              7.011227130889893,
              7.011227130889893,
              0.0,
              0.0,
              1.0,
              0.0,
              0.0,
              0.0,
              14.00008773803711,
              14.00008773803711,
              8.005624771118164,
              8.005624771118164,
              1.0,
          ],
          rtol=1e-5,
          atol=1e-5,
      )
      # Group 2.2: Buffer B
      assert_allclose(
          actual=f[75:93],
          desired=[
              0.0,
              1.0,
              0.0,
              13.000176429748535,
              13.000176429748535,
              7.011227130889893,
              7.011227130889893,
              0.0,
              0.0,
              1.0,
              0.0,
              0.0,
              0.0,
              14.00008773803711,
              14.00008773803711,
              8.005624771118164,
              8.005624771118164,
              1.0,
          ],
          rtol=1e-5,
          atol=1e-5,
      )
      # Group 2.3: Dummy padding
      assert_allclose(
          actual=f[93:111],
          desired=[0.0] * 18,
          rtol=1e-5,
          atol=1e-5,
      )
      # Group 2.4: Dummy padding
      assert_allclose(
          actual=f[111:129],
          desired=[0.0] * 18,
          rtol=1e-5,
          atol=1e-5,
      )
      # Group 2.5: Dummy padding
      assert_allclose(
          actual=f[129:147],
          desired=[0.0] * 18,
          rtol=1e-5,
          atol=1e-5,
      )
      # Group 3: Arithmetic intensity
      assert_allclose(
          actual=f[147:157],
          desired=[0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0],
          rtol=1e-5,
          atol=1e-5,
      )
      # Group 4 & 5
      assert_allclose(
          actual=f[157:164],
          desired=[
              13.000176,
              11.000703811645508,
              1.0,
              11.000703811645508,
              11.000703811645508,
              1.5849624872207642,
              0.0,
          ],
          rtol=1e-5,
          atol=1e-5,
      )
      ## Features for BufferStore(C)
      f = feature[1]
      # Group 1.1: arith
      assert_allclose(
          actual=f[0:16],
          # fmt: off
          desired=[0.0] * 16,
          # fmt: on
          rtol=1e-5,
          atol=1e-5,
      )
      # Group 1.2: vectorize
      assert_allclose(
          actual=f[16:27],
          desired=[
              0.0,
              0.0,
              0.0,
              1.0,
              0.0,
              0.0,
              0.0,
              0.0,
              0.0,
              0.0,
              0.0,
          ],
          rtol=1e-5,
          atol=1e-5,
      )
      # Group 1.3: unroll
      assert_allclose(
          actual=f[27:38],
          desired=[
              0.0,
              0.0,
              0.0,
              1.0,
              0.0,
              0.0,
              0.0,
              0.0,
              0.0,
              0.0,
              0.0,
          ],
          rtol=1e-5,
          atol=1e-5,
      )
      # Group 1.4: parallel
      assert_allclose(
          actual=f[38:49],
          desired=[
              0.0,
              0.0,
              0.0,
              1.0,
              0.0,
              0.0,
              0.0,
              0.0,
              0.0,
              0.0,
              0.0,
          ],
          rtol=1e-5,
          atol=1e-5,
      )
      # Group 1.5: is_gpu, blockIdx.x/y/z, threadIdx.x/y/z, vthread
      assert_allclose(
          actual=f[49:57],
          desired=[0.0, 1.0, 1.0, 1.0, 1.0, 1.0, 1.0, 1.0],
          rtol=1e-5,
          atol=1e-5,
      )
      # Group 2.1: Buffer B
      assert_allclose(
          actual=f[57:75],
          desired=[
              1.0,
              0.0,
              0.0,
              13.000176429748535,
              13.000176429748535,
              7.011227130889893,
              7.011227130889893,
              0.0,
              1.0,
              0.0,
              1.0,
              4.087462902069092,
              1.0,
              13.000176429748535,
              13.000176429748535,
              7.011227130889893,
              7.011227130889893,
              1.0,
          ],
          rtol=1e-5,
          atol=1e-5,
      )
      # Group 2.2: Buffer C
      assert_allclose(
          actual=f[75:93],
          desired=[
              0.0,
              1.0,
              0.0,
              13.000176429748535,
              13.000176429748535,
              7.011227130889893,
              7.011227130889893,
              0.0,
              0.0,
              1.0,
              0.0,
              0.0,
              0.0,
              14.00008773803711,
              14.00008773803711,
              8.005624771118164,
              8.005624771118164,
              1.0,
          ],
          rtol=1e-5,
          atol=1e-5,
      )
      # Group 2.3: Dummy padding
      assert_allclose(
          actual=f[93:111],
          desired=[0.0] * 18,
          rtol=1e-5,
          atol=1e-5,
      )
      # Group 2.4: Dummy padding
      assert_allclose(
          actual=f[111:129],
          desired=[0.0] * 18,
          rtol=1e-5,
          atol=1e-5,
      )
      # Group 2.5: Dummy padding
      assert_allclose(
          actual=f[129:147],
          desired=[0.0] * 18,
          rtol=1e-5,
          atol=1e-5,
      )
      # Group 3: Arithmetic intensity
      assert_allclose(
          actual=f[147:157],
          desired=[0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0],
          rtol=1e-5,
          atol=1e-5,
      )
      # Group 4 & 5
      assert_allclose(
          actual=f[157:164],
          desired=[
              13.000176429748535,
              11.000703811645508,
              1.0,
              11.000703811645508,
              11.000703811645508,
              1.5849624872207642,
              0.0,
          ],
          rtol=1e-5,
          atol=1e-5,
      )
  
imports: import sys
  from typing import Callable, List
  import pytest
  import tvm
  import tvm.testing
  from numpy.testing import assert_allclose
  from tvm import meta_schedule
  from tvm import te, tir
  from tvm.script import tir
label_desc: Use image node_id 8 for calling a global function with example usage: test_cpu_fusion() and returns: tir