This is a deep(ish) dive into the Vortex open-source GPU. This is a very interesting codebase with both software and hardware components to get things up and running.
Hello-world Link to heading
Starting with installation, if needed, modify the configure script to add your distribution. In my case, I was running Mint, so I had to add this branch in the configure script:
git clone --depth=1 --recursive https://github.com/vortexgpgpu/vortex.git
cd vortex
diff --git a/configure b/configure
index fbcd3f1..3d5898b 100755
--- a/configure
+++ b/configure
@@ -37,6 +37,7 @@ detect_osversion() {
# Add new versions as needed
esac
;;
+ linuxmint) osversion="ubuntu/focal";;
esac
fi
echo "$osversion"
Then build dependencies and toolchain with install_dependencies.sh:
sudo ./ci/install_dependencies.sh
mkdir build; cd build
../configure --xlen=64 --tooldir=$PWD/tools
make -s
./ci/toolchain_install.sh --all
Finally, to run Vortex, we need to source ./ci/toolchain_env.sh and use ./ci/blackbox.sh to run simulation target. In this example, I run vecadd and it generates the following output:
source ./ci/toolchain_env.sh
./ci/blackbox.sh --cores=2 --app=vecadd
Workload size=64
Create context
Allocate device buffers
Create program from kernel source
Upload source buffers
Execute the kernel
Elapsed time: 51 ms
Download destination buffer
Verify result
PASSED!
PERF: core0: instrs=5596, cycles=13524, IPC=0.413783
PERF: core1: instrs=5596, cycles=13523, IPC=0.413814
PERF: instrs=11192, cycles=13524, IPC=0.827566
vecadd walkthrough Link to heading
To understand how things work, I am going through the vecadd source code.
main.cc Link to heading
Starting with tests/opencl/vecadd/main.cc, starting from the top where context is created with clCreateContext:
CL_CHECK(clGetPlatformIDs(1, &platform_id, NULL));
CL_CHECK(clGetDeviceIDs(platform_id, CL_DEVICE_TYPE_DEFAULT, 1, &device_id, NULL));
printf("Create context\n");
context = CL_CHECK2(clCreateContext(NULL, 1, &device_id, NULL, NULL, &_err));
And kernel.cl is loaded into program to create kernel:
printf("Create program from kernel source\n");
if (0 != read_kernel_file("kernel.cl", &kernel_bin, &kernel_size))
return -1;
program = CL_CHECK2(clCreateProgramWithSource(
context, 1, (const char**)&kernel_bin, &kernel_size, &_err));
// Build program
CL_CHECK(clBuildProgram(program, 1, &device_id, NULL, NULL, NULL));
// Create kernel
kernel = CL_CHECK2(clCreateKernel(program, KERNEL_NAME, &_err));
printf("Allocate device buffers\n");
size_t nbytes = size * sizeof(TYPE);
a_memobj = CL_CHECK2(clCreateBuffer(context, CL_MEM_READ_ONLY, nbytes, NULL, &_err));
b_memobj = CL_CHECK2(clCreateBuffer(context, CL_MEM_READ_ONLY, nbytes, NULL, &_err));
c_memobj = CL_CHECK2(clCreateBuffer(context, CL_MEM_WRITE_ONLY, nbytes, NULL, &_err));
...
...
// Set kernel arguments
CL_CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *)&a_memobj));
CL_CHECK(clSetKernelArg(kernel, 1, sizeof(cl_mem), (void *)&b_memobj));
CL_CHECK(clSetKernelArg(kernel, 2, sizeof(cl_mem), (void *)&c_memobj));
...
...
// Creating command queue
commandQueue = CL_CHECK2(clCreateCommandQueue(context, device_id, 0, &_err));
printf("Upload source buffers\n");
CL_CHECK(clEnqueueWriteBuffer(commandQueue, a_memobj, CL_TRUE, 0, nbytes, h_a.data(), 0, NULL, NULL));
CL_CHECK(clEnqueueWriteBuffer(commandQueue, b_memobj, CL_TRUE, 0, nbytes, h_b.data(), 0, NULL, NULL));
Finally, kernel is executed and return buffer is copied back
printf("Execute the kernel\n");
size_t global_work_size[1] = {size};
size_t local_work_size[1] = {1};
auto time_start = std::chrono::high_resolution_clock::now();
CL_CHECK(clEnqueueNDRangeKernel(commandQueue, kernel, 1, NULL, global_work_size, local_work_size, 0, NULL, NULL));
CL_CHECK(clFinish(commandQueue));
auto time_end = std::chrono::high_resolution_clock::now();
double elapsed = std::chrono::duration_cast<std::chrono::milliseconds>(time_end - time_start).count();
printf("Elapsed time: %lg ms\n", elapsed);
printf("Download destination buffer\n");
CL_CHECK(clEnqueueReadBuffer(commandQueue, c_memobj, CL_TRUE, 0, nbytes, h_c.data(), 0, NULL, NULL));
The example calls vecadd_cpu to compare with GPU output
std::vector<TYPE> h_a(size);
std::vector<TYPE> h_b(size);
std::vector<TYPE> h_c(size);
// Generate input values
for (uint32_t i = 0; i < size; ++i) {
h_a[i] = Comparator<TYPE>::generate();
h_b[i] = Comparator<TYPE>::generate();
}
std::vector<TYPE> h_ref(size);
vecadd_cpu(h_ref.data(), h_a.data(), h_b.data(), size);
static void vecadd_cpu(TYPE *C, const TYPE* A, const TYPE *B, int N) {
for (int i = 0; i < N; ++i) {
C[i] = A[i] + B[i];
}
}
kernel.cl Link to heading
The kernel loaded in vecadd above is just addition of A and B
tests/opencl/vecadd/kernel.cl
#include "common.h"
__kernel void vecadd (__global const TYPE *A,
__global const TYPE *B,
__global TYPE *C)
{
int gid = get_global_id(0);
C[gid] = A[gid] + B[gid];
}
Vortex Simulation Drivers Link to heading
Vortex comes with different drivers, These are the most interesting ones
- rtlsim
- simx
- xrt
./ci/blackbox.sh --driver=rtlsim --clusters=2 --cores=2 --warps=2 --threads=4 --app=basic
./ci/blackbox.sh --driver=simx --clusters=2 --cores=2 --warps=2 --threads=4 --app=basic
./ci/blackbox.sh --driver=xrt --clusters=2 --cores=2 --warps=2 --threads=4 --app=basic
I will look at rtlsim and xrt in little more depth
Vortex rtlsim Link to heading
RTL sim is cycle accurate simulation using rtlsim_shim in sim/rtlsim/Makefile
RTL_PKGS = $(RTL_DIR)/VX_gpu_pkg.sv $(RTL_DIR)/fpu/VX_fpu_pkg.sv
TOP = rtlsim_shim
VL_FLAGS = --exe
VL_FLAGS += --language 1800-2009 --assert -Wall -Wpedantic
VL_FLAGS += -Wno-DECLFILENAME -Wno-REDEFMACRO
VL_FLAGS += --x-initial unique --x-assign unique
VL_FLAGS += verilator.vlt
VL_FLAGS += -DSIMULATION -DSV_DPI
VL_FLAGS += -DXLEN_$(XLEN)
VL_FLAGS += $(CONFIGS)
VL_FLAGS += $(RTL_INCLUDE)
VL_FLAGS += $(RTL_PKGS)
VL_FLAGS += --cc $(TOP) --top-module $(TOP)
PROJECT := rtlsim
SRCS = $(COMMON_DIR)/util.cpp $(COMMON_DIR)/mem.cpp $(COMMON_DIR)/softfloat_ext.cpp $(COMMON_DIR)/rvfloats.cpp $(COMMON_DIR)/dram_sim.cpp
SRCS += $(DPI_DIR)/util_dpi.cpp $(DPI_DIR)/float_dpi.cpp
SRCS += $(SRC_DIR)/processor.cpp
$(DESTDIR)/$(PROJECT): $(SRCS) $(SRC_DIR)/main.cpp
verilator --build $(VL_FLAGS) $^ -CFLAGS '$(CXXFLAGS) -DSTARTUP_ADDR=0x80000000' -LDFLAGS '$(LDFLAGS)' --Mdir $@.obj_dir -o $@
$(DESTDIR)/lib$(PROJECT).so: $(SRCS)
verilator --build $(VL_FLAGS) $^ -CFLAGS '$(CXXFLAGS)' -LDFLAGS '-shared $(LDFLAGS)' --Mdir $@.obj_dir -o $@
The top rtlsim_shim uses Vortex block in sim/rtlsim/rtlsim_shim.sv
Vortex vortex (
`SCOPE_IO_BIND (0)
.clk (clk),
.reset (reset),
.mem_req_valid (vx_mem_req_valid),
.mem_req_rw (vx_mem_req_rw),
.mem_req_byteen (vx_mem_req_byteen),
.mem_req_addr (vx_mem_req_addr),
.mem_req_data (vx_mem_req_data),
.mem_req_tag (vx_mem_req_tag),
.mem_req_ready (vx_mem_req_ready),
.mem_rsp_valid (vx_mem_rsp_valid),
.mem_rsp_data (vx_mem_rsp_data),
.mem_rsp_tag (vx_mem_rsp_tag),
.mem_rsp_ready (vx_mem_rsp_ready),
.dcr_wr_valid (dcr_wr_valid),
.dcr_wr_addr (dcr_wr_addr),
.dcr_wr_data (dcr_wr_data),
.busy (busy)
);
Vortex xrt Link to heading
Xilinx Runtime (XRT) is implemented as as a combination of userspace and kernel driver components.
The Makefile used vortex_afu_shim top to build verilator executable
SRCS = $(COMMON_DIR)/util.cpp $(COMMON_DIR)/mem.cpp $(COMMON_DIR)/softfloat_ext.cpp $(COMMON_DIR)/rvfloats.cpp $(COMMON_DIR)/dram_sim.cpp
SRCS += $(DPI_DIR)/util_dpi.cpp $(DPI_DIR)/float_dpi.cpp
SRCS += $(SRC_DIR)/xrt.cpp $(SRC_DIR)/xrt_sim.cpp
RTL_PKGS += $(RTL_DIR)/VX_gpu_pkg.sv $(RTL_DIR)/fpu/VX_fpu_pkg.sv
FPU_INCLUDE = -I$(RTL_DIR)/fpu
ifneq (,$(findstring FPU_FPNEW,$(CONFIGS)))
RTL_PKGS += $(THIRD_PARTY_DIR)/cvfpu/src/fpnew_pkg.sv $(THIRD_PARTY_DIR)/cvfpu/src/common_cells/src/cf_math_pkg $(THIRD_PARTY_DIR)/cvfpu/src/fpu_div_sqrt_mvp/hdl/defs_div_sqrt_mvp.sv
FPU_INCLUDE += -I$(THIRD_PARTY_DIR)/cvfpu/src/common_cells/include -I$(THIRD_PARTY_DIR)/cvfpu/src/common_cells/src -I$(THIRD_PARTY_DIR)/cvfpu/src/fpu_div_sqrt_mvp/hdl -I$(THIRD_PARTY_DIR)/cvfpu/src
endif
RTL_INCLUDE = -I$(SRC_DIR) -I$(RTL_DIR) -I$(DPI_DIR) -I$(RTL_DIR)/libs -I$(RTL_DIR)/interfaces -I$(RTL_DIR)/core -I$(RTL_DIR)/mem -I$(RTL_DIR)/cache $(FPU_INCLUDE)
RTL_INCLUDE += -I$(AFU_DIR)
TOP = vortex_afu_shim
VL_FLAGS += --language 1800-2009 --assert -Wall -Wpedantic
VL_FLAGS += -Wno-DECLFILENAME -Wno-REDEFMACRO
VL_FLAGS += --x-initial unique --x-assign unique
VL_FLAGS += -DSIMULATION -DSV_DPI
VL_FLAGS += -DXLEN_$(XLEN)
VL_FLAGS += $(CONFIGS)
VL_FLAGS += verilator.vlt
VL_FLAGS += $(RTL_INCLUDE)
VL_FLAGS += $(RTL_PKGS)
$(DESTDIR)/vortex.xml:
verilator --xml-only -O0 $(VL_FLAGS) $(TOP) --xml-output $@
$(DESTDIR)/scope.json: $(DESTDIR)/vortex.xml
$(SCRIPT_DIR)/scope.py $^ -o $@
$(DESTDIR)/$(PROJECT): $(SRCS) $(SCOPE_JSON)
verilator --build --exe $(VL_FLAGS) --cc $(TOP) --top-module $(TOP) $(SRCS) -CFLAGS '$(CXXFLAGS)' -LDFLAGS '$(LDFLAGS)' --Mdir $@.obj_dir -o $@
This is chain of modules instances starting from vortex_afu_shim.sv down to Vortex
sim/xrtsim/vortex_afu_shim.sv
VX_afu_wrap #(
.C_S_AXI_CTRL_ADDR_WIDTH (C_S_AXI_CTRL_ADDR_WIDTH),
.C_S_AXI_CTRL_DATA_WIDTH (C_S_AXI_CTRL_DATA_WIDTH),
.C_M_AXI_MEM_ID_WIDTH (C_M_AXI_MEM_ID_WIDTH),
.C_M_AXI_MEM_DATA_WIDTH (C_M_AXI_MEM_DATA_WIDTH),
.C_M_AXI_MEM_ADDR_WIDTH (C_M_AXI_MEM_ADDR_WIDTH),
.C_M_AXI_MEM_NUM_BANKS (C_M_AXI_MEM_NUM_BANKS)
) afu_wrap (
.clk (ap_clk),
.reset (~ap_rst_n),
`REPEAT (`PLATFORM_MEMORY_BANKS, AXI_MEM_ARGS, REPEAT_COMMA),
.s_axi_ctrl_awvalid (s_axi_ctrl_awvalid),
.s_axi_ctrl_awready (s_axi_ctrl_awready),
.s_axi_ctrl_awaddr (s_axi_ctrl_awaddr),
.s_axi_ctrl_wvalid (s_axi_ctrl_wvalid),
.s_axi_ctrl_wready (s_axi_ctrl_wready),
.s_axi_ctrl_wdata (s_axi_ctrl_wdata),
.s_axi_ctrl_wstrb (s_axi_ctrl_wstrb),
.s_axi_ctrl_arvalid (s_axi_ctrl_arvalid),
.s_axi_ctrl_arready (s_axi_ctrl_arready),
.s_axi_ctrl_araddr (s_axi_ctrl_araddr),
.s_axi_ctrl_rvalid (s_axi_ctrl_rvalid),
.s_axi_ctrl_rready (s_axi_ctrl_rready),
.s_axi_ctrl_rdata (s_axi_ctrl_rdata),
.s_axi_ctrl_rresp (s_axi_ctrl_rresp),
.s_axi_ctrl_bvalid (s_axi_ctrl_bvalid),
.s_axi_ctrl_bready (s_axi_ctrl_bready),
.s_axi_ctrl_bresp (s_axi_ctrl_bresp),
.interrupt (interrupt)
);
hw/rtl/afu/xrt/VX_afu_wrap.sv
Vortex_axi #(
.AXI_DATA_WIDTH (C_M_AXI_MEM_DATA_WIDTH),
.AXI_ADDR_WIDTH (M_AXI_MEM_ADDR_WIDTH),
.AXI_TID_WIDTH (C_M_AXI_MEM_ID_WIDTH),
.AXI_NUM_BANKS (C_M_AXI_MEM_NUM_BANKS)
) vortex_axi (
`SCOPE_IO_BIND (1)
.clk (clk),
.reset (vx_reset),
.m_axi_awvalid (m_axi_mem_awvalid_a),
.m_axi_awready (m_axi_mem_awready_a),
.m_axi_awaddr (m_axi_mem_awaddr_u),
.m_axi_awid (m_axi_mem_awid_a),
.m_axi_awlen (m_axi_mem_awlen_a),
`UNUSED_PIN (m_axi_awsize),
`UNUSED_PIN (m_axi_awburst),
`UNUSED_PIN (m_axi_awlock),
`UNUSED_PIN (m_axi_awcache),
`UNUSED_PIN (m_axi_awprot),
`UNUSED_PIN (m_axi_awqos),
`UNUSED_PIN (m_axi_awregion),
.m_axi_wvalid (m_axi_mem_wvalid_a),
.m_axi_wready (m_axi_mem_wready_a),
.m_axi_wdata (m_axi_mem_wdata_a),
.m_axi_wstrb (m_axi_mem_wstrb_a),
.m_axi_wlast (m_axi_mem_wlast_a),
.m_axi_bvalid (m_axi_mem_bvalid_a),
.m_axi_bready (m_axi_mem_bready_a),
.m_axi_bid (m_axi_mem_bid_a),
.m_axi_bresp (m_axi_mem_bresp_a),
.m_axi_arvalid (m_axi_mem_arvalid_a),
.m_axi_arready (m_axi_mem_arready_a),
.m_axi_araddr (m_axi_mem_araddr_u),
.m_axi_arid (m_axi_mem_arid_a),
.m_axi_arlen (m_axi_mem_arlen_a),
`UNUSED_PIN (m_axi_arsize),
`UNUSED_PIN (m_axi_arburst),
`UNUSED_PIN (m_axi_arlock),
`UNUSED_PIN (m_axi_arcache),
`UNUSED_PIN (m_axi_arprot),
`UNUSED_PIN (m_axi_arqos),
`UNUSED_PIN (m_axi_arregion),
.m_axi_rvalid (m_axi_mem_rvalid_a),
.m_axi_rready (m_axi_mem_rready_a),
.m_axi_rdata (m_axi_mem_rdata_a),
.m_axi_rlast (m_axi_mem_rlast_a),
.m_axi_rid (m_axi_mem_rid_a),
.m_axi_rresp (m_axi_mem_rresp_a),
.dcr_wr_valid (dcr_wr_valid),
.dcr_wr_addr (dcr_wr_addr),
.dcr_wr_data (dcr_wr_data),
.busy (vx_busy)
);
hw/rtl/Vortex_axi.sv
Vortex vortex (
`SCOPE_IO_BIND (0)
.clk (clk),
.reset (reset),
.mem_req_valid (mem_req_valid),
.mem_req_rw (mem_req_rw),
.mem_req_byteen (mem_req_byteen),
.mem_req_addr (mem_req_addr),
.mem_req_data (mem_req_data),
.mem_req_tag (mem_req_tag),
.mem_req_ready (mem_req_ready),
.mem_rsp_valid (mem_rsp_valid),
.mem_rsp_data (mem_rsp_data),
.mem_rsp_tag (mem_rsp_tag),
.mem_rsp_ready (mem_rsp_ready),
.dcr_wr_valid (dcr_wr_valid),
.dcr_wr_addr (dcr_wr_addr),
.dcr_wr_data (dcr_wr_data),
.busy (busy)
);