Refactor counters to scoped-counters

[PGZXB]

Following #4561, in C++ codebase of Taichi, many counters are used to identify a instance. But the counters are global now, which causes some problems. Using Taichi AST string generated by irpass::print as offline-cache-key (#4401) cannot identify a kernel completely because of global counters. We need scoped-counters.

An example in which I found problem

ti.init(arch=ti.cpu, print_preprocessed_ir=True)

@ti.kernel
def f1():
  for i in range(100):
    pass

@ti.kernel
def f2():
  for i in range(100):
    pass

1st: f1(); f2(), Ouput:

[I 03/19/22 10:47:49.763 17728] [kernel.cpp:lower@79] [f1_c56_0] Preprocessed IR:
kernel {
  $0 : for @tmp0 in range((cast_value<i32> 0), (cast_value<i32> 100)) block_dim=adaptive {
  }
}
[I 03/19/22 10:47:49.792 17728] [kernel.cpp:lower@79] [f2_c58_0] Preprocessed IR:
kernel {
  $0 : for @tmp2 in range((cast_value<i32> 0), (cast_value<i32> 100)) block_dim=adaptive {
  }
}

2nd: f2(); f1(), Output:

[I 03/19/22 10:49:09.890 17800] [kernel.cpp:lower@79] [f2_c58_0] Preprocessed IR:
kernel {
  $0 : for @tmp0 in range((cast_value<i32> 0), (cast_value<i32> 100)) block_dim=adaptive {
  }
}
[I 03/19/22 10:49:09.919 17800] [kernel.cpp:lower@79] [f1_c56_0] Preprocessed IR:
kernel {
  $0 : for @tmp2 in range((cast_value<i32> 0), (cast_value<i32> 100)) block_dim=adaptive {
  }
}

One kernel's ir/ast-string is related with calling order, because the id: Identifier which is member of IdExpression has auto incremental id: int (depending on static int variable id_counter). Global id_counter causes the problem.
I think we can remove this counters, and make them scoped.

Counters in C++ Side

(Simply using reg "static.*int.*counter" to search)

• Identifier::id_counter, global scope and kernel/func scope
• SNode::counter, program(global) scope
• Stmt::instance_id_counter, kernel/func scope
• TaskLaunchRecord::task_counter, maybe kernel/func scope