wu-arch/vortex
1
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

update

Browse Source
This commit is contained in:
Blaise Tine
2019-11-22 00:17:19 -05:00
parent 9ebf806c6d d4f6a7e3b2
commit 4eeab73ce6
25 changed files with 225 additions and 430016 deletions
Download Patch File Download Diff File Expand all files Collapse all files

8
benchmarks/opencl/aes/Makefile
View File

@@ -34,7 +34,11 @@ PROJECT=saxpy
all: $(PROJECT).dump $(PROJECT).hex all: $(PROJECT).dump $(PROJECT).hex
lib$(PROJECT).a: kernel.cl lib$(PROJECT).a: kernel.cl
<<<<<<< HEAD
POCL_DEBUG=all POCL_DEBUG_LLVM_PASSES=1 LD_LIBRARY_PATH=$(RISCV_TOOL_PATH)/lib:$(POCL_CC_PATH)/lib $(POCL_CC_PATH)/bin/poclcc -o lib$(PROJECT).a kernel.cl POCL_DEBUG=all POCL_DEBUG_LLVM_PASSES=1 LD_LIBRARY_PATH=$(RISCV_TOOL_PATH)/lib:$(POCL_CC_PATH)/lib $(POCL_CC_PATH)/bin/poclcc -o lib$(PROJECT).a kernel.cl
=======
POCL_DEBUG=all POCL_DEBUG_LLVM_PASSES=1 LD_LIBRARY_PATH=$(RISCV_TOOL_PATH)/lib:$(POCL_CC_PATH)/lib $(POCL_CC_PATH)/bin/poclcc -o kernel.pocl kernel.cl
>>>>>>> d4f6a7e3b221ae64441558037b40b87dbf432798
$(PROJECT).elf: main.cc lib$(PROJECT).a $(PROJECT).elf: main.cc lib$(PROJECT).a
$(CXX) $(CXXFLAGS) -I$(POCL_RT_PATH)/include -L$(POCL_RT_PATH)/lib/static -L. $(VX_SRCS) main.cc -Wl,--whole-archive -l$(PROJECT) -Wl,--no-whole-archive $(LIBS) -o $(PROJECT).elf $(CXX) $(CXXFLAGS) -I$(POCL_RT_PATH)/include -L$(POCL_RT_PATH)/lib/static -L. $(VX_SRCS) main.cc -Wl,--whole-archive -l$(PROJECT) -Wl,--no-whole-archive $(LIBS) -o $(PROJECT).elf
@@ -49,4 +53,8 @@ run:
$(VX_SIMX_PATH)/Vcache_simX -E -a rv32i --core $(PROJECT).hex -s -b 1> emulator.debug $(VX_SIMX_PATH)/Vcache_simX -E -a rv32i --core $(PROJECT).hex -s -b 1> emulator.debug
clean: clean:
<<<<<<< HEAD
rm -rf *.elf *.dump *.hex *.a *.pocl *.o rm -rf *.elf *.dump *.hex *.a *.pocl *.o
=======
rm -rf *.elf *.dump *.hex *.a *.pocl
>>>>>>> d4f6a7e3b221ae64441558037b40b87dbf432798

4
benchmarks/opencl/bfs/Makefile
View File

@@ -37,7 +37,11 @@ lib$(PROJECT).a: kernel.cl
POCL_DEBUG=all POCL_DEBUG_LLVM_PASSES=1 LD_LIBRARY_PATH=$(RISCV_TOOL_PATH)/lib:$(POCL_CC_PATH)/lib $(POCL_CC_PATH)/bin/poclcc -o lib$(PROJECT).a kernel.cl POCL_DEBUG=all POCL_DEBUG_LLVM_PASSES=1 LD_LIBRARY_PATH=$(RISCV_TOOL_PATH)/lib:$(POCL_CC_PATH)/lib $(POCL_CC_PATH)/bin/poclcc -o lib$(PROJECT).a kernel.cl
$(PROJECT).elf: main.cc lib$(PROJECT).a $(PROJECT).elf: main.cc lib$(PROJECT).a
<<<<<<< HEAD
$(CXX) $(CXXFLAGS) -I$(POCL_RT_PATH)/include -L$(POCL_RT_PATH)/lib/static -L. $(VX_SRCS) main.cc timer.cc -Wl,--whole-archive -l$(PROJECT) -Wl,--no-whole-archive $(LIBS) -o $(PROJECT).elf $(CXX) $(CXXFLAGS) -I$(POCL_RT_PATH)/include -L$(POCL_RT_PATH)/lib/static -L. $(VX_SRCS) main.cc timer.cc -Wl,--whole-archive -l$(PROJECT) -Wl,--no-whole-archive $(LIBS) -o $(PROJECT).elf
=======
$(CXX) $(CXXFLAGS) -I$(POCL_RT_PATH)/include -L$(POCL_RT_PATH)/lib/static -L. $(VX_SRCS) main.cc -Wl,--whole-archive -l$(PROJECT) -Wl,--no-whole-archive $(LIBS) -o $(PROJECT).elf
>>>>>>> d4f6a7e3b221ae64441558037b40b87dbf432798
$(PROJECT).hex: $(PROJECT).elf $(PROJECT).hex: $(PROJECT).elf
$(HEX) -O ihex $(PROJECT).elf $(PROJECT).hex $(HEX) -O ihex $(PROJECT).elf $(PROJECT).hex

8
benchmarks/opencl/fft/Makefile
View File

@@ -34,7 +34,11 @@ PROJECT=saxpy
all: $(PROJECT).dump $(PROJECT).hex all: $(PROJECT).dump $(PROJECT).hex
lib$(PROJECT).a: kernel.cl lib$(PROJECT).a: kernel.cl
<<<<<<< HEAD
POCL_DEBUG=all POCL_DEBUG_LLVM_PASSES=1 LD_LIBRARY_PATH=$(RISCV_TOOL_PATH)/lib:$(POCL_CC_PATH)/lib $(POCL_CC_PATH)/bin/poclcc -o lib$(PROJECT).a kernel.cl POCL_DEBUG=all POCL_DEBUG_LLVM_PASSES=1 LD_LIBRARY_PATH=$(RISCV_TOOL_PATH)/lib:$(POCL_CC_PATH)/lib $(POCL_CC_PATH)/bin/poclcc -o lib$(PROJECT).a kernel.cl
=======
POCL_DEBUG=all POCL_DEBUG_LLVM_PASSES=1 LD_LIBRARY_PATH=$(RISCV_TOOL_PATH)/lib:$(POCL_CC_PATH)/lib $(POCL_CC_PATH)/bin/poclcc -o kernel.pocl kernel.cl
>>>>>>> d4f6a7e3b221ae64441558037b40b87dbf432798
$(PROJECT).elf: main.cc lib$(PROJECT).a $(PROJECT).elf: main.cc lib$(PROJECT).a
$(CXX) $(CXXFLAGS) -I$(POCL_RT_PATH)/include -L$(POCL_RT_PATH)/lib/static -L. $(VX_SRCS) main.cc -Wl,--whole-archive -l$(PROJECT) -Wl,--no-whole-archive $(LIBS) -o $(PROJECT).elf $(CXX) $(CXXFLAGS) -I$(POCL_RT_PATH)/include -L$(POCL_RT_PATH)/lib/static -L. $(VX_SRCS) main.cc -Wl,--whole-archive -l$(PROJECT) -Wl,--no-whole-archive $(LIBS) -o $(PROJECT).elf
@@ -49,4 +53,8 @@ run:
$(VX_SIMX_PATH)/Vcache_simX -E -a rv32i --core $(PROJECT).hex -s -b 1> emulator.debug $(VX_SIMX_PATH)/Vcache_simX -E -a rv32i --core $(PROJECT).hex -s -b 1> emulator.debug
clean: clean:
<<<<<<< HEAD
rm -rf *.elf *.dump *.hex *.a *.pocl *.o rm -rf *.elf *.dump *.hex *.a *.pocl *.o
=======
rm -rf *.elf *.dump *.hex *.a *.pocl
>>>>>>> d4f6a7e3b221ae64441558037b40b87dbf432798

8
benchmarks/opencl/hotspot/Makefile
View File

@@ -34,7 +34,11 @@ PROJECT=saxpy
all: $(PROJECT).dump $(PROJECT).hex all: $(PROJECT).dump $(PROJECT).hex
lib$(PROJECT).a: kernel.cl lib$(PROJECT).a: kernel.cl
<<<<<<< HEAD
POCL_DEBUG=all POCL_DEBUG_LLVM_PASSES=1 LD_LIBRARY_PATH=$(RISCV_TOOL_PATH)/lib:$(POCL_CC_PATH)/lib $(POCL_CC_PATH)/bin/poclcc -o lib$(PROJECT).a kernel.cl POCL_DEBUG=all POCL_DEBUG_LLVM_PASSES=1 LD_LIBRARY_PATH=$(RISCV_TOOL_PATH)/lib:$(POCL_CC_PATH)/lib $(POCL_CC_PATH)/bin/poclcc -o lib$(PROJECT).a kernel.cl
=======
POCL_DEBUG=all POCL_DEBUG_LLVM_PASSES=1 LD_LIBRARY_PATH=$(RISCV_TOOL_PATH)/lib:$(POCL_CC_PATH)/lib $(POCL_CC_PATH)/bin/poclcc -o kernel.pocl kernel.cl
>>>>>>> d4f6a7e3b221ae64441558037b40b87dbf432798
$(PROJECT).elf: main.cc lib$(PROJECT).a $(PROJECT).elf: main.cc lib$(PROJECT).a
$(CXX) $(CXXFLAGS) -I$(POCL_RT_PATH)/include -L$(POCL_RT_PATH)/lib/static -L. $(VX_SRCS) main.cc -Wl,--whole-archive -l$(PROJECT) -Wl,--no-whole-archive $(LIBS) -o $(PROJECT).elf $(CXX) $(CXXFLAGS) -I$(POCL_RT_PATH)/include -L$(POCL_RT_PATH)/lib/static -L. $(VX_SRCS) main.cc -Wl,--whole-archive -l$(PROJECT) -Wl,--no-whole-archive $(LIBS) -o $(PROJECT).elf
@@ -49,4 +53,8 @@ run:
$(VX_SIMX_PATH)/Vcache_simX -E -a rv32i --core $(PROJECT).hex -s -b 1> emulator.debug $(VX_SIMX_PATH)/Vcache_simX -E -a rv32i --core $(PROJECT).hex -s -b 1> emulator.debug
clean: clean:
<<<<<<< HEAD
rm -rf *.elf *.dump *.hex *.a *.pocl *.o rm -rf *.elf *.dump *.hex *.a *.pocl *.o
=======
rm -rf *.elf *.dump *.hex *.a *.pocl
>>>>>>> d4f6a7e3b221ae64441558037b40b87dbf432798

11
benchmarks/opencl/kmeans/Makefile
View File

@@ -29,15 +29,26 @@ CXXFLAGS += -Wl,--gc-sections # enable garbage collection of unused input sectio
LIBS = -lOpenCL LIBS = -lOpenCL
<<<<<<< HEAD
PROJECT=kmeans PROJECT=kmeans
=======
PROJECT=saxpy
>>>>>>> d4f6a7e3b221ae64441558037b40b87dbf432798
all: $(PROJECT).dump $(PROJECT).hex all: $(PROJECT).dump $(PROJECT).hex
lib$(PROJECT).a: kernel.cl lib$(PROJECT).a: kernel.cl
<<<<<<< HEAD
POCL_DEBUG=all POCL_DEBUG_LLVM_PASSES=1 LD_LIBRARY_PATH=$(RISCV_TOOL_PATH)/lib:$(POCL_CC_PATH)/lib $(POCL_CC_PATH)/bin/poclcc -o lib$(PROJECT).a kernel.cl POCL_DEBUG=all POCL_DEBUG_LLVM_PASSES=1 LD_LIBRARY_PATH=$(RISCV_TOOL_PATH)/lib:$(POCL_CC_PATH)/lib $(POCL_CC_PATH)/bin/poclcc -o lib$(PROJECT).a kernel.cl
$(PROJECT).elf: main.cc lib$(PROJECT).a $(PROJECT).elf: main.cc lib$(PROJECT).a
$(CXX) $(CXXFLAGS) -I$(POCL_RT_PATH)/include -L$(POCL_RT_PATH)/lib/static -L. $(VX_SRCS) main.cc rmse.c read_input.c cluster.c kmeans_clustering.c -Wl,--whole-archive -l$(PROJECT) -Wl,--no-whole-archive $(LIBS) -o $(PROJECT).elf $(CXX) $(CXXFLAGS) -I$(POCL_RT_PATH)/include -L$(POCL_RT_PATH)/lib/static -L. $(VX_SRCS) main.cc rmse.c read_input.c cluster.c kmeans_clustering.c -Wl,--whole-archive -l$(PROJECT) -Wl,--no-whole-archive $(LIBS) -o $(PROJECT).elf
=======
POCL_DEBUG=all POCL_DEBUG_LLVM_PASSES=1 LD_LIBRARY_PATH=$(RISCV_TOOL_PATH)/lib:$(POCL_CC_PATH)/lib $(POCL_CC_PATH)/bin/poclcc -o kernel.pocl kernel.cl
$(PROJECT).elf: main.cc lib$(PROJECT).a
$(CXX) $(CXXFLAGS) -I$(POCL_RT_PATH)/include -L$(POCL_RT_PATH)/lib/static -L. $(VX_SRCS) main.cc -Wl,--whole-archive -l$(PROJECT) -Wl,--no-whole-archive $(LIBS) -o $(PROJECT).elf
>>>>>>> d4f6a7e3b221ae64441558037b40b87dbf432798
$(PROJECT).hex: $(PROJECT).elf $(PROJECT).hex: $(PROJECT).elf
$(HEX) -O ihex $(PROJECT).elf $(PROJECT).hex $(HEX) -O ihex $(PROJECT).elf $(PROJECT).hex

8
benchmarks/opencl/sort/Makefile
View File

@@ -34,7 +34,11 @@ PROJECT=saxpy
all: $(PROJECT).dump $(PROJECT).hex all: $(PROJECT).dump $(PROJECT).hex
lib$(PROJECT).a: kernel.cl lib$(PROJECT).a: kernel.cl
<<<<<<< HEAD
POCL_DEBUG=all POCL_DEBUG_LLVM_PASSES=1 LD_LIBRARY_PATH=$(RISCV_TOOL_PATH)/lib:$(POCL_CC_PATH)/lib $(POCL_CC_PATH)/bin/poclcc -o lib$(PROJECT).a kernel.cl POCL_DEBUG=all POCL_DEBUG_LLVM_PASSES=1 LD_LIBRARY_PATH=$(RISCV_TOOL_PATH)/lib:$(POCL_CC_PATH)/lib $(POCL_CC_PATH)/bin/poclcc -o lib$(PROJECT).a kernel.cl
=======
POCL_DEBUG=all POCL_DEBUG_LLVM_PASSES=1 LD_LIBRARY_PATH=$(RISCV_TOOL_PATH)/lib:$(POCL_CC_PATH)/lib $(POCL_CC_PATH)/bin/poclcc -o kernel.pocl kernel.cl
>>>>>>> d4f6a7e3b221ae64441558037b40b87dbf432798
$(PROJECT).elf: main.cc lib$(PROJECT).a $(PROJECT).elf: main.cc lib$(PROJECT).a
$(CXX) $(CXXFLAGS) -I$(POCL_RT_PATH)/include -L$(POCL_RT_PATH)/lib/static -L. $(VX_SRCS) main.cc -Wl,--whole-archive -l$(PROJECT) -Wl,--no-whole-archive $(LIBS) -o $(PROJECT).elf $(CXX) $(CXXFLAGS) -I$(POCL_RT_PATH)/include -L$(POCL_RT_PATH)/lib/static -L. $(VX_SRCS) main.cc -Wl,--whole-archive -l$(PROJECT) -Wl,--no-whole-archive $(LIBS) -o $(PROJECT).elf
@@ -49,4 +53,8 @@ run:
$(VX_SIMX_PATH)/Vcache_simX -E -a rv32i --core $(PROJECT).hex -s -b 1> emulator.debug $(VX_SIMX_PATH)/Vcache_simX -E -a rv32i --core $(PROJECT).hex -s -b 1> emulator.debug
clean: clean:
<<<<<<< HEAD
rm -rf *.elf *.dump *.hex *.a *.pocl *.o rm -rf *.elf *.dump *.hex *.a *.pocl *.o
=======
rm -rf *.elf *.dump *.hex *.a *.pocl
>>>>>>> d4f6a7e3b221ae64441558037b40b87dbf432798

4
benchmarks/opencl/spmv/Makefile
View File

@@ -34,7 +34,11 @@ PROJECT=saxpy
all: $(PROJECT).dump $(PROJECT).hex all: $(PROJECT).dump $(PROJECT).hex
lib$(PROJECT).a: kernel.cl lib$(PROJECT).a: kernel.cl
<<<<<<< HEAD
POCL_DEBUG=all POCL_DEBUG_LLVM_PASSES=1 LD_LIBRARY_PATH=$(RISCV_TOOL_PATH)/lib:$(POCL_CC_PATH)/lib $(POCL_CC_PATH)/bin/poclcc -o lib$(PROJECT).a kernel.cl POCL_DEBUG=all POCL_DEBUG_LLVM_PASSES=1 LD_LIBRARY_PATH=$(RISCV_TOOL_PATH)/lib:$(POCL_CC_PATH)/lib $(POCL_CC_PATH)/bin/poclcc -o lib$(PROJECT).a kernel.cl
=======
POCL_DEBUG=all POCL_DEBUG_LLVM_PASSES=1 LD_LIBRARY_PATH=$(RISCV_TOOL_PATH)/lib:$(POCL_CC_PATH)/lib $(POCL_CC_PATH)/bin/poclcc -o kernel.pocl kernel.cl
>>>>>>> d4f6a7e3b221ae64441558037b40b87dbf432798
$(PROJECT).elf: main.cc lib$(PROJECT).a $(PROJECT).elf: main.cc lib$(PROJECT).a
$(CXX) $(CXXFLAGS) -I$(POCL_RT_PATH)/include -L$(POCL_RT_PATH)/lib/static -L. $(VX_SRCS) main.cc -Wl,--whole-archive -l$(PROJECT) -Wl,--no-whole-archive $(LIBS) -o $(PROJECT).elf $(CXX) $(CXXFLAGS) -I$(POCL_RT_PATH)/include -L$(POCL_RT_PATH)/lib/static -L. $(VX_SRCS) main.cc -Wl,--whole-archive -l$(PROJECT) -Wl,--no-whole-archive $(LIBS) -o $(PROJECT).elf

429896
emulator/emulator.debug
View File

File diff suppressed because it is too large Load Diff

BIN
emulator/enc.o
View File

Binary file not shown.

BIN
emulator/harptool
View File

Binary file not shown.

BIN
emulator/harptool.o
View File

Binary file not shown.

34
emulator/include/mem.h
View File

@@ -53,32 +53,18 @@ namespace Harp {
}; };
class Core; class Core;
// class ConsoleMemDevice : public MemDevice { class ConsoleMemDevice : public MemDevice {
// public: public:
// ConsoleMemDevice(Size wS, std::ostream &o, Core &core, bool batch = false); ConsoleMemDevice(Size wS, std::ostream &o, Core &core, bool batch = false) {}
// ~ConsoleMemDevice() {} ~ConsoleMemDevice() {}
// //virtual Size wordSize() const { return wordSize; } //virtual Size wordSize() const { return wordSize; }
// virtual Size size() const { return wordSize; } virtual Size size() const { return 1; }
// virtual Word read(Addr) { pthread_mutex_lock(&cBufLock); virtual Word read(Addr) { Word(5); }
// char c = cBuf.front(); virtual void write(Addr a, Word w) { }
// cBuf.pop();
// pthread_mutex_unlock(&cBufLock);
// return Word(c); }
// virtual void write(Addr a, Word w) { output << char(w); }
// void poll(); void poll() {}
};
// friend void *Harp::consoleInputThread(void *);
// private:
// std::ostream &output;
// Size wordSize;
// Core &core;
// std::queue<char> cBuf;
// pthread_mutex_t cBufLock;
// };
class DiskControllerMemDevice : public MemDevice { class DiskControllerMemDevice : public MemDevice {
public: public:

10
emulator/instruction.cpp
View File

@@ -578,9 +578,15 @@ void Instruction::executeOn(Warp &c) {
reg[rdest] = ((immsrc << 12) & 0xfffff000) + (c.pc - 4); reg[rdest] = ((immsrc << 12) & 0xfffff000) + (c.pc - 4);
break; break;
case JAL_INST: case JAL_INST:
//std::cout << "JAL_INST\n"; std::cout << "JAL_INST\n";
if (!pcSet) nextPc = (c.pc - 4) + immsrc; if (!pcSet) nextPc = (c.pc - 4) + immsrc;
if (!pcSet) {/*std::cout << "JAL... SETTING PC: " << nextPc << "\n"; */}
if (!pcSet)
{
std::cout << "JAL... immsrc: " << hex << immsrc << "\n";
std::cout << "JAL... pc base: " << hex << (c.pc - 4) << "\n";
std::cout << "JAL... SETTING PC: " << nextPc << "\n";
}
if (rdest != 0) if (rdest != 0)
{ {
reg[rdest] = c.pc; reg[rdest] = c.pc;

BIN
emulator/instruction.o
View File

Binary file not shown.

BIN
emulator/libharplib.a
View File

Binary file not shown.

BIN
emulator/libharplib.so
View File

Binary file not shown.

BIN
emulator/mem.o
View File

Binary file not shown.

1
emulator/test_riscv.sh
View File

@@ -2,4 +2,5 @@ echo start > results.txt
# echo ../kernel/vortex_test.hex # echo ../kernel/vortex_test.hex
./harptool -E -a rv32i --core ../runtime/mains/simple/vx_simple_main.hex -s -b 1> emulator.debug ./harptool -E -a rv32i --core ../runtime/mains/simple/vx_simple_main.hex -s -b 1> emulator.debug
# ./harptool -E -a rv32i --core ../benchmarks/opencl/sgemm/sgemm.hex -s -b 1> emulator.debug
# ./harptool -E -a rv32i --core ../runtime/mains/vector_test/vx_vector_main.hex -s -b 1> emulator.debug # ./harptool -E -a rv32i --core ../runtime/mains/vector_test/vx_vector_main.hex -s -b 1> emulator.debug

5
rtl/VX_define.v
View File

@@ -11,8 +11,13 @@
// `define ONLY // `define ONLY
// `define SYN 1 // `define SYN 1
<<<<<<< HEAD
//`define ASIC 1 //`define ASIC 1
//`define SYN_FUNC 1 //`define SYN_FUNC 1
=======
// `define ASIC 1
// `define SYN_FUNC 1
>>>>>>> d4f6a7e3b221ae64441558037b40b87dbf432798
`define NUM_BARRIERS 4 `define NUM_BARRIERS 4

156
rtl/VX_gpr.v
View File

@@ -85,83 +85,87 @@ module VX_gpr (
wire[`NT_M1:0][31:0] to_write = (VX_writeback_inter.rd != 0) ? VX_writeback_inter.write_data : 0; wire[`NT_M1:0][31:0] to_write = (VX_writeback_inter.rd != 0) ? VX_writeback_inter.write_data : 0;
/* verilator lint_off PINCONNECTEMPTY */ genvar curr_base_thread;
rf2_32x128_wm1 first_ram ( for (curr_base_thread = 0; curr_base_thread < 'NT; curr_base_thread=curr_base_thread+4)
.CENYA(), begin
.AYA(), /* verilator lint_off PINCONNECTEMPTY */
.CENYB(), rf2_32x128_wm1 first_ram (
.WENYB(), .CENYA(),
.AYB(), .AYA(),
.QA(temp_a), .CENYB(),
.SOA(), .WENYB(),
.SOB(), .AYB(),
.CLKA(clk), .QA(temp_a[(curr_base_thread+3):(curr_base_thread)]),
.CENA(cena_1), .SOA(),
.AA(VX_gpr_read.rs1), .SOB(),
.CLKB(clk), .CLKA(clk),
.CENB(cenb), .CENA(cena_1),
.WENB(write_bit_mask), .AA(VX_gpr_read.rs1[(curr_base_thread+3):(curr_base_thread)]),
.AB(VX_writeback_inter.rd), .CLKB(clk),
.DB(to_write), .CENB(cenb),
.EMAA(3'b011), .WENB(write_bit_mask[(curr_base_thread+3):(curr_base_thread)]),
.EMASA(1'b0), .AB(VX_writeback_inter.rd[(curr_base_thread+3):(curr_base_thread)]),
.EMAB(3'b011), .DB(to_write[(curr_base_thread+3):(curr_base_thread)]),
.TENA(1'b1), .EMAA(3'b011),
.TCENA(1'b0), .EMASA(1'b0),
.TAA(5'b0), .EMAB(3'b011),
.TENB(1'b1), .TENA(1'b1),
.TCENB(1'b0), .TCENA(1'b0),
.TWENB(128'b0), .TAA(5'b0),
.TAB(5'b0), .TENB(1'b1),
.TDB(128'b0), .TCENB(1'b0),
.RET1N(1'b1), .TWENB(128'b0),
.SIA(2'b0), .TAB(5'b0),
.SEA(1'b0), .TDB(128'b0),
.DFTRAMBYP(1'b0), .RET1N(1'b1),
.SIB(2'b0), .SIA(2'b0),
.SEB(1'b0), .SEA(1'b0),
.COLLDISN(1'b1) .DFTRAMBYP(1'b0),
); .SIB(2'b0),
/* verilator lint_on PINCONNECTEMPTY */ .SEB(1'b0),
.COLLDISN(1'b1)
);
/* verilator lint_on PINCONNECTEMPTY */
/* verilator lint_off PINCONNECTEMPTY */ /* verilator lint_off PINCONNECTEMPTY */
rf2_32x128_wm1 second_ram ( rf2_32x128_wm1 second_ram (
.CENYA(), .CENYA(),
.AYA(), .AYA(),
.CENYB(), .CENYB(),
.WENYB(), .WENYB(),
.AYB(), .AYB(),
.QA(temp_b), .QA(temp_b[(curr_base_thread+3):(curr_base_thread)]),
.SOA(), .SOA(),
.SOB(), .SOB(),
.CLKA(clk), .CLKA(clk),
.CENA(cena_2), .CENA(cena_2),
.AA(VX_gpr_read.rs2), .AA(VX_gpr_read.rs2[(curr_base_thread+3):(curr_base_thread)]),
.CLKB(clk), .CLKB(clk),
.CENB(cenb), .CENB(cenb),
.WENB(write_bit_mask), .WENB(write_bit_mask[(curr_base_thread+3):(curr_base_thread)]),
.AB(VX_writeback_inter.rd), .AB(VX_writeback_inter.rd[(curr_base_thread+3):(curr_base_thread)]),
.DB(to_write), .DB(to_write[(curr_base_thread+3):(curr_base_thread)]),
.EMAA(3'b011), .EMAA(3'b011),
.EMASA(1'b0), .EMASA(1'b0),
.EMAB(3'b011), .EMAB(3'b011),
.TENA(1'b1), .TENA(1'b1),
.TCENA(1'b0), .TCENA(1'b0),
.TAA(5'b0), .TAA(5'b0),
.TENB(1'b1), .TENB(1'b1),
.TCENB(1'b0), .TCENB(1'b0),
.TWENB(128'b0), .TWENB(128'b0),
.TAB(5'b0), .TAB(5'b0),
.TDB(128'b0), .TDB(128'b0),
.RET1N(1'b1), .RET1N(1'b1),
.SIA(2'b0), .SIA(2'b0),
.SEA(1'b0), .SEA(1'b0),
.DFTRAMBYP(1'b0), .DFTRAMBYP(1'b0),
.SIB(2'b0), .SIB(2'b0),
.SEB(1'b0), .SEB(1'b0),
.COLLDISN(1'b1) .COLLDISN(1'b1)
); );
/* verilator lint_on PINCONNECTEMPTY */ /* verilator lint_on PINCONNECTEMPTY */
end
`endif `endif

28
rtl/VX_writeback.v
View File

@@ -63,14 +63,40 @@ module VX_writeback (
wire zero = 0; wire zero = 0;
wire[`NT-1:0][31:0] use_wb_data;
reg prev_is_mem;
always @(posedge clk, posedge reset) begin
if (reset)
begin
prev_is_mem = 0;
end begin
prev_is_mem = mem_wb && !no_slot_mem;
end
end
VX_generic_register #(.N(39 + `NW_M1 + 1 + `NT*33)) wb_register( VX_generic_register #(.N(39 + `NW_M1 + 1 + `NT*33)) wb_register(
.clk (clk), .clk (clk),
.reset(reset), .reset(reset),
.stall(zero), .stall(zero),
.flush(zero), .flush(zero),
.in ({VX_writeback_tempp.write_data, VX_writeback_tempp.wb_valid, VX_writeback_tempp.rd, VX_writeback_tempp.wb, VX_writeback_tempp.wb_warp_num, VX_writeback_tempp.wb_pc}), .in ({VX_writeback_tempp.write_data, VX_writeback_tempp.wb_valid, VX_writeback_tempp.rd, VX_writeback_tempp.wb, VX_writeback_tempp.wb_warp_num, VX_writeback_tempp.wb_pc}),
.out ({VX_writeback_inter.write_data, VX_writeback_inter.wb_valid, VX_writeback_inter.rd, VX_writeback_inter.wb, VX_writeback_inter.wb_warp_num, VX_writeback_inter.wb_pc}) .out ({use_wb_data , VX_writeback_inter.wb_valid, VX_writeback_inter.rd, VX_writeback_inter.wb, VX_writeback_inter.wb_warp_num, VX_writeback_inter.wb_pc})
); );
`ifdef SYN
assign VX_writeback_inter.write_data = prev_is_mem ? VX_writeback_tempp.write_data : use_wb_data;
`else
assign VX_writeback_inter.write_data = use_wb_data;
`endif
endmodule // VX_writeback endmodule // VX_writeback

4
rtl/modelsim/Makefile
View File

@@ -79,7 +79,9 @@ SRC = \
../../models/memory/cln28hpm/rf2_128x128_wm1/rf2_128x128_wm1.v \ ../../models/memory/cln28hpm/rf2_128x128_wm1/rf2_128x128_wm1.v \
../../models/memory/cln28hpm/rf2_256x128_wm1/rf2_256x128_wm1.v \ ../../models/memory/cln28hpm/rf2_256x128_wm1/rf2_256x128_wm1.v \
../../models/memory/cln28hpm/rf2_256x19_wm0/rf2_256x19_wm0.v \ ../../models/memory/cln28hpm/rf2_256x19_wm0/rf2_256x19_wm0.v \
../../models/memory/cln28hpm/rf2_32x128_wm1/rf2_32x128_wm1.v ../../models/memory/cln28hpm/rf2_32x128_wm1/rf2_32x128_wm1.v \
../../models/memory/cln28hpm/rf2_32x19_wm0/rf2_32x19_wm0.v
# ../../models/memory/cln28hpc/rf2_32x128_wm1/rf2_32x128_wm1.v # ../../models/memory/cln28hpc/rf2_32x128_wm1/rf2_32x128_wm1.v
# vortex_dpi.h # vortex_dpi.h

20
runtime/mains/vecadd/vecadd.cl Normal file
View File

@@ -0,0 +1,20 @@
#include "../../intrinsics/vx_intrinsics.h"
kernel void
vecadd (__global const int *a,
__global const int *b,
__global int *c)
{
int gid = get_global_id(0);
__if (gid < 2)
{
c[gid] = a[gid] + b[gid];
}
__else
{
c[gid] = b[gid] - a[gid];
}
__endif
}

30
simX/core.cpp
View File

@@ -405,11 +405,11 @@ void Core::fetch()
printTrace(&inst_in_fetch, "Fetch"); printTrace(&inst_in_fetch, "Fetch");
// #ifdef PRINT_ACTIVE_THREADS // #ifdef PRINT_ACTIVE_THREADS
// for (unsigned j = 0; j < w[schedule_w].tmask.size(); ++j) { for (unsigned j = 0; j < w[schedule_w].tmask.size(); ++j) {
// if (w[schedule_w].activeThreads > j && w[schedule_w].tmask[j]) cout << " 1"; if (w[schedule_w].activeThreads > j && w[schedule_w].tmask[j]) cout << " 1";
// else cout << " 0"; else cout << " 0";
// if (j != w[schedule_w].tmask.size()-1 || schedule_w != w.size()-1) cout << ','; if (j != w[schedule_w].tmask.size()-1 || schedule_w != w.size()-1) cout << ',';
// } }
// #endif // #endif
@@ -430,7 +430,7 @@ void Core::decode()
INIT_TRACE(inst_in_fetch); INIT_TRACE(inst_in_fetch);
} }
printTrace(&inst_in_decode, "Decode"); //printTrace(&inst_in_decode, "Decode");
} }
void Core::scheduler() void Core::scheduler()
@@ -442,7 +442,7 @@ void Core::scheduler()
INIT_TRACE(inst_in_decode); INIT_TRACE(inst_in_decode);
} }
printTrace(&inst_in_scheduler, "scheduler"); //printTrace(&inst_in_scheduler, "scheduler");
} }
void Core::load_store() void Core::load_store()
@@ -496,7 +496,7 @@ void Core::load_store()
if (inst_in_lsu.mem_stall_cycles > 0) inst_in_lsu.mem_stall_cycles--; if (inst_in_lsu.mem_stall_cycles > 0) inst_in_lsu.mem_stall_cycles--;
printTrace(&inst_in_lsu, "LSU"); //printTrace(&inst_in_lsu, "LSU");
} }
void Core::execute_unit() void Core::execute_unit()
@@ -548,7 +548,7 @@ void Core::execute_unit()
// } // }
printTrace(&inst_in_exe, "execute_unit"); //printTrace(&inst_in_exe, "execute_unit");
// INIT_TRACE(inst_in_exe); // INIT_TRACE(inst_in_exe);
} }
@@ -604,7 +604,7 @@ void Core::writeback()
// if (!serviced_exe && !serviced_mem) INIT_TRACE(inst_in_wb); // if (!serviced_exe && !serviced_mem) INIT_TRACE(inst_in_wb);
printTrace(&inst_in_wb, "Writeback"); //printTrace(&inst_in_wb, "Writeback");
} }
@@ -712,12 +712,12 @@ void Warp::step(trace_inst_t * trace_inst) {
// At Debug Level 3, print debug info after each instruction. // At Debug Level 3, print debug info after each instruction.
#ifdef USE_DEBUG // #ifdef USE_DEBUG
if (USE_DEBUG >= 3) { // if (USE_DEBUG >= 3) {
D(3, "Register state:"); D(3, "Register state:");
for (unsigned i = 0; i < reg[0].size(); ++i) { for (unsigned i = 0; i < reg[0].size(); ++i) {
D_RAW(" %r" << setfill(' ') << setw(2) << dec << i << ':'); D_RAW(" %r" << setfill(' ') << setw(2) << dec << i << ':');
for (unsigned j = 0; j < reg.size(); ++j) for (unsigned j = 0; j < (this->activeThreads); ++j)
D_RAW(' ' << setfill('0') << setw(8) << hex << reg[j][i] << setfill(' ') << ' '); D_RAW(' ' << setfill('0') << setw(8) << hex << reg[j][i] << setfill(' ') << ' ');
D_RAW('(' << shadowReg[i] << ')' << endl); D_RAW('(' << shadowReg[i] << ')' << endl);
} }
@@ -729,8 +729,8 @@ void Warp::step(trace_inst_t * trace_inst) {
D_RAW(endl); D_RAW(endl);
D_RAW(endl); D_RAW(endl);
D_RAW(endl); D_RAW(endl);
} // }
#endif // #endif
// Clean up. // Clean up.
delete inst; delete inst;

4
syn/esyn.tcl
View File

@@ -37,7 +37,11 @@ uniquify
define_name_rules verilog -remove_internal_net_bus -remove_port_bus define_name_rules verilog -remove_internal_net_bus -remove_port_bus
change_names -rule verilog -hierarchy change_names -rule verilog -hierarchy
<<<<<<< HEAD
report_qor report_qor
=======
# report_qor
>>>>>>> d4f6a7e3b221ae64441558037b40b87dbf432798
report_area report_area
report_hierarchy report_hierarchy
report_cell report_cell