[Tensor Core 编程教程]对Warp 指令的疑问

[hbsun2113]

我没有任何GPU编程基础，所以问题可能过于基础，请见谅。
在Tensor Core 编程教程中有如下优化：

我的问题：

1. 对于“使 threadIdx.x extent=32”：
   • threadIdx.x extent是什么？
   • 我在上图code中没有找到与32有关的绑定，请问是如何实现的？
2. “threadIdx.y”和“threadIdx.x”的区别是什么，为什么“warp-level operators”对x有限制，而对y无限制？
3. 对于“确保一个 warp 中的所有线程都可以同时调用 Tensor Core”，请问是如何确保的？

[Hzfengsy]

1. threadIdx.x extent是32，通过GPU的其他部分schedule生成的（比如coop fetch），因为单个kernel内threads和blocks数量相同。
2. warp-level operators 要求是一个warp中的线程同时工作，并不会直接对threadIdx.x有要求（原文提到是“是决此问题的最简单方法之一”，并非唯一方法）
3. 这个是由cuda编程模型保证的，TVM只负责生成cuda代码，具体后续编译执行由NVIDIA相关软硬件保证