Ventus GPGPU is based on RISCV RV32IMAZfinxZve32f ISA with fully redefined concept of V-extension.
The Ventus GPGPU OpenCL compiler based on LLVM is developed by Terapines Technology (Wuhan) Co., Ltd
承影GPGPU OpenCL编译器由Terapines(兆松科技)负责开发
For more architecture detail, please refer to Ventus GPGPU Arch
Download all the repositories firstly and place them in the same path.
- llvm-ventus : git clone https://github.com/THU-DSP-LAB/llvm-project.git
- pocl : git clone https://github.com/THU-DSP-LAB/pocl.git
- ocl-icd : git clone https://github.com/OCL-dev/ocl-icd.git
- isa-simulator(spike) : git clone https://github.com/THU-DSP-LAB/ventus-gpgpu-isa-simulator.git
- driver : git clone https://github.com/THU-DSP-LAB/ventus-driver.git
ATTENTION: Remember to check branch for every repository, cause the project are under development, if you get any build errors, feel free to give an issue or just contact authors
Our program is based on LLVM, so the need packages to build ventus are almost the same as what are needed to build LLVM, you can refer to official website for detailed llvm building guidance, we just list most important needed packages here.
- ccache
- cmake
- ninja
- clang
If you see any packages missing information, just install them
The following packages are needed for other repositories:
- device-tree-compiler
- bsdmainutils
Run ./build-ventus.sh to automatically build all the programs, but we need to run firstly
- For developers who want to build
Debugversion for llvm,export BUILD_TYPE=Debug, since it's set default to be 'Release' export POCL_DIR=<path-to-pocl-dir>, default folder path will be set to be<llvm-ventus-parentFolder>/poclexport OCL_ICD_DIR=<path-to-ocl-icd-dir>, default folder path will be set to be<llvm-ventus-parentFolder>/ocl-icd
You can dive into build-ventus.sh file to see the detailed information about build process
Run export VENTUS_INSTALL_PREFIX=<path_to_install> to set VENTUS_INSTALL_PREFIX environment variable(system environment variable recommended)
Run export LD_LIBRARY_PATH=${VENTUS_INSTALL_PREFIX}/lib to tell OpenCL application to use your own built libOpenCL.so, also to correctly locate LLVM shared libraries
Run export OCL_ICD_VENDORS=${VENTUS_INSTALL_PREFIX}/lib/libpocl.so to tell ocl icd loader where the icd driver is.
Finally, run export POCL_DEVICES="ventus" to tell pocl driver which device is available(should we set ventus as default device?).
You will see Ventus GPGPU device is found if your setup is correct.
NOTE: OpenCL host side program should be linked with icd loader -lOpenCL.
$ <pocl-install-dir>/bin/poclcc -l
LIST OF DEVICES:
0:
Vendor: THU
Name: Ventus GPGPU device
Version: 2.2 HSTR: THU-ventus-gpgpu
Also, you can try to set POCL_DEBUG=all and run example under <pocl-build-dir> to see the full OpenCL software stack execution pipeline. For example(Work in progress).
we can now use our built compiler to generate an ELF file, and using spike to complete the isa simulation.
Cause the address space requirement in spike, we use a customized linker script for our compiler
Take vecadd.cl below as an example :
__kernel void vectorAdd(__global float* A, __global float* B) {
unsigned tid = get_global_id(0);
A[tid] += B[tid];
}
Remember to build libclc too because we need the libclc library
Use command line under the root directory of llvm-ventus
./install/bin/clang -cl-std=CL2.0 -target riscv32 -mcpu=ventus-gpgpu vecadd.cl ./install/lib/crt0.o -L./install/lib -lworkitem -I./libclc/generic/include -nodefaultlibs ./libclc/riscv32/lib/workitem/get_global_id.cl -O1 -cl-std=CL2.0 -Wl,-T,utils/ldscripts/ventus/elf32lriscv.ld -o vecadd.riscv
- Compile OpenCL code to LLVM IR assembly (.ll file):
./install/bin/clang -S -cl-std=CL2.0 -target riscv32 -mcpu=ventus-gpgpu vecadd.cl -emit-llvm -o vecadd.ll- Compile LLVM IR to RISC-V assembly or object file:
./install/bin/llc -mtriple=riscv32 -mcpu=ventus-gpgpu vecadd.ll -o vecadd.s./install/bin/llc -mtriple=riscv32 -mcpu=ventus-gpgpu --filetype=obj vecadd.ll -o vecadd.o- Link essential library:
Linking
crt0andlibclcAll the libclc workitem functions' implementation is included inriscv32clc.o
./install/bin/ld.lld -o vecadd.riscv -T utils/ldscripts/ventus/elf32lriscv.ld vecadd.o ./install/lib/crt0.o ./install/lib/riscv32clc.o -L./install/lib -lworkitem --gc-sectionsTake custome instructions custome.s as an example :
vftta.vv v0, v0, v1
vfexp v0, v1
vadd12.vi v0, v1, 8./install/bin/clang -c -target riscv32 -mcpu=ventus-gpgpu custom.s -o custom.o./install/bin/llvm-objdump -d --mattr=+v,+zfinx vecadd.riscv >& vecadd.txt
you will see output like below, 0x80000000 is the space address required by spike for _start function, this is the reason why we use a customized linker script
vecadd.riscv: file format elf32-littleriscv
Disassembly of section .text:
80000000 <_start>:
80000000: 97 21 00 00 auipc gp, 2
80000004: 93 81 01 80 addi gp, gp, -2048
80000008: 93 0e 00 02 li t4, 32
8000000c: d7 fe 0e 0d vsetvli t4, t4, e32, m1, ta, ma
80000010: b7 2e 00 00 lui t4, 2
80000014: f3 ae 0e 30 csrrs t4, mstatus, t4
80000018: 93 0e 00 00 li t4, 0
8000001c: 73 21 60 80 csrr sp, 2054
80000020: 73 22 70 80 csrr tp, 2055
80000024 <.Lpcrel_hi1>:
80000024: 17 15 00 00 auipc a0, 1
80000028: 13 05 85 fe addi a0, a0, -24
....
....
....
or you can check encoding of custom instructions
./install/bin/llvm-objdump -d --mattr=+v,+zfinx custom.o >& custom.txt
custom.o: file format elf32-littleriscv
Disassembly of section .text:
00000000 <.text>:
0: 0b c0 00 0e vftta.vv v0, v0, v1
4: 0b 60 10 0a vfexp v0, v1
8: 0b 80 80 00 vadd12.vi v0, v1, 8
We need to run the isa simulator to verify our compiler
Use spike from THU and follow the README.md
Accordingly, after all the building process, you can change directory to <llvm-ventus-parentFolder>/pocl/build/examples/vecadd directory, then export variables as what Bridge icd loader does, finally just execute the file vecadd
the workflow file is .github/workflows/ventus-build.yml, including below jobs
- Build llvm
- Build ocl-icd
- Build libclc
- Build isa-simulator
- Build sumulator-driver
- Build pocl
- Isa simulation test
- GPU-rodinia testsuite
- Pocl testing
If the user needs to build the toolchain of the Ventus project in an environment other than Ubuntu, such as the CentOS system, we provide the Dockerfile for building the CentOS image. The file is under '.github/workflows/containers/dockerfiles'.
Note: When using build-ventus.sh to build the instantiated centos container, the following modifications are required, which are different from the above "2: Build all the programs":
--- a/build-ventus.sh
+++ b/build-ventus.sh
@@ -119,6 +119,8 @@ build_llvm() {
-DLLVM_CCACHE_BUILD=ON \
-DLLVM_OPTIMIZED_TABLEGEN=ON \
-DLLVM_PARALLEL_LINK_JOBS=12 \
+ -DCMAKE_C_COMPILER=clang \
+ -DCMAKE_CXX_COMPILER=clang++ \
-DCMAKE_BUILD_TYPE=${BUILD_TYPE} \
-DLLVM_ENABLE_PROJECTS="clang;lld;libclc" \
-DLLVM_TARGETS_TO_BUILD="AMDGPU;X86;RISCV" \
@@ -232,7 +234,7 @@ export_elements() {
export SPIKE_TARGET_DIR=${VENTUS_INSTALL_PREFIX}
export VENTUS_INSTALL_PREFIX=${VENTUS_INSTALL_PREFIX}
export POCL_DEVICES="ventus"
- export OCL_ICD_VENDORS=${VENTUS_INSTALL_PREFIX}/lib/libpocl.so
+ export OCL_ICD_VENDORS=${VENTUS_INSTALL_PREFIX}/lib64/libpocl.so
}
React
Typescript
Django
Laravel
Facebook
Microsoft
Google
Alibaba
D3
Tencent