CUDA中的三维网格、线程块如何映射到硬件？ about ventus-gpgpu HOT 6 CLOSED

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

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

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

from ventus-gpgpu.

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。

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

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

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

from ventus-gpgpu.

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

from ventus-gpgpu.

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。

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

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

from ventus-gpgpu.

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

from ventus-gpgpu.

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

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

from ventus-gpgpu.