[MetaSchedule] Fix Dynamic Loop from AutoBinding

include/tvm/meta_schedule/schedule/cuda/thread_bind.h

@@ -45,6 +45,7 @@ std::function<tir::ExprRV(int64_t)> MakeFactorSampler(tir::Schedule sch,
  * \param max_threadblocks The maximum number of threadblocks allowed.
  * \param max_threads_per_block The maximum number of threads allowed.
  * \param get_factor A function that returns the tiling factor.
+ * \return The binded loops in the order of blockIdx.x, threadIdx.x, and the rest.
  */
 Array<tir::LoopRV> BindSpatialLoop(tir::Schedule sch, tir::LoopRV loop,  //
                                    int64_t max_threadblocks, int64_t max_threads_per_block,

src/meta_schedule/schedule/cuda/thread_bind.cc

@@ -80,7 +80,7 @@ Array<LoopRV> BindSpatialLoop(Schedule sch, LoopRV loop, int64_t max_threadblock
     sch->Reorder({splits[1], splits[2], splits[0]});
     sch->Bind(splits[1], "blockIdx.x");
     sch->Bind(splits[2], "threadIdx.x");
-    return {splits[1], splits[2]};
+    return {splits[1], splits[2], splits[0]};
   }
 }
 

src/meta_schedule/schedule/cuda/winograd.cc

@@ -117,7 +117,8 @@ TVM_REGISTER_GLOBAL("meta_schedule.cuda.conv2d_nchw_winograd_data_pack")
         sch->Unroll(loops[4]);
         sch->Unroll(loops[5]);
         outer = BindSpatialLoop(sch, sch->Fuse({loops[2], loops[3]}), max_threadblocks,
-                                max_threads_per_block)[1];
+                                max_threads_per_block, /*get_factor=*/nullptr)
+                    .back();
       }
       {
         BlockRV data_pack_local = sch->CacheWrite(data_pack, 0, "local");