[7_GPU_and_Specialized_Hardware]一个thread block里有多少个thread?

[hbsun2113]

在7_GPU_and_Specialized_Hardware发现对于不同模型设置的thread num是不同的。我的疑问：为了提供并行度，不是应该尽可能把一个thread block(e.g. stream multi-processors)里所有的thread都用上吗? 为什么thread num会不同？

对于MyModuleWindowSum设定的thread num是128：

对于blocking设定的thread num是8*8=64：

如果使用meta_schedule进行调优，对于这两个模型使用的thread num也是不同的：
MyModuleWindowSum:

from tvm import meta_schedule as ms

sch_tuned_MyModuleWindowSum = ms.tune_tir(
    mod=MyModuleWindowSum,
    target="nvidia/tesla-p100",
    config=ms.TuneConfig(
      max_trials_global=1000,
      num_trials_per_iter=1000,
    ),
    work_dir="./tune_tmp_MyModuleWindowSum",
    task_name="main"
)
sch_tuned_MyModuleWindowSum.mod.show()
print(sch_tuned_MyModuleWindowSum.trace)

MyModuleMatmul:

from tvm import meta_schedule as ms

sch_tuned = ms.tune_tir(
    mod=MyModuleMatmul,
    target="nvidia/tesla-p100",
    config=ms.TuneConfig(
      max_trials_global=64,
      num_trials_per_iter=64,
    ),
    work_dir="./tune_tmp",
    task_name="main"
)

sch_tuned.mod.show()

[Hzfengsy]

对于GPU而言，一个warp的线程数是32，参考 (https://docs.nvidia.com/cuda/cuda-c-best-practices-guide/index.html#differences-between-host-and-device) ，超过32个thread会被串行执行

[hbsun2113]

所以32个thread应该是最好的？ 那为什么自动调优给出的结果是64呢，是考虑了其他因素吗？

[Hzfengsy]

thread数量保持32的倍数都是可以接受的。如果是64近似于

for i in range(2):
    for j in T.thread_binding(32, thread="threadIdx.x")

当然硬件会有更加复杂的scheduler，所以可能有时候32比较好，也有可能64会更好。
比较明确的是，如果小于32，肯定没有充分利用计算资源；如果不是32的整数倍（如48），那么最后一次线性执行（48 % 32 = 16线程）也没有充分利用计算资源

[yzh119]

一个thread block内可以有多个warp，gpu的warp scheduler可以调度这些warp的执行（可以参考https://www.cs.ucr.edu/~nael/217-f19/lectures/WarpScheduling.pptx）。
使用多个warp的好处：

• 如果某一个warp发射了一个在全局内存里读取数据的指令（可能需要比较长的时钟周期）停顿住了，warp scheduler可以转向执行下一个无需等待可以立即执行的warp，以此减小等待的时间，覆盖掉部分数据读写的开销。