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
2: Build all the programs
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
NOTE: If you see any packages missing information, just install them.
The following packages are needed for other repositories:
device-tree-compiler
bsdmainutils
ATTENTION: In addition, we also provide Dockerfiles for Ubuntu and CentOS in .github/workflows/containers/dockerfiles. You can use them directly if needed. The following “6: Docker image” has the corresponding usage.
Before running ./build-ventus.sh to automatically build all the programs, we need to set the following commands:
For developers who want to build Debug version 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 <path-to-llvm-ventus>/../pocl.
export OCL_ICD_DIR=<path-to-ocl-icd-dir>, default folder path will be set to be <path-to-llvm-ventus>/../ocl-icd.
You can dive into build-ventus.sh file to see the detailed information about build process.
3: Bridge icd loader
Run export VENTUS_INSTALL_PREFIX=<path_to_install> to set VENTUS_INSTALL_PREFIX environment variable(system environment variable recommended), default folder path will be set to be <path-to-llvm-ventus>/install.
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:
$ <pocl-install-dir>/bin/poclcc -l
// The following output should be shown:
LIST OF DEVICES:
0:
Vendor: THU
Name: Ventus GPGPU device
Version: 2.2 HSTR: THU-ventus-gpgpu
NOTE: OpenCL host side program should be linked with icd loader -lOpenCL.
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:
./<pocl-install-dir>/examples/vecadd/vecadd
You will see that the program runs correctly.
4: Compiler using example
We can now use our built compiler to generate an ELF file, and using spike to complete the isa simulation.
NOTE: Cause the address space requirement in spike, we use a customized linker script for our compiler.
First, name the following program vecadd.cl, and place it under <path-to-llvm-ventus>:
__kernel void vectorAdd(__global float* A, __global float* B) {
unsigned tid = get_global_id(0);
A[tid] += B[tid];
}
Then, run the commands listed as follows under the same directory.
NOTE: Remember to build libclc too because we need the libclc library.
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:
We need to run the isa simulator to verify our compiler. Use spike from THU and follow the README.md.
4.4: Driver using example
Accordingly, after all the building process, you can change directory to <path-to-llvm-ventus>/../pocl/build/examples/vecadd directory, then export variables as what 3: Bridge icd loader does, finally just execute the file vecadd.
5: Github actions
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
6: Docker image
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 2: Build all the programs:
This is the Ventus GPGPU port of LLVM Compiler Infrastructure
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
Getting Started
1: Programs related repositories
Download all the repositories firstly and place them in the same path.
ventus_readme.md
.)2: Build all the programs
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.
The following packages are needed for other repositories:
Before running
./build-ventus.sh
to automatically build all the programs, we need to set the following commands:Debug
version 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<path-to-llvm-ventus>/../pocl
.export OCL_ICD_DIR=<path-to-ocl-icd-dir>
, default folder path will be set to be<path-to-llvm-ventus>/../ocl-icd
.You can dive into
build-ventus.sh
file to see the detailed information about build process.3: Bridge icd loader
Run
export VENTUS_INSTALL_PREFIX=<path_to_install>
to setVENTUS_INSTALL_PREFIX
environment variable(system environment variable recommended), default folder path will be set to be<path-to-llvm-ventus>/install
.Run
export LD_LIBRARY_PATH=${VENTUS_INSTALL_PREFIX}/lib
to tell OpenCL application to use your own builtlibOpenCL.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: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:You will see that the program runs correctly.
4: Compiler using example
We can now use our built compiler to generate an ELF file, and using spike to complete the isa simulation.
First, name the following program
vecadd.cl
, and place it under<path-to-llvm-ventus>
:Then, run the commands listed as follows under the same directory.
4.1: Generate ELF file
4.1.1 Compile directly
4.1.2 Compile step-by-step
crt0
andlibclc
All the libclc workitem functions’ implementation is included inriscv32clc.o
4.1.3 Compile assembly code to object file (
.s
to.o
)Take custome instructions
custome.s
as an example :4.2: Dump file
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:or you can check encoding of custom instructions:
4.3: Running in spike
We need to run the isa simulator to verify our compiler. Use spike from THU and follow the
README.md
.4.4: Driver using example
Accordingly, after all the building process, you can change directory to
<path-to-llvm-ventus>/../pocl/build/examples/vecadd
directory, then export variables as what 3: Bridge icd loader does, finally just execute the filevecadd
.5: Github actions
the workflow file is
.github/workflows/ventus-build.yml
, including below jobs:6: Docker image
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 2: Build all the programs: