THU-DSP-LAB/ventus-gpgpu

CUDA中的三维网格、线程块如何映射到硬件？

kaitoukito opened this issue · 6 comments

kaitoukito commented

我在阅读源码时，有一个疑惑，想请教一下：

class host2CTA_data extends Bundle{
  val host_wg_id            = (UInt(WG_ID_WIDTH.W))
  val host_num_wf           = (UInt(WF_COUNT_WIDTH.W))
  val host_wf_size          = (UInt(WAVE_ITEM_WIDTH.W))
  val host_start_pc         = (UInt(MEM_ADDR_WIDTH.W))
  val host_vgpr_size_total  = (UInt((VGPR_ID_WIDTH + 1).W))
  val host_sgpr_size_total  = (UInt((SGPR_ID_WIDTH + 1).W))
  val host_lds_size_total   = (UInt((LDS_ID_WIDTH + 1).W))
  val host_gds_size_total   = (UInt((GDS_ID_WIDTH + 1).W))
  val host_vgpr_size_per_wf = (UInt((VGPR_ID_WIDTH + 1).W))
  val host_sgpr_size_per_wf = (UInt((SGPR_ID_WIDTH + 1).W))
  val host_gds_baseaddr = UInt(MEM_ADDR_WIDTH.W)
}

貌似host2CTA_data所包含的信息中，并没有一个CUDA中三维网格、线程块ID的概念，那么它们是如何映射到硬件上的呢？即：

  • blockIdx.x
  • blockIdx.y
  • blockIdx.z
  • threadIdx.x
  • threadIdx.y
  • threadIdx.z
yangkex commented

blockIdx和threadIdx在硬件上都是一维的，可以通过CUDA中的三维坐标乘以blockDim等来计算出其一维的Idx
至于具体的一维Idx在我们的代码中：

  • class host2CTA_data是host向GPGPU发送的任务，以block为单位，其blockIdx是host_wg_id
  • class CTAreqData是CTA_scheduler将任务按warp拆分后发送到SM上的，此时threadIdx的基值为dispatch2cu_wf_tag_dispatch，还需要在汇编中加上vid.v指令返回的值才能得到threadIdx
kaitoukito commented

blockIdx和threadIdx在硬件上都是一维的，可以通过CUDA中的三维坐标乘以blockDim等来计算出其一维的Idx 至于具体的一维Idx在我们的代码中：

  • class host2CTA_data是host向GPGPU发送的任务，以block为单位，其blockIdx是host_wg_id
  • class CTAreqData是CTA_scheduler将任务按warp拆分后发送到SM上的，此时threadIdx的基值为dispatch2cu_wf_tag_dispatch，还需要在汇编中加上vid.v指令返回的值才能得到threadIdx

谢谢。我再追问一下：举一个二维的例子，如下图所示，假设在一个block内，我逻辑上划分好了二维的thread。
image

但我实际的数据行与行之间存在padding，并不是完全连续的。那此时其实是需要同时用到 threadIdx.x、threadIdx.y。
即：address = base_address + threadIdx.y * stride + threadIdx.x
image

那么是否应该分离地把 threadIdx.x、threadIdx.y 告诉硬件、而不是把它们计算成一维再告诉硬件呢？
还是说有其他方法处理这种情况？

yangkex commented

若您在CUDA编程中做了padding、使用了threadIdx.x、threadIdx.y等来计算address，编译器会对应转换成一维形式的地址，硬件上也是按照此接口来的。
“承影”目前没有配套的编译器工具链，所以如果您想在其上测试自己的程序，只能在rvv汇编程序里手动完成该过程。address的计算可以参考我们的一些示例程序，通过csrrs+vid.v的方式取得。

kaitoukito commented

若您在CUDA编程中做了padding、使用了threadIdx.x、threadIdx.y等来计算address，编译器会对应转换成一维形式的地址，硬件上也是按照此接口来的。 “承影”目前没有配套的编译器工具链，所以如果您想在其上测试自己的程序，只能在rvv汇编程序里手动完成该过程。address的计算可以参考我们的一些示例程序，通过csrrs+vid.v的方式取得。

好的，谢谢。
想请教一下你们后续的技术路线，假如准备增加编译器的支持，那么高层次代码是写成vector风格（并增加自动向量化的支持以映射到底层rvv汇编）、还是CUDA风格呢？

yangkex commented

若您在CUDA编程中做了padding、使用了threadIdx.x、threadIdx.y等来计算address，编译器会对应转换成一维形式的地址，硬件上也是按照此接口来的。 “承影”目前没有配套的编译器工具链，所以如果您想在其上测试自己的程序，只能在rvv汇编程序里手动完成该过程。address的计算可以参考我们的一些示例程序，通过csrrs+vid.v的方式取得。

好的，谢谢。 想请教一下你们后续的技术路线，假如准备增加编译器的支持，那么高层次代码是写成vector风格（并增加自动向量化的支持以映射到底层rvv汇编）、还是CUDA风格呢？

我们今日进行了第一次开发者活动，目前考虑是优先支持OpenCL。

  • 👍1
yangkex commented

blockIdx和threadIdx在硬件上都是一维的，可以通过CUDA中的三维坐标乘以blockDim等来计算出其一维的Idx 至于具体的一维Idx在我们的代码中：

  • class host2CTA_data是host向GPGPU发送的任务，以block为单位，其blockIdx是host_wg_id
  • class CTAreqData是CTA_scheduler将任务按warp拆分后发送到SM上的，此时threadIdx的基值为dispatch2cu_wf_tag_dispatch，还需要在汇编中加上vid.v指令返回的值才能得到threadIdx

谢谢。我再追问一下：举一个二维的例子，如下图所示，假设在一个block内，我逻辑上划分好了二维的thread。 image

但我实际的数据行与行之间存在padding，并不是完全连续的。那此时其实是需要同时用到 threadIdx.x、threadIdx.y。 即：address = base_address + threadIdx.y * stride + threadIdx.x image

那么是否应该分离地把 threadIdx.x、threadIdx.y 告诉硬件、而不是把它们计算成一维再告诉硬件呢？ 还是说有其他方法处理这种情况？

我们之前的方案里是编译器计算好后直接传递threadIdx，期望能在此层次上进行一些简化，但按照PTX（和SASS）的做法是x y z均有各自的Special Register，OpenCL中也有get_global_id()获取各维度坐标和大小。
分离地把threadIdx.x、threadIdx.y告诉硬件，再由软件计算，应该是个对编译器实现更友好的方案。感谢您的提议！

  • 👍1