Skip to content
vortexgpgpu
Volt
Public

About

No description, website, or topics provided.

Resources

Readme
Apache-2.0 license
Activity
Custom properties

Stars

23 stars

Watchers

0 watching

Forks

4 forks
Report repository
BranchesTags

Repository files navigation

Vortex-Optimized Light-weight Toolchain (VOLT)

VOLT is an end-to-end, open-source compiler framework for the Vortex open-source GPU and its variants. It supports multiple front-end languages, including CUDA and OpenCL, by extending two open-source projects, PoCL and CuPBoP. VOLT also extends LLVM to enable SIMT-aware IR-level optimizations for portability across GPU variants, and augments LLVM-RISC-V to provide full support for the Vortex ISA.

Volt Overview

The figure shows an overview of VOLT, with the modules extended or introduced by VOLT highlighted in blue.

Vortex 3.0 Support (Volt v1.1, Vortex Tutorial 2026)

Volt v1.1 targets Vortex 3.0 (release notes):

Submodule Branch Commit
vortex master (Vortex 3.0) 50cb17454
pocl vortex_3.x (PoCL 7.0) 5e06f2f44
cupbop vortex3.0 d289d365f
llvm vortex_3.x (LLVM 20.1.8 + Volt divergence levels -vortex-branch-opt) ca6791f23

Build and test everything (OpenCL through PoCL, CUDA through CuPBoP, both on SimX):

git clone --recurse-submodules https://github.com/vortexgpgpu/Volt.git
cd Volt/tests/vortex26_tutorial
./build.sh          # LLVM, Vortex 3.0 + toolchain, PoCL, CuPBoP
./run_opencl.sh     # 33 OpenCL benchmarks  -> total=33 pass=33 fail=0
./run_cuda.sh       # 24 CUDA examples      -> total=24 pass=24 fail=0
  • tests/vortex26_tutorial/README.md: dependencies, options, and the divergence levels (VX_BRANCH_OPT / VORTEX_DIVERGENCE_OPT_LEVEL).
  • Documentation: updated for Vortex 3.0. The tutorials (section 7: kernel/host function extension, shared-memory lowering, and the new tensor-core / wgmma / DXA tutorial) come with runnable scripts in tests/vortex26_tutorial/tutorials.

CC'26 artifact (Vortex 2.x, Volt LLVM 18): use the v1.0 tag (git clone --recurse-submodules --branch v1.0 https://github.com/vortexgpgpu/Volt.git) and follow tests/cc26_ae/README.md. The "Build Toolchain (Vortex v2)" section below describes that version.

Citation and Presentation Materials

Vortex Tutorials

VOLT paper (link)

The VOLT paper was accepted in the 35th ACM SIGPLAN International Conference on Compiler Construction (CC' 2026).

BibTeX
@inproceedings{10.1145/3771775.3786275,
author = {Jeong, Shinnung and Ahn, Chihyo and Pu, Huanzhi and Zhao, Jisheng and Kim, Hyesoon and Tine, Blaise},
title = {Inside VOLT: Designing an Open-Source GPU Compiler (Tool)},
year = {2026},
isbn = {9798400722745},
url = {https://doi.org/10.1145/3771775.3786275},
doi = {10.1145/3771775.3786275},
booktitle = {Proceedings of the 35th ACM SIGPLAN International Conference on Compiler Construction},
keywords = {Compiler, Divergence Management, Open GPU},
location = {Sydney, NSW, Australia},
series = {CC '26}
}

Build Toolchain (Vortex v2, 64-bit)

This section and "Test Toolchain" below apply to the CC'26 version (v1.0 tag: Vortex 2.x, PoCL volt, Volt LLVM 18). For Vortex 3.0, see Vortex 3.0 Support.

1. Folder structure

<VOLT>
├── tools # location of libraries and bin files
│   ├── libc32
│   ├── libc64
│   ├── libcrt32
│   ├── libcrt64
│   ├── llvm-vortex
│   ├── pocl
│   ├── riscv32-gnu-toolchain
│   ├── riscv64-gnu-toolchain
│   ├── sv2v
│   ├── verilator
│   ├── vortex
│   └── yosys
├── pocl 
├── llvm-vortex
└── vortex

2. Set up folders

cd Volt
mkdir tools

3. Make a config file (add to the bash file)

  export PRJ=<path-to-volt-folder>
  export TOOL_DIR=$PRJ/tools
  export VORTEX_PREFIX=$TOOL_DIR/vortex
  export LLVM_PREFIX=$TOOL_DIR/llvm-vortex
  export POCL_PREFIX=$TOOL_DIR/pocl
  source $PRJ/vortex/build/ci/toolchain_env.sh

4. Install vortex

cd $PRJ/vortex
sudo ./ci/install_dependencies.sh
mkdir build && cd build

../configure --xlen=64 --tooldir=$TOOL_DIR --prefix=$VORTEX_PREFIX #setup 32bit
./ci/toolchain_install.sh --all # install prebuilt toolchain
source ./ci/toolchain_env.sh
make -s
make install
./ci/blackbox.sh --cores=2 --app=opencl/vecadd # test

5. Install llvm-vortex

cd $PRJ
mv $LLVM_PREFIX $TOOL_DIR/llvm-vortex-prebuilt 
mkdir $TOOL_DIR/llvm-vortex

cd llvm
mkdir build && cd build

export RISCV_TOOLCHAIN_ELF_PATH=$PRJ/tools/riscv64-gnu-toolchain/riscv64-unknown-elf

cmake -G "Unix Makefiles" -DLLVM_ABI_BREAKING_CHECKS=FORCE_OFF -DLLVM_INCLUDE_EXAMPLES=OFF -DLLVM_INCLUDE_TESTS=OFF -DBUILD_SHARED_LIBS=True -DCMAKE_BUILD_TYPE=RelWithDebInfo -DLLVM_ENABLE_ASSERTIONS=ON -DLLVM_ENABLE_PROJECTS="clang" -DCMAKE_INSTALL_PREFIX=$LLVM_PREFIX -DDEFAULT_SYSROOT=$RISCV_TOOLCHAIN_ELF_PATH -DLLVM_DEFAULT_TARGET_TRIPLE="riscv64-unknown-elf" -DLLVM_TARGETS_TO_BUILD="X86;RISCV;NVPTX" ../llvm
make -j`nproc`
make install

6. Install POCL-Vortex

cd $PRJ
mv $POCL_PREFIX $TOOL_DIR/pocl-vortex-prebuilt 
mkdir $TOOL_DIR/pocl-vortex

cd $PRJ/pocl
mkdir build && cd build

cmake -G "Unix Makefiles" \
  -DCMAKE_BUILD_TYPE=Debug \
  -DCMAKE_INSTALL_PREFIX=$POCL_PREFIX \
  -DWITH_LLVM_CONFIG=$LLVM_PREFIX/bin/llvm-config \
  -DVORTEX_HOME=$VORTEX_HOME \
  -DVORTEX_BUILD=$VORTEX_PREFIX \
  -DENABLE_VORTEX=ON \
  -DKERNEL_CACHE_DEFAULT=OFF \
  -DENABLE_HOST_CPU_DEVICES=OFF \
  -DENABLE_TESTS=OFF \
  -DPOCL_DEBUG_MESSAGES=ON \
  -DENABLE_ICD=OFF ..

make -j`nproc`
make install

7. Install CuPBoP-Vortex

cd $PRJ/cupbop
export CuPBoP_PATH=pwd
export LD_LIBRARY_PATH=$CuPBoP_PATH/build/runtime:$CuPBoP_PATH/build/runtime/threadPool:$LD_LIBRARY_PATH

wget "https://dl.dropboxusercontent.com/scl/fi/m9ap1tiybau4zk720t2z7/cuda-header.tar.gz?rlkey=zmdpst5l66t48ywrbtkj426nu&st=luao6zy7" -O cuda-header.tar.gz
tar -xzf 'cuda-header.tar.gz'
cp -r include/* runtime/threadPool/include/

wget https://developer.download.nvidia.com/compute/cuda/12.1.1/local_installers/cuda_12.1.1_530.30.02_linux.run
mkdir -p cuda-12.1
sh cuda_12.1.1_530.30.02_linux.run --silent --toolkit --toolkitpath=$CuPBoP_PATH/cuda-12.1

$CuPBoP_PATH/CuPBoP_env_setup.sh
export VORTEX_ARCHITECTUE = 64                  # 64 if using 64-bit vortex
mkdir build && cd build
cmake .. -DLLVM_CONFIG_PATH=`which llvm-config` # need path to llvm-config
make

Extend the config file (add to the bash file)

export RISCV_TOOLCHAIN=${TOOL_DIR}/riscv64-gnu-toolchain # Assume Vortex is setup with 64-bit RISC-V
export VORTEX_HOME= # VORTEX repo path
export CuPBoP_PATH= # CUPBOP repo path
export VORTEX_PATH=$VORTEX_HOME/build
export PATH=$LLVM_PREFIX/bin:$PATH
export LD_LIBRARY_PATH=$CuPBoP_PATH/build/runtime:$CuPBoP_PATH/build/runtime/threadPool:$LD_LIBRARY_PATH
export LD_LIBRARY_PATH=$CuPBoP_PATH/cuda-12.1/lib64:$LD_LIBRARY_PATH
export LD_LIBRARY_PATH=${LLVM_PREFIX}/lib:${LD_LIBRARY_PATH}
export LD_LIBRARY_PATH=${VORTEX_PATH}/runtime/:${LD_LIBRARY_PATH}

CuPBoP Running Test

cd ./exampless/cfd
./kjrun_llvm18.sh

Test Toolchain

0. Test Benchmark Location for Each Pipeline

  • OpenCL to Vortex pipeline is in vortex/tests/opencl
  • CUDA to Vortex pipeline is in cupbop/examples

1. Test OpenCL Benchmarks with Thread Divergence Optimizations

cd $PRJ/tests/test_opencl_w_divergence_opt
./run_evaluation.sh

2. Test CuPBoP with ISA extensions

cd $PRJ/tests/test_cupbop
./run_isa_extensions.sh

3. Test CuPBoP with Memory Hierarchy Support and API Support

cd $PRJ/tests/test_cupbop
./run_mem_hierarchy.sh

About

No description, website, or topics provided.

Resources

Readme
Apache-2.0 license
Activity
Custom properties

Stars

23 stars

Watchers

0 watching

Forks

4 forks
Report repository

Releases

Packages

Contributors

Languages