Comments (6)
若您在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.
Related Issues (20)
- 你好,请问当前版本实现了MMU了吗? HOT 2
- make verilog编译报错:1 targets failed HOT 1
- make verilog编译报错 HOT 5
- make tests出错
- support FPGA? HOT 11
- When run make test, it occur with the following message. Mill version is 0.10.8. HOT 1
- When we use command "make compile", here report the following errors: val config = chipsalliance.rocketchip.config HOT 1
- Where can I configure the number of cores? HOT 1
- 【RFC】关于ventus的gnu工具链支持 HOT 1
- Chisel 3.6+ Support HOT 3
- make test报错 commit: 19dc44a7529198411a7ed5b17d19443d545bddd2 HOT 1
- saxpy2 test error HOT 4
- Cannot use the released version "VENTUS-2.0.2" of software part HOT 2
- fpga_test HOT 1
- done
- 关于承影的正确性验证 HOT 2
- 生成Verilog时出现的问题 HOT 1
- fpga_test support and documentation HOT 4
- How to generate .data file HOT 2
- 承影有SIMT-STACK的退栈指令吗? HOT 1
Recommend Projects
-
React
A declarative, efficient, and flexible JavaScript library for building user interfaces.
-
Vue.js
🖖 Vue.js is a progressive, incrementally-adoptable JavaScript framework for building UI on the web.
-
Typescript
TypeScript is a superset of JavaScript that compiles to clean JavaScript output.
-
TensorFlow
An Open Source Machine Learning Framework for Everyone
-
Django
The Web framework for perfectionists with deadlines.
-
Laravel
A PHP framework for web artisans
-
D3
Bring data to life with SVG, Canvas and HTML. 📊📈🎉
-
Recommend Topics
-
javascript
JavaScript (JS) is a lightweight interpreted programming language with first-class functions.
-
web
Some thing interesting about web. New door for the world.
-
server
A server is a program made to process requests and deliver data to clients.
-
Machine learning
Machine learning is a way of modeling and interpreting data that allows a piece of software to respond intelligently.
-
Visualization
Some thing interesting about visualization, use data art
-
Game
Some thing interesting about game, make everyone happy.
Recommend Org
-
Facebook
We are working to build community through open source technology. NB: members must have two-factor auth.
-
Microsoft
Open source projects and samples from Microsoft.
-
Google
Google ❤️ Open Source for everyone.
-
Alibaba
Alibaba Open Source for everyone
-
D3
Data-Driven Documents codes.
-
Tencent
China tencent open source team.
from ventus-gpgpu.