[TIR] Enable HoistIfThenElse in the default lowering procedure

[roastduck]

Enabling the HoistIfThenElse pass (#3865) in the default lowering procedure. HoistIfThenElse can be very helpful for sparse applications, since LoopPartition cannot eliminate their if statements with dynamic (unknown at compile time) conditions.

Changes:

• Move HoistIfThenElse from src/tir/pass to src/tir/transform.
• Added it into lower.
• Fixed a breaking test.

@tqchen Requesting for a review.

[roastduck]

Well, a test for TensorCore fails.

HoistIfThenElse transforms

for (n.inner, 0, 2) {
for (o.inner, 0, 2) {
if ((((threadIdx.y*2) + n.inner) < 2)) {
if ((((threadIdx.z*2) + o.inner) < 4)) {
if ((threadIdx.y < 1)) {
if ((threadIdx.z < 2)) {
tvm_store_matrix_sync(Conv.wmma.accumulator, 16, 16, 16, ((n.inner*2) + o.inner), tvm_access_ptr(type_annotation(), Conv, (((((threadIdx.y*401408) + (n.inner*200704)) + (blockIdx.z*1024)) + (threadIdx.z*512)) + (o.inner*256)), 256, 2), 16, "row_major")
}
}
}
}
}
}

into

if ((((threadIdx.y*2) + n.inner) < 2)) {
if ((threadIdx.y < 1)) {
if ((threadIdx.z < 2)) {
for (n.inner, 0, 2) {
for (o.inner, 0, 2) {
if ((((threadIdx.z*2) + o.inner) < 4)) {
tvm_store_matrix_sync(Conv.wmma.accumulator, 16, 16, 16, ((n.inner*2) + o.inner), tvm_access_ptr(type_annotation(), Conv, (((((threadIdx.y*401408) + (n.inner*200704)) + (blockIdx.z*1024)) + (threadIdx.z*512)) + (o.inner*256)), 256, 2), 16, "row_major")
}
}
}
}
}
}

where the if containing n.inner is wrongly hoisted, seeming like a bug of HoistIfThenElse. I will try to dig it out.

[roastduck]

Fixed and added an extra test.

[roastduck]

Renamed Vars in hoisted else_cases, to be different from those in then_cases, so that SSA verification won't panic.

[roastduck]

Fixed a bug where if nodes containing thread indices could be hoisted over the definition of the indices. This would happen when Attr node for thread_extent is scheduled into the body of a For node, using a compute_at command.

[roastduck]

Since there are too many bugs in HoistIfThenElse, let's track them in another issue, and enable the pass after they are all fixed. Closing for now.