Merge branch 'master' into graphics

This commit is contained in:
Blaise Tine
2021-10-15 19:32:11 -07:00
542 changed files with 124552 additions and 18682 deletions

View File

@@ -5,7 +5,7 @@ all:
$(MAKE) -C sfilter
$(MAKE) -C nearn
$(MAKE) -C guassian
$(MAKE) -C printf
$(MAKE) -C oclprintf
$(MAKE) -C psort
run-simx:
@@ -15,8 +15,8 @@ run-simx:
$(MAKE) -C sfilter run-simx
$(MAKE) -C nearn run-simx
$(MAKE) -C guassian run-simx
$(MAKE) -C printf run-simx
#$(MAKE) -C psort run-simx
$(MAKE) -C oclprintf run-simx
$(MAKE) -C psort run-simx
run-rtlsim:
$(MAKE) -C vecadd run-rtlsim
@@ -25,8 +25,8 @@ run-rtlsim:
$(MAKE) -C sfilter run-rtlsim
$(MAKE) -C nearn run-rtlsim
$(MAKE) -C guassian run-rtlsim
$(MAKE) -C printf run-rtlsim
#$(MAKE) -C psort run-rtlsim
$(MAKE) -C oclprintf run-rtlsim
$(MAKE) -C psort run-rtlsim
run-vlsim:
$(MAKE) -C vecadd run-vlsim
@@ -35,8 +35,8 @@ run-vlsim:
$(MAKE) -C sfilter run-vlsim
$(MAKE) -C nearn run-vlsim
$(MAKE) -C guassian run-vlsim
$(MAKE) -C printf run-vlsim
#$(MAKE) -C psort run-vlsim
$(MAKE) -C oclprintf run-vlsim
$(MAKE) -C psort run-vlsim
clean:
$(MAKE) -C vecadd clean
@@ -45,7 +45,7 @@ clean:
$(MAKE) -C sfilter clean
$(MAKE) -C nearn clean
$(MAKE) -C guassian clean
$(MAKE) -C printf clean
$(MAKE) -C oclprintf clean
$(MAKE) -C psort clean
clean-all:
@@ -55,5 +55,5 @@ clean-all:
$(MAKE) -C sfilter clean-all
$(MAKE) -C nearn clean-all
$(MAKE) -C guassian clean-all
$(MAKE) -C printf clean-all
$(MAKE) -C oclprintf clean-all
$(MAKE) -C psort clean-all

View File

@@ -7,8 +7,8 @@ POCL_RT_PATH ?= /opt/pocl/runtime
VORTEX_DRV_PATH ?= $(realpath ../../../driver)
VORTEX_RT_PATH ?= $(realpath ../../../runtime)
K_LLCFLAGS += "-O3 -march=riscv32 -target-abi=ilp32f -mcpu=generic-rv32 -mattr=+m,+f -float-abi=hard -code-model=small"
K_CFLAGS += "-v -O3 -Wstack-usage=1024 --sysroot=$(SYSROOT) --gcc-toolchain=$(RISCV_TOOLCHAIN_PATH) -march=rv32imf -mabi=ilp32f -I$(VORTEX_RT_PATH)/include -fno-rtti -fno-exceptions -ffreestanding -nostartfiles -fdata-sections -ffunction-sections"
K_LLCFLAGS += "-O3 -march=riscv32 -target-abi=ilp32f -mcpu=generic-rv32 -mattr=+m,+f -mattr=+vortex -float-abi=hard -code-model=small"
K_CFLAGS += "-v -O3 -Wstack-usage=1024 --sysroot=$(SYSROOT) --gcc-toolchain=$(RISCV_TOOLCHAIN_PATH) -march=rv32imf -mabi=ilp32f -Xclang -target-feature -Xclang +vortex -I$(VORTEX_RT_PATH)/include -fno-rtti -fno-exceptions -ffreestanding -nostartfiles -fdata-sections -ffunction-sections"
K_LDFLAGS += "-Wl,-Bstatic,-T$(VORTEX_RT_PATH)/linker/vx_link.ld -Wl,--gc-sections $(VORTEX_RT_PATH)/libvortexrt.a -lm"
CXXFLAGS += -std=c++11 -O2 -Wall -Wextra -pedantic -Wfatal-errors
@@ -31,13 +31,13 @@ $(PROJECT): $(SRCS)
$(CXX) $(CXXFLAGS) $^ $(LDFLAGS) -o $@
run-fpga: $(PROJECT) kernel.pocl
LD_LIBRARY_PATH=$(POCL_RT_PATH)/lib:$(VORTEX_DRV_PATH)/opae:$(LD_LIBRARY_PATH) ./$(PROJECT)
LD_LIBRARY_PATH=$(POCL_RT_PATH)/lib:$(VORTEX_DRV_PATH)/fpga:$(LD_LIBRARY_PATH) ./$(PROJECT)
run-asesim: $(PROJECT) kernel.pocl
LD_LIBRARY_PATH=$(POCL_RT_PATH)/lib:$(VORTEX_DRV_PATH)/opae/ase:$(LD_LIBRARY_PATH) ./$(PROJECT)
LD_LIBRARY_PATH=$(POCL_RT_PATH)/lib:$(VORTEX_DRV_PATH)/asesim:$(LD_LIBRARY_PATH) ./$(PROJECT)
run-vlsim: $(PROJECT) kernel.pocl
LD_LIBRARY_PATH=$(POCL_RT_PATH)/lib:$(VORTEX_DRV_PATH)/opae/vlsim:$(LD_LIBRARY_PATH) ./$(PROJECT)
LD_LIBRARY_PATH=$(POCL_RT_PATH)/lib:$(VORTEX_DRV_PATH)/vlsim:$(LD_LIBRARY_PATH) ./$(PROJECT)
run-simx: $(PROJECT) kernel.pocl
LD_LIBRARY_PATH=$(POCL_RT_PATH)/lib:$(VORTEX_DRV_PATH)/simx:$(LD_LIBRARY_PATH) ./$(PROJECT)

View File

@@ -7,8 +7,8 @@ POCL_RT_PATH ?= /opt/pocl/runtime
VORTEX_DRV_PATH ?= $(realpath ../../../driver)
VORTEX_RT_PATH ?= $(realpath ../../../runtime)
K_LLCFLAGS += "-O3 -march=riscv32 -target-abi=ilp32f -mcpu=generic-rv32 -mattr=+m,+f -float-abi=hard -code-model=small"
K_CFLAGS += "-v -O3 --sysroot=$(SYSROOT) --gcc-toolchain=$(RISCV_TOOLCHAIN_PATH) -march=rv32imf -mabi=ilp32f -I$(VORTEX_RT_PATH)/include -fno-rtti -fno-exceptions -ffreestanding -nostartfiles -fdata-sections -ffunction-sections"
K_LLCFLAGS += "-O3 -march=riscv32 -target-abi=ilp32f -mcpu=generic-rv32 -mattr=+m,+f -mattr=+vortex -float-abi=hard -code-model=small"
K_CFLAGS += "-v -O3 --sysroot=$(SYSROOT) --gcc-toolchain=$(RISCV_TOOLCHAIN_PATH) -march=rv32imf -mabi=ilp32f -Xclang -target-feature -Xclang +vortex -I$(VORTEX_RT_PATH)/include -fno-rtti -fno-exceptions -ffreestanding -nostartfiles -fdata-sections -ffunction-sections"
K_LDFLAGS += "-Wl,-Bstatic,-T$(VORTEX_RT_PATH)/linker/vx_link.ld -Wl,--gc-sections $(VORTEX_RT_PATH)/libvortexrt.a -lm"
CXXFLAGS += -std=c++11 -O2 -Wall -Wextra -pedantic -Wfatal-errors
@@ -31,13 +31,13 @@ $(PROJECT): $(SRCS)
$(CXX) $(CXXFLAGS) $^ $(LDFLAGS) -o $@
run-fpga: $(PROJECT) kernel.pocl
LD_LIBRARY_PATH=$(POCL_RT_PATH)/lib:$(VORTEX_DRV_PATH)/opae:$(LD_LIBRARY_PATH) ./$(PROJECT)
LD_LIBRARY_PATH=$(POCL_RT_PATH)/lib:$(VORTEX_DRV_PATH)/fpga:$(LD_LIBRARY_PATH) ./$(PROJECT)
run-asesim: $(PROJECT) kernel.pocl
LD_LIBRARY_PATH=$(POCL_RT_PATH)/lib:$(VORTEX_DRV_PATH)/opae/ase:$(LD_LIBRARY_PATH) ./$(PROJECT)
LD_LIBRARY_PATH=$(POCL_RT_PATH)/lib:$(VORTEX_DRV_PATH)/asesim:$(LD_LIBRARY_PATH) ./$(PROJECT)
run-vlsim: $(PROJECT) kernel.pocl
LD_LIBRARY_PATH=$(POCL_RT_PATH)/lib:$(VORTEX_DRV_PATH)/opae/vlsim:$(LD_LIBRARY_PATH) ./$(PROJECT)
LD_LIBRARY_PATH=$(POCL_RT_PATH)/lib:$(VORTEX_DRV_PATH)/vlsim:$(LD_LIBRARY_PATH) ./$(PROJECT)
run-simx: $(PROJECT) kernel.pocl
LD_LIBRARY_PATH=$(POCL_RT_PATH)/lib:$(VORTEX_DRV_PATH)/simx:$(LD_LIBRARY_PATH) ./$(PROJECT)

File diff suppressed because it is too large Load Diff

File diff suppressed because it is too large Load Diff

View File

@@ -7,8 +7,8 @@ POCL_RT_PATH ?= /opt/pocl/runtime
VORTEX_DRV_PATH ?= $(realpath ../../../driver)
VORTEX_RT_PATH ?= $(realpath ../../../runtime)
K_LLCFLAGS += "-O3 -march=riscv32 -target-abi=ilp32f -mcpu=generic-rv32 -mattr=+m,+f -float-abi=hard -code-model=small"
K_CFLAGS += "-v -O3 --sysroot=$(SYSROOT) --gcc-toolchain=$(RISCV_TOOLCHAIN_PATH) -march=rv32imf -mabi=ilp32f -I$(VORTEX_RT_PATH)/include -fno-rtti -fno-exceptions -ffreestanding -nostartfiles -fdata-sections -ffunction-sections"
K_LLCFLAGS += "-O3 -march=riscv32 -target-abi=ilp32f -mcpu=generic-rv32 -mattr=+m,+f -mattr=+vortex -float-abi=hard -code-model=small"
K_CFLAGS += "-v -O3 --sysroot=$(SYSROOT) --gcc-toolchain=$(RISCV_TOOLCHAIN_PATH) -march=rv32imf -mabi=ilp32f -Xclang -target-feature -Xclang +vortex -I$(VORTEX_RT_PATH)/include -fno-rtti -fno-exceptions -ffreestanding -nostartfiles -fdata-sections -ffunction-sections"
K_LDFLAGS += "-Wl,-Bstatic,-T$(VORTEX_RT_PATH)/linker/vx_link.ld -Wl,--gc-sections $(VORTEX_RT_PATH)/libvortexrt.a -lm"
CXXFLAGS += -std=c++11 -O2 -Wall -Wextra -Wfatal-errors
@@ -34,13 +34,13 @@ $(PROJECT): $(SRCS)
$(CXX) $(CXXFLAGS) $^ $(LDFLAGS) -o $@
run-fpga: $(PROJECT) kernel.pocl
LD_LIBRARY_PATH=$(POCL_RT_PATH)/lib:$(VORTEX_DRV_PATH)/opae:$(LD_LIBRARY_PATH) ./$(PROJECT)
LD_LIBRARY_PATH=$(POCL_RT_PATH)/lib:$(VORTEX_DRV_PATH)/fpga:$(LD_LIBRARY_PATH) ./$(PROJECT)
run-asesim: $(PROJECT) kernel.pocl
LD_LIBRARY_PATH=$(POCL_RT_PATH)/lib:$(VORTEX_DRV_PATH)/opae/ase:$(LD_LIBRARY_PATH) ./$(PROJECT)
LD_LIBRARY_PATH=$(POCL_RT_PATH)/lib:$(VORTEX_DRV_PATH)/asesim:$(LD_LIBRARY_PATH) ./$(PROJECT)
run-vlsim: $(PROJECT) kernel.pocl
LD_LIBRARY_PATH=$(POCL_RT_PATH)/lib:$(VORTEX_DRV_PATH)/opae/vlsim:$(LD_LIBRARY_PATH) ./$(PROJECT)
LD_LIBRARY_PATH=$(POCL_RT_PATH)/lib:$(VORTEX_DRV_PATH)/vlsim:$(LD_LIBRARY_PATH) ./$(PROJECT)
run-simx: $(PROJECT) kernel.pocl
LD_LIBRARY_PATH=$(POCL_RT_PATH)/lib:$(VORTEX_DRV_PATH)/simx:$(LD_LIBRARY_PATH) ./$(PROJECT)

Binary file not shown.

View File

@@ -7,8 +7,8 @@ POCL_RT_PATH ?= /opt/pocl/runtime
VORTEX_DRV_PATH ?= $(realpath ../../../driver)
VORTEX_RT_PATH ?= $(realpath ../../../runtime)
K_LLCFLAGS += "-O3 -march=riscv32 -target-abi=ilp32f -mcpu=generic-rv32 -mattr=+m,+f -float-abi=hard -code-model=small"
K_CFLAGS += "-v -O3 --sysroot=$(SYSROOT) --gcc-toolchain=$(RISCV_TOOLCHAIN_PATH) -march=rv32imf -mabi=ilp32f -I$(VORTEX_RT_PATH)/include -fno-rtti -fno-exceptions -ffreestanding -nostartfiles -fdata-sections -ffunction-sections"
K_LLCFLAGS += "-O3 -march=riscv32 -target-abi=ilp32f -mcpu=generic-rv32 -mattr=+m,+f -mattr=+vortex -float-abi=hard -code-model=small"
K_CFLAGS += "-v -O3 --sysroot=$(SYSROOT) --gcc-toolchain=$(RISCV_TOOLCHAIN_PATH) -march=rv32imf -mabi=ilp32f -Xclang -target-feature -Xclang +vortex -I$(VORTEX_RT_PATH)/include -fno-rtti -fno-exceptions -ffreestanding -nostartfiles -fdata-sections -ffunction-sections"
K_LDFLAGS += "-Wl,-Bstatic,-T$(VORTEX_RT_PATH)/linker/vx_link.ld -Wl,--gc-sections $(VORTEX_RT_PATH)/libvortexrt.a -lm"
CXXFLAGS += -std=c++11 -O2 -Wall -Wextra -pedantic -Wfatal-errors
@@ -31,13 +31,13 @@ $(PROJECT): $(SRCS)
$(CXX) $(CXXFLAGS) $^ $(LDFLAGS) -o $@
run-fpga: $(PROJECT) kernel.pocl
LD_LIBRARY_PATH=$(POCL_RT_PATH)/lib:$(VORTEX_DRV_PATH)/opae:$(LD_LIBRARY_PATH) ./$(PROJECT)
LD_LIBRARY_PATH=$(POCL_RT_PATH)/lib:$(VORTEX_DRV_PATH)/fpga:$(LD_LIBRARY_PATH) ./$(PROJECT)
run-asesim: $(PROJECT) kernel.pocl
LD_LIBRARY_PATH=$(POCL_RT_PATH)/lib:$(VORTEX_DRV_PATH)/opae/ase:$(LD_LIBRARY_PATH) ./$(PROJECT)
LD_LIBRARY_PATH=$(POCL_RT_PATH)/lib:$(VORTEX_DRV_PATH)/asesim:$(LD_LIBRARY_PATH) ./$(PROJECT)
run-vlsim: $(PROJECT) kernel.pocl
LD_LIBRARY_PATH=$(POCL_RT_PATH)/lib:$(VORTEX_DRV_PATH)/opae/vlsim:$(LD_LIBRARY_PATH) ./$(PROJECT)
LD_LIBRARY_PATH=$(POCL_RT_PATH)/lib:$(VORTEX_DRV_PATH)/vlsim:$(LD_LIBRARY_PATH) ./$(PROJECT)
run-simx: $(PROJECT) kernel.pocl
LD_LIBRARY_PATH=$(POCL_RT_PATH)/lib:$(VORTEX_DRV_PATH)/simx:$(LD_LIBRARY_PATH) ./$(PROJECT)

View File

@@ -9,8 +9,8 @@ OPTS ?= filelist.txt
VORTEX_DRV_PATH ?= $(realpath ../../../driver)
VORTEX_RT_PATH ?= $(realpath ../../../runtime)
K_LLCFLAGS += "-O3 -march=riscv32 -target-abi=ilp32f -mcpu=generic-rv32 -mattr=+m,+f -float-abi=hard -code-model=small"
K_CFLAGS += "-v -O3 --sysroot=$(SYSROOT) --gcc-toolchain=$(RISCV_TOOLCHAIN_PATH) -march=rv32imf -mabi=ilp32f -I$(VORTEX_RT_PATH)/include -fno-rtti -fno-exceptions -ffreestanding -nostartfiles -fdata-sections -ffunction-sections"
K_LLCFLAGS += "-O3 -march=riscv32 -target-abi=ilp32f -mcpu=generic-rv32 -mattr=+m,+f -mattr=+vortex -float-abi=hard -code-model=small"
K_CFLAGS += "-v -O3 --sysroot=$(SYSROOT) --gcc-toolchain=$(RISCV_TOOLCHAIN_PATH) -march=rv32imf -mabi=ilp32f -Xclang -target-feature -Xclang +vortex -I$(VORTEX_RT_PATH)/include -fno-rtti -fno-exceptions -ffreestanding -nostartfiles -fdata-sections -ffunction-sections"
K_LDFLAGS += "-Wl,-Bstatic,-T$(VORTEX_RT_PATH)/linker/vx_link.ld -Wl,--gc-sections $(VORTEX_RT_PATH)/libvortexrt.a -lm"
CXXFLAGS += -std=c++11 -O2 -Wall -Wextra -Wfatal-errors
@@ -36,13 +36,13 @@ $(PROJECT): $(SRCS)
$(CXX) $(CXXFLAGS) $^ $(LDFLAGS) -o $@
run-fpga: $(PROJECT) kernel.pocl
LD_LIBRARY_PATH=$(POCL_RT_PATH)/lib:$(VORTEX_DRV_PATH)/opae:$(LD_LIBRARY_PATH) ./$(PROJECT) $(OPTS)
LD_LIBRARY_PATH=$(POCL_RT_PATH)/lib:$(VORTEX_DRV_PATH)/fpga:$(LD_LIBRARY_PATH) ./$(PROJECT) $(OPTS)
run-asesim: $(PROJECT) kernel.pocl
LD_LIBRARY_PATH=$(POCL_RT_PATH)/lib:$(VORTEX_DRV_PATH)/opae/ase:$(LD_LIBRARY_PATH) ./$(PROJECT) $(OPTS)
LD_LIBRARY_PATH=$(POCL_RT_PATH)/lib:$(VORTEX_DRV_PATH)/asesim:$(LD_LIBRARY_PATH) ./$(PROJECT) $(OPTS)
run-vlsim: $(PROJECT) kernel.pocl
LD_LIBRARY_PATH=$(POCL_RT_PATH)/lib:$(VORTEX_DRV_PATH)/opae/vlsim:$(LD_LIBRARY_PATH) ./$(PROJECT) $(OPTS)
LD_LIBRARY_PATH=$(POCL_RT_PATH)/lib:$(VORTEX_DRV_PATH)/vlsim:$(LD_LIBRARY_PATH) ./$(PROJECT) $(OPTS)
run-simx: $(PROJECT) kernel.pocl
LD_LIBRARY_PATH=$(POCL_RT_PATH)/lib:$(VORTEX_DRV_PATH)/simx:$(LD_LIBRARY_PATH) ./$(PROJECT) $(OPTS)

File diff suppressed because it is too large Load Diff

Binary file not shown.

View File

@@ -9,8 +9,8 @@ OPTS ?= -n1
VORTEX_DRV_PATH ?= $(realpath ../../../driver)
VORTEX_RT_PATH ?= $(realpath ../../../runtime)
K_LLCFLAGS += "-O3 -march=riscv32 -target-abi=ilp32f -mcpu=generic-rv32 -mattr=+m,+f -float-abi=hard -code-model=small"
K_CFLAGS += "-v -O3 --sysroot=$(SYSROOT) --gcc-toolchain=$(RISCV_TOOLCHAIN_PATH) -march=rv32imf -mabi=ilp32f -I$(VORTEX_RT_PATH)/include -fno-rtti -fno-exceptions -ffreestanding -nostartfiles -fdata-sections -ffunction-sections"
K_LLCFLAGS += "-O3 -march=riscv32 -target-abi=ilp32f -mcpu=generic-rv32 -mattr=+m,+f -mattr=+vortex -float-abi=hard -code-model=small"
K_CFLAGS += "-v -O3 --sysroot=$(SYSROOT) --gcc-toolchain=$(RISCV_TOOLCHAIN_PATH) -march=rv32imf -mabi=ilp32f -Xclang -target-feature -Xclang +vortex -I$(VORTEX_RT_PATH)/include -fno-rtti -fno-exceptions -ffreestanding -nostartfiles -fdata-sections -ffunction-sections"
K_LDFLAGS += "-Wl,-Bstatic,-T$(VORTEX_RT_PATH)/linker/vx_link.ld -Wl,--gc-sections $(VORTEX_RT_PATH)/libvortexrt.a -lm"
CXXFLAGS += -std=c++11 -O2 -Wall -Wextra -Wfatal-errors
@@ -35,13 +35,13 @@ $(PROJECT): $(SRCS)
$(CXX) $(CXXFLAGS) $^ $(LDFLAGS) -o $@
run-fpga: $(PROJECT) kernel.pocl
LD_LIBRARY_PATH=$(POCL_RT_PATH)/lib:$(VORTEX_DRV_PATH)/opae:$(LD_LIBRARY_PATH) ./$(PROJECT) $(OPTS)
LD_LIBRARY_PATH=$(POCL_RT_PATH)/lib:$(VORTEX_DRV_PATH)/fpga:$(LD_LIBRARY_PATH) ./$(PROJECT) $(OPTS)
run-asesim: $(PROJECT) kernel.pocl
LD_LIBRARY_PATH=$(POCL_RT_PATH)/lib:$(VORTEX_DRV_PATH)/opae/ase:$(LD_LIBRARY_PATH) ./$(PROJECT) $(OPTS)
LD_LIBRARY_PATH=$(POCL_RT_PATH)/lib:$(VORTEX_DRV_PATH)/asesim:$(LD_LIBRARY_PATH) ./$(PROJECT) $(OPTS)
run-vlsim: $(PROJECT) kernel.pocl
LD_LIBRARY_PATH=$(POCL_RT_PATH)/lib:$(VORTEX_DRV_PATH)/opae/vlsim:$(LD_LIBRARY_PATH) ./$(PROJECT) $(OPTS)
LD_LIBRARY_PATH=$(POCL_RT_PATH)/lib:$(VORTEX_DRV_PATH)/vlsim:$(LD_LIBRARY_PATH) ./$(PROJECT) $(OPTS)
run-simx: $(PROJECT) kernel.pocl
LD_LIBRARY_PATH=$(POCL_RT_PATH)/lib:$(VORTEX_DRV_PATH)/simx:$(LD_LIBRARY_PATH) ./$(PROJECT) $(OPTS)

View File

@@ -1,6 +1,6 @@
__kernel void test_printf (__global const int *A)
__kernel void oclprintf (__global const int *A)
{
int gid = get_global_id(0);
int value = A[gid];
printf("Print Test! value[%d]=%d\n", gid, value);
}
}

View File

@@ -7,7 +7,7 @@
#include <string.h>
#include <chrono>
#define KERNEL_NAME "test_printf"
#define KERNEL_NAME "oclprintf"
#define CL_CHECK(_expr) \
do { \

File diff suppressed because it is too large Load Diff

View File

@@ -4,13 +4,13 @@ SYSROOT ?= $(RISCV_TOOLCHAIN_PATH)/riscv32-unknown-elf
POCL_CC_PATH ?= /opt/pocl/compiler
POCL_RT_PATH ?= /opt/pocl/runtime
OPTS ?= -n32
OPTS ?= -f -n16
VORTEX_DRV_PATH ?= $(realpath ../../../driver)
VORTEX_RT_PATH ?= $(realpath ../../../runtime)
K_LLCFLAGS += "-O3 -march=riscv32 -target-abi=ilp32f -mcpu=generic-rv32 -mattr=+m,+f -float-abi=hard -code-model=small"
K_CFLAGS += "-v -O3 --sysroot=$(SYSROOT) --gcc-toolchain=$(RISCV_TOOLCHAIN_PATH) -march=rv32imf -mabi=ilp32f -I$(VORTEX_RT_PATH)/include -fno-rtti -fno-exceptions -ffreestanding -nostartfiles -fdata-sections -ffunction-sections"
K_LLCFLAGS += "-O3 -march=riscv32 -target-abi=ilp32f -mcpu=generic-rv32 -mattr=+m,+f -mattr=+vortex -float-abi=hard -code-model=small"
K_CFLAGS += "-v -O3 --sysroot=$(SYSROOT) --gcc-toolchain=$(RISCV_TOOLCHAIN_PATH) -march=rv32imf -mabi=ilp32f -Xclang -target-feature -Xclang +vortex -I$(VORTEX_RT_PATH)/include -fno-rtti -fno-exceptions -ffreestanding -nostartfiles -fdata-sections -ffunction-sections"
K_LDFLAGS += "-Wl,-Bstatic,-T$(VORTEX_RT_PATH)/linker/vx_link.ld -Wl,--gc-sections $(VORTEX_RT_PATH)/libvortexrt.a -lm"
CXXFLAGS += -std=c++11 -O2 -Wall -Wextra -Wfatal-errors
@@ -35,13 +35,13 @@ $(PROJECT): $(SRCS)
$(CXX) $(CXXFLAGS) $^ $(LDFLAGS) -o $@
run-fpga: $(PROJECT) kernel.pocl
LD_LIBRARY_PATH=$(POCL_RT_PATH)/lib:$(VORTEX_DRV_PATH)/opae:$(LD_LIBRARY_PATH) ./$(PROJECT) $(OPTS)
LD_LIBRARY_PATH=$(POCL_RT_PATH)/lib:$(VORTEX_DRV_PATH)/fpga:$(LD_LIBRARY_PATH) ./$(PROJECT) $(OPTS)
run-asesim: $(PROJECT) kernel.pocl
LD_LIBRARY_PATH=$(POCL_RT_PATH)/lib:$(VORTEX_DRV_PATH)/opae/ase:$(LD_LIBRARY_PATH) ./$(PROJECT) $(OPTS)
LD_LIBRARY_PATH=$(POCL_RT_PATH)/lib:$(VORTEX_DRV_PATH)/asesim:$(LD_LIBRARY_PATH) ./$(PROJECT) $(OPTS)
run-vlsim: $(PROJECT) kernel.pocl
LD_LIBRARY_PATH=$(POCL_RT_PATH)/lib:$(VORTEX_DRV_PATH)/opae/vlsim:$(LD_LIBRARY_PATH) ./$(PROJECT) $(OPTS)
LD_LIBRARY_PATH=$(POCL_RT_PATH)/lib:$(VORTEX_DRV_PATH)/vlsim:$(LD_LIBRARY_PATH) ./$(PROJECT) $(OPTS)
run-simx: $(PROJECT) kernel.pocl
LD_LIBRARY_PATH=$(POCL_RT_PATH)/lib:$(VORTEX_DRV_PATH)/simx:$(LD_LIBRARY_PATH) ./$(PROJECT) $(OPTS)

View File

@@ -1,4 +1,19 @@
__kernel void psort (__global const float *in, __global float *out)
__kernel void psorti (__global const int *in, __global int *out)
{
int gid = get_global_id(0);
int n = get_global_size(0);
int ref = in[gid];
int pos = 0;
for (int i = 0; i < n; ++i) {
int cur = in[i];
pos += (cur < ref) || ((cur == ref) && (i < gid));
}
out[pos] = ref;
}
__kernel void psortf (__global const float *in, __global float *out)
{
int gid = get_global_id(0);
int n = get_global_size(0);
@@ -8,7 +23,13 @@ __kernel void psort (__global const float *in, __global float *out)
int pos = 0;
for (int i = 0; i < n; ++i) {
float cur = in[i];
pos += (cur < ref) || (cur == ref && i < gid);
pos += (cur < ref) || ((cur == ref) && (i < gid));
/*int cl = (cur < ref);
int ce = (cur == ref);
int ls = (i < gid);
int x = ce && ls;
int y = cl || x;
pos += y;*/
}
out[pos] = ref;
}

Binary file not shown.

View File

@@ -7,7 +7,8 @@
#include <string.h>
#include <chrono>
#define KERNEL_NAME "psort"
#define KERNEL0_NAME "psorti"
#define KERNEL1_NAME "psortf"
#define CL_CHECK(_expr) \
do { \
@@ -52,14 +53,6 @@ static int read_kernel_file(const char* filename, uint8_t** data, size_t* size)
return 0;
}
static bool almost_equal(float a, float b, int ulp = 4) {
union fi_t { int i; float f; };
fi_t fa, fb;
fa.f = a;
fb.f = b;
return std::abs(fa.i - fb.i) <= ulp;
}
cl_device_id device_id = NULL;
cl_context context = NULL;
cl_command_queue commandQueue = NULL;
@@ -67,8 +60,8 @@ cl_program program = NULL;
cl_kernel kernel = NULL;
cl_mem a_memobj = NULL;
cl_mem c_memobj = NULL;
float *h_a = NULL;
float *h_c = NULL;
int *h_a = NULL;
int *h_c = NULL;
uint8_t *kernel_bin = NULL;
static void cleanup() {
@@ -86,15 +79,19 @@ static void cleanup() {
}
int size = 64;
bool float_enable = false;
static void show_usage() {
printf("Usage: [-n size] [-h: help]\n");
printf("Usage: [-f] [-n size] [-h: help]\n");
}
static void parse_args(int argc, char **argv) {
int c;
while ((c = getopt(argc, argv, "n:h?")) != -1) {
while ((c = getopt(argc, argv, "fn:h?")) != -1) {
switch (c) {
case 'f':
float_enable = 1;
break;
case 'n':
size = atoi(optarg);
break;
@@ -132,7 +129,7 @@ int main (int argc, char **argv) {
context = CL_CHECK2(clCreateContext(NULL, 1, &device_id, NULL, NULL, &_err));
printf("Allocate device buffers\n");
size_t nbytes = size * sizeof(float);
size_t nbytes = size * sizeof(int);
a_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));
@@ -148,21 +145,28 @@ int main (int argc, char **argv) {
CL_CHECK(clBuildProgram(program, 1, &device_id, NULL, NULL, NULL));
// Create kernel
kernel = CL_CHECK2(clCreateKernel(program, KERNEL_NAME, &_err));
kernel = CL_CHECK2(clCreateKernel(program, (float_enable ? KERNEL1_NAME : KERNEL0_NAME), &_err));
// Set kernel arguments
CL_CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *)&a_memobj));
CL_CHECK(clSetKernelArg(kernel, 1, sizeof(cl_mem), (void *)&c_memobj));
// Allocate memories for input arrays and output arrays.
h_a = (float*)malloc(nbytes);
h_c = (float*)malloc(nbytes);
h_a = (int*)malloc(nbytes);
h_c = (int*)malloc(nbytes);
// Initialize values for array members.
for (int i = 0; i < size; ++i) {
h_a[i] = sinf(i)*sinf(i);
h_c[i] = 0xdeadbeef;
printf("*** [%d]: h_a=%f\n", i, h_a[i]);
if (float_enable) {
float value = sinf(i)*sinf(i);
h_a[i] = *(int*)&value;
printf("*** [%d]: h_a=%f\n", i, value);
} else {
int value = size*sinf(i);
h_a[i] = value;
printf("*** [%d]: h_a=%d\n", i, value);
}
}
// Creating command queue
@@ -185,17 +189,37 @@ int main (int argc, char **argv) {
CL_CHECK(clEnqueueReadBuffer(commandQueue, c_memobj, CL_TRUE, 0, nbytes, h_c, 0, NULL, NULL));
printf("Verify result\n");
for (int i = 0; i < size; ++i) {
int value = h_c[i];
if (float_enable) {
printf("*** [%d]: h_a=%f\n", i, *(float*)&value);
} else {
printf("*** [%d]: h_a=%d\n", i, value);
}
}
int errors = 0;
for (int i = 0; i < size; ++i) {
float ref = h_a[i];
int ref = h_a[i];
float ref_f = *(float*)&ref;
int pos = 0;
for (int j = 0; j < size; ++j) {
float cur = h_a[j];
pos += (cur < ref) || (cur == ref && j < i);
int cur = h_a[j];
if (float_enable) {
float cur_f = *(float*)&cur;
pos += (cur_f < ref_f) || (cur_f == ref_f && j < i);
} else {
pos += (cur < ref) || (cur == ref && j < i);
}
}
if (!almost_equal(h_c[pos], ref)) {
if (errors < 100)
printf("*** error: [%d] expected=%f, actual=%f\n", pos, ref, h_c[pos]);
int value = h_c[pos];
if (value != ref) {
if (errors < 100) {
if (float_enable) {
printf("*** error: [%d] expected=%f, actual=%f\n", pos, ref_f, *(float*)&value);
} else {
printf("*** error: [%d] expected=%d, actual=%d\n", pos, ref, value);
}
}
++errors;
}
}

File diff suppressed because it is too large Load Diff

File diff suppressed because it is too large Load Diff

File diff suppressed because it is too large Load Diff

View File

@@ -9,8 +9,8 @@ VORTEX_RT_PATH ?= $(realpath ../../../runtime)
OPTS ?= -n1024
K_LLCFLAGS += "-O3 -march=riscv32 -target-abi=ilp32f -mcpu=generic-rv32 -mattr=+m,+f -float-abi=hard -code-model=small"
K_CFLAGS += "-v -O3 --sysroot=$(SYSROOT) --gcc-toolchain=$(RISCV_TOOLCHAIN_PATH) -march=rv32imf -mabi=ilp32f -I$(VORTEX_RT_PATH)/include -fno-rtti -fno-exceptions -ffreestanding -nostartfiles -fdata-sections -ffunction-sections"
K_LLCFLAGS += "-O3 -march=riscv32 -target-abi=ilp32f -mcpu=generic-rv32 -mattr=+m,+f -mattr=+vortex -float-abi=hard -code-model=small"
K_CFLAGS += "-v -O3 --sysroot=$(SYSROOT) --gcc-toolchain=$(RISCV_TOOLCHAIN_PATH) -march=rv32imf -mabi=ilp32f -Xclang -target-feature -Xclang +vortex -I$(VORTEX_RT_PATH)/include -fno-rtti -fno-exceptions -ffreestanding -nostartfiles -fdata-sections -ffunction-sections"
K_LDFLAGS += "-Wl,-Bstatic,-T$(VORTEX_RT_PATH)/linker/vx_link.ld -Wl,--gc-sections $(VORTEX_RT_PATH)/libvortexrt.a -lm"
CXXFLAGS += -std=c++11 -O2 -Wall -Wextra -Wfatal-errors
@@ -35,13 +35,13 @@ $(PROJECT): $(SRCS)
$(CXX) $(CXXFLAGS) $^ $(LDFLAGS) -o $@
run-fpga: $(PROJECT) kernel.pocl
LD_LIBRARY_PATH=$(POCL_RT_PATH)/lib:$(VORTEX_DRV_PATH)/opae:$(LD_LIBRARY_PATH) ./$(PROJECT) $(OPTS)
LD_LIBRARY_PATH=$(POCL_RT_PATH)/lib:$(VORTEX_DRV_PATH)/fpga:$(LD_LIBRARY_PATH) ./$(PROJECT) $(OPTS)
run-asesim: $(PROJECT) kernel.pocl
LD_LIBRARY_PATH=$(POCL_RT_PATH)/lib:$(VORTEX_DRV_PATH)/opae/ase:$(LD_LIBRARY_PATH) ./$(PROJECT) $(OPTS)
LD_LIBRARY_PATH=$(POCL_RT_PATH)/lib:$(VORTEX_DRV_PATH)/asesim:$(LD_LIBRARY_PATH) ./$(PROJECT) $(OPTS)
run-vlsim: $(PROJECT) kernel.pocl
LD_LIBRARY_PATH=$(POCL_RT_PATH)/lib:$(VORTEX_DRV_PATH)/opae/vlsim:$(LD_LIBRARY_PATH) ./$(PROJECT) $(OPTS)
LD_LIBRARY_PATH=$(POCL_RT_PATH)/lib:$(VORTEX_DRV_PATH)/vlsim:$(LD_LIBRARY_PATH) ./$(PROJECT) $(OPTS)
run-simx: $(PROJECT) kernel.pocl
LD_LIBRARY_PATH=$(POCL_RT_PATH)/lib:$(VORTEX_DRV_PATH)/simx:$(LD_LIBRARY_PATH) ./$(PROJECT) $(OPTS)

Binary file not shown.

File diff suppressed because it is too large Load Diff

View File

@@ -9,8 +9,8 @@ VORTEX_RT_PATH ?= $(realpath ../../../runtime)
OPTS ?= -n16
K_LLCFLAGS += "-O3 -march=riscv32 -target-abi=ilp32f -mcpu=generic-rv32 -mattr=+m,+f -float-abi=hard -code-model=small"
K_CFLAGS += "-v -O3 --sysroot=$(SYSROOT) --gcc-toolchain=$(RISCV_TOOLCHAIN_PATH) -march=rv32imf -mabi=ilp32f -I$(VORTEX_RT_PATH)/include -fno-rtti -fno-exceptions -ffreestanding -nostartfiles -fdata-sections -ffunction-sections"
K_LLCFLAGS += "-O3 -march=riscv32 -target-abi=ilp32f -mcpu=generic-rv32 -mattr=+m,+f -mattr=+vortex -float-abi=hard -code-model=small"
K_CFLAGS += "-v -O3 --sysroot=$(SYSROOT) --gcc-toolchain=$(RISCV_TOOLCHAIN_PATH) -march=rv32imf -mabi=ilp32f -Xclang -target-feature -Xclang +vortex -I$(VORTEX_RT_PATH)/include -fno-rtti -fno-exceptions -ffreestanding -nostartfiles -fdata-sections -ffunction-sections"
K_LDFLAGS += "-Wl,-Bstatic,-T$(VORTEX_RT_PATH)/linker/vx_link.ld -Wl,--gc-sections $(VORTEX_RT_PATH)/libvortexrt.a -lm"
CXXFLAGS += -std=c++11 -O2 -Wall -Wextra -Wfatal-errors
@@ -35,13 +35,13 @@ $(PROJECT): $(SRCS)
$(CXX) $(CXXFLAGS) $^ $(LDFLAGS) -o $@
run-fpga: $(PROJECT) kernel.pocl
LD_LIBRARY_PATH=$(POCL_RT_PATH)/lib:$(VORTEX_DRV_PATH)/opae:$(LD_LIBRARY_PATH) ./$(PROJECT) $(OPTS)
LD_LIBRARY_PATH=$(POCL_RT_PATH)/lib:$(VORTEX_DRV_PATH)/fpga:$(LD_LIBRARY_PATH) ./$(PROJECT) $(OPTS)
run-asesim: $(PROJECT) kernel.pocl
LD_LIBRARY_PATH=$(POCL_RT_PATH)/lib:$(VORTEX_DRV_PATH)/opae/ase:$(LD_LIBRARY_PATH) ./$(PROJECT) $(OPTS)
LD_LIBRARY_PATH=$(POCL_RT_PATH)/lib:$(VORTEX_DRV_PATH)/asesim:$(LD_LIBRARY_PATH) ./$(PROJECT) $(OPTS)
run-vlsim: $(PROJECT) kernel.pocl
LD_LIBRARY_PATH=$(POCL_RT_PATH)/lib:$(VORTEX_DRV_PATH)/opae/vlsim:$(LD_LIBRARY_PATH) ./$(PROJECT) $(OPTS)
LD_LIBRARY_PATH=$(POCL_RT_PATH)/lib:$(VORTEX_DRV_PATH)/vlsim:$(LD_LIBRARY_PATH) ./$(PROJECT) $(OPTS)
run-simx: $(PROJECT) kernel.pocl
LD_LIBRARY_PATH=$(POCL_RT_PATH)/lib:$(VORTEX_DRV_PATH)/simx:$(LD_LIBRARY_PATH) ./$(PROJECT) $(OPTS)

Binary file not shown.

File diff suppressed because it is too large Load Diff

View File

@@ -9,8 +9,8 @@ OPTS ?= -n32
VORTEX_DRV_PATH ?= $(realpath ../../../driver)
VORTEX_RT_PATH ?= $(realpath ../../../runtime)
K_LLCFLAGS += "-O3 -march=riscv32 -target-abi=ilp32f -mcpu=generic-rv32 -mattr=+m,+f -float-abi=hard -code-model=small"
K_CFLAGS += "-v -O3 --sysroot=$(SYSROOT) --gcc-toolchain=$(RISCV_TOOLCHAIN_PATH) -march=rv32imf -mabi=ilp32f -I$(VORTEX_RT_PATH)/include -fno-rtti -fno-exceptions -ffreestanding -nostartfiles -fdata-sections -ffunction-sections"
K_LLCFLAGS += "-O3 -march=riscv32 -target-abi=ilp32f -mcpu=generic-rv32 -mattr=+m,+f -mattr=+vortex -float-abi=hard -code-model=small"
K_CFLAGS += "-v -O3 --sysroot=$(SYSROOT) --gcc-toolchain=$(RISCV_TOOLCHAIN_PATH) -march=rv32imf -mabi=ilp32f -Xclang -target-feature -Xclang +vortex -I$(VORTEX_RT_PATH)/include -fno-rtti -fno-exceptions -ffreestanding -nostartfiles -fdata-sections -ffunction-sections"
K_LDFLAGS += "-Wl,-Bstatic,-T$(VORTEX_RT_PATH)/linker/vx_link.ld -Wl,--gc-sections $(VORTEX_RT_PATH)/libvortexrt.a -lm"
CXXFLAGS += -std=c++11 -O2 -Wall -Wextra -Wfatal-errors
@@ -35,13 +35,13 @@ $(PROJECT): $(SRCS)
$(CXX) $(CXXFLAGS) $^ $(LDFLAGS) -o $@
run-fpga: $(PROJECT) kernel.pocl
LD_LIBRARY_PATH=$(POCL_RT_PATH)/lib:$(VORTEX_DRV_PATH)/opae:$(LD_LIBRARY_PATH) ./$(PROJECT) $(OPTS)
LD_LIBRARY_PATH=$(POCL_RT_PATH)/lib:$(VORTEX_DRV_PATH)/fpga:$(LD_LIBRARY_PATH) ./$(PROJECT) $(OPTS)
run-asesim: $(PROJECT) kernel.pocl
LD_LIBRARY_PATH=$(POCL_RT_PATH)/lib:$(VORTEX_DRV_PATH)/opae/ase:$(LD_LIBRARY_PATH) ./$(PROJECT) $(OPTS)
LD_LIBRARY_PATH=$(POCL_RT_PATH)/lib:$(VORTEX_DRV_PATH)/asesim:$(LD_LIBRARY_PATH) ./$(PROJECT) $(OPTS)
run-vlsim: $(PROJECT) kernel.pocl
LD_LIBRARY_PATH=$(POCL_RT_PATH)/lib:$(VORTEX_DRV_PATH)/opae/vlsim:$(LD_LIBRARY_PATH) ./$(PROJECT) $(OPTS)
LD_LIBRARY_PATH=$(POCL_RT_PATH)/lib:$(VORTEX_DRV_PATH)/vlsim:$(LD_LIBRARY_PATH) ./$(PROJECT) $(OPTS)
run-simx: $(PROJECT) kernel.pocl
LD_LIBRARY_PATH=$(POCL_RT_PATH)/lib:$(VORTEX_DRV_PATH)/simx:$(LD_LIBRARY_PATH) ./$(PROJECT) $(OPTS)

Binary file not shown.

File diff suppressed because it is too large Load Diff

View File

@@ -7,8 +7,8 @@ POCL_RT_PATH ?= /opt/pocl/runtime
VORTEX_DRV_PATH ?= $(realpath ../../../driver)
VORTEX_RT_PATH ?= $(realpath ../../../runtime)
K_LLCFLAGS += "-O3 -march=riscv32 -target-abi=ilp32f -mcpu=generic-rv32 -mattr=+m,+f -float-abi=hard -code-model=small"
K_CFLAGS += "-v -O3 --sysroot=$(SYSROOT) --gcc-toolchain=$(RISCV_TOOLCHAIN_PATH) -march=rv32imf -mabi=ilp32f -I$(VORTEX_RT_PATH)/include -fno-rtti -fno-exceptions -ffreestanding -nostartfiles -fdata-sections -ffunction-sections"
K_LLCFLAGS += "-O3 -march=riscv32 -target-abi=ilp32f -mcpu=generic-rv32 -mattr=+m,+f -mattr=+vortex -float-abi=hard -code-model=small"
K_CFLAGS += "-v -O3 --sysroot=$(SYSROOT) --gcc-toolchain=$(RISCV_TOOLCHAIN_PATH) -march=rv32imf -mabi=ilp32f -Xclang -target-feature -Xclang +vortex -I$(VORTEX_RT_PATH)/include -fno-rtti -fno-exceptions -ffreestanding -nostartfiles -fdata-sections -ffunction-sections"
K_LDFLAGS += "-Wl,-Bstatic,-T$(VORTEX_RT_PATH)/linker/vx_link.ld -Wl,--gc-sections $(VORTEX_RT_PATH)/libvortexrt.a -lm"
CXXFLAGS += -std=c++11 -O2 -Wall -Wextra -pedantic -Wfatal-errors
@@ -31,13 +31,13 @@ $(PROJECT): $(SRCS)
$(CXX) $(CXXFLAGS) $^ $(LDFLAGS) -o $@
run-fpga: $(PROJECT) kernel.pocl
LD_LIBRARY_PATH=$(POCL_RT_PATH)/lib:$(VORTEX_DRV_PATH)/opae:$(LD_LIBRARY_PATH) ./$(PROJECT)
LD_LIBRARY_PATH=$(POCL_RT_PATH)/lib:$(VORTEX_DRV_PATH)/fpga:$(LD_LIBRARY_PATH) ./$(PROJECT)
run-asesim: $(PROJECT) kernel.pocl
LD_LIBRARY_PATH=$(POCL_RT_PATH)/lib:$(VORTEX_DRV_PATH)/opae/ase:$(LD_LIBRARY_PATH) ./$(PROJECT)
LD_LIBRARY_PATH=$(POCL_RT_PATH)/lib:$(VORTEX_DRV_PATH)/asesim:$(LD_LIBRARY_PATH) ./$(PROJECT)
run-vlsim: $(PROJECT) kernel.pocl
LD_LIBRARY_PATH=$(POCL_RT_PATH)/lib:$(VORTEX_DRV_PATH)/opae/vlsim:$(LD_LIBRARY_PATH) ./$(PROJECT)
LD_LIBRARY_PATH=$(POCL_RT_PATH)/lib:$(VORTEX_DRV_PATH)/vlsim:$(LD_LIBRARY_PATH) ./$(PROJECT)
run-simx: $(PROJECT) kernel.pocl
LD_LIBRARY_PATH=$(POCL_RT_PATH)/lib:$(VORTEX_DRV_PATH)/simx:$(LD_LIBRARY_PATH) ./$(PROJECT)

View File

@@ -9,8 +9,8 @@ OPTS ?= -n64
VORTEX_DRV_PATH ?= $(realpath ../../../driver)
VORTEX_RT_PATH ?= $(realpath ../../../runtime)
K_LLCFLAGS += "-O3 -march=riscv32 -target-abi=ilp32f -mcpu=generic-rv32 -mattr=+m,+f -float-abi=hard -code-model=small"
K_CFLAGS += "-v -O3 --sysroot=$(SYSROOT) --gcc-toolchain=$(RISCV_TOOLCHAIN_PATH) -march=rv32imf -mabi=ilp32f -I$(VORTEX_RT_PATH)/include -fno-rtti -fno-exceptions -ffreestanding -nostartfiles -fdata-sections -ffunction-sections"
K_LLCFLAGS += "-O3 -march=riscv32 -target-abi=ilp32f -mcpu=generic-rv32 -mattr=+m,+f -mattr=+vortex -float-abi=hard -code-model=small"
K_CFLAGS += "-v -O3 --sysroot=$(SYSROOT) --gcc-toolchain=$(RISCV_TOOLCHAIN_PATH) -march=rv32imf -mabi=ilp32f -Xclang -target-feature -Xclang +vortex -I$(VORTEX_RT_PATH)/include -fno-rtti -fno-exceptions -ffreestanding -nostartfiles -fdata-sections -ffunction-sections"
K_LDFLAGS += "-Wl,-Bstatic,-T$(VORTEX_RT_PATH)/linker/vx_link.ld -Wl,--gc-sections $(VORTEX_RT_PATH)/libvortexrt.a -lm"
CXXFLAGS += -std=c++11 -O2 -Wall -Wextra -Wfatal-errors
@@ -35,13 +35,13 @@ $(PROJECT): $(SRCS)
$(CXX) $(CXXFLAGS) $^ $(LDFLAGS) -o $@
run-fpga: $(PROJECT) kernel.pocl
LD_LIBRARY_PATH=$(POCL_RT_PATH)/lib:$(VORTEX_DRV_PATH)/opae:$(LD_LIBRARY_PATH) ./$(PROJECT) $(OPTS)
LD_LIBRARY_PATH=$(POCL_RT_PATH)/lib:$(VORTEX_DRV_PATH)/fpga:$(LD_LIBRARY_PATH) ./$(PROJECT) $(OPTS)
run-asesim: $(PROJECT) kernel.pocl
LD_LIBRARY_PATH=$(POCL_RT_PATH)/lib:$(VORTEX_DRV_PATH)/opae/ase:$(LD_LIBRARY_PATH) ./$(PROJECT) $(OPTS)
LD_LIBRARY_PATH=$(POCL_RT_PATH)/lib:$(VORTEX_DRV_PATH)/asesim:$(LD_LIBRARY_PATH) ./$(PROJECT) $(OPTS)
run-vlsim: $(PROJECT) kernel.pocl
LD_LIBRARY_PATH=$(POCL_RT_PATH)/lib:$(VORTEX_DRV_PATH)/opae/vlsim:$(LD_LIBRARY_PATH) ./$(PROJECT) $(OPTS)
LD_LIBRARY_PATH=$(POCL_RT_PATH)/lib:$(VORTEX_DRV_PATH)/vlsim:$(LD_LIBRARY_PATH) ./$(PROJECT) $(OPTS)
run-simx: $(PROJECT) kernel.pocl
LD_LIBRARY_PATH=$(POCL_RT_PATH)/lib:$(VORTEX_DRV_PATH)/simx:$(LD_LIBRARY_PATH) ./$(PROJECT) $(OPTS)

Binary file not shown.

File diff suppressed because it is too large Load Diff

View File

@@ -45,13 +45,13 @@ run-simx: $(PROJECT) kernel.bin
LD_LIBRARY_PATH=$(POCL_RT_PATH)/lib:$(VORTEX_DRV_PATH)/simx:$(LD_LIBRARY_PATH) ./$(PROJECT) $(OPTS)
run-fpga: $(PROJECT) kernel.bin
LD_LIBRARY_PATH=$(POCL_RT_PATH)/lib:$(VORTEX_DRV_PATH)/opae:$(LD_LIBRARY_PATH) ./$(PROJECT) $(OPTS)
LD_LIBRARY_PATH=$(POCL_RT_PATH)/lib:$(VORTEX_DRV_PATH)/fpga:$(LD_LIBRARY_PATH) ./$(PROJECT) $(OPTS)
run-asesim: $(PROJECT) kernel.bin
LD_LIBRARY_PATH=$(POCL_RT_PATH)/lib:$(VORTEX_DRV_PATH)/opae/ase:$(LD_LIBRARY_PATH) ./$(PROJECT) $(OPTS)
LD_LIBRARY_PATH=$(POCL_RT_PATH)/lib:$(VORTEX_DRV_PATH)/asesim:$(LD_LIBRARY_PATH) ./$(PROJECT) $(OPTS)
run-vlsim: $(PROJECT) kernel.bin
LD_LIBRARY_PATH=$(POCL_RT_PATH)/lib:$(VORTEX_DRV_PATH)/opae/vlsim:$(LD_LIBRARY_PATH) ./$(PROJECT) $(OPTS)
LD_LIBRARY_PATH=$(POCL_RT_PATH)/lib:$(VORTEX_DRV_PATH)/vlsim:$(LD_LIBRARY_PATH) ./$(PROJECT) $(OPTS)
run-rtlsim: $(PROJECT) kernel.bin
LD_LIBRARY_PATH=$(POCL_RT_PATH)/lib:$(VORTEX_DRV_PATH)/rtlsim:$(LD_LIBRARY_PATH) ./$(PROJECT) $(OPTS)

View File

@@ -3,10 +3,10 @@
#define KERNEL_ARG_DEV_MEM_ADDR 0x7ffff000
struct kernel_arg_t {
typedef struct {
uint32_t count;
uint32_t src_ptr;
uint32_t dst_ptr;
};
} kernel_arg_t;
#endif

View File

@@ -3,7 +3,7 @@
#include "common.h"
void main() {
struct kernel_arg_t* arg = (struct kernel_arg_t*)KERNEL_ARG_DEV_MEM_ADDR;
const kernel_arg_t* arg = (const kernel_arg_t*)KERNEL_ARG_DEV_MEM_ADDR;
uint32_t count = arg->count;
int32_t* src_ptr = (int32_t*)arg->src_ptr;
int32_t* dst_ptr = (int32_t*)arg->dst_ptr;

View File

@@ -45,13 +45,13 @@ run-simx: $(PROJECT) kernel.bin
LD_LIBRARY_PATH=$(POCL_RT_PATH)/lib:$(VORTEX_DRV_PATH)/simx:$(LD_LIBRARY_PATH) ./$(PROJECT) $(OPTS)
run-fpga: $(PROJECT) kernel.bin
LD_LIBRARY_PATH=$(POCL_RT_PATH)/lib:$(VORTEX_DRV_PATH)/opae:$(LD_LIBRARY_PATH) ./$(PROJECT) $(OPTS)
LD_LIBRARY_PATH=$(POCL_RT_PATH)/lib:$(VORTEX_DRV_PATH)/fpga:$(LD_LIBRARY_PATH) ./$(PROJECT) $(OPTS)
run-asesim: $(PROJECT) kernel.bin
LD_LIBRARY_PATH=$(POCL_RT_PATH)/lib:$(VORTEX_DRV_PATH)/opae/ase:$(LD_LIBRARY_PATH) ./$(PROJECT) $(OPTS)
LD_LIBRARY_PATH=$(POCL_RT_PATH)/lib:$(VORTEX_DRV_PATH)/asesim:$(LD_LIBRARY_PATH) ./$(PROJECT) $(OPTS)
run-vlsim: $(PROJECT) kernel.bin
LD_LIBRARY_PATH=$(POCL_RT_PATH)/lib:$(VORTEX_DRV_PATH)/opae/vlsim:$(LD_LIBRARY_PATH) ./$(PROJECT) $(OPTS)
LD_LIBRARY_PATH=$(POCL_RT_PATH)/lib:$(VORTEX_DRV_PATH)/vlsim:$(LD_LIBRARY_PATH) ./$(PROJECT) $(OPTS)
run-rtlsim: $(PROJECT) kernel.bin
LD_LIBRARY_PATH=$(POCL_RT_PATH)/lib:$(VORTEX_DRV_PATH)/rtlsim:$(LD_LIBRARY_PATH) ./$(PROJECT) $(OPTS)

View File

@@ -3,12 +3,12 @@
#define KERNEL_ARG_DEV_MEM_ADDR 0x7ffff000
struct kernel_arg_t {
typedef struct {
uint32_t num_tasks;
uint32_t task_size;
uint32_t src0_ptr;
uint32_t src1_ptr;
uint32_t dst_ptr;
};
} kernel_arg_t;
#endif

View File

@@ -3,12 +3,11 @@
#include <vx_spawn.h>
#include "common.h"
void kernel_body(int task_id, void* arg) {
struct kernel_arg_t* _arg = (struct kernel_arg_t*)(arg);
uint32_t count = _arg->task_size;
int32_t* src0_ptr = (int32_t*)_arg->src0_ptr;
int32_t* src1_ptr = (int32_t*)_arg->src1_ptr;
int32_t* dst_ptr = (int32_t*)_arg->dst_ptr;
void kernel_body(int task_id, const kernel_arg_t* arg) {
uint32_t count = arg->task_size;
int32_t* src0_ptr = (int32_t*)arg->src0_ptr;
int32_t* src1_ptr = (int32_t*)arg->src1_ptr;
int32_t* dst_ptr = (int32_t*)arg->dst_ptr;
uint32_t offset = task_id * count;
@@ -18,6 +17,6 @@ void kernel_body(int task_id, void* arg) {
}
void main() {
struct kernel_arg_t* arg = (struct kernel_arg_t*)KERNEL_ARG_DEV_MEM_ADDR;
const kernel_arg_t* arg = (const kernel_arg_t*)KERNEL_ARG_DEV_MEM_ADDR;
vx_spawn_tasks(arg->num_tasks, kernel_body, arg);
}

View File

@@ -45,13 +45,13 @@ run-simx: $(PROJECT) kernel.bin
LD_LIBRARY_PATH=$(POCL_RT_PATH)/lib:$(VORTEX_DRV_PATH)/simx:$(LD_LIBRARY_PATH) ./$(PROJECT) $(OPTS)
run-fpga: $(PROJECT) kernel.bin
LD_LIBRARY_PATH=$(POCL_RT_PATH)/lib:$(VORTEX_DRV_PATH)/opae:$(LD_LIBRARY_PATH) ./$(PROJECT) $(OPTS)
LD_LIBRARY_PATH=$(POCL_RT_PATH)/lib:$(VORTEX_DRV_PATH)/fpga:$(LD_LIBRARY_PATH) ./$(PROJECT) $(OPTS)
run-asesim: $(PROJECT) kernel.bin
LD_LIBRARY_PATH=$(POCL_RT_PATH)/lib:$(VORTEX_DRV_PATH)/opae/ase:$(LD_LIBRARY_PATH) ./$(PROJECT) $(OPTS)
LD_LIBRARY_PATH=$(POCL_RT_PATH)/lib:$(VORTEX_DRV_PATH)/asesim:$(LD_LIBRARY_PATH) ./$(PROJECT) $(OPTS)
run-vlsim: $(PROJECT) kernel.bin
LD_LIBRARY_PATH=$(POCL_RT_PATH)/lib:$(VORTEX_DRV_PATH)/opae/vlsim:$(LD_LIBRARY_PATH) ./$(PROJECT) $(OPTS)
LD_LIBRARY_PATH=$(POCL_RT_PATH)/lib:$(VORTEX_DRV_PATH)/vlsim:$(LD_LIBRARY_PATH) ./$(PROJECT) $(OPTS)
run-rtlsim: $(PROJECT) kernel.bin
LD_LIBRARY_PATH=$(POCL_RT_PATH)/lib:$(VORTEX_DRV_PATH)/rtlsim:$(LD_LIBRARY_PATH) ./$(PROJECT) $(OPTS)

View File

@@ -3,10 +3,10 @@
#define KERNEL_ARG_DEV_MEM_ADDR 0x7ffff000
struct kernel_arg_t {
typedef struct {
uint32_t num_points;
uint32_t src_ptr;
uint32_t dst_ptr;
};
} kernel_arg_t;
#endif

View File

@@ -5,10 +5,9 @@
// Parallel Selection sort
void kernel_body(int task_id, void* arg) {
struct kernel_arg_t* _arg = (struct kernel_arg_t*)(arg);
int32_t* src_ptr = (int32_t*)_arg->src_ptr;
int32_t* dst_ptr = (int32_t*)_arg->dst_ptr;
void kernel_body(int task_id, const kernel_arg_t* arg) {
int32_t* src_ptr = (int32_t*)arg->src_ptr;
int32_t* dst_ptr = (int32_t*)arg->dst_ptr;
int value = src_ptr[task_id];
@@ -45,6 +44,6 @@ void kernel_body(int task_id, void* arg) {
}
void main() {
struct kernel_arg_t* arg = (struct kernel_arg_t*)KERNEL_ARG_DEV_MEM_ADDR;
const kernel_arg_t* arg = (const kernel_arg_t*)KERNEL_ARG_DEV_MEM_ADDR;
vx_spawn_tasks(arg->num_points, kernel_body, arg);
}

View File

@@ -21,7 +21,7 @@ VX_SRCS = kernel.c
#CXXFLAGS += -std=c++11 -O2 -Wall -Wextra -pedantic -Wfatal-errors
CXXFLAGS += -std=c++11 -O0 -g -Wall -Wextra -pedantic -Wfatal-errors
CXXFLAGS += -I$(VORTEX_DRV_PATH)/include
CXXFLAGS += -I$(VORTEX_DRV_PATH)/include -I$(VORTEX_RT_PATH)/../hw
LDFLAGS += -L$(VORTEX_DRV_PATH)/stub -lvortex
@@ -47,13 +47,13 @@ run-simx: $(PROJECT) kernel.bin
LD_LIBRARY_PATH=$(POCL_RT_PATH)/lib:$(VORTEX_DRV_PATH)/simx:$(LD_LIBRARY_PATH) ./$(PROJECT) $(OPTS)
run-fpga: $(PROJECT) kernel.bin
LD_LIBRARY_PATH=$(POCL_RT_PATH)/lib:$(VORTEX_DRV_PATH)/opae:$(LD_LIBRARY_PATH) ./$(PROJECT) $(OPTS)
LD_LIBRARY_PATH=$(POCL_RT_PATH)/lib:$(VORTEX_DRV_PATH)/fpga:$(LD_LIBRARY_PATH) ./$(PROJECT) $(OPTS)
run-asesim: $(PROJECT) kernel.bin
LD_LIBRARY_PATH=$(POCL_RT_PATH)/lib:$(VORTEX_DRV_PATH)/opae/ase:$(LD_LIBRARY_PATH) ./$(PROJECT) $(OPTS)
LD_LIBRARY_PATH=$(POCL_RT_PATH)/lib:$(VORTEX_DRV_PATH)/asesim:$(LD_LIBRARY_PATH) ./$(PROJECT) $(OPTS)
run-vlsim: $(PROJECT) kernel.bin
LD_LIBRARY_PATH=$(POCL_RT_PATH)/lib:$(VORTEX_DRV_PATH)/opae/vlsim:$(LD_LIBRARY_PATH) ./$(PROJECT) $(OPTS)
LD_LIBRARY_PATH=$(POCL_RT_PATH)/lib:$(VORTEX_DRV_PATH)/vlsim:$(LD_LIBRARY_PATH) ./$(PROJECT) $(OPTS)
run-rtlsim: $(PROJECT) kernel.bin
LD_LIBRARY_PATH=$(POCL_RT_PATH)/lib:$(VORTEX_DRV_PATH)/rtlsim:$(LD_LIBRARY_PATH) ./$(PROJECT) $(OPTS)

View File

@@ -3,13 +3,13 @@
#define KERNEL_ARG_DEV_MEM_ADDR 0x7ffff000
struct kernel_arg_t {
typedef struct {
uint32_t testid;
uint32_t num_tasks;
uint32_t task_size;
uint32_t src0_ptr;
uint32_t src1_ptr;
uint32_t dst_ptr;
};
} kernel_arg_t;
#endif

View File

@@ -4,19 +4,18 @@
#include <vx_spawn.h>
#include "common.h"
typedef void (*PFN_Kernel)(int task_id, void* arg);
typedef void (*PFN_Kernel)(int task_id, const kernel_arg_t* arg);
inline float __ieee754_sqrtf (float x) {
asm ("fsqrt.s %0, %1" : "=f" (x) : "f" (x));
return x;
}
void kernel_iadd(int task_id, void* arg) {
struct kernel_arg_t* _arg = (struct kernel_arg_t*)(arg);
uint32_t count = _arg->task_size;
int32_t* src0_ptr = (int32_t*)_arg->src0_ptr;
int32_t* src1_ptr = (int32_t*)_arg->src1_ptr;
int32_t* dst_ptr = (int32_t*)_arg->dst_ptr;
void kernel_iadd(int task_id, const kernel_arg_t* arg) {
uint32_t count = arg->task_size;
int32_t* src0_ptr = (int32_t*)arg->src0_ptr;
int32_t* src1_ptr = (int32_t*)arg->src1_ptr;
int32_t* dst_ptr = (int32_t*)arg->dst_ptr;
uint32_t offset = task_id * count;
for (uint32_t i = 0; i < count; ++i) {
@@ -27,12 +26,11 @@ void kernel_iadd(int task_id, void* arg) {
}
}
void kernel_imul(int task_id, void* arg) {
struct kernel_arg_t* _arg = (struct kernel_arg_t*)(arg);
uint32_t count = _arg->task_size;
int32_t* src0_ptr = (int32_t*)_arg->src0_ptr;
int32_t* src1_ptr = (int32_t*)_arg->src1_ptr;
int32_t* dst_ptr = (int32_t*)_arg->dst_ptr;
void kernel_imul(int task_id, const kernel_arg_t* arg) {
uint32_t count = arg->task_size;
int32_t* src0_ptr = (int32_t*)arg->src0_ptr;
int32_t* src1_ptr = (int32_t*)arg->src1_ptr;
int32_t* dst_ptr = (int32_t*)arg->dst_ptr;
uint32_t offset = task_id * count;
for (uint32_t i = 0; i < count; ++i) {
@@ -43,12 +41,11 @@ void kernel_imul(int task_id, void* arg) {
}
}
void kernel_idiv(int task_id, void* arg) {
struct kernel_arg_t* _arg = (struct kernel_arg_t*)(arg);
uint32_t count = _arg->task_size;
int32_t* src0_ptr = (int32_t*)_arg->src0_ptr;
int32_t* src1_ptr = (int32_t*)_arg->src1_ptr;
int32_t* dst_ptr = (int32_t*)_arg->dst_ptr;
void kernel_idiv(int task_id, const kernel_arg_t* arg) {
uint32_t count = arg->task_size;
int32_t* src0_ptr = (int32_t*)arg->src0_ptr;
int32_t* src1_ptr = (int32_t*)arg->src1_ptr;
int32_t* dst_ptr = (int32_t*)arg->dst_ptr;
uint32_t offset = task_id * count;
for (uint32_t i = 0; i < count; ++i) {
@@ -59,12 +56,11 @@ void kernel_idiv(int task_id, void* arg) {
}
}
void kernel_idiv_mul(int task_id, void* arg) {
struct kernel_arg_t* _arg = (struct kernel_arg_t*)(arg);
uint32_t count = _arg->task_size;
int32_t* src0_ptr = (int32_t*)_arg->src0_ptr;
int32_t* src1_ptr = (int32_t*)_arg->src1_ptr;
int32_t* dst_ptr = (int32_t*)_arg->dst_ptr;
void kernel_idiv_mul(int task_id, const kernel_arg_t* arg) {
uint32_t count = arg->task_size;
int32_t* src0_ptr = (int32_t*)arg->src0_ptr;
int32_t* src1_ptr = (int32_t*)arg->src1_ptr;
int32_t* dst_ptr = (int32_t*)arg->dst_ptr;
uint32_t offset = task_id * count;
for (uint32_t i = 0; i < count; ++i) {
@@ -77,12 +73,11 @@ void kernel_idiv_mul(int task_id, void* arg) {
}
}
void kernel_fadd(int task_id, void* arg) {
struct kernel_arg_t* _arg = (struct kernel_arg_t*)(arg);
uint32_t count = _arg->task_size;
float* src0_ptr = (float*)_arg->src0_ptr;
float* src1_ptr = (float*)_arg->src1_ptr;
float* dst_ptr = (float*)_arg->dst_ptr;
void kernel_fadd(int task_id, const kernel_arg_t* arg) {
uint32_t count = arg->task_size;
float* src0_ptr = (float*)arg->src0_ptr;
float* src1_ptr = (float*)arg->src1_ptr;
float* dst_ptr = (float*)arg->dst_ptr;
uint32_t offset = task_id * count;
for (uint32_t i = 0; i < count; ++i) {
@@ -93,12 +88,11 @@ void kernel_fadd(int task_id, void* arg) {
}
}
void kernel_fsub(int task_id, void* arg) {
struct kernel_arg_t* _arg = (struct kernel_arg_t*)(arg);
uint32_t count = _arg->task_size;
float* src0_ptr = (float*)_arg->src0_ptr;
float* src1_ptr = (float*)_arg->src1_ptr;
float* dst_ptr = (float*)_arg->dst_ptr;
void kernel_fsub(int task_id, const kernel_arg_t* arg) {
uint32_t count = arg->task_size;
float* src0_ptr = (float*)arg->src0_ptr;
float* src1_ptr = (float*)arg->src1_ptr;
float* dst_ptr = (float*)arg->dst_ptr;
uint32_t offset = task_id * count;
for (uint32_t i = 0; i < count; ++i) {
@@ -109,12 +103,11 @@ void kernel_fsub(int task_id, void* arg) {
}
}
void kernel_fmul(int task_id, void* arg) {
struct kernel_arg_t* _arg = (struct kernel_arg_t*)(arg);
uint32_t count = _arg->task_size;
float* src0_ptr = (float*)_arg->src0_ptr;
float* src1_ptr = (float*)_arg->src1_ptr;
float* dst_ptr = (float*)_arg->dst_ptr;
void kernel_fmul(int task_id, const kernel_arg_t* arg) {
uint32_t count = arg->task_size;
float* src0_ptr = (float*)arg->src0_ptr;
float* src1_ptr = (float*)arg->src1_ptr;
float* dst_ptr = (float*)arg->dst_ptr;
uint32_t offset = task_id * count;
for (uint32_t i = 0; i < count; ++i) {
@@ -125,12 +118,11 @@ void kernel_fmul(int task_id, void* arg) {
}
}
void kernel_fmadd(int task_id, void* arg) {
struct kernel_arg_t* _arg = (struct kernel_arg_t*)(arg);
uint32_t count = _arg->task_size;
float* src0_ptr = (float*)_arg->src0_ptr;
float* src1_ptr = (float*)_arg->src1_ptr;
float* dst_ptr = (float*)_arg->dst_ptr;
void kernel_fmadd(int task_id, const kernel_arg_t* arg) {
uint32_t count = arg->task_size;
float* src0_ptr = (float*)arg->src0_ptr;
float* src1_ptr = (float*)arg->src1_ptr;
float* dst_ptr = (float*)arg->dst_ptr;
uint32_t offset = task_id * count;
for (uint32_t i = 0; i < count; ++i) {
@@ -141,12 +133,11 @@ void kernel_fmadd(int task_id, void* arg) {
}
}
void kernel_fmsub(int task_id, void* arg) {
struct kernel_arg_t* _arg = (struct kernel_arg_t*)(arg);
uint32_t count = _arg->task_size;
float* src0_ptr = (float*)_arg->src0_ptr;
float* src1_ptr = (float*)_arg->src1_ptr;
float* dst_ptr = (float*)_arg->dst_ptr;
void kernel_fmsub(int task_id, const kernel_arg_t* arg) {
uint32_t count = arg->task_size;
float* src0_ptr = (float*)arg->src0_ptr;
float* src1_ptr = (float*)arg->src1_ptr;
float* dst_ptr = (float*)arg->dst_ptr;
uint32_t offset = task_id * count;
for (uint32_t i = 0; i < count; ++i) {
@@ -157,12 +148,11 @@ void kernel_fmsub(int task_id, void* arg) {
}
}
void kernel_fnmadd(int task_id, void* arg) {
struct kernel_arg_t* _arg = (struct kernel_arg_t*)(arg);
uint32_t count = _arg->task_size;
float* src0_ptr = (float*)_arg->src0_ptr;
float* src1_ptr = (float*)_arg->src1_ptr;
float* dst_ptr = (float*)_arg->dst_ptr;
void kernel_fnmadd(int task_id, const kernel_arg_t* arg) {
uint32_t count = arg->task_size;
float* src0_ptr = (float*)arg->src0_ptr;
float* src1_ptr = (float*)arg->src1_ptr;
float* dst_ptr = (float*)arg->dst_ptr;
uint32_t offset = task_id * count;
for (uint32_t i = 0; i < count; ++i) {
@@ -173,12 +163,11 @@ void kernel_fnmadd(int task_id, void* arg) {
}
}
void kernel_fnmsub(int task_id, void* arg) {
struct kernel_arg_t* _arg = (struct kernel_arg_t*)(arg);
uint32_t count = _arg->task_size;
float* src0_ptr = (float*)_arg->src0_ptr;
float* src1_ptr = (float*)_arg->src1_ptr;
float* dst_ptr = (float*)_arg->dst_ptr;
void kernel_fnmsub(int task_id, const kernel_arg_t* arg) {
uint32_t count = arg->task_size;
float* src0_ptr = (float*)arg->src0_ptr;
float* src1_ptr = (float*)arg->src1_ptr;
float* dst_ptr = (float*)arg->dst_ptr;
uint32_t offset = task_id * count;
for (uint32_t i = 0; i < count; ++i) {
@@ -189,12 +178,11 @@ void kernel_fnmsub(int task_id, void* arg) {
}
}
void kernel_fnmadd_madd(int task_id, void* arg) {
struct kernel_arg_t* _arg = (struct kernel_arg_t*)(arg);
uint32_t count = _arg->task_size;
float* src0_ptr = (float*)_arg->src0_ptr;
float* src1_ptr = (float*)_arg->src1_ptr;
float* dst_ptr = (float*)_arg->dst_ptr;
void kernel_fnmadd_madd(int task_id, const kernel_arg_t* arg) {
uint32_t count = arg->task_size;
float* src0_ptr = (float*)arg->src0_ptr;
float* src1_ptr = (float*)arg->src1_ptr;
float* dst_ptr = (float*)arg->dst_ptr;
uint32_t offset = task_id * count;
for (uint32_t i = 0; i < count; ++i) {
@@ -207,12 +195,11 @@ void kernel_fnmadd_madd(int task_id, void* arg) {
}
}
void kernel_fdiv(int task_id, void* arg) {
struct kernel_arg_t* _arg = (struct kernel_arg_t*)(arg);
uint32_t count = _arg->task_size;
float* src0_ptr = (float*)_arg->src0_ptr;
float* src1_ptr = (float*)_arg->src1_ptr;
float* dst_ptr = (float*)_arg->dst_ptr;
void kernel_fdiv(int task_id, const kernel_arg_t* arg) {
uint32_t count = arg->task_size;
float* src0_ptr = (float*)arg->src0_ptr;
float* src1_ptr = (float*)arg->src1_ptr;
float* dst_ptr = (float*)arg->dst_ptr;
uint32_t offset = task_id * count;
for (uint32_t i = 0; i < count; ++i) {
@@ -223,12 +210,11 @@ void kernel_fdiv(int task_id, void* arg) {
}
}
void kernel_fdiv2(int task_id, void* arg) {
struct kernel_arg_t* _arg = (struct kernel_arg_t*)(arg);
uint32_t count = _arg->task_size;
float* src0_ptr = (float*)_arg->src0_ptr;
float* src1_ptr = (float*)_arg->src1_ptr;
float* dst_ptr = (float*)_arg->dst_ptr;
void kernel_fdiv2(int task_id, const kernel_arg_t* arg) {
uint32_t count = arg->task_size;
float* src0_ptr = (float*)arg->src0_ptr;
float* src1_ptr = (float*)arg->src1_ptr;
float* dst_ptr = (float*)arg->dst_ptr;
uint32_t offset = task_id * count;
for (uint32_t i = 0; i < count; ++i) {
@@ -241,12 +227,11 @@ void kernel_fdiv2(int task_id, void* arg) {
}
}
void kernel_fsqrt(int task_id, void* arg) {
struct kernel_arg_t* _arg = (struct kernel_arg_t*)(arg);
uint32_t count = _arg->task_size;
float* src0_ptr = (float*)_arg->src0_ptr;
float* src1_ptr = (float*)_arg->src1_ptr;
float* dst_ptr = (float*)_arg->dst_ptr;
void kernel_fsqrt(int task_id, const kernel_arg_t* arg) {
uint32_t count = arg->task_size;
float* src0_ptr = (float*)arg->src0_ptr;
float* src1_ptr = (float*)arg->src1_ptr;
float* dst_ptr = (float*)arg->dst_ptr;
uint32_t offset = task_id * count;
for (uint32_t i = 0; i < count; ++i) {
@@ -257,12 +242,11 @@ void kernel_fsqrt(int task_id, void* arg) {
}
}
void kernel_ftoi(int task_id, void* arg) {
struct kernel_arg_t* _arg = (struct kernel_arg_t*)(arg);
uint32_t count = _arg->task_size;
float* src0_ptr = (float*)_arg->src0_ptr;
float* src1_ptr = (float*)_arg->src1_ptr;
int32_t* dst_ptr = (int32_t*)_arg->dst_ptr;
void kernel_ftoi(int task_id, const kernel_arg_t* arg) {
uint32_t count = arg->task_size;
float* src0_ptr = (float*)arg->src0_ptr;
float* src1_ptr = (float*)arg->src1_ptr;
int32_t* dst_ptr = (int32_t*)arg->dst_ptr;
uint32_t offset = task_id * count;
for (uint32_t i = 0; i < count; ++i) {
@@ -274,12 +258,11 @@ void kernel_ftoi(int task_id, void* arg) {
}
}
void kernel_ftou(int task_id, void* arg) {
struct kernel_arg_t* _arg = (struct kernel_arg_t*)(arg);
uint32_t count = _arg->task_size;
float* src0_ptr = (float*)_arg->src0_ptr;
float* src1_ptr = (float*)_arg->src1_ptr;
uint32_t* dst_ptr = (uint32_t*)_arg->dst_ptr;
void kernel_ftou(int task_id, const kernel_arg_t* arg) {
uint32_t count = arg->task_size;
float* src0_ptr = (float*)arg->src0_ptr;
float* src1_ptr = (float*)arg->src1_ptr;
uint32_t* dst_ptr = (uint32_t*)arg->dst_ptr;
uint32_t offset = task_id * count;
for (uint32_t i = 0; i < count; ++i) {
@@ -291,12 +274,11 @@ void kernel_ftou(int task_id, void* arg) {
}
}
void kernel_itof(int task_id, void* arg) {
struct kernel_arg_t* _arg = (struct kernel_arg_t*)(arg);
uint32_t count = _arg->task_size;
int32_t* src0_ptr = (int32_t*)_arg->src0_ptr;
int32_t* src1_ptr = (int32_t*)_arg->src1_ptr;
float* dst_ptr = (float*)_arg->dst_ptr;
void kernel_itof(int task_id, const kernel_arg_t* arg) {
uint32_t count = arg->task_size;
int32_t* src0_ptr = (int32_t*)arg->src0_ptr;
int32_t* src1_ptr = (int32_t*)arg->src1_ptr;
float* dst_ptr = (float*)arg->dst_ptr;
uint32_t offset = task_id * count;
for (uint32_t i = 0; i < count; ++i) {
@@ -308,12 +290,11 @@ void kernel_itof(int task_id, void* arg) {
}
}
void kernel_utof(int task_id, void* arg) {
struct kernel_arg_t* _arg = (struct kernel_arg_t*)(arg);
uint32_t count = _arg->task_size;
int32_t* src0_ptr = (int32_t*)_arg->src0_ptr;
int32_t* src1_ptr = (int32_t*)_arg->src1_ptr;
float* dst_ptr = (float*)_arg->dst_ptr;
void kernel_utof(int task_id, const kernel_arg_t* arg) {
uint32_t count = arg->task_size;
int32_t* src0_ptr = (int32_t*)arg->src0_ptr;
int32_t* src1_ptr = (int32_t*)arg->src1_ptr;
float* dst_ptr = (float*)arg->dst_ptr;
uint32_t offset = task_id * count;
for (uint32_t i = 0; i < count; ++i) {
@@ -348,6 +329,6 @@ static const PFN_Kernel sc_tests[] = {
};
void main() {
struct kernel_arg_t* arg = (struct kernel_arg_t*)KERNEL_ARG_DEV_MEM_ADDR;
const kernel_arg_t* arg = (const kernel_arg_t*)KERNEL_ARG_DEV_MEM_ADDR;
vx_spawn_tasks(arg->num_tasks, sc_tests[arg->testid], arg);
}

View File

@@ -3,6 +3,7 @@
#include <unistd.h>
#include <string.h>
#include <vortex.h>
#include <VX_config.h>
#include "testcases.h"
#include "common.h"

View File

@@ -45,13 +45,13 @@ run-simx: $(PROJECT) kernel.bin
LD_LIBRARY_PATH=$(POCL_RT_PATH)/lib:$(VORTEX_DRV_PATH)/simx:$(LD_LIBRARY_PATH) ./$(PROJECT) $(OPTS)
run-fpga: $(PROJECT) kernel.bin
LD_LIBRARY_PATH=$(POCL_RT_PATH)/lib:$(VORTEX_DRV_PATH)/opae:$(LD_LIBRARY_PATH) ./$(PROJECT) $(OPTS)
LD_LIBRARY_PATH=$(POCL_RT_PATH)/lib:$(VORTEX_DRV_PATH)/fpga:$(LD_LIBRARY_PATH) ./$(PROJECT) $(OPTS)
run-asesim: $(PROJECT) kernel.bin
LD_LIBRARY_PATH=$(POCL_RT_PATH)/lib:$(VORTEX_DRV_PATH)/opae/ase:$(LD_LIBRARY_PATH) ./$(PROJECT) $(OPTS)
LD_LIBRARY_PATH=$(POCL_RT_PATH)/lib:$(VORTEX_DRV_PATH)/asesim:$(LD_LIBRARY_PATH) ./$(PROJECT) $(OPTS)
run-vlsim: $(PROJECT) kernel.bin
LD_LIBRARY_PATH=$(POCL_RT_PATH)/lib:$(VORTEX_DRV_PATH)/opae/vlsim:$(LD_LIBRARY_PATH) ./$(PROJECT) $(OPTS)
LD_LIBRARY_PATH=$(POCL_RT_PATH)/lib:$(VORTEX_DRV_PATH)/vlsim:$(LD_LIBRARY_PATH) ./$(PROJECT) $(OPTS)
run-rtlsim: $(PROJECT) kernel.bin
LD_LIBRARY_PATH=$(POCL_RT_PATH)/lib:$(VORTEX_DRV_PATH)/rtlsim:$(LD_LIBRARY_PATH) ./$(PROJECT) $(OPTS)

View File

@@ -3,12 +3,12 @@
#define KERNEL_ARG_DEV_MEM_ADDR 0x7ffff000
struct kernel_arg_t {
typedef struct {
uint32_t num_tasks;
uint32_t task_size;
uint32_t src0_ptr;
uint32_t src1_ptr;
uint32_t dst_ptr;
};
} kernel_arg_t;
#endif

View File

@@ -3,12 +3,11 @@
#include <vx_spawn.h>
#include "common.h"
void kernel_body(int task_id, void* arg) {
struct kernel_arg_t* _arg = (struct kernel_arg_t*)(arg);
uint32_t count = _arg->task_size;
int32_t* src0_ptr = (int32_t*)_arg->src0_ptr;
int32_t* src1_ptr = (int32_t*)_arg->src1_ptr;
int32_t* dst_ptr = (int32_t*)_arg->dst_ptr;
void kernel_body(int task_id, const kernel_arg_t* arg) {
uint32_t count = arg->task_size;
int32_t* src0_ptr = (int32_t*)arg->src0_ptr;
int32_t* src1_ptr = (int32_t*)arg->src1_ptr;
int32_t* dst_ptr = (int32_t*)arg->dst_ptr;
uint32_t offset = task_id * count;
@@ -20,6 +19,6 @@ void kernel_body(int task_id, void* arg) {
}
void main() {
struct kernel_arg_t* arg = (struct kernel_arg_t*)KERNEL_ARG_DEV_MEM_ADDR;
const kernel_arg_t* arg = (const kernel_arg_t*)KERNEL_ARG_DEV_MEM_ADDR;
vx_spawn_tasks(arg->num_tasks, kernel_body, arg);
}

View File

@@ -45,13 +45,13 @@ run-simx: $(PROJECT) kernel.bin
LD_LIBRARY_PATH=$(POCL_RT_PATH)/lib:$(VORTEX_DRV_PATH)/simx:$(LD_LIBRARY_PATH) ./$(PROJECT) $(OPTS)
run-fpga: $(PROJECT) kernel.bin
LD_LIBRARY_PATH=$(POCL_RT_PATH)/lib:$(VORTEX_DRV_PATH)/opae:$(LD_LIBRARY_PATH) ./$(PROJECT) $(OPTS)
LD_LIBRARY_PATH=$(POCL_RT_PATH)/lib:$(VORTEX_DRV_PATH)/fpga:$(LD_LIBRARY_PATH) ./$(PROJECT) $(OPTS)
run-asesim: $(PROJECT) kernel.bin
LD_LIBRARY_PATH=$(POCL_RT_PATH)/lib:$(VORTEX_DRV_PATH)/opae/ase:$(LD_LIBRARY_PATH) ./$(PROJECT) $(OPTS)
LD_LIBRARY_PATH=$(POCL_RT_PATH)/lib:$(VORTEX_DRV_PATH)/asesim:$(LD_LIBRARY_PATH) ./$(PROJECT) $(OPTS)
run-vlsim: $(PROJECT) kernel.bin
LD_LIBRARY_PATH=$(POCL_RT_PATH)/lib:$(VORTEX_DRV_PATH)/opae/vlsim:$(LD_LIBRARY_PATH) ./$(PROJECT) $(OPTS)
LD_LIBRARY_PATH=$(POCL_RT_PATH)/lib:$(VORTEX_DRV_PATH)/vlsim:$(LD_LIBRARY_PATH) ./$(PROJECT) $(OPTS)
run-rtlsim: $(PROJECT) kernel.bin
LD_LIBRARY_PATH=$(POCL_RT_PATH)/lib:$(VORTEX_DRV_PATH)/rtlsim:$(LD_LIBRARY_PATH) ./$(PROJECT) $(OPTS)

View File

@@ -3,10 +3,10 @@
#define KERNEL_ARG_DEV_MEM_ADDR 0x7ffff000
struct kernel_arg_t {
typedef struct {
uint32_t num_points;
uint32_t src_ptr;
uint32_t dst_ptr;
};
} kernel_arg_t;
#endif

View File

@@ -3,10 +3,9 @@
#include <vx_spawn.h>
#include "common.h"
void kernel_body(int task_id, void* arg) {
struct kernel_arg_t* _arg = (struct kernel_arg_t*)(arg);
uint32_t* src_ptr = (uint32_t*)_arg->src_ptr;
uint32_t* dst_ptr = (uint32_t*)_arg->dst_ptr;
void kernel_body(int task_id, const kernel_arg_t* arg) {
uint32_t* src_ptr = (uint32_t*)arg->src_ptr;
uint32_t* dst_ptr = (uint32_t*)arg->dst_ptr;
int32_t* addr_ptr = (int32_t*)(src_ptr[task_id]);
@@ -14,6 +13,6 @@ void kernel_body(int task_id, void* arg) {
}
void main() {
struct kernel_arg_t* arg = (struct kernel_arg_t*)KERNEL_ARG_DEV_MEM_ADDR;
const kernel_arg_t* arg = (const kernel_arg_t*)KERNEL_ARG_DEV_MEM_ADDR;
vx_spawn_tasks(arg->num_points, kernel_body, arg);
}

View File

@@ -45,13 +45,13 @@ run-simx: $(PROJECT) kernel.bin
LD_LIBRARY_PATH=$(POCL_RT_PATH)/lib:$(VORTEX_DRV_PATH)/simx:$(LD_LIBRARY_PATH) ./$(PROJECT) $(OPTS)
run-fpga: $(PROJECT) kernel.bin
LD_LIBRARY_PATH=$(POCL_RT_PATH)/lib:$(VORTEX_DRV_PATH)/opae:$(LD_LIBRARY_PATH) ./$(PROJECT) $(OPTS)
LD_LIBRARY_PATH=$(POCL_RT_PATH)/lib:$(VORTEX_DRV_PATH)/fpga:$(LD_LIBRARY_PATH) ./$(PROJECT) $(OPTS)
run-asesim: $(PROJECT) kernel.bin
LD_LIBRARY_PATH=$(POCL_RT_PATH)/lib:$(VORTEX_DRV_PATH)/opae/ase:$(LD_LIBRARY_PATH) ./$(PROJECT) $(OPTS)
LD_LIBRARY_PATH=$(POCL_RT_PATH)/lib:$(VORTEX_DRV_PATH)/asesim:$(LD_LIBRARY_PATH) ./$(PROJECT) $(OPTS)
run-vlsim: $(PROJECT) kernel.bin
LD_LIBRARY_PATH=$(POCL_RT_PATH)/lib:$(VORTEX_DRV_PATH)/opae/vlsim:$(LD_LIBRARY_PATH) ./$(PROJECT) $(OPTS)
LD_LIBRARY_PATH=$(POCL_RT_PATH)/lib:$(VORTEX_DRV_PATH)/vlsim:$(LD_LIBRARY_PATH) ./$(PROJECT) $(OPTS)
run-rtlsim: $(PROJECT) kernel.bin
LD_LIBRARY_PATH=$(POCL_RT_PATH)/lib:$(VORTEX_DRV_PATH)/rtlsim:$(LD_LIBRARY_PATH) ./$(PROJECT) $(OPTS)

View File

@@ -5,13 +5,13 @@
#define NUM_LOADS 8
struct kernel_arg_t {
typedef struct {
uint32_t num_tasks;
uint32_t size;
uint32_t stride;
uint32_t addr_ptr;
uint32_t src_ptr;
uint32_t dst_ptr;
};
} kernel_arg_t;
#endif

View File

@@ -3,12 +3,11 @@
#include <vx_spawn.h>
#include "common.h"
void kernel_body(int task_id, void* arg) {
struct kernel_arg_t* _arg = (struct kernel_arg_t*)(arg);
uint32_t stride = _arg->stride;
uint32_t* addr_ptr = (uint32_t*)_arg->addr_ptr;
float* src_ptr = (float*)_arg->src_ptr;
float* dst_ptr = (float*)_arg->dst_ptr;
void kernel_body(int task_id, const kernel_arg_t* arg) {
uint32_t stride = arg->stride;
uint32_t* addr_ptr = (uint32_t*)arg->addr_ptr;
float* src_ptr = (float*)arg->src_ptr;
float* dst_ptr = (float*)arg->dst_ptr;
uint32_t offset = task_id * stride;
@@ -24,6 +23,6 @@ void kernel_body(int task_id, void* arg) {
}
void main() {
struct kernel_arg_t* arg = (struct kernel_arg_t*)KERNEL_ARG_DEV_MEM_ADDR;
const kernel_arg_t* arg = (const kernel_arg_t*)KERNEL_ARG_DEV_MEM_ADDR;
vx_spawn_tasks(arg->num_tasks, kernel_body, arg);
}

View File

@@ -45,13 +45,13 @@ run-simx: $(PROJECT) kernel.bin
LD_LIBRARY_PATH=$(POCL_RT_PATH)/lib:$(VORTEX_DRV_PATH)/simx:$(LD_LIBRARY_PATH) ./$(PROJECT) $(OPTS)
run-fpga: $(PROJECT) kernel.bin
LD_LIBRARY_PATH=$(POCL_RT_PATH)/lib:$(VORTEX_DRV_PATH)/opae:$(LD_LIBRARY_PATH) ./$(PROJECT) $(OPTS)
LD_LIBRARY_PATH=$(POCL_RT_PATH)/lib:$(VORTEX_DRV_PATH)/fpga:$(LD_LIBRARY_PATH) ./$(PROJECT) $(OPTS)
run-asesim: $(PROJECT) kernel.bin
LD_LIBRARY_PATH=$(POCL_RT_PATH)/lib:$(VORTEX_DRV_PATH)/opae/ase:$(LD_LIBRARY_PATH) ./$(PROJECT) $(OPTS)
LD_LIBRARY_PATH=$(POCL_RT_PATH)/lib:$(VORTEX_DRV_PATH)/asesim:$(LD_LIBRARY_PATH) ./$(PROJECT) $(OPTS)
run-vlsim: $(PROJECT) kernel.bin
LD_LIBRARY_PATH=$(POCL_RT_PATH)/lib:$(VORTEX_DRV_PATH)/opae/vlsim:$(LD_LIBRARY_PATH) ./$(PROJECT) $(OPTS)
LD_LIBRARY_PATH=$(POCL_RT_PATH)/lib:$(VORTEX_DRV_PATH)/vlsim:$(LD_LIBRARY_PATH) ./$(PROJECT) $(OPTS)
run-rtlsim: $(PROJECT) kernel.bin
LD_LIBRARY_PATH=$(POCL_RT_PATH)/lib:$(VORTEX_DRV_PATH)/rtlsim:$(LD_LIBRARY_PATH) ./$(PROJECT) $(OPTS)

View File

@@ -3,10 +3,10 @@
#define KERNEL_ARG_DEV_MEM_ADDR 0x7ffff000
struct kernel_arg_t {
typedef struct {
uint32_t size;
uint32_t src_ptr;
uint32_t dst_ptr;
};
} kernel_arg_t;
#endif

View File

@@ -4,7 +4,7 @@
#include "common.h"
void main() {
struct kernel_arg_t* arg = (struct kernel_arg_t*)KERNEL_ARG_DEV_MEM_ADDR;
const kernel_arg_t* arg = (const kernel_arg_t*)KERNEL_ARG_DEV_MEM_ADDR;
uint32_t size = arg->size;
int32_t* src_ptr = (int32_t*)arg->src_ptr;

View File

@@ -47,13 +47,13 @@ run-simx: $(PROJECT) kernel.bin
LD_LIBRARY_PATH=$(POCL_RT_PATH)/lib:$(VORTEX_DRV_PATH)/simx:$(LD_LIBRARY_PATH) ./$(PROJECT) $(OPTS)
run-fpga: $(PROJECT) kernel.bin
LD_LIBRARY_PATH=$(POCL_RT_PATH)/lib:$(VORTEX_DRV_PATH)/opae:$(LD_LIBRARY_PATH) ./$(PROJECT) $(OPTS)
LD_LIBRARY_PATH=$(POCL_RT_PATH)/lib:$(VORTEX_DRV_PATH)/fpga:$(LD_LIBRARY_PATH) ./$(PROJECT) $(OPTS)
run-asesim: $(PROJECT) kernel.bin
LD_LIBRARY_PATH=$(POCL_RT_PATH)/lib:$(VORTEX_DRV_PATH)/opae/ase:$(LD_LIBRARY_PATH) ./$(PROJECT) $(OPTS)
LD_LIBRARY_PATH=$(POCL_RT_PATH)/lib:$(VORTEX_DRV_PATH)/asesim:$(LD_LIBRARY_PATH) ./$(PROJECT) $(OPTS)
run-vlsim: $(PROJECT) kernel.bin
LD_LIBRARY_PATH=$(POCL_RT_PATH)/lib:$(VORTEX_DRV_PATH)/opae/vlsim:$(LD_LIBRARY_PATH) ./$(PROJECT) $(OPTS)
LD_LIBRARY_PATH=$(POCL_RT_PATH)/lib:$(VORTEX_DRV_PATH)/vlsim:$(LD_LIBRARY_PATH) ./$(PROJECT) $(OPTS)
run-rtlsim: $(PROJECT) kernel.bin
LD_LIBRARY_PATH=$(POCL_RT_PATH)/lib:$(VORTEX_DRV_PATH)/rtlsim:$(LD_LIBRARY_PATH) ./$(PROJECT) $(OPTS)

View File

@@ -3,10 +3,10 @@
#define KERNEL_ARG_DEV_MEM_ADDR 0x7ffff000
struct kernel_arg_t {
typedef struct {
uint32_t size;
uint32_t src_ptr;
uint32_t dst_ptr;
};
} kernel_arg_t;
#endif

View File

@@ -4,7 +4,7 @@
#include "common.h"
void main() {
struct kernel_arg_t* arg = (struct kernel_arg_t*)KERNEL_ARG_DEV_MEM_ADDR;
const kernel_arg_t* arg = (const kernel_arg_t*)KERNEL_ARG_DEV_MEM_ADDR;
uint32_t size = arg->size;
int32_t* src_ptr = (int32_t*)arg->src_ptr;

View File

@@ -45,13 +45,13 @@ run-simx: $(PROJECT) kernel.bin
LD_LIBRARY_PATH=$(POCL_RT_PATH)/lib:$(VORTEX_DRV_PATH)/simx:$(LD_LIBRARY_PATH) ./$(PROJECT) $(OPTS)
run-fpga: $(PROJECT) kernel.bin
LD_LIBRARY_PATH=$(POCL_RT_PATH)/lib:$(VORTEX_DRV_PATH)/opae:$(LD_LIBRARY_PATH) ./$(PROJECT) $(OPTS)
LD_LIBRARY_PATH=$(POCL_RT_PATH)/lib:$(VORTEX_DRV_PATH)/fpga:$(LD_LIBRARY_PATH) ./$(PROJECT) $(OPTS)
run-asesim: $(PROJECT) kernel.bin
LD_LIBRARY_PATH=$(POCL_RT_PATH)/lib:$(VORTEX_DRV_PATH)/opae/ase:$(LD_LIBRARY_PATH) ./$(PROJECT) $(OPTS)
LD_LIBRARY_PATH=$(POCL_RT_PATH)/lib:$(VORTEX_DRV_PATH)/asesim:$(LD_LIBRARY_PATH) ./$(PROJECT) $(OPTS)
run-vlsim: $(PROJECT) kernel.bin
LD_LIBRARY_PATH=$(POCL_RT_PATH)/lib:$(VORTEX_DRV_PATH)/opae/vlsim:$(LD_LIBRARY_PATH) ./$(PROJECT) $(OPTS)
LD_LIBRARY_PATH=$(POCL_RT_PATH)/lib:$(VORTEX_DRV_PATH)/vlsim:$(LD_LIBRARY_PATH) ./$(PROJECT) $(OPTS)
run-rtlsim: $(PROJECT) kernel.bin
LD_LIBRARY_PATH=$(POCL_RT_PATH)/lib:$(VORTEX_DRV_PATH)/rtlsim:$(LD_LIBRARY_PATH) ./$(PROJECT) $(OPTS)

View File

@@ -3,9 +3,9 @@
#define KERNEL_ARG_DEV_MEM_ADDR 0x7ffff000
struct kernel_arg_t {
typedef struct {
uint32_t num_points;
uint32_t src_ptr;
};
} kernel_arg_t;
#endif

View File

@@ -4,13 +4,12 @@
#include <vx_spawn.h>
#include "common.h"
void kernel_body(int task_id, void* arg) {
struct kernel_arg_t* _arg = (struct kernel_arg_t*)(arg);
int* src_ptr = (int*)_arg->src_ptr;
void kernel_body(int task_id, const kernel_arg_t* arg) {
int* src_ptr = (int*)arg->src_ptr;
vx_printf("task=%d, value=%d\n", task_id, src_ptr[task_id]);
}
void main() {
struct kernel_arg_t* arg = (struct kernel_arg_t*)KERNEL_ARG_DEV_MEM_ADDR;
const kernel_arg_t* arg = (const kernel_arg_t*)KERNEL_ARG_DEV_MEM_ADDR;
vx_spawn_tasks(arg->num_points, kernel_body, arg);
}

View File

@@ -58,9 +58,7 @@ void cleanup() {
}
}
int run_test(const kernel_arg_t& kernel_arg,
uint32_t buf_size,
uint32_t num_points) {
int run_test() {
// start device
std::cout << "start device" << std::endl;
RT_CHECK(vx_start(device));
@@ -137,7 +135,7 @@ int main(int argc, char *argv[]) {
// run tests
std::cout << "run tests" << std::endl;
RT_CHECK(run_test(kernel_arg, buf_size, num_points));
RT_CHECK(run_test());
// cleanup
std::cout << "cleanup" << std::endl;

View File

@@ -45,13 +45,13 @@ run-simx: $(PROJECT) kernel.bin
LD_LIBRARY_PATH=$(POCL_RT_PATH)/lib:$(VORTEX_DRV_PATH)/simx:$(LD_LIBRARY_PATH) ./$(PROJECT) $(OPTS)
run-fpga: $(PROJECT) kernel.bin
LD_LIBRARY_PATH=$(POCL_RT_PATH)/lib:$(VORTEX_DRV_PATH)/opae:$(LD_LIBRARY_PATH) ./$(PROJECT) $(OPTS)
LD_LIBRARY_PATH=$(POCL_RT_PATH)/lib:$(VORTEX_DRV_PATH)/fpga:$(LD_LIBRARY_PATH) ./$(PROJECT) $(OPTS)
run-asesim: $(PROJECT) kernel.bin
LD_LIBRARY_PATH=$(POCL_RT_PATH)/lib:$(VORTEX_DRV_PATH)/opae/ase:$(LD_LIBRARY_PATH) ./$(PROJECT) $(OPTS)
LD_LIBRARY_PATH=$(POCL_RT_PATH)/lib:$(VORTEX_DRV_PATH)/asesim:$(LD_LIBRARY_PATH) ./$(PROJECT) $(OPTS)
run-vlsim: $(PROJECT) kernel.bin
LD_LIBRARY_PATH=$(POCL_RT_PATH)/lib:$(VORTEX_DRV_PATH)/opae/vlsim:$(LD_LIBRARY_PATH) ./$(PROJECT) $(OPTS)
LD_LIBRARY_PATH=$(POCL_RT_PATH)/lib:$(VORTEX_DRV_PATH)/vlsim:$(LD_LIBRARY_PATH) ./$(PROJECT) $(OPTS)
run-rtlsim: $(PROJECT) kernel.bin
LD_LIBRARY_PATH=$(POCL_RT_PATH)/lib:$(VORTEX_DRV_PATH)/rtlsim:$(LD_LIBRARY_PATH) ./$(PROJECT) $(OPTS)

View File

@@ -3,10 +3,10 @@
#define KERNEL_ARG_DEV_MEM_ADDR 0x7ffff000
struct kernel_arg_t {
typedef struct {
uint32_t num_points;
uint32_t src_ptr;
uint32_t dst_ptr;
};
} kernel_arg_t;
#endif

View File

@@ -20,11 +20,10 @@ int __attribute__((noinline)) __smaller(int index, int tid, int32_t cur_value, i
return ret;
}
void kernel_body(int task_id, void* arg) {
struct kernel_arg_t* _arg = (struct kernel_arg_t*)(arg);
uint32_t num_points = _arg->num_points;
int32_t* src_ptr = (int32_t*)_arg->src_ptr;
int32_t* dst_ptr = (int32_t*)_arg->dst_ptr;
void kernel_body(int task_id, const kernel_arg_t* arg) {
uint32_t num_points = arg->num_points;
int32_t* src_ptr = (int32_t*)arg->src_ptr;
int32_t* dst_ptr = (int32_t*)arg->dst_ptr;
int32_t ref_value = src_ptr[task_id];
@@ -38,6 +37,6 @@ void kernel_body(int task_id, void* arg) {
}
void main() {
struct kernel_arg_t* arg = (struct kernel_arg_t*)KERNEL_ARG_DEV_MEM_ADDR;
const kernel_arg_t* arg = (const kernel_arg_t*)KERNEL_ARG_DEV_MEM_ADDR;
vx_spawn_tasks(arg->num_points, kernel_body, arg);
}

View File

@@ -1,13 +1,18 @@
ALL_TESTS := $(wildcard *.hex)
D_TESTS := $(wildcard *ud-p-*.hex)
V_TESTS := $(wildcard *-v-*.hex)
EXCLUDED_TESTS := $(V_TESTS) rv32si-p-scall.hex rv32si-p-sbreak.hex rv32mi-p-breakpoint.hex rv32ud-p-fclass.hex rv32ua-p-amomax_w.hex rv32ua-p-amoxor_w.hex rv32ud-p-ldst.hex rv32ua-p-amoor_w.hex rv32mi-p-ma_addr.hex rv32ud-p-fdiv.hex rv32ud-p-fcmp.hex rv32mi-p-mcsr.hex rv32ua-p-amoswap_w.hex rv32mi-p-ma_fetch.hex rv32mi-p-csr.hex rv32ua-p-amoadd_w.hex rv32si-p-dirty.hex rv32ud-p-fcvt.hex rv32ui-p-fence_i.hex rv32si-p-csr.hex rv32mi-p-shamt.hex rv32ua-p-amomin_w.hex rv32ua-p-lrsc.hex rv32ud-p-fmadd.hex rv32ud-p-fadd.hex rv32si-p-wfi.hex rv32ua-p-amomaxu_w.hex rv32si-p-ma_fetch.hex rv32ud-p-fmin.hex rv32mi-p-illegal.hex rv32uc-p-rvc.hex rv32mi-p-sbreak.hex rv32ua-p-amominu_w.hex rv32ua-p-amoand_w.hex
EXCLUDED_TESTS := $(V_TESTS) $(D_TESTS) rv32si-p-scall.hex rv32si-p-sbreak.hex rv32mi-p-breakpoint.hex rv32ua-p-amomax_w.hex rv32ua-p-amoxor_w.hex rv32ua-p-amoor_w.hex rv32mi-p-ma_addr.hex rv32mi-p-mcsr.hex rv32ua-p-amoswap_w.hex rv32mi-p-ma_fetch.hex rv32mi-p-csr.hex rv32ua-p-amoadd_w.hex rv32si-p-dirty.hex rv32ui-p-fence_i.hex rv32si-p-csr.hex rv32mi-p-shamt.hex rv32ua-p-amomin_w.hex rv32ua-p-lrsc.hex rv32si-p-wfi.hex rv32ua-p-amomaxu_w.hex rv32si-p-ma_fetch.hex rv32mi-p-illegal.hex rv32uc-p-rvc.hex rv32mi-p-sbreak.hex rv32ua-p-amominu_w.hex rv32ua-p-amoand_w.hex
TESTS := $(filter-out $(EXCLUDED_TESTS), $(ALL_TESTS))
all:
run-simx:
$(foreach test, $(TESTS), ../../../simX/simX -r -a rv32i -c 1 -i $(test) || exit;)
$(foreach test, $(TESTS), ../../../sim/simX/simX -r -a rv32i -c 1 -i $(test) || exit;)
run-rtlsim:
$(foreach test, $(TESTS), ../../../hw/simulate/obj_dir/VVortex -r $(test) || exit;)
$(foreach test, $(TESTS), ../../../sim/rtlsim/rtlsim -r $(test) || exit;)
clean:

View File

@@ -15,25 +15,25 @@ PROJECT = fibonacci
SRCS = main.cpp
all: $(PROJECT).elf $(PROJECT).hex $(PROJECT).dump
all: $(PROJECT).elf $(PROJECT).bin $(PROJECT).dump
$(PROJECT).dump: $(PROJECT).elf
$(DP) -D $(PROJECT).elf > $(PROJECT).dump
$(PROJECT).hex: $(PROJECT).elf
$(CP) -O ihex $(PROJECT).elf $(PROJECT).hex
$(PROJECT).bin: $(PROJECT).elf
$(CP) -O binary $(PROJECT).elf $(PROJECT).bin
$(PROJECT).elf: $(SRCS)
$(CC) $(CFLAGS) $(SRCS) $(LDFLAGS) -o $(PROJECT).elf
run-rtlsim: $(PROJECT).hex
../../../hw/simulate/obj_dir/VVortex $(PROJECT).hex
run-rtlsim: $(PROJECT).bin
../../../sim/rtlsim/rtlsim $(PROJECT).bin
run-simx: $(PROJECT).hex
../../../simX/simX -a rv32i -c 1 -i $(PROJECT).hex
run-simx: $(PROJECT).bin
../../../sim/simX/simX -a rv32i -c 1 -i $(PROJECT).bin
.depend: $(SRCS)
$(CC) $(CFLAGS) -MM $^ > .depend;
clean:
rm -rf *.elf *.hex *.dump .depend
rm -rf *.elf *.bin *.dump .depend

View File

@@ -15,25 +15,25 @@ PROJECT = hello
SRCS = main.cpp
all: $(PROJECT).elf $(PROJECT).hex $(PROJECT).dump
all: $(PROJECT).elf $(PROJECT).bin $(PROJECT).dump
$(PROJECT).dump: $(PROJECT).elf
$(DP) -D $(PROJECT).elf > $(PROJECT).dump
$(PROJECT).hex: $(PROJECT).elf
$(CP) -O ihex $(PROJECT).elf $(PROJECT).hex
$(PROJECT).bin: $(PROJECT).elf
$(CP) -O binary $(PROJECT).elf $(PROJECT).bin
$(PROJECT).elf: $(SRCS)
$(CC) $(CFLAGS) $(SRCS) $(LDFLAGS) -o $(PROJECT).elf
run-rtlsim: $(PROJECT).hex
../../../hw/simulate/obj_dir/VVortex $(PROJECT).hex
run-rtlsim: $(PROJECT).bin
../../../sim/rtlsim/rtlsim $(PROJECT).bin
run-simx: $(PROJECT).hex
../../../simX/simX -a rv32i -c 1 -i $(PROJECT).hex
run-simx: $(PROJECT).bin
../../../sim/simX/simX -a rv32i -c 1 -i $(PROJECT).bin
.depend: $(SRCS)
$(CC) $(CFLAGS) -MM $^ > .depend;
clean:
rm -rf *.elf *.hex *.dump .depend
rm -rf *.elf *.bin *.dump .depend

View File

@@ -6,7 +6,7 @@ AR = $(RISCV_TOOLCHAIN_PATH)/bin/riscv32-unknown-elf-gcc-ar
DP = $(RISCV_TOOLCHAIN_PATH)/bin/riscv32-unknown-elf-objdump
CP = $(RISCV_TOOLCHAIN_PATH)/bin/riscv32-unknown-elf-objcopy
CFLAGS += -march=rv32imf -mabi=ilp32f -O3 -Wstack-usage=1024 -ffreestanding -nostartfiles -fdata-sections -ffunction-sections
CFLAGS += -march=rv32imf -mabi=ilp32f -O3 -Wstack-usage=1024 -ffreestanding -nostartfiles -fdata-sections -ffunction-sections -fpermissive
CFLAGS += -I$(VORTEX_RT_PATH)/include -I$(VORTEX_RT_PATH)/../hw
LDFLAGS += -Wl,-Bstatic,-T,$(VORTEX_RT_PATH)/linker/vx_link.ld -Wl,--gc-sections $(VORTEX_RT_PATH)/libvortexrt.a
@@ -15,25 +15,25 @@ PROJECT = simple
SRCS = main.cpp tests.cpp
all: $(PROJECT).elf $(PROJECT).hex $(PROJECT).dump
all: $(PROJECT).elf $(PROJECT).bin $(PROJECT).dump
$(PROJECT).dump: $(PROJECT).elf
$(DP) -D $(PROJECT).elf > $(PROJECT).dump
$(PROJECT).hex: $(PROJECT).elf
$(CP) -O ihex $(PROJECT).elf $(PROJECT).hex
$(PROJECT).bin: $(PROJECT).elf
$(CP) -O binary $(PROJECT).elf $(PROJECT).bin
$(PROJECT).elf: $(SRCS)
$(CC) $(CFLAGS) $(SRCS) $(LDFLAGS) -o $(PROJECT).elf
run-rtlsim: $(PROJECT).hex
../../../hw/simulate/obj_dir/VVortex $(PROJECT).hex
run-rtlsim: $(PROJECT).bin
../../../sim/rtlsim/rtlsim $(PROJECT).bin
run-simx: $(PROJECT).hex
../../../simX/simX -a rv32i -c 1 -i $(PROJECT).hex
run-simx: $(PROJECT).bin
../../../sim/simX/simX -a rv32i -c 1 -i $(PROJECT).bin
.depend: $(SRCS)
$(CC) $(CFLAGS) -MM $^ > .depend;
clean:
rm -rf *.elf *.hex *.dump .depend
rm -rf *.elf *.bin *.dump .depend

View File

@@ -4,8 +4,6 @@
int main() {
int errors = 0;
vx_printf("Simple Test\n");
errors += test_global_memory();
errors += test_stack_memory();
@@ -14,12 +12,20 @@ int main() {
errors += test_tmc();
errors += test_pred();
errors += test_divergence();
errors += test_wsapwn();
errors += test_spawn_tasks();
errors += test_serial();
errors += test_tmask();
errors += test_barrier();
if (0 == errors) {
vx_printf("Passed!\n");
} else {

View File

@@ -1,12 +1,13 @@
#include "tests.h"
#include <stdio.h>
#include <algorithm>
#include <vx_intrinsics.h>
#include <vx_print.h>
#include <vx_spawn.h>
int check_error(const int* buffer, int size) {
int __attribute__ ((noinline)) check_error(const int* buffer, int offset, int size) {
int errors = 0;
for (int i = 0; i < size; i++) {
for (int i = offset; i < size; i++) {
int value = buffer[i];
int ref_value = 65 + i;
if (value == ref_value) {
@@ -19,37 +20,42 @@ int check_error(const int* buffer, int size) {
return errors;
}
int __attribute__ ((noinline)) make_select_tmask(int tid) {
return (1 << tid);
}
int __attribute__ ((noinline)) make_full_tmask(int num_threads) {
return (1 << num_threads) - 1;
}
///////////////////////////////////////////////////////////////////////////////
#define GLOBAL_MEM_SZ 8
int global_buffer[GLOBAL_MEM_SZ];
int test_global_memory() {
int errors = 0;
vx_printf("Global Memory test\n");
int test_global_memory() {
vx_printf("Global Memory Test\n");
for (int i = 0; i < GLOBAL_MEM_SZ; i++) {
global_buffer[i] = 65 + i;
}
return check_error(global_buffer, GLOBAL_MEM_SZ);
return check_error(global_buffer, 0, GLOBAL_MEM_SZ);
}
///////////////////////////////////////////////////////////////////////////////
int test_stack_memory() {
vx_printf("Stack Memory Test\n");
static const int STACK_MEM_SZ = 8;
int stack_buffer[STACK_MEM_SZ];
int errors = 0;
vx_printf("Stack Memory test\n");
for (int i = 0; i < STACK_MEM_SZ; i++) {
stack_buffer[i] = 65 + i;
}
return check_error(stack_buffer, STACK_MEM_SZ);
return check_error(stack_buffer, 0, STACK_MEM_SZ);
}
///////////////////////////////////////////////////////////////////////////////
@@ -57,92 +63,123 @@ int test_stack_memory() {
int test_shared_memory() {
static const int SHARED_MEM_SZ = 8;
int* shared_buffer = (int*)(SMEM_BASE_ADDR-(SMEM_SIZE-SHARED_MEM_SZ-4));
int errors = 0;
vx_printf("Shared Memory test\n");
vx_printf("Shared Memory Test\n");
for (int i = 0; i < SHARED_MEM_SZ; i++) {
shared_buffer[i] = 65 + i;
}
return check_error(shared_buffer, SHARED_MEM_SZ);
return check_error(shared_buffer, 0, SHARED_MEM_SZ);
}
///////////////////////////////////////////////////////////////////////////////
int tmc_buffer[NUM_THREADS];
int tmc_buffer[8];
int test_tmc() {
int errors = 0;
vx_printf("Thread mask test\n");
vx_tmc(NUM_THREADS);
void __attribute__ ((noinline)) do_tmc() {
unsigned tid = vx_thread_id();
tmc_buffer[tid] = 65 + tid;
}
int test_tmc() {
vx_printf("TMC Test\n");
int num_threads = std::min(vx_num_threads(), 8);
int tmask = make_full_tmask(num_threads);
vx_tmc(tmask);
do_tmc();
vx_tmc(1);
return check_error(tmc_buffer, NUM_THREADS);
return check_error(tmc_buffer, 0, num_threads);
}
///////////////////////////////////////////////////////////////////////////////
int wspawn_buffer[NUM_WARPS];
int pred_buffer[8];
void simple_kernel() {
void __attribute__ ((noinline)) do_pred() {
unsigned tid = vx_thread_id();
pred_buffer[tid] = 65 + tid;
}
int test_pred() {
vx_printf("PRED Test\n");
int num_threads = std::min(vx_num_threads(), 8);
int tmask = make_full_tmask(num_threads);
for (int i = 0; i < num_threads; i++) {
pred_buffer[i] = 0;
}
vx_pred(~1);
do_pred();
vx_tmc(1);
int status_n0 = (0 == tmc_buffer[0]);
int status_n1 = check_error(tmc_buffer, 1, num_threads);
return status_n0 && status_n1;
}
///////////////////////////////////////////////////////////////////////////////
int wspawn_buffer[8];
void wspawn_kernel() {
unsigned wid = vx_warp_id();
wspawn_buffer[wid] = 65 + wid;
vx_tmc(0 == wid);
}
int test_wsapwn() {
vx_printf("test_wspawn\n");
vx_wspawn(NUM_WARPS, simple_kernel);
simple_kernel();
vx_printf("Wspawn Test\n");
int num_warps = std::min(vx_num_warps(), 8);
vx_wspawn(num_warps, wspawn_kernel);
wspawn_kernel();
return check_error(wspawn_buffer, NUM_WARPS);
return check_error(wspawn_buffer, 0, num_warps);
}
///////////////////////////////////////////////////////////////////////////////
#define DIV_BUF_SZ ((NUM_THREADS > 4) ? 4 : NUM_THREADS)
int div_buffer[DIV_BUF_SZ];
int dvg_buffer[4];
int test_divergence() {
int errors = 0;
vx_printf("Control divergence test\n");
vx_tmc(DIV_BUF_SZ);
void __attribute__ ((noinline)) do_divergence() {
unsigned tid = vx_thread_id();
bool b = tid < 2;
__if (b) {
bool c = tid < 1;
__if (c) {
div_buffer[tid] = 65;
__if (tid < 2) {
__if (tid < 1) {
dvg_buffer[tid] = 65;
}
__else {
div_buffer[tid] = 66;
dvg_buffer[tid] = 66;
}
__endif
}
__else {
bool c = tid < 3;
__if (c) {
div_buffer[tid] = 67;
__if (tid < 3) {
dvg_buffer[tid] = 67;
}
__else {
div_buffer[tid] = 68;
dvg_buffer[tid] = 68;
}
__endif
}
__endif
}
int test_divergence() {
vx_printf("Control Divergence Test\n");
int num_threads = std::min(vx_num_threads(), 4);
int tmask = make_full_tmask(num_threads);
vx_tmc(tmask);
do_divergence();
vx_tmc(1);
return check_error(div_buffer, DIV_BUF_SZ);
return check_error(dvg_buffer, 0, num_threads);
}
///////////////////////////////////////////////////////////////////////////////
@@ -156,25 +193,111 @@ typedef struct {
int st_buffer_src[ST_BUF_SZ];
int st_buffer_dst[ST_BUF_SZ];
void st_kernel(int task_id, void * arg) {
st_args_t * arguments = (st_args_t *) arg;
arguments->dst[task_id] = arguments->src[task_id];
void st_kernel(int task_id, const st_args_t * arg) {
arg->dst[task_id] = arg->src[task_id];
}
int test_spawn_tasks() {
int error = 0;
vx_printf("SpawnTasks Test\n");
st_args_t arg;
arg.src = st_buffer_src;
arg.dst = st_buffer_dst;
vx_printf("spawning %d tasks\n", ST_BUF_SZ);
for (int i = 0; i < ST_BUF_SZ; i++) {
st_buffer_src[i] = 65 + i;
}
vx_spawn_tasks(ST_BUF_SZ, st_kernel, &arg);
return check_error(st_buffer_dst, ST_BUF_SZ);
return check_error(st_buffer_dst, 0, ST_BUF_SZ);
}
///////////////////////////////////////////////////////////////////////////////
#define SR_BUF_SZ 8
typedef struct {
int * buf;
} sr_args_t;
int sr_buffer[SR_BUF_SZ];
void sr_kernel(const sr_args_t * arg) {
int tid = vx_thread_id();
arg->buf[tid] = 65 + tid;
}
void __attribute__ ((noinline)) do_serial() {
sr_args_t arg;
arg.buf = sr_buffer;
vx_serial(sr_kernel, &arg);
}
int test_serial() {
vx_printf("Serial Test\n");
int num_threads = std::min(vx_num_threads(), 8);
int tmask = make_full_tmask(num_threads);
vx_tmc(tmask);
do_serial();
vx_tmc(1);
return check_error(sr_buffer, 0, num_threads);
}
///////////////////////////////////////////////////////////////////////////////
int tmask_buffer[8];
int __attribute__ ((noinline)) do_tmask() {
int tid = vx_thread_id();
int tmask = make_select_tmask(tid);
int cur_tmask = vx_thread_mask();
tmask_buffer[tid] = (cur_tmask == tmask) ? (65 + tid) : 0;
return tid + 1;
}
int test_tmask() {
vx_printf("Thread Mask Test\n");
// activate all thread to populate shared variables
vx_tmc(-1);
int num_threads = std::min(vx_num_threads(), 8);
int tid = 0;
l_start:
int tmask = make_select_tmask(tid);
vx_tmc(tmask);
tid = do_tmask();
if (tid < num_threads)
goto l_start;
vx_tmc(1);
return check_error(tmask_buffer, 0, num_threads);
}
///////////////////////////////////////////////////////////////////////////////
int barrier_buffer[8];
volatile int barrier_ctr;
volatile int barrier_stall;
void barrier_kernel() {
unsigned wid = vx_warp_id();
for (int i = 0; i <= (wid * 256); ++i) {
++barrier_stall;
}
barrier_buffer[wid] = 65 + wid;
vx_barrier(0, barrier_ctr);
vx_tmc(0 == wid);
}
int test_barrier() {
vx_printf("Barrier Test\n");
int num_warps = std::min(vx_num_warps(), 8);
barrier_ctr = num_warps;
barrier_stall = 0;
vx_wspawn(num_warps, barrier_kernel);
barrier_kernel();
return check_error(barrier_buffer, 0, num_warps);
}

View File

@@ -9,10 +9,18 @@ int test_shared_memory();
int test_tmc();
int test_pred();
int test_divergence();
int test_wsapwn();
int test_spawn_tasks();
int test_serial();
int test_tmask();
int test_barrier();
#endif