update
This commit is contained in:
@@ -34,7 +34,11 @@ PROJECT=saxpy
|
||||
all: $(PROJECT).dump $(PROJECT).hex
|
||||
|
||||
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 kernel.pocl kernel.cl
|
||||
>>>>>>> d4f6a7e3b221ae64441558037b40b87dbf432798
|
||||
|
||||
$(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
|
||||
@@ -49,4 +53,8 @@ run:
|
||||
$(VX_SIMX_PATH)/Vcache_simX -E -a rv32i --core $(PROJECT).hex -s -b 1> emulator.debug
|
||||
|
||||
clean:
|
||||
<<<<<<< HEAD
|
||||
rm -rf *.elf *.dump *.hex *.a *.pocl *.o
|
||||
=======
|
||||
rm -rf *.elf *.dump *.hex *.a *.pocl
|
||||
>>>>>>> d4f6a7e3b221ae64441558037b40b87dbf432798
|
||||
|
||||
@@ -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
|
||||
|
||||
$(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 -Wl,--whole-archive -l$(PROJECT) -Wl,--no-whole-archive $(LIBS) -o $(PROJECT).elf
|
||||
>>>>>>> d4f6a7e3b221ae64441558037b40b87dbf432798
|
||||
|
||||
$(PROJECT).hex: $(PROJECT).elf
|
||||
$(HEX) -O ihex $(PROJECT).elf $(PROJECT).hex
|
||||
|
||||
@@ -34,7 +34,11 @@ PROJECT=saxpy
|
||||
all: $(PROJECT).dump $(PROJECT).hex
|
||||
|
||||
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 kernel.pocl kernel.cl
|
||||
>>>>>>> d4f6a7e3b221ae64441558037b40b87dbf432798
|
||||
|
||||
$(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
|
||||
@@ -49,4 +53,8 @@ run:
|
||||
$(VX_SIMX_PATH)/Vcache_simX -E -a rv32i --core $(PROJECT).hex -s -b 1> emulator.debug
|
||||
|
||||
clean:
|
||||
<<<<<<< HEAD
|
||||
rm -rf *.elf *.dump *.hex *.a *.pocl *.o
|
||||
=======
|
||||
rm -rf *.elf *.dump *.hex *.a *.pocl
|
||||
>>>>>>> d4f6a7e3b221ae64441558037b40b87dbf432798
|
||||
|
||||
@@ -34,7 +34,11 @@ PROJECT=saxpy
|
||||
all: $(PROJECT).dump $(PROJECT).hex
|
||||
|
||||
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 kernel.pocl kernel.cl
|
||||
>>>>>>> d4f6a7e3b221ae64441558037b40b87dbf432798
|
||||
|
||||
$(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
|
||||
@@ -49,4 +53,8 @@ run:
|
||||
$(VX_SIMX_PATH)/Vcache_simX -E -a rv32i --core $(PROJECT).hex -s -b 1> emulator.debug
|
||||
|
||||
clean:
|
||||
<<<<<<< HEAD
|
||||
rm -rf *.elf *.dump *.hex *.a *.pocl *.o
|
||||
=======
|
||||
rm -rf *.elf *.dump *.hex *.a *.pocl
|
||||
>>>>>>> d4f6a7e3b221ae64441558037b40b87dbf432798
|
||||
|
||||
@@ -29,15 +29,26 @@ CXXFLAGS += -Wl,--gc-sections # enable garbage collection of unused input sectio
|
||||
|
||||
LIBS = -lOpenCL
|
||||
|
||||
<<<<<<< HEAD
|
||||
PROJECT=kmeans
|
||||
=======
|
||||
PROJECT=saxpy
|
||||
>>>>>>> d4f6a7e3b221ae64441558037b40b87dbf432798
|
||||
|
||||
all: $(PROJECT).dump $(PROJECT).hex
|
||||
|
||||
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
|
||||
|
||||
$(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
|
||||
=======
|
||||
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
|
||||
$(HEX) -O ihex $(PROJECT).elf $(PROJECT).hex
|
||||
|
||||
@@ -34,7 +34,11 @@ PROJECT=saxpy
|
||||
all: $(PROJECT).dump $(PROJECT).hex
|
||||
|
||||
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 kernel.pocl kernel.cl
|
||||
>>>>>>> d4f6a7e3b221ae64441558037b40b87dbf432798
|
||||
|
||||
$(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
|
||||
@@ -49,4 +53,8 @@ run:
|
||||
$(VX_SIMX_PATH)/Vcache_simX -E -a rv32i --core $(PROJECT).hex -s -b 1> emulator.debug
|
||||
|
||||
clean:
|
||||
<<<<<<< HEAD
|
||||
rm -rf *.elf *.dump *.hex *.a *.pocl *.o
|
||||
=======
|
||||
rm -rf *.elf *.dump *.hex *.a *.pocl
|
||||
>>>>>>> d4f6a7e3b221ae64441558037b40b87dbf432798
|
||||
|
||||
@@ -34,7 +34,11 @@ PROJECT=saxpy
|
||||
all: $(PROJECT).dump $(PROJECT).hex
|
||||
|
||||
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 kernel.pocl kernel.cl
|
||||
>>>>>>> d4f6a7e3b221ae64441558037b40b87dbf432798
|
||||
|
||||
$(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
|
||||
|
||||
429896
emulator/emulator.debug
429896
emulator/emulator.debug
File diff suppressed because it is too large
Load Diff
BIN
emulator/enc.o
BIN
emulator/enc.o
Binary file not shown.
Binary file not shown.
Binary file not shown.
@@ -53,32 +53,18 @@ namespace Harp {
|
||||
};
|
||||
|
||||
class Core;
|
||||
// class ConsoleMemDevice : public MemDevice {
|
||||
// public:
|
||||
// ConsoleMemDevice(Size wS, std::ostream &o, Core &core, bool batch = false);
|
||||
// ~ConsoleMemDevice() {}
|
||||
class ConsoleMemDevice : public MemDevice {
|
||||
public:
|
||||
ConsoleMemDevice(Size wS, std::ostream &o, Core &core, bool batch = false) {}
|
||||
~ConsoleMemDevice() {}
|
||||
|
||||
// //virtual Size wordSize() const { return wordSize; }
|
||||
// virtual Size size() const { return wordSize; }
|
||||
// virtual Word read(Addr) { pthread_mutex_lock(&cBufLock);
|
||||
// char c = cBuf.front();
|
||||
// cBuf.pop();
|
||||
// pthread_mutex_unlock(&cBufLock);
|
||||
// return Word(c); }
|
||||
// virtual void write(Addr a, Word w) { output << char(w); }
|
||||
//virtual Size wordSize() const { return wordSize; }
|
||||
virtual Size size() const { return 1; }
|
||||
virtual Word read(Addr) { Word(5); }
|
||||
virtual void write(Addr a, Word w) { }
|
||||
|
||||
// void poll();
|
||||
|
||||
// friend void *Harp::consoleInputThread(void *);
|
||||
|
||||
// private:
|
||||
// std::ostream &output;
|
||||
// Size wordSize;
|
||||
// Core &core;
|
||||
|
||||
// std::queue<char> cBuf;
|
||||
// pthread_mutex_t cBufLock;
|
||||
// };
|
||||
void poll() {}
|
||||
};
|
||||
|
||||
class DiskControllerMemDevice : public MemDevice {
|
||||
public:
|
||||
|
||||
@@ -578,9 +578,15 @@ void Instruction::executeOn(Warp &c) {
|
||||
reg[rdest] = ((immsrc << 12) & 0xfffff000) + (c.pc - 4);
|
||||
break;
|
||||
case JAL_INST:
|
||||
//std::cout << "JAL_INST\n";
|
||||
std::cout << "JAL_INST\n";
|
||||
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)
|
||||
{
|
||||
reg[rdest] = c.pc;
|
||||
|
||||
Binary file not shown.
Binary file not shown.
Binary file not shown.
BIN
emulator/mem.o
BIN
emulator/mem.o
Binary file not shown.
@@ -2,4 +2,5 @@ echo start > results.txt
|
||||
|
||||
# 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 ../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
|
||||
|
||||
@@ -11,8 +11,13 @@
|
||||
// `define ONLY
|
||||
|
||||
// `define SYN 1
|
||||
<<<<<<< HEAD
|
||||
//`define ASIC 1
|
||||
//`define SYN_FUNC 1
|
||||
=======
|
||||
// `define ASIC 1
|
||||
// `define SYN_FUNC 1
|
||||
>>>>>>> d4f6a7e3b221ae64441558037b40b87dbf432798
|
||||
|
||||
`define NUM_BARRIERS 4
|
||||
|
||||
|
||||
156
rtl/VX_gpr.v
156
rtl/VX_gpr.v
@@ -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;
|
||||
|
||||
/* verilator lint_off PINCONNECTEMPTY */
|
||||
rf2_32x128_wm1 first_ram (
|
||||
.CENYA(),
|
||||
.AYA(),
|
||||
.CENYB(),
|
||||
.WENYB(),
|
||||
.AYB(),
|
||||
.QA(temp_a),
|
||||
.SOA(),
|
||||
.SOB(),
|
||||
.CLKA(clk),
|
||||
.CENA(cena_1),
|
||||
.AA(VX_gpr_read.rs1),
|
||||
.CLKB(clk),
|
||||
.CENB(cenb),
|
||||
.WENB(write_bit_mask),
|
||||
.AB(VX_writeback_inter.rd),
|
||||
.DB(to_write),
|
||||
.EMAA(3'b011),
|
||||
.EMASA(1'b0),
|
||||
.EMAB(3'b011),
|
||||
.TENA(1'b1),
|
||||
.TCENA(1'b0),
|
||||
.TAA(5'b0),
|
||||
.TENB(1'b1),
|
||||
.TCENB(1'b0),
|
||||
.TWENB(128'b0),
|
||||
.TAB(5'b0),
|
||||
.TDB(128'b0),
|
||||
.RET1N(1'b1),
|
||||
.SIA(2'b0),
|
||||
.SEA(1'b0),
|
||||
.DFTRAMBYP(1'b0),
|
||||
.SIB(2'b0),
|
||||
.SEB(1'b0),
|
||||
.COLLDISN(1'b1)
|
||||
);
|
||||
/* verilator lint_on PINCONNECTEMPTY */
|
||||
genvar curr_base_thread;
|
||||
for (curr_base_thread = 0; curr_base_thread < 'NT; curr_base_thread=curr_base_thread+4)
|
||||
begin
|
||||
/* verilator lint_off PINCONNECTEMPTY */
|
||||
rf2_32x128_wm1 first_ram (
|
||||
.CENYA(),
|
||||
.AYA(),
|
||||
.CENYB(),
|
||||
.WENYB(),
|
||||
.AYB(),
|
||||
.QA(temp_a[(curr_base_thread+3):(curr_base_thread)]),
|
||||
.SOA(),
|
||||
.SOB(),
|
||||
.CLKA(clk),
|
||||
.CENA(cena_1),
|
||||
.AA(VX_gpr_read.rs1[(curr_base_thread+3):(curr_base_thread)]),
|
||||
.CLKB(clk),
|
||||
.CENB(cenb),
|
||||
.WENB(write_bit_mask[(curr_base_thread+3):(curr_base_thread)]),
|
||||
.AB(VX_writeback_inter.rd[(curr_base_thread+3):(curr_base_thread)]),
|
||||
.DB(to_write[(curr_base_thread+3):(curr_base_thread)]),
|
||||
.EMAA(3'b011),
|
||||
.EMASA(1'b0),
|
||||
.EMAB(3'b011),
|
||||
.TENA(1'b1),
|
||||
.TCENA(1'b0),
|
||||
.TAA(5'b0),
|
||||
.TENB(1'b1),
|
||||
.TCENB(1'b0),
|
||||
.TWENB(128'b0),
|
||||
.TAB(5'b0),
|
||||
.TDB(128'b0),
|
||||
.RET1N(1'b1),
|
||||
.SIA(2'b0),
|
||||
.SEA(1'b0),
|
||||
.DFTRAMBYP(1'b0),
|
||||
.SIB(2'b0),
|
||||
.SEB(1'b0),
|
||||
.COLLDISN(1'b1)
|
||||
);
|
||||
/* verilator lint_on PINCONNECTEMPTY */
|
||||
|
||||
/* verilator lint_off PINCONNECTEMPTY */
|
||||
rf2_32x128_wm1 second_ram (
|
||||
.CENYA(),
|
||||
.AYA(),
|
||||
.CENYB(),
|
||||
.WENYB(),
|
||||
.AYB(),
|
||||
.QA(temp_b),
|
||||
.SOA(),
|
||||
.SOB(),
|
||||
.CLKA(clk),
|
||||
.CENA(cena_2),
|
||||
.AA(VX_gpr_read.rs2),
|
||||
.CLKB(clk),
|
||||
.CENB(cenb),
|
||||
.WENB(write_bit_mask),
|
||||
.AB(VX_writeback_inter.rd),
|
||||
.DB(to_write),
|
||||
.EMAA(3'b011),
|
||||
.EMASA(1'b0),
|
||||
.EMAB(3'b011),
|
||||
.TENA(1'b1),
|
||||
.TCENA(1'b0),
|
||||
.TAA(5'b0),
|
||||
.TENB(1'b1),
|
||||
.TCENB(1'b0),
|
||||
.TWENB(128'b0),
|
||||
.TAB(5'b0),
|
||||
.TDB(128'b0),
|
||||
.RET1N(1'b1),
|
||||
.SIA(2'b0),
|
||||
.SEA(1'b0),
|
||||
.DFTRAMBYP(1'b0),
|
||||
.SIB(2'b0),
|
||||
.SEB(1'b0),
|
||||
.COLLDISN(1'b1)
|
||||
);
|
||||
/* verilator lint_on PINCONNECTEMPTY */
|
||||
/* verilator lint_off PINCONNECTEMPTY */
|
||||
rf2_32x128_wm1 second_ram (
|
||||
.CENYA(),
|
||||
.AYA(),
|
||||
.CENYB(),
|
||||
.WENYB(),
|
||||
.AYB(),
|
||||
.QA(temp_b[(curr_base_thread+3):(curr_base_thread)]),
|
||||
.SOA(),
|
||||
.SOB(),
|
||||
.CLKA(clk),
|
||||
.CENA(cena_2),
|
||||
.AA(VX_gpr_read.rs2[(curr_base_thread+3):(curr_base_thread)]),
|
||||
.CLKB(clk),
|
||||
.CENB(cenb),
|
||||
.WENB(write_bit_mask[(curr_base_thread+3):(curr_base_thread)]),
|
||||
.AB(VX_writeback_inter.rd[(curr_base_thread+3):(curr_base_thread)]),
|
||||
.DB(to_write[(curr_base_thread+3):(curr_base_thread)]),
|
||||
.EMAA(3'b011),
|
||||
.EMASA(1'b0),
|
||||
.EMAB(3'b011),
|
||||
.TENA(1'b1),
|
||||
.TCENA(1'b0),
|
||||
.TAA(5'b0),
|
||||
.TENB(1'b1),
|
||||
.TCENB(1'b0),
|
||||
.TWENB(128'b0),
|
||||
.TAB(5'b0),
|
||||
.TDB(128'b0),
|
||||
.RET1N(1'b1),
|
||||
.SIA(2'b0),
|
||||
.SEA(1'b0),
|
||||
.DFTRAMBYP(1'b0),
|
||||
.SIB(2'b0),
|
||||
.SEB(1'b0),
|
||||
.COLLDISN(1'b1)
|
||||
);
|
||||
/* verilator lint_on PINCONNECTEMPTY */
|
||||
end
|
||||
|
||||
`endif
|
||||
|
||||
|
||||
@@ -63,14 +63,40 @@ module VX_writeback (
|
||||
|
||||
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(
|
||||
.clk (clk),
|
||||
.reset(reset),
|
||||
.stall(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}),
|
||||
.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
|
||||
@@ -79,7 +79,9 @@ SRC = \
|
||||
../../models/memory/cln28hpm/rf2_128x128_wm1/rf2_128x128_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_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
|
||||
|
||||
# vortex_dpi.h
|
||||
|
||||
20
runtime/mains/vecadd/vecadd.cl
Normal file
20
runtime/mains/vecadd/vecadd.cl
Normal 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
|
||||
}
|
||||
@@ -405,11 +405,11 @@ void Core::fetch()
|
||||
printTrace(&inst_in_fetch, "Fetch");
|
||||
|
||||
// #ifdef PRINT_ACTIVE_THREADS
|
||||
// for (unsigned j = 0; j < w[schedule_w].tmask.size(); ++j) {
|
||||
// if (w[schedule_w].activeThreads > j && w[schedule_w].tmask[j]) cout << " 1";
|
||||
// else cout << " 0";
|
||||
// if (j != w[schedule_w].tmask.size()-1 || schedule_w != w.size()-1) cout << ',';
|
||||
// }
|
||||
for (unsigned j = 0; j < w[schedule_w].tmask.size(); ++j) {
|
||||
if (w[schedule_w].activeThreads > j && w[schedule_w].tmask[j]) cout << " 1";
|
||||
else cout << " 0";
|
||||
if (j != w[schedule_w].tmask.size()-1 || schedule_w != w.size()-1) cout << ',';
|
||||
}
|
||||
// #endif
|
||||
|
||||
|
||||
@@ -430,7 +430,7 @@ void Core::decode()
|
||||
INIT_TRACE(inst_in_fetch);
|
||||
}
|
||||
|
||||
printTrace(&inst_in_decode, "Decode");
|
||||
//printTrace(&inst_in_decode, "Decode");
|
||||
}
|
||||
|
||||
void Core::scheduler()
|
||||
@@ -442,7 +442,7 @@ void Core::scheduler()
|
||||
INIT_TRACE(inst_in_decode);
|
||||
}
|
||||
|
||||
printTrace(&inst_in_scheduler, "scheduler");
|
||||
//printTrace(&inst_in_scheduler, "scheduler");
|
||||
}
|
||||
|
||||
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--;
|
||||
|
||||
printTrace(&inst_in_lsu, "LSU");
|
||||
//printTrace(&inst_in_lsu, "LSU");
|
||||
}
|
||||
|
||||
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);
|
||||
}
|
||||
|
||||
@@ -604,7 +604,7 @@ void Core::writeback()
|
||||
|
||||
// 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.
|
||||
#ifdef USE_DEBUG
|
||||
if (USE_DEBUG >= 3) {
|
||||
// #ifdef USE_DEBUG
|
||||
// if (USE_DEBUG >= 3) {
|
||||
D(3, "Register state:");
|
||||
for (unsigned i = 0; i < reg[0].size(); ++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('(' << shadowReg[i] << ')' << endl);
|
||||
}
|
||||
@@ -729,8 +729,8 @@ void Warp::step(trace_inst_t * trace_inst) {
|
||||
D_RAW(endl);
|
||||
D_RAW(endl);
|
||||
D_RAW(endl);
|
||||
}
|
||||
#endif
|
||||
// }
|
||||
// #endif
|
||||
|
||||
// Clean up.
|
||||
delete inst;
|
||||
|
||||
@@ -37,7 +37,11 @@ uniquify
|
||||
define_name_rules verilog -remove_internal_net_bus -remove_port_bus
|
||||
change_names -rule verilog -hierarchy
|
||||
|
||||
<<<<<<< HEAD
|
||||
report_qor
|
||||
=======
|
||||
# report_qor
|
||||
>>>>>>> d4f6a7e3b221ae64441558037b40b87dbf432798
|
||||
report_area
|
||||
report_hierarchy
|
||||
report_cell
|
||||
|
||||
Reference in New Issue
Block a user