[TIR][Transform] Add LiftThreadBinding Pass

[junrushao]

This PR adds a pass LiftThreadBinding to TIR.

Previously, during GPU cross-thread reduction, a temporary local buffer will be created in the RF buffer, as a concrete example:

rf_local = T.alloc_buffer(..., scope="local")

// Step 1. Data parallel RF block
for tx in T.thread_binding(..., thread="threadIdx.x")
  rf_local[tx, ...] =

// Step 2. Cross-thread reduction to accumuate rf_local
for ...:
  for tx' in T.thread_binding(..., thread="threadIdx.x"):
    ... += rf_local[tx', ...]

In this case, the buffer rf_local will only be accessed by a single point tx or tx', but during the pass CompactBuffeRegion, the two variables as thread bindings are treated as two separate variables, i.e. the information that tx and tx' are always identical to each other is discarded, which means the accessed region on rf_local are estimated as Union({tx}, {tx'}) as opposed to {tx}, leading over allocation of local registers.

This pass is introduced to address this issue by lifting thread bindings to their LCAs.

[tvm-bot]

Thanks for contributing to TVM! Please refer to the contributing guidelines https://tvm.apache.org/docs/contribute/ for useful information and tips. Please request code reviews from Reviewers by @-ing them in a comment.

• cc @Hzfengsy, @quic-sanirudh, @shingjan _{See #10317 for details}

_{Generated by tvm-bot}