From e841e5f21b9d86910a6cc10de3af016912c43ce0 Mon Sep 17 00:00:00 2001 From: sspenst Date: Wed, 6 Jul 2016 15:08:11 -0700 Subject: Added the ability to load from sstarr memory after data has been stored in it --- src/gpgpu-sim/gpu-sim.cc | 3 +++ src/gpgpu-sim/shader.cc | 1 + src/gpgpu-sim/shader.h | 1 + 3 files changed, 5 insertions(+) (limited to 'src/gpgpu-sim') diff --git a/src/gpgpu-sim/gpu-sim.cc b/src/gpgpu-sim/gpu-sim.cc index eafb909..e7ba8e5 100644 --- a/src/gpgpu-sim/gpu-sim.cc +++ b/src/gpgpu-sim/gpu-sim.cc @@ -1049,6 +1049,9 @@ void shader_core_ctx::mem_instruction_stats(const warp_inst_t &inst) case shared_space: m_stats->gpgpu_n_shmem_insn += active_count; break; + case sstarr_space: + m_stats->gpgpu_n_sstarr_insn += active_count; + break; case const_space: m_stats->gpgpu_n_const_insn += active_count; break; diff --git a/src/gpgpu-sim/shader.cc b/src/gpgpu-sim/shader.cc index ff2fac7..9c3f816 100644 --- a/src/gpgpu-sim/shader.cc +++ b/src/gpgpu-sim/shader.cc @@ -386,6 +386,7 @@ void shader_core_stats::print( FILE* fout ) const fprintf(fout, "gpgpu_n_load_insn = %d\n", gpgpu_n_load_insn); fprintf(fout, "gpgpu_n_store_insn = %d\n", gpgpu_n_store_insn); fprintf(fout, "gpgpu_n_shmem_insn = %d\n", gpgpu_n_shmem_insn); + fprintf(fout, "gpgpu_n_shmem_insn = %d\n", gpgpu_n_sstarr_insn); fprintf(fout, "gpgpu_n_tex_insn = %d\n", gpgpu_n_tex_insn); fprintf(fout, "gpgpu_n_const_mem_insn = %d\n", gpgpu_n_const_insn); fprintf(fout, "gpgpu_n_param_mem_insn = %d\n", gpgpu_n_param_insn); diff --git a/src/gpgpu-sim/shader.h b/src/gpgpu-sim/shader.h index 38d09e9..3b9859f 100644 --- a/src/gpgpu-sim/shader.h +++ b/src/gpgpu-sim/shader.h @@ -1372,6 +1372,7 @@ struct shader_core_stats_pod { unsigned gpgpu_n_load_insn; unsigned gpgpu_n_store_insn; unsigned gpgpu_n_shmem_insn; + unsigned gpgpu_n_sstarr_insn; unsigned gpgpu_n_tex_insn; unsigned gpgpu_n_const_insn; unsigned gpgpu_n_param_insn; -- cgit v1.3 From 68674d4ba230df0d3bf9f4e5b035f4cf9cfc185b Mon Sep 17 00:00:00 2001 From: negargoli93 Date: Sat, 12 May 2018 16:09:04 -0700 Subject: commit for eece527project --- src/abstract_hardware_model.h | 6 +- src/cuda-sim/cuda-sim.cc | 11 ++- src/cuda-sim/instructions.cc | 213 +++++++++++++++++++----------------------- src/cuda-sim/opcodes.def | 2 +- src/cuda-sim/ptx.l | 2 +- src/cuda-sim/ptx_ir.h | 25 +++++ src/gpgpu-sim/scoreboard.cc | 4 + src/gpgpusim_entrypoint.cc | 1 - 8 files changed, 136 insertions(+), 128 deletions(-) (limited to 'src/gpgpu-sim') diff --git a/src/abstract_hardware_model.h b/src/abstract_hardware_model.h index cdd9cf3..9dc58d4 100644 --- a/src/abstract_hardware_model.h +++ b/src/abstract_hardware_model.h @@ -750,7 +750,7 @@ public: }; // the maximum number of destination, source, or address uarch operands in a instruction -#define MAX_REG_OPERANDS 8 +#define MAX_REG_OPERANDS 32 struct dram_callback_t { dram_callback_t() { function=NULL; instruction=NULL; thread=NULL; } @@ -825,8 +825,8 @@ public: address_type reconvergence_pc; // -1 => not a branch, -2 => use function return address - unsigned out[4]; - unsigned in[4]; + unsigned out[8]; + unsigned in[8]; unsigned char is_vectorin; unsigned char is_vectorout; int pred; // predicate register number diff --git a/src/cuda-sim/cuda-sim.cc b/src/cuda-sim/cuda-sim.cc index 54d8796..006738a 100644 --- a/src/cuda-sim/cuda-sim.cc +++ b/src/cuda-sim/cuda-sim.cc @@ -792,9 +792,9 @@ void ptx_instruction::set_opcode_and_latency() initiation_interval = dp_init[2]; op = SFU_OP; break; - case BSMAD_OP: - latency = int_precision/int_lane_width; - initiation_interval = int_init_precision/int_init_lane_width; + case MMA_OP: + latency = 64; + initiation_interval = 64; break; case SHFL_OP: latency = 32; @@ -1301,6 +1301,7 @@ void ptx_thread_info::ptx_exec_inst( warp_inst_t &inst, unsigned lane_id) *((warp_inst_t*)pJ) = inst; // copy active mask information pI = pJ; } + if((pI->get_opcode()!=MMA_OP)||((pI->get_opcode()==MMA_OP)&&(lane_id==0))){ switch ( pI->get_opcode() ) { #define OP_DEF(OP,FUNC,STR,DST,CLASSIFICATION) case OP: FUNC(pI,this); op_classification = CLASSIFICATION; break; #define OP_W_DEF(OP,FUNC,STR,DST,CLASSIFICATION) case OP: FUNC(pI,get_core(),inst); op_classification = CLASSIFICATION; break; @@ -1308,7 +1309,7 @@ void ptx_thread_info::ptx_exec_inst( warp_inst_t &inst, unsigned lane_id) #undef OP_DEF #undef OP_W_DEF default: printf( "Execution error: Invalid opcode (0x%x)\n", pI->get_opcode() ); break; - } + }} delete pJ; pI = pI_saved; @@ -1930,7 +1931,7 @@ void functionalCoreSim::executeWarp(unsigned i, bool &allAtBarrier, bool & someO { if(!m_warpAtBarrier[i] && m_liveThreadCount[i]!=0){ warp_inst_t inst =getExecuteWarp(i); - execute_warp_inst_t(inst,i); + execute_warp_inst_t(inst,i); if(inst.isatomic()) inst.do_atomic(true); if(inst.op==BARRIER_OP || inst.op==MEMORY_BARRIER_OP ) m_warpAtBarrier[i]=true; updateSIMTStack( i, &inst ); diff --git a/src/cuda-sim/instructions.cc b/src/cuda-sim/instructions.cc index 493e307..7903343 100644 --- a/src/cuda-sim/instructions.cc +++ b/src/cuda-sim/instructions.cc @@ -771,6 +771,17 @@ void add_impl( const ptx_instruction *pI, ptx_thread_info *thread ) unsigned i_type = pI->get_type(); src1_data = thread->get_operand_value(src1, dst, i_type, thread, 1); src2_data = thread->get_operand_value(src2, dst, i_type, thread, 1); + //unsigned warpId_aa,warp_size_aa; + //warpId_aa = pI->warp_id(); + //warp_size_aa=32; + //dim3 t=thread->get_tid(); + //unsigned tid_aa=warp_size_aa*warpId_aa; + + ptx_thread_info *thread2; + thread2=thread; + src1_data = thread2->get_operand_value(src1, dst, i_type, thread2, 1); + src2_data = thread2->get_operand_value(src2, dst, i_type, thread2, 1); + unsigned rounding_mode = pI->rounding_mode(); int orig_rm = fegetround(); @@ -1483,135 +1494,102 @@ unsigned trunc(unsigned num, unsigned precision) { return num; } -void bsmad_impl( const ptx_instruction *pI, core_t *core, warp_inst_t inst ) +void mma_impl( const ptx_instruction *pI, core_t *core, warp_inst_t inst ) { - // operands: - // 0 = output - // 1 = input precision - // 2 = output precision - // 3 = buffer0 - // 4 = buffer1 - // 5 = buffer2 - // 6 = buffer3 - // 7 = synapse value - // 8 = output value + int i,j,k,thrd; + int row,offset; + printf("mmaWorld\n"); + ptx_reg_t matrix_a[16][16]; + ptx_reg_t matrix_b[16][16]; + ptx_reg_t matrix_c[16][16]; + ptx_reg_t matrix_d[16][16]; + ptx_reg_t src_data; + ptx_thread_info *thread; - const operand_info &dst = pI->dst(); - const operand_info &src1 = pI->src1(); - const operand_info &src2 = pI->src2(); unsigned type = pI->get_type(); int tid = inst.warp_id_func() * core->get_warp_size(); - ptx_thread_info *thread = core->get_thread_info()[tid]; - const int ip = (thread->get_operand_value(src1, dst, type, thread, 1)).u32; - const int op = (thread->get_operand_value(src2, dst, type, thread, 1)).u32; - const int THREADS = inst.active_count(); - const int INBUFFERS = 4; - const int OUTBUFFERS = (((32/ip)*INBUFFERS) / (32/op)) + ((((32/ip)*INBUFFERS) % (32/op)) != 0); - if (OUTBUFFERS > THREADS) { - printf("GPGPU-Sim PTX: BSMAD ERROR - Number of output registers required (%d) is greater than the number available (%d)\n", OUTBUFFERS, THREADS); - abort(); - } - ptx_warp_info *warp_info = thread->m_warp_info; - warp_info->inc_done_threads(); + const operand_info &dst = pI->operand_lookup(0); + +//NOT WOR thread = core->get_thread_info()[tid]; +//NOT WOR const operand_info &src_a= pI->operand_lookup(1); +//NOT WOR src_data= (thread->get_operand_value(src_a, dst, type, thread, 1)); +//NOT WOR thread->set_operand_value(dst, src_data, type, thread, pI); + for (thrd=0; thrd < core->get_warp_size(); thrd++){ + row=thrd/2; + offset=8*(thrd%2); + thread = core->get_thread_info()[tid+thrd]; + printf("thread=%d:",thrd); + for(i=8;i<=31;i++){ + const operand_info &src_a= pI->operand_lookup(i); + src_data= (thread->get_operand_value(src_a, dst, type, thread, 1)); + printf("%f ",src_data.f32); + if(i<=15) + matrix_a[row][offset+(i)%8]=src_data; + else if((i>15)&&(i<=23)) + matrix_b[row][offset+(i)%8]=src_data; + else if(i>23) + matrix_c[row][offset+(i)%8]=src_data; + - // threads within the warp are executed sequentially by the simulator, store output in first four registers - if (warp_info->get_done_threads() <= OUTBUFFERS) { - unsigned buffer[THREADS][INBUFFERS]; - unsigned synapse[THREADS]; - unsigned output; - - // loop through all threads in the warp and get all data - for (unsigned i = 0, j = 0; i < core->get_warp_size(); i++) { - if (inst.active(i)) { - const operand_info &src3 = pI->operand_lookup(3); - const operand_info &src4 = pI->operand_lookup(4); - const operand_info &src5 = pI->operand_lookup(5); - const operand_info &src6 = pI->operand_lookup(6); - const operand_info &src7 = pI->operand_lookup(7); - const operand_info &src8 = pI->operand_lookup(8); - - thread = core->get_thread_info()[tid+i]; - // get buffer data and synapse data from each thread - buffer[j][0] = (thread->get_operand_value(src3, dst, type, thread, 1)).u32; - buffer[j][1] = (thread->get_operand_value(src4, dst, type, thread, 1)).u32; - buffer[j][2] = (thread->get_operand_value(src5, dst, type, thread, 1)).u32; - buffer[j][3] = (thread->get_operand_value(src6, dst, type, thread, 1)).u32; - synapse[j] = (thread->get_operand_value(src7, dst, type, thread, 1)).u32; - j++; - // get output data from the first 4 threads - if (j == warp_info->get_done_threads()) { - output = (thread->get_operand_value(src8, dst, type, thread, 1)).u32; - } - } } + printf("\n"); + } - // unpack registers, compute enough outputs to fill an output register - unsigned *unpacked_output = (unsigned*)calloc(32/op,sizeof(unsigned)); - unsigned buffer_data_start = (32/op)*(warp_info->get_done_threads()-1); - for (unsigned i = buffer_data_start; i < (32/op + buffer_data_start) && i < (32/ip)*INBUFFERS; i++) { - unsigned buf = i/(32/ip); - unsigned pos = i%(32/ip); - // sum values from the buffers - int sum = 0; - unsigned mask = (unsigned)(pow(2,ip)-1) << (pos*ip); - for (int j = 0; j < THREADS; j++) { - //sum += ((mask & buffer[j][buf]) >> (pos*ip)) * synapse[j]; - sum += trunc(((mask & buffer[j][buf]) >> (pos*ip)) * synapse[j], op); - } - // get the previous output - mask = (unsigned)(pow(2,op)-1) << (op*(i-buffer_data_start)); - int past_output = (mask & output) >> (op*(i-buffer_data_start)); - unpacked_output[i-buffer_data_start] = trunc(trunc(sum,op) + past_output,op); - // truncate sum, truncate (truncated sum + past_output) + printf("MATRIX_A\n"); + for (i=0;i<16;i++){ + for(j=0;j<16;j++){ + printf("%f ",matrix_a[i][j].f32); } - - // truncate output - /*for (unsigned i = 0; i < 32/op; i++) { - int mask = 1, latest_one = -1; - unsigned data = unpacked_output[i]; - for (unsigned j = 0; j < sizeof(unsigned)*8; j++) { - int bit = data & mask; - if (bit == 1) latest_one = j; - data >>= 1; - } - if (latest_one >= op) { - // round_up is 1 if the most significant truncated digit is a 1, otherwise it is 0 - int round_up = (unpacked_output[i] & (1 << (latest_one-op))) >> (latest_one-op); - unsigned shifted_output = unpacked_output[i] >> (latest_one-op+1); - // if shifted_output is a number like 1111, don't round up - if (shifted_output == (pow(2,op)-1)) round_up = 0; - unpacked_output[i] = shifted_output + round_up; - } - }*/ - - // pack the outputs into one register - unsigned mask = pow(2,op)-1; - unsigned output_data = 0; - for (int i = 0; i < 32/op; i++) { - output_data |= (unpacked_output[i] & mask) << (op*i); + printf("\n"); + } + printf("MATRIX_B\n"); + for (i=0;i<16;i++){ + for(j=0;j<16;j++){ + printf("%f ",matrix_b[i][j].f32); } - - // store the result in the correct thread's output register - for (unsigned i = 0, j = 0; i < core->get_warp_size(); i++) { - if (inst.active(i)) j++; - if (j == warp_info->get_done_threads()) { - thread = core->get_thread_info()[tid+i]; - ptx_reg_t data; - data.u32 = output_data; - thread->set_operand_value(dst, data, type, thread, pI); - break; + printf("\n"); + } + printf("MATRIX_C\n"); + for (i=0;i<16;i++){ + for(j=0;j<16;j++){ + printf("%f ",matrix_c[i][j].f32); + } + printf("\n"); + } + for (i=0;i<16;i++){ + for(j=0;j<16;j++){ + matrix_d[i][j].f32=0; + } + } + + for (i=0;i<16;i++){ + for(j=0;j<16;j++){ + for(k=0;k<16;k++){ + matrix_d[i][j].f32=matrix_d[i][j].f32+matrix_a[i][k].f32*matrix_b[j][k].f32; } + matrix_d[i][j].f32+=matrix_c[i][j].f32; } } - - // once the warp has finished, set the number of completed threads back to 0 for the next warp - if (warp_info->get_done_threads() == THREADS) { - warp_info->reset_done_threads(); + printf("MATRIX_D\n"); + for (i=0;i<16;i++){ + for(j=0;j<16;j++){ + printf("%f ",matrix_d[i][j].f32); + } + printf("\n"); + } + for (thrd=0; thrd < core->get_warp_size(); thrd++){ + thread = core->get_thread_info()[tid+thrd]; + row=thrd/2; + offset=8*(thrd%2); + for(i=0;i<8;i++){ + const operand_info &dst = pI->operand_lookup(i); + const symbol *r2; + r2=dst.get_symbol(); + printf("thrd=%d,i=%d,register%s, data=%f\n",thrd,i,(r2->name()).c_str(),matrix_d[row][offset+i].f32); + thread->set_operand_value(dst, matrix_d[row][offset+i], type, thread, pI); + } } - - // set the latency assuming 4 bits of each input get processed every cycle - // mutable latency variable??? - //pI->latency = (ip+3)/4; + } void call_impl( const ptx_instruction *pI, ptx_thread_info *thread ) @@ -4098,7 +4076,8 @@ void st_impl( const ptx_instruction *pI, ptx_thread_info *thread ) if (!vector_spec) { data = thread->get_operand_value(src1, dst, type, thread, 1); mem->write(addr,size/8,&data.s64,thread,pI); - } else { + printf("addr=%d data=%d\n",addr,data.s64); + } else { if (vector_spec == V2_TYPE) { ptx_reg_t* ptx_regs = new ptx_reg_t[2]; thread->get_vector_operand_values(src1, ptx_regs, 2); diff --git a/src/cuda-sim/opcodes.def b/src/cuda-sim/opcodes.def index 41f2f22..a3cc83f 100644 --- a/src/cuda-sim/opcodes.def +++ b/src/cuda-sim/opcodes.def @@ -52,7 +52,7 @@ OP_DEF(BRA_OP,bra_impl,"bra",0,3) OP_DEF(BRX_OP,brx_impl,"brx",0,3) OP_DEF(BREV_OP,brev_impl,"brev",1,1) OP_DEF(BRKPT_OP,brkpt_impl,"brkpt",1,9) -OP_W_DEF(BSMAD_OP,bsmad_impl,"bsmad",1,1) +OP_W_DEF(MMA_OP,mma_impl,"mma",1,1) OP_DEF(CALL_OP,call_impl,"call",1,3) OP_DEF(CALLP_OP,callp_impl,"callp",1,3) OP_DEF(CLZ_OP,clz_impl,"clz",1,1) diff --git a/src/cuda-sim/ptx.l b/src/cuda-sim/ptx.l index 7620134..e07e339 100644 --- a/src/cuda-sim/ptx.l +++ b/src/cuda-sim/ptx.l @@ -68,7 +68,7 @@ bra TC; ptx_lval.int_value = BRA_OP; return OPCODE; brx TC; ptx_lval.int_value = BRX_OP; return OPCODE; brev TC; ptx_lval.int_value = BREV_OP; return OPCODE; brkpt TC; ptx_lval.int_value = BRKPT_OP; return OPCODE; -bsmad TC; ptx_lval.int_value = BSMAD_OP; return OPCODE; +mma TC; ptx_lval.int_value = MMA_OP; return OPCODE; call TC; BEGIN(NOT_OPCODE); ptx_lval.int_value = CALL_OP; return OPCODE; // blocking opcode token in case the callee has the same name as an opcode callp TC; BEGIN(NOT_OPCODE); ptx_lval.int_value = CALLP_OP; return OPCODE; clz TC; ptx_lval.int_value = CLZ_OP; return OPCODE; diff --git a/src/cuda-sim/ptx_ir.h b/src/cuda-sim/ptx_ir.h index 4c10373..0601b97 100644 --- a/src/cuda-sim/ptx_ir.h +++ b/src/cuda-sim/ptx_ir.h @@ -948,6 +948,31 @@ public: assert( m_operands.size() > 3 ); return m_operands[3]; } + const operand_info &src4() const + { + assert( m_operands.size() > 4 ); + return m_operands[4]; + } + const operand_info &src5() const + { + assert( m_operands.size() > 5 ); + return m_operands[5]; + } + const operand_info &src6() const + { + assert( m_operands.size() > 6 ); + return m_operands[6]; + } + const operand_info &src7() const + { + assert( m_operands.size() > 7 ); + return m_operands[7]; + } + const operand_info &src8() const + { + assert( m_operands.size() > 8 ); + return m_operands[8]; + } const operand_info &operand_lookup( unsigned n ) const { diff --git a/src/gpgpu-sim/scoreboard.cc b/src/gpgpu-sim/scoreboard.cc index f412054..b538fdf 100644 --- a/src/gpgpu-sim/scoreboard.cc +++ b/src/gpgpu-sim/scoreboard.cc @@ -146,6 +146,10 @@ bool Scoreboard::checkCollision( unsigned wid, const class inst_t *inst ) const if(inst->in[1] > 0) inst_regs.insert(inst->in[1]); if(inst->in[2] > 0) inst_regs.insert(inst->in[2]); if(inst->in[3] > 0) inst_regs.insert(inst->in[3]); + if(inst->in[3] > 0) inst_regs.insert(inst->in[4]); + if(inst->in[3] > 0) inst_regs.insert(inst->in[5]); + if(inst->in[3] > 0) inst_regs.insert(inst->in[6]); + if(inst->in[3] > 0) inst_regs.insert(inst->in[7]); if(inst->pred > 0) inst_regs.insert(inst->pred); if(inst->ar1 > 0) inst_regs.insert(inst->ar1); if(inst->ar2 > 0) inst_regs.insert(inst->ar2); diff --git a/src/gpgpusim_entrypoint.cc b/src/gpgpusim_entrypoint.cc index 04845e7..a6d7eb4 100644 --- a/src/gpgpusim_entrypoint.cc +++ b/src/gpgpusim_entrypoint.cc @@ -134,7 +134,6 @@ void *gpgpu_sim_thread_concurrent(void*) gpgpu_cuda_ptx_sim_main_func(*kernel); g_the_gpu->finish_functional_sim(kernel); } - //performance simulation if( g_the_gpu->active() ) { g_the_gpu->cycle(); -- cgit v1.3 From fa0089a5d3a86ef348fae9a83a862f5219892bab Mon Sep 17 00:00:00 2001 From: aamir Date: Wed, 30 May 2018 23:07:44 -0700 Subject: adding code for wmma_ld_impl, error at decode space --- cuda-kernels/.tensor_core_ptx.swp | Bin 0 -> 16384 bytes cuda-kernels/Makefile | 3 +- cuda-kernels/tensor_core | Bin 48541 -> 2750968 bytes cuda-kernels/tensor_core_ptx | 171 +++++++++++ src/Makefile | 2 +- src/cuda-sim/.ptx.y.swp | Bin 0 -> 36864 bytes src/cuda-sim/Makefile | 2 +- src/cuda-sim/instructions.cc | 72 ++++- src/cuda-sim/ptx.y~ | 608 ++++++++++++++++++++++++++++++++++++++ src/cuda-sim/ptx_ir.h | 4 +- src/cuda-sim/ptx_sim.h | 9 + src/gpgpu-sim/Makefile | 2 +- src/intersim2/Makefile | 2 +- 13 files changed, 865 insertions(+), 10 deletions(-) create mode 100644 cuda-kernels/.tensor_core_ptx.swp create mode 100644 cuda-kernels/tensor_core_ptx create mode 100644 src/cuda-sim/.ptx.y.swp create mode 100644 src/cuda-sim/ptx.y~ (limited to 'src/gpgpu-sim') diff --git a/cuda-kernels/.tensor_core_ptx.swp b/cuda-kernels/.tensor_core_ptx.swp new file mode 100644 index 0000000..6d7bad4 Binary files /dev/null and b/cuda-kernels/.tensor_core_ptx.swp differ diff --git a/cuda-kernels/Makefile b/cuda-kernels/Makefile index 51a7760..673460f 100755 --- a/cuda-kernels/Makefile +++ b/cuda-kernels/Makefile @@ -1,5 +1,6 @@ all: tensor_core.cu - nvcc -arch=sm_70 -lcudart -g -o tensor_core tensor_core.cu + nvcc --gpu-architecture=compute_70 --gpu-code=compute_70 -lcudart -g -o tensor_core tensor_core.cu +# nvcc -arch=sm_70 -lcudart -g -o tensor_core tensor_core.cu .PHONY: clean: diff --git a/cuda-kernels/tensor_core b/cuda-kernels/tensor_core index b25f3d9..cb53851 100755 Binary files a/cuda-kernels/tensor_core and b/cuda-kernels/tensor_core differ diff --git a/cuda-kernels/tensor_core_ptx b/cuda-kernels/tensor_core_ptx new file mode 100644 index 0000000..36074cb --- /dev/null +++ b/cuda-kernels/tensor_core_ptx @@ -0,0 +1,171 @@ +// +// Generated by NVIDIA NVVM Compiler +// +// Compiler Build ID: CL-22781540 +// Cuda compilation tools, release 9.0, V9.0.176 +// Based on LLVM 3.4svn +// + +.version 6.0 +.target sm_70 +.address_size 64 + + // .globl _Z12wmma_exampleP6__halfS0_Pfiiiff +.extern .func (.param .b32 func_retval0) vprintf +( + .param .b64 vprintf_param_0, + .param .b64 vprintf_param_1 +) +; +.global .align 16 .b8 $str[9] = {99, 108, 111, 99, 107, 61, 37, 100, 0}; + +.visible .entry _Z12wmma_exampleP6__halfS0_Pfiiiff( + .param .u64 _Z12wmma_exampleP6__halfS0_Pfiiiff_param_0, + .param .u64 _Z12wmma_exampleP6__halfS0_Pfiiiff_param_1, + .param .u64 _Z12wmma_exampleP6__halfS0_Pfiiiff_param_2, + .param .u32 _Z12wmma_exampleP6__halfS0_Pfiiiff_param_3, + .param .u32 _Z12wmma_exampleP6__halfS0_Pfiiiff_param_4, + .param .u32 _Z12wmma_exampleP6__halfS0_Pfiiiff_param_5, + .param .f32 _Z12wmma_exampleP6__halfS0_Pfiiiff_param_6, + .param .f32 _Z12wmma_exampleP6__halfS0_Pfiiiff_param_7 +) +{ + .local .align 8 .b8 __local_depot0[8]; + .reg .b64 %SP; + .reg .b64 %SPL; + .reg .pred %p<6>; + .reg .f32 %f<34>; + .reg .b32 %r<38>; + .reg .b64 %rd<18>; + + + mov.u64 %rd17, __local_depot0; + cvta.local.u64 %SP, %rd17; + ld.param.u64 %rd1, [_Z12wmma_exampleP6__halfS0_Pfiiiff_param_0]; + ld.param.u64 %rd2, [_Z12wmma_exampleP6__halfS0_Pfiiiff_param_1]; + ld.param.u64 %rd3, [_Z12wmma_exampleP6__halfS0_Pfiiiff_param_2]; + ld.param.u32 %r4, [_Z12wmma_exampleP6__halfS0_Pfiiiff_param_3]; + ld.param.u32 %r7, [_Z12wmma_exampleP6__halfS0_Pfiiiff_param_4]; + ld.param.u32 %r5, [_Z12wmma_exampleP6__halfS0_Pfiiiff_param_5]; + // inline asm + mov.u32 %r6, %clock; + // inline asm + mov.u32 %r8, %ntid.x; + mov.u32 %r9, %ctaid.x; + mov.u32 %r10, %tid.x; + mad.lo.s32 %r11, %r8, %r9, %r10; + mov.u32 %r12, WARP_SZ; + div.u32 %r13, %r11, %r12; + mov.u32 %r14, %ntid.y; + mov.u32 %r15, %ctaid.y; + mov.u32 %r16, %tid.y; + mad.lo.s32 %r17, %r14, %r15, %r16; + shl.b32 %r2, %r13, 4; + shl.b32 %r3, %r17, 4; + setp.lt.s32 %p1, %r2, %r4; + setp.gt.s32 %p2, %r5, 0; + and.pred %p3, %p1, %p2; + setp.lt.s32 %p4, %r3, %r7; + and.pred %p5, %p3, %p4; + mov.f32 %f26, 0f00000000; + mov.f32 %f27, %f26; + mov.f32 %f28, %f26; + mov.f32 %f29, %f26; + mov.f32 %f30, %f26; + mov.f32 %f31, %f26; + mov.f32 %f32, %f26; + mov.f32 %f33, %f26; + @!%p5 bra BB0_2; + bra.uni BB0_1; + +BB0_1: + mul.wide.s32 %rd4, %r2, 2; + add.s64 %rd5, %rd1, %rd4; + wmma.load.a.sync.row.m16n16k16.f16 {%r18, %r19, %r20, %r21, %r22, %r23, %r24, %r25}, [%rd5], %r4; + mul.wide.s32 %rd6, %r3, 2; + add.s64 %rd7, %rd2, %rd6; + wmma.load.b.sync.col.m16n16k16.f16 {%r26, %r27, %r28, %r29, %r30, %r31, %r32, %r33}, [%rd7], %r5; + mov.f32 %f25, 0f00000000; + wmma.mma.sync.row.col.m16n16k16.f32.f32 {%f33, %f32, %f31, %f30, %f29, %f28, %f27, %f26}, {%r18, %r19, %r20, %r21, %r22, %r23, %r24, %r25}, {%r26, %r27, %r28, %r29, %r30, %r31, %r32, %r33}, {%f25, %f25, %f25, %f25, %f25, %f25, %f25, %f25}; + +BB0_2: + add.u64 %rd8, %SP, 0; + cvta.to.local.u64 %rd9, %rd8; + mul.lo.s32 %r35, %r3, %r4; + cvt.s64.s32 %rd10, %r35; + cvt.s64.s32 %rd11, %r2; + add.s64 %rd12, %rd10, %rd11; + shl.b64 %rd13, %rd12, 2; + add.s64 %rd14, %rd3, %rd13; + wmma.store.d.sync.col.m16n16k16.f32 [%rd14], {%f33, %f32, %f31, %f30, %f29, %f28, %f27, %f26}, %r4; + // inline asm + mov.u32 %r34, %clock; + // inline asm + sub.s32 %r36, %r34, %r6; + st.local.u32 [%rd9], %r36; + mov.u64 %rd15, $str; + cvta.global.u64 %rd16, %rd15; + // Callseq Start 0 + { + .reg .b32 temp_param_reg; + // } + .param .b64 param0; + st.param.b64 [param0+0], %rd16; + .param .b64 param1; + st.param.b64 [param1+0], %rd8; + .param .b32 retval0; + call.uni (retval0), + vprintf, + ( + param0, + param1 + ); + ld.param.b32 %r37, [retval0+0]; + + //{ + }// Callseq End 0 + ret; +} + + // .globl _Z17convertFp32ToFp16P6__halfPfi +.visible .entry _Z17convertFp32ToFp16P6__halfPfi( + .param .u64 _Z17convertFp32ToFp16P6__halfPfi_param_0, + .param .u64 _Z17convertFp32ToFp16P6__halfPfi_param_1, + .param .u32 _Z17convertFp32ToFp16P6__halfPfi_param_2 +) +{ + .reg .pred %p<2>; + .reg .b16 %rs<2>; + .reg .f32 %f<2>; + .reg .b32 %r<6>; + .reg .b64 %rd<9>; + + + ld.param.u64 %rd1, [_Z17convertFp32ToFp16P6__halfPfi_param_0]; + ld.param.u64 %rd2, [_Z17convertFp32ToFp16P6__halfPfi_param_1]; + ld.param.u32 %r2, [_Z17convertFp32ToFp16P6__halfPfi_param_2]; + mov.u32 %r3, %ntid.x; + mov.u32 %r4, %ctaid.x; + mov.u32 %r5, %tid.x; + mad.lo.s32 %r1, %r4, %r3, %r5; + setp.ge.s32 %p1, %r1, %r2; + @%p1 bra BB1_2; + + cvta.to.global.u64 %rd3, %rd2; + mul.wide.s32 %rd4, %r1, 4; + add.s64 %rd5, %rd3, %rd4; + ld.global.f32 %f1, [%rd5]; + // inline asm + { cvt.rn.f16.f32 %rs1, %f1;} + + // inline asm + cvta.to.global.u64 %rd6, %rd1; + mul.wide.s32 %rd7, %r1, 2; + add.s64 %rd8, %rd6, %rd7; + st.global.u16 [%rd8], %rs1; + +BB1_2: + ret; +} + + diff --git a/src/Makefile b/src/Makefile index 6001669..09194f3 100644 --- a/src/Makefile +++ b/src/Makefile @@ -46,7 +46,7 @@ ifeq ($(TRACE),1) endif ifneq ($(DEBUG),1) - OPTFLAGS += -O3 + OPTFLAGS += -O0 else CXXFLAGS += endif diff --git a/src/cuda-sim/.ptx.y.swp b/src/cuda-sim/.ptx.y.swp new file mode 100644 index 0000000..c8a83b5 Binary files /dev/null and b/src/cuda-sim/.ptx.y.swp differ diff --git a/src/cuda-sim/Makefile b/src/cuda-sim/Makefile index 999dad7..a65e8e1 100644 --- a/src/cuda-sim/Makefile +++ b/src/cuda-sim/Makefile @@ -42,7 +42,7 @@ include ../../version_detection.mk OUTPUT_DIR=$(SIM_OBJ_FILES_DIR)/cuda-sim -OPT := -O3 -g3 -Wall -Wno-unused-function -Wno-sign-compare +OPT := -O0 -g3 -Wall -Wno-unused-function -Wno-sign-compare ifeq ($(DEBUG),1) OPT := -g3 -Wall -Wno-unused-function -Wno-sign-compare endif diff --git a/src/cuda-sim/instructions.cc b/src/cuda-sim/instructions.cc index 446cdbf..16f33c6 100644 --- a/src/cuda-sim/instructions.cc +++ b/src/cuda-sim/instructions.cc @@ -643,6 +643,33 @@ void ptx_thread_info::set_vector_operand_values( const operand_info &dst, m_last_set_operand_value = data1; } +void ptx_thread_info::set_wmma_vector_operand_values( const operand_info &dst, + const ptx_reg_t &data1, + const ptx_reg_t &data2, + const ptx_reg_t &data3, + const ptx_reg_t &data4, + const ptx_reg_t &data5, + const ptx_reg_t &data6, + const ptx_reg_t &data7, + const ptx_reg_t &data8 ) +{ + unsigned num_elements = dst.get_vect_nelem(); + if (num_elements > 7) { + set_reg(dst.vec_symbol(0), data1); + set_reg(dst.vec_symbol(1), data2); + set_reg(dst.vec_symbol(2), data3); + set_reg(dst.vec_symbol(3), data4); + set_reg(dst.vec_symbol(4), data5); + set_reg(dst.vec_symbol(5), data6); + set_reg(dst.vec_symbol(6), data7); + set_reg(dst.vec_symbol(7), data8); + } + else{ + printf("error:set_wmma_vector_operands"); + } + + m_last_set_operand_value = data1; +} #define my_abs(a) (((a)<0)?(-a):(a)) @@ -1493,9 +1520,6 @@ unsigned trunc(unsigned num, unsigned precision) { } return num; } -void mma_ld_impl( const ptx_instruction *pI, core_t *core, warp_inst_t inst ) -{ -} void mma_st_impl( const ptx_instruction *pI, core_t *core, warp_inst_t inst ) { } @@ -2595,6 +2619,48 @@ void ldu_impl( const ptx_instruction *pI, ptx_thread_info *thread ) { ld_exec(pI,thread); } +void mma_ld_impl( const ptx_instruction *pI, core_t *core, warp_inst_t inst ) +{ + const operand_info &dst = pI->dst(); + const operand_info &src1 = pI->src1(); + const operand_info &src2 = pI->src2(); + + unsigned type = pI->get_type(); + + int tid = inst.warp_id_func() * core->get_warp_size(); + int thrd; + ptx_thread_info *thread; + thread = core->get_thread_info()[tid]; + + ptx_reg_t src1_data = thread->get_operand_value(src1, dst, type, thread, 1); + + ptx_reg_t data; + memory_space_t space = pI->get_space(); + + memory_space *mem = NULL; + addr_t addr = src1_data.u32; + + decode_space(space,thread,src1,mem,addr); + + size_t size; + int t; + data.u64=0; + type_info_key::type_decode(type,size,t); + ptx_reg_t data1, data2, data3, data4; + ptx_reg_t data5, data6, data7, data8; + mem->read(addr,size/8,&data1.s64); + mem->read(addr+size/8,size/8,&data2.s64); + mem->read(addr+2*size/8,size/8,&data3.s64); + mem->read(addr+3*size/8,size/8,&data4.s64); + mem->read(addr+4*size/8,size/8,&data5.s64); + mem->read(addr+5*size/8,size/8,&data6.s64); + mem->read(addr+6*size/8,size/8,&data7.s64); + mem->read(addr+7*size/8,size/8,&data8.s64); + thread->set_wmma_vector_operand_values(dst,data1,data2,data3,data4,data5,data6,data7,data8); + + thread->m_last_effective_address = addr; + thread->m_last_memory_space = space; +} void lg2_impl( const ptx_instruction *pI, ptx_thread_info *thread ) { diff --git a/src/cuda-sim/ptx.y~ b/src/cuda-sim/ptx.y~ new file mode 100644 index 0000000..0710ecd --- /dev/null +++ b/src/cuda-sim/ptx.y~ @@ -0,0 +1,608 @@ +/* +Copyright (c) 2009-2011, Tor M. Aamodt +The University of British Columbia +All rights reserved. + +Redistribution and use in source and binary forms, with or without +modification, are permitted provided that the following conditions are met: + +Redistributions of source code must retain the above copyright notice, this +list of conditions and the following disclaimer. +Redistributions in binary form must reproduce the above copyright notice, this +list of conditions and the following disclaimer in the documentation and/or +other materials provided with the distribution. +Neither the name of The University of British Columbia nor the names of its +contributors may be used to endorse or promote products derived from this +software without specific prior written permission. + +THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND +ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED +WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE +DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE +FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL +DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR +SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER +CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, +OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE +OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. +*/ + +%union { + double double_value; + float float_value; + int int_value; + char * string_value; + void * ptr_value; +} + +%token STRING +%token OPCODE +%token WMMA_DIRECTIVE +%token LAYOUT +%token CONFIGURATION +%token ALIGN_DIRECTIVE +%token BRANCHTARGETS_DIRECTIVE +%token BYTE_DIRECTIVE +%token CALLPROTOTYPE_DIRECTIVE +%token CALLTARGETS_DIRECTIVE +%token CONST_DIRECTIVE +%token CONSTPTR_DIRECTIVE +%token PTR_DIRECTIVE +%token ENTRY_DIRECTIVE +%token EXTERN_DIRECTIVE +%token WEAK_DIRECTIVE +%token FILE_DIRECTIVE +%token FUNC_DIRECTIVE +%token GLOBAL_DIRECTIVE +%token LOCAL_DIRECTIVE +%token LOC_DIRECTIVE +%token MAXNCTAPERSM_DIRECTIVE +%token MAXNNREG_DIRECTIVE +%token MAXNTID_DIRECTIVE +%token MINNCTAPERSM_DIRECTIVE +%token PARAM_DIRECTIVE +%token PRAGMA_DIRECTIVE +%token REG_DIRECTIVE +%token REQNTID_DIRECTIVE +%token SECTION_DIRECTIVE +%token SHARED_DIRECTIVE +%token SREG_DIRECTIVE +%token SSTARR_DIRECTIVE +%token STRUCT_DIRECTIVE +%token SURF_DIRECTIVE +%token TARGET_DIRECTIVE +%token TEX_DIRECTIVE +%token UNION_DIRECTIVE +%token VERSION_DIRECTIVE +%token ADDRESS_SIZE_DIRECTIVE +%token VISIBLE_DIRECTIVE +%token WEAK_DIRECTIVE +%token IDENTIFIER +%token INT_OPERAND +%token FLOAT_OPERAND +%token DOUBLE_OPERAND +%token S8_TYPE +%token S16_TYPE +%token S32_TYPE +%token S64_TYPE +%token U8_TYPE +%token U16_TYPE +%token U32_TYPE +%token U64_TYPE +%token F16_TYPE +%token F32_TYPE +%token F64_TYPE +%token FF64_TYPE +%token B8_TYPE +%token B16_TYPE +%token B32_TYPE +%token B64_TYPE +%token BB64_TYPE +%token BB128_TYPE +%token PRED_TYPE +%token TEXREF_TYPE +%token SAMPLERREF_TYPE +%token SURFREF_TYPE +%token V2_TYPE +%token V3_TYPE +%token V4_TYPE +%token COMMA +%token PRED +%token HALF_OPTION +%token EXTP_OPTION +%token EQ_OPTION +%token NE_OPTION +%token LT_OPTION +%token LE_OPTION +%token GT_OPTION +%token GE_OPTION +%token LO_OPTION +%token LS_OPTION +%token HI_OPTION +%token HS_OPTION +%token EQU_OPTION +%token NEU_OPTION +%token LTU_OPTION +%token LEU_OPTION +%token GTU_OPTION +%token GEU_OPTION +%token NUM_OPTION +%token NAN_OPTION +%token CF_OPTION +%token SF_OPTION +%token NSF_OPTION +%token LEFT_SQUARE_BRACKET +%token RIGHT_SQUARE_BRACKET +%token WIDE_OPTION +%token SPECIAL_REGISTER +%token MINUS +%token PLUS +%token COLON +%token SEMI_COLON +%token EXCLAMATION +%token PIPE +%token RIGHT_BRACE +%token LEFT_BRACE +%token EQUALS +%token PERIOD +%token BACKSLASH +%token DIMENSION_MODIFIER +%token RN_OPTION +%token RZ_OPTION +%token RM_OPTION +%token RP_OPTION +%token RNI_OPTION +%token RZI_OPTION +%token RMI_OPTION +%token RPI_OPTION +%token UNI_OPTION +%token GEOM_MODIFIER_1D +%token GEOM_MODIFIER_2D +%token GEOM_MODIFIER_3D +%token SAT_OPTION +%token FTZ_OPTION +%token NEG_OPTION +%token SYNC_OPTION +%token RED_OPTION +%token ARRIVE_OPTION +%token ATOMIC_POPC +%token ATOMIC_AND +%token ATOMIC_OR +%token ATOMIC_XOR +%token ATOMIC_CAS +%token ATOMIC_EXCH +%token ATOMIC_ADD +%token ATOMIC_INC +%token ATOMIC_DEC +%token ATOMIC_MIN +%token ATOMIC_MAX +%token LEFT_ANGLE_BRACKET +%token RIGHT_ANGLE_BRACKET +%token LEFT_PAREN +%token RIGHT_PAREN +%token APPROX_OPTION +%token FULL_OPTION +%token ANY_OPTION +%token ALL_OPTION +%token BALLOT_OPTION +%token GLOBAL_OPTION +%token CTA_OPTION +%token SYS_OPTION +%token EXIT_OPTION +%token ABS_OPTION +%token TO_OPTION +%token CA_OPTION; +%token CG_OPTION; +%token CS_OPTION; +%token LU_OPTION; +%token CV_OPTION; +%token WB_OPTION; +%token WT_OPTION; +%token NC_OPTION; +%token UP_OPTION; +%token DOWN_OPTION; +%token BFLY_OPTION; +%token IDX_OPTION; + +%type function_decl_header +%type function_decl + +%{ + #include "ptx_parser.h" + #include + #include + #include + void syntax_not_implemented(); + extern int g_func_decl; + int ptx_lex(void); + int ptx_error(const char *); +%} + +%% + +input: /* empty */ + | input directive_statement + | input function_defn + | input function_decl + ; + +function_defn: function_decl { set_symtab($1); func_header(".skip"); } statement_block { end_function(); } + | function_decl { set_symtab($1); } block_spec_list { func_header(".skip"); } statement_block { end_function(); } + ; + +block_spec: MAXNTID_DIRECTIVE INT_OPERAND COMMA INT_OPERAND COMMA INT_OPERAND {func_header_info_int(".maxntid", $2); + func_header_info_int(",", $4); + func_header_info_int(",", $6); } + | MINNCTAPERSM_DIRECTIVE INT_OPERAND { func_header_info_int(".minnctapersm", $2); printf("GPGPU-Sim: Warning: .minnctapersm ignored. \n"); } + | MAXNCTAPERSM_DIRECTIVE INT_OPERAND { func_header_info_int(".maxnctapersm", $2); printf("GPGPU-Sim: Warning: .maxnctapersm ignored. \n"); } + ; + +block_spec_list: block_spec + | block_spec_list block_spec + ; + +function_decl: function_decl_header LEFT_PAREN { start_function($1); func_header_info("(");} param_entry RIGHT_PAREN {func_header_info(")");} function_ident_param { $$ = reset_symtab(); } + | function_decl_header { start_function($1); } function_ident_param { $$ = reset_symtab(); } + | function_decl_header { start_function($1); add_function_name(""); g_func_decl=0; $$ = reset_symtab(); } + ; + +function_ident_param: IDENTIFIER { add_function_name($1); } LEFT_PAREN {func_header_info("(");} param_list RIGHT_PAREN { g_func_decl=0; func_header_info(")"); } + | IDENTIFIER { add_function_name($1); g_func_decl=0; } + ; + +function_decl_header: ENTRY_DIRECTIVE { $$ = 1; g_func_decl=1; func_header(".entry"); } + | VISIBLE_DIRECTIVE ENTRY_DIRECTIVE { $$ = 1; g_func_decl=1; func_header(".entry"); } + | WEAK_DIRECTIVE ENTRY_DIRECTIVE { $$ = 1; g_func_decl=1; func_header(".entry"); } + | FUNC_DIRECTIVE { $$ = 0; g_func_decl=1; func_header(".func"); } + | VISIBLE_DIRECTIVE FUNC_DIRECTIVE { $$ = 0; g_func_decl=1; func_header(".func"); } + | WEAK_DIRECTIVE FUNC_DIRECTIVE { $$ = 0; g_func_decl=1; func_header(".func"); } + | EXTERN_DIRECTIVE FUNC_DIRECTIVE { $$ = 2; g_func_decl=1; func_header(".func"); } + | WEAK_DIRECTIVE FUNC_DIRECTIVE { $$ = 0; g_func_decl=1; func_header(".func"); } + ; + +param_list: /*empty*/ + | param_entry { add_directive(); } + | param_list COMMA {func_header_info(",");} param_entry { add_directive(); } + +param_entry: PARAM_DIRECTIVE { add_space_spec(param_space_unclassified,0); } variable_spec ptr_spec identifier_spec { add_function_arg(); } + | REG_DIRECTIVE { add_space_spec(reg_space,0); } variable_spec identifier_spec { add_function_arg(); } + +ptr_spec: /*empty*/ + | PTR_DIRECTIVE ptr_space_spec ptr_align_spec + | PTR_DIRECTIVE ptr_align_spec + +ptr_space_spec: GLOBAL_DIRECTIVE { add_ptr_spec(global_space); } + | LOCAL_DIRECTIVE { add_ptr_spec(local_space); } + | SHARED_DIRECTIVE { add_ptr_spec(shared_space); } + +ptr_align_spec: ALIGN_DIRECTIVE INT_OPERAND + +statement_block: LEFT_BRACE statement_list RIGHT_BRACE + +statement_list: directive_statement { add_directive(); } + | instruction_statement { add_instruction(); } + | statement_list directive_statement { add_directive(); } + | statement_list instruction_statement { add_instruction(); } + | statement_list {start_inst_group();} statement_block {end_inst_group();} + | {start_inst_group();} statement_block {end_inst_group();} + ; + +directive_statement: variable_declaration SEMI_COLON + | VERSION_DIRECTIVE DOUBLE_OPERAND { add_version_info($2, 0); } + | VERSION_DIRECTIVE DOUBLE_OPERAND PLUS { add_version_info($2,1); } + | ADDRESS_SIZE_DIRECTIVE INT_OPERAND {/*Do nothing*/} + | TARGET_DIRECTIVE IDENTIFIER COMMA IDENTIFIER { target_header2($2,$4); } + | TARGET_DIRECTIVE IDENTIFIER COMMA IDENTIFIER COMMA IDENTIFIER { target_header3($2,$4,$6); } + | TARGET_DIRECTIVE IDENTIFIER { target_header($2); } + | FILE_DIRECTIVE INT_OPERAND STRING { add_file($2,$3); } + | FILE_DIRECTIVE INT_OPERAND STRING COMMA INT_OPERAND COMMA INT_OPERAND { add_file($2,$3); } + | LOC_DIRECTIVE INT_OPERAND INT_OPERAND INT_OPERAND + | PRAGMA_DIRECTIVE STRING SEMI_COLON { add_pragma($2); } + | function_decl SEMI_COLON {/*Do nothing*/} + ; + +variable_declaration: variable_spec identifier_list { add_variables(); } + | variable_spec identifier_spec EQUALS initializer_list { add_variables(); } + | variable_spec identifier_spec EQUALS literal_operand { add_variables(); } + | CONSTPTR_DIRECTIVE IDENTIFIER COMMA IDENTIFIER COMMA INT_OPERAND { add_constptr($2, $4, $6); } + ; + +variable_spec: var_spec_list { set_variable_type(); } + +identifier_list: identifier_spec + | identifier_list COMMA identifier_spec; + +identifier_spec: IDENTIFIER { add_identifier($1,0,NON_ARRAY_IDENTIFIER); func_header_info($1);} + | IDENTIFIER LEFT_ANGLE_BRACKET INT_OPERAND RIGHT_ANGLE_BRACKET { func_header_info($1); func_header_info_int("<", $3); func_header_info(">"); + int i,lbase,l; + char *id = NULL; + lbase = strlen($1); + for( i=0; i < $3; i++ ) { + l = lbase + (int)log10(i+1)+10; + id = (char*) malloc(l); + snprintf(id,l,"%s%u",$1,i); + add_identifier(id,0,NON_ARRAY_IDENTIFIER); + } + free($1); + } + | IDENTIFIER LEFT_SQUARE_BRACKET RIGHT_SQUARE_BRACKET { add_identifier($1,0,ARRAY_IDENTIFIER_NO_DIM); func_header_info($1); func_header_info("["); func_header_info("]");} + | IDENTIFIER LEFT_SQUARE_BRACKET INT_OPERAND RIGHT_SQUARE_BRACKET { add_identifier($1,$3,ARRAY_IDENTIFIER); func_header_info($1); func_header_info_int("[",$3); func_header_info("]");} + ; + +var_spec_list: var_spec + | var_spec_list var_spec; + +var_spec: space_spec + | type_spec + | align_spec + | EXTERN_DIRECTIVE { add_extern_spec(); } + | WEAK_DIRECTIVE + ; + +align_spec: ALIGN_DIRECTIVE INT_OPERAND { add_alignment_spec($2); } + +space_spec: REG_DIRECTIVE { add_space_spec(reg_space,0); } + | SREG_DIRECTIVE { add_space_spec(reg_space,0); } + | addressable_spec + ; + +addressable_spec: CONST_DIRECTIVE { add_space_spec(const_space,$1); } + | GLOBAL_DIRECTIVE { add_space_spec(global_space,0); } + | LOCAL_DIRECTIVE { add_space_spec(local_space,0); } + | PARAM_DIRECTIVE { add_space_spec(param_space_unclassified,0); } + | SHARED_DIRECTIVE { add_space_spec(shared_space,0); } + | SSTARR_DIRECTIVE { add_space_spec(sstarr_space,0); } + | SURF_DIRECTIVE { add_space_spec(surf_space,0); } + | TEX_DIRECTIVE { add_space_spec(tex_space,0); } + ; + +type_spec: scalar_type + | vector_spec scalar_type + ; + +vector_spec: V2_TYPE { add_option(V2_TYPE); func_header_info(".v2");} + | V3_TYPE { add_option(V3_TYPE); func_header_info(".v3");} + | V4_TYPE { add_option(V4_TYPE); func_header_info(".v4");} + ; + +scalar_type: S8_TYPE { add_scalar_type_spec( S8_TYPE ); } + | S16_TYPE { add_scalar_type_spec( S16_TYPE ); } + | S32_TYPE { add_scalar_type_spec( S32_TYPE ); } + | S64_TYPE { add_scalar_type_spec( S64_TYPE ); } + | U8_TYPE { add_scalar_type_spec( U8_TYPE ); } + | U16_TYPE { add_scalar_type_spec( U16_TYPE ); } + | U32_TYPE { add_scalar_type_spec( U32_TYPE ); } + | U64_TYPE { add_scalar_type_spec( U64_TYPE ); } + | F16_TYPE { add_scalar_type_spec( F16_TYPE ); } + | F32_TYPE { add_scalar_type_spec( F32_TYPE ); } + | F64_TYPE { add_scalar_type_spec( F64_TYPE ); } + | FF64_TYPE { add_scalar_type_spec( FF64_TYPE ); } + | B8_TYPE { add_scalar_type_spec( B8_TYPE ); } + | B16_TYPE { add_scalar_type_spec( B16_TYPE ); } + | B32_TYPE { add_scalar_type_spec( B32_TYPE ); } + | B64_TYPE { add_scalar_type_spec( B64_TYPE ); } + | BB64_TYPE { add_scalar_type_spec( BB64_TYPE ); } + | BB128_TYPE { add_scalar_type_spec( BB128_TYPE ); } + | PRED_TYPE { add_scalar_type_spec( PRED_TYPE ); } + | TEXREF_TYPE { add_scalar_type_spec( TEXREF_TYPE ); } + | SAMPLERREF_TYPE { add_scalar_type_spec( SAMPLERREF_TYPE ); } + | SURFREF_TYPE { add_scalar_type_spec( SURFREF_TYPE ); } + ; + +initializer_list: LEFT_BRACE literal_list RIGHT_BRACE { add_array_initializer(); } + | LEFT_BRACE initializer_list RIGHT_BRACE { syntax_not_implemented(); } + +literal_list: literal_operand + | literal_list COMMA literal_operand; + +instruction_statement: instruction SEMI_COLON + | IDENTIFIER COLON { add_label($1); } + | pred_spec instruction SEMI_COLON; + +instruction: opcode_spec LEFT_PAREN operand RIGHT_PAREN { set_return(); } COMMA operand COMMA LEFT_PAREN operand_list RIGHT_PAREN + | opcode_spec operand COMMA LEFT_PAREN operand_list RIGHT_PAREN + | opcode_spec operand COMMA LEFT_PAREN RIGHT_PAREN + | opcode_spec operand_list + | opcode_spec + ; + +opcode_spec: OPCODE { add_opcode($1); } option_list + | OPCODE { add_opcode($1); } + +pred_spec: PRED IDENTIFIER { add_pred($2,0, -1); } + | PRED EXCLAMATION IDENTIFIER { add_pred($3,1, -1); } + | PRED IDENTIFIER LT_OPTION { add_pred($2,0,1); } + | PRED IDENTIFIER EQ_OPTION { add_pred($2,0,2); } + | PRED IDENTIFIER LE_OPTION { add_pred($2,0,3); } + | PRED IDENTIFIER NE_OPTION { add_pred($2,0,5); } + | PRED IDENTIFIER GE_OPTION { add_pred($2,0,6); } + | PRED IDENTIFIER EQU_OPTION { add_pred($2,0,10); } + | PRED IDENTIFIER GTU_OPTION { add_pred($2,0,12); } + | PRED IDENTIFIER NEU_OPTION { add_pred($2,0,13); } + | PRED IDENTIFIER CF_OPTION { add_pred($2,0,17); } + | PRED IDENTIFIER SF_OPTION { add_pred($2,0,19); } + | PRED IDENTIFIER NSF_OPTION { add_pred($2,0,28); } + ; + +option_list: option + | option option_list ; + +option: type_spec + | compare_spec + | addressable_spec + | rounding_mode + | wmma_spec + | SYNC_OPTION { add_option(SYNC_OPTION); } + | ARRIVE_OPTION { add_option(ARRIVE_OPTION); } + | RED_OPTION { add_option(RED_OPTION); } + | UNI_OPTION { add_option(UNI_OPTION); } + | WIDE_OPTION { add_option(WIDE_OPTION); } + | ANY_OPTION { add_option(ANY_OPTION); } + | ALL_OPTION { add_option(ALL_OPTION); } + | BALLOT_OPTION { add_option(BALLOT_OPTION); } + | GLOBAL_OPTION { add_option(GLOBAL_OPTION); } + | CTA_OPTION { add_option(CTA_OPTION); } + | SYS_OPTION { add_option(SYS_OPTION); } + | GEOM_MODIFIER_1D { add_option(GEOM_MODIFIER_1D); } + | GEOM_MODIFIER_2D { add_option(GEOM_MODIFIER_2D); } + | GEOM_MODIFIER_3D { add_option(GEOM_MODIFIER_3D); } + | SAT_OPTION { add_option(SAT_OPTION); } + | FTZ_OPTION { add_option(FTZ_OPTION); } + | NEG_OPTION { add_option(NEG_OPTION); } + | APPROX_OPTION { add_option(APPROX_OPTION); } + | FULL_OPTION { add_option(FULL_OPTION); } + | EXIT_OPTION { add_option(EXIT_OPTION); } + | ABS_OPTION { add_option(ABS_OPTION); } + | atomic_operation_spec ; + | TO_OPTION { add_option(TO_OPTION); } + | HALF_OPTION { add_option(HALF_OPTION); } + | EXTP_OPTION { add_option(EXTP_OPTION); } + | CA_OPTION { add_option(CA_OPTION); } + | CG_OPTION { add_option(CG_OPTION); } + | CS_OPTION { add_option(CS_OPTION); } + | LU_OPTION { add_option(LU_OPTION); } + | CV_OPTION { add_option(CV_OPTION); } + | WB_OPTION { add_option(WB_OPTION); } + | WT_OPTION { add_option(WT_OPTION); } + | NC_OPTION { add_option(NC_OPTION); } + | UP_OPTION { add_option(UP_OPTION); } + | DOWN_OPTION { add_option(DOWN_OPTION); } + | BFLY_OPTION { add_option(BFLY_OPTION); } + | IDX_OPTION { add_option(IDX_OPTION); } + ; + +atomic_operation_spec: ATOMIC_AND { add_option(ATOMIC_AND); } + | ATOMIC_POPC { add_option(ATOMIC_POPC); } + | ATOMIC_OR { add_option(ATOMIC_OR); } + | ATOMIC_XOR { add_option(ATOMIC_XOR); } + | ATOMIC_CAS { add_option(ATOMIC_CAS); } + | ATOMIC_EXCH { add_option(ATOMIC_EXCH); } + | ATOMIC_ADD { add_option(ATOMIC_ADD); } + | ATOMIC_INC { add_option(ATOMIC_INC); } + | ATOMIC_DEC { add_option(ATOMIC_DEC); } + | ATOMIC_MIN { add_option(ATOMIC_MIN); } + | ATOMIC_MAX { add_option(ATOMIC_MAX); } + ; + +rounding_mode: floating_point_rounding_mode + | integer_rounding_mode; + + +floating_point_rounding_mode: RN_OPTION { add_option(RN_OPTION); } + | RZ_OPTION { add_option(RZ_OPTION); } + | RM_OPTION { add_option(RM_OPTION); } + | RP_OPTION { add_option(RP_OPTION); } + ; + +integer_rounding_mode: RNI_OPTION { add_option(RNI_OPTION); } + | RZI_OPTION { add_option(RZI_OPTION); } + | RMI_OPTION { add_option(RMI_OPTION); } + | RPI_OPTION { add_option(RPI_OPTION); } + ; + +compare_spec:EQ_OPTION { add_option(EQ_OPTION); } + | NE_OPTION { add_option(NE_OPTION); } + | LT_OPTION { add_option(LT_OPTION); } + | LE_OPTION { add_option(LE_OPTION); } + | GT_OPTION { add_option(GT_OPTION); } + | GE_OPTION { add_option(GE_OPTION); } + | LO_OPTION { add_option(LO_OPTION); } + | LS_OPTION { add_option(LS_OPTION); } + | HI_OPTION { add_option(HI_OPTION); } + | HS_OPTION { add_option(HS_OPTION); } + | EQU_OPTION { add_option(EQU_OPTION); } + | NEU_OPTION { add_option(NEU_OPTION); } + | LTU_OPTION { add_option(LTU_OPTION); } + | LEU_OPTION { add_option(LEU_OPTION); } + | GTU_OPTION { add_option(GTU_OPTION); } + | GEU_OPTION { add_option(GEU_OPTION); } + | NUM_OPTION { add_option(NUM_OPTION); } + | NAN_OPTION { add_option(NAN_OPTION); } + ; + +wmma_spec: WMMA_DIRECTIVE LAYOUT CONFIGURATION{add_wmma_option($1);add_wmma_option($2);add_wmma_option($3);} + | WMMA_DIRECTIVE LAYOUT LAYOUT CONFIGURATION{add_wmma_option($1);add_wmma_option($2),add_wmma_option($3),add_wmma_option($4)} + | WMMA_DIRECTIVE LAYOUT CONFIGURATION ptr_space_spec{add_wmma_option($1);add_wmma_option($2),add_wmma_option($3),add_wmma_option($4)} + | WMMA_DIRECTIVE LAYOUT LAYOUT CONFIGURATION ptr_space_spec{add_wmma_option($1);add_wmma_option($2),add_wmma_option($3),add_wmma_option($4)} + ; + +operand_list: operand + | operand COMMA operand_list; + +operand: IDENTIFIER { add_scalar_operand( $1 ); } + | EXCLAMATION IDENTIFIER { add_neg_pred_operand( $2 ); } + | MINUS IDENTIFIER { add_scalar_operand( $2 ); change_operand_neg(); } + | memory_operand + | literal_operand + | builtin_operand + | vector_operand + | MINUS vector_operand { change_operand_neg(); } + | tex_operand + | IDENTIFIER PLUS INT_OPERAND { add_address_operand($1,$3); } + | IDENTIFIER LO_OPTION { add_scalar_operand( $1 ); change_operand_lohi(1);} + | MINUS IDENTIFIER LO_OPTION { add_scalar_operand( $2 ); change_operand_lohi(1); change_operand_neg();} + | IDENTIFIER HI_OPTION { add_scalar_operand( $1 ); change_operand_lohi(2);} + | MINUS IDENTIFIER HI_OPTION { add_scalar_operand( $2 ); change_operand_lohi(2); change_operand_neg();} + | IDENTIFIER PIPE IDENTIFIER { add_2vector_operand($1,$3); change_double_operand_type(-1);} + | IDENTIFIER PIPE IDENTIFIER LO_OPTION { add_2vector_operand($1,$3); change_double_operand_type(-1); change_operand_lohi(1);} + | IDENTIFIER PIPE IDENTIFIER HI_OPTION { add_2vector_operand($1,$3); change_double_operand_type(-1); change_operand_lohi(2);} + | IDENTIFIER BACKSLASH IDENTIFIER { add_2vector_operand($1,$3); change_double_operand_type(-3);} + | IDENTIFIER BACKSLASH IDENTIFIER LO_OPTION { add_2vector_operand($1,$3); change_double_operand_type(-3); change_operand_lohi(1);} + | IDENTIFIER BACKSLASH IDENTIFIER HI_OPTION { add_2vector_operand($1,$3); change_double_operand_type(-3); change_operand_lohi(2);} + ; + +vector_operand: LEFT_BRACE IDENTIFIER COMMA IDENTIFIER RIGHT_BRACE { add_2vector_operand($2,$4); } + | LEFT_BRACE IDENTIFIER COMMA IDENTIFIER COMMA IDENTIFIER RIGHT_BRACE { add_3vector_operand($2,$4,$6); } + | LEFT_BRACE IDENTIFIER COMMA IDENTIFIER COMMA IDENTIFIER COMMA IDENTIFIER RIGHT_BRACE { add_4vector_operand($2,$4,$6,$8); } + | LEFT_BRACE IDENTIFIER COMMA IDENTIFIER COMMA IDENTIFIER COMMA IDENTIFIER COMMA IDENTIFIER COMMA IDENTIFIER COMMA IDENTIFIER COMMA IDENTIFIER RIGHT_BRACE { add_8vector_operand($2,$4,$6,$8,$10,$12,$14,$16); } + | LEFT_BRACE IDENTIFIER RIGHT_BRACE { add_1vector_operand($2); } + ; + +tex_operand: LEFT_SQUARE_BRACKET IDENTIFIER COMMA { add_scalar_operand($2); } + vector_operand + RIGHT_SQUARE_BRACKET + ; + +builtin_operand: SPECIAL_REGISTER DIMENSION_MODIFIER { add_builtin_operand($1,$2); } + | SPECIAL_REGISTER { add_builtin_operand($1,-1); } + ; + +memory_operand : LEFT_SQUARE_BRACKET address_expression RIGHT_SQUARE_BRACKET { add_memory_operand(); } + | IDENTIFIER LEFT_SQUARE_BRACKET address_expression RIGHT_SQUARE_BRACKET { add_memory_operand(); change_memory_addr_space($1); } + | IDENTIFIER LEFT_SQUARE_BRACKET literal_operand RIGHT_SQUARE_BRACKET { change_memory_addr_space($1); } + | IDENTIFIER LEFT_SQUARE_BRACKET twin_operand RIGHT_SQUARE_BRACKET { change_memory_addr_space($1); add_memory_operand();} + | MINUS memory_operand { change_operand_neg(); } + ; + +twin_operand : IDENTIFIER PLUS IDENTIFIER { add_double_operand($1,$3); change_double_operand_type(1); } + | IDENTIFIER PLUS IDENTIFIER LO_OPTION { add_double_operand($1,$3); change_double_operand_type(1); change_operand_lohi(1); } + | IDENTIFIER PLUS IDENTIFIER HI_OPTION { add_double_operand($1,$3); change_double_operand_type(1); change_operand_lohi(2); } + | IDENTIFIER PLUS EQUALS IDENTIFIER { add_double_operand($1,$4); change_double_operand_type(2); } + | IDENTIFIER PLUS EQUALS IDENTIFIER LO_OPTION { add_double_operand($1,$4); change_double_operand_type(2); change_operand_lohi(1); } + | IDENTIFIER PLUS EQUALS IDENTIFIER HI_OPTION { add_double_operand($1,$4); change_double_operand_type(2); change_operand_lohi(2); } + | IDENTIFIER PLUS EQUALS INT_OPERAND { add_address_operand($1,$4); change_double_operand_type(3); } + ; + +literal_operand : INT_OPERAND { add_literal_int($1); } + | FLOAT_OPERAND { add_literal_float($1); } + | DOUBLE_OPERAND { add_literal_double($1); } + ; + +address_expression: IDENTIFIER { add_address_operand($1,0); } + | IDENTIFIER LO_OPTION { add_address_operand($1,0); change_operand_lohi(1);} + | IDENTIFIER HI_OPTION { add_address_operand($1,0); change_operand_lohi(2); } + | IDENTIFIER PLUS INT_OPERAND { add_address_operand($1,$3); } + | INT_OPERAND { add_address_operand2($1); } + ; + +%% + +extern int ptx_lineno; +extern const char *g_filename; + +void syntax_not_implemented() +{ + printf("Parse error (%s:%u): this syntax is not (yet) implemented:\n",g_filename,ptx_lineno); + ptx_error(NULL); + abort(); +} diff --git a/src/cuda-sim/ptx_ir.h b/src/cuda-sim/ptx_ir.h index 833f175..16cc975 100644 --- a/src/cuda-sim/ptx_ir.h +++ b/src/cuda-sim/ptx_ir.h @@ -1078,7 +1078,7 @@ public: int membar_level() const { return m_membar_level; } bool has_memory_read() const { - if( m_opcode == LD_OP || m_opcode == LDU_OP || m_opcode == TEX_OP ) + if( m_opcode == LD_OP || m_opcode == LDU_OP || m_opcode == TEX_OP|| m_opcode==MMA_LD_OP) return true; // Check PTXPlus operand type below // Source operands are memory operands @@ -1090,7 +1090,7 @@ public: return false; } bool has_memory_write() const { - if( m_opcode == ST_OP ) return true; + if( m_opcode == ST_OP || m_opcode==MMA_ST_OP ) return true; // Check PTXPlus operand type below // Destination operand is a memory operand ptx_instruction::const_iterator op=op_iter_begin(); diff --git a/src/cuda-sim/ptx_sim.h b/src/cuda-sim/ptx_sim.h index 05acf20..403ce5b 100644 --- a/src/cuda-sim/ptx_sim.h +++ b/src/cuda-sim/ptx_sim.h @@ -303,6 +303,15 @@ public: const ptx_reg_t &data2, const ptx_reg_t &data3, const ptx_reg_t &data4 ); + void set_wmma_vector_operand_values( const operand_info &dst, + const ptx_reg_t &data1, + const ptx_reg_t &data2, + const ptx_reg_t &data3, + const ptx_reg_t &data4, + const ptx_reg_t &data5, + const ptx_reg_t &data6, + const ptx_reg_t &data7, + const ptx_reg_t &data8 ); function_info *func_info() { diff --git a/src/gpgpu-sim/Makefile b/src/gpgpu-sim/Makefile index f10a8a4..4f77699 100644 --- a/src/gpgpu-sim/Makefile +++ b/src/gpgpu-sim/Makefile @@ -48,7 +48,7 @@ ifeq ($(GNUC_CPP0X), 1) endif ifneq ($(DEBUG),1) - OPTFLAGS += -O3 + OPTFLAGS += -O0 else CXXFLAGS += endif diff --git a/src/intersim2/Makefile b/src/intersim2/Makefile index bd42000..4ef21ac 100644 --- a/src/intersim2/Makefile +++ b/src/intersim2/Makefile @@ -44,7 +44,7 @@ endif CPPFLAGS += -Wall $(INCPATH) $(DEFINE) ifneq ($(DEBUG),1) -CPPFLAGS += -O3 +CPPFLAGS += -O0 endif CPPFLAGS += -g CPPFLAGS += -fPIC -- cgit v1.3 From 774e28201d239433357db4cbd00c0a62273773cc Mon Sep 17 00:00:00 2001 From: Aamir Raihan Date: Fri, 8 Jun 2018 13:56:54 -0700 Subject: regression passed --- src/Makefile | 2 +- src/cuda-sim/Makefile | 2 +- src/cuda-sim/instructions.cc | 11 ----------- src/cuda-sim/ptx.l | 4 ++-- src/gpgpu-sim/Makefile | 2 +- src/intersim2/Makefile | 2 +- 6 files changed, 6 insertions(+), 17 deletions(-) (limited to 'src/gpgpu-sim') diff --git a/src/Makefile b/src/Makefile index 09194f3..6001669 100644 --- a/src/Makefile +++ b/src/Makefile @@ -46,7 +46,7 @@ ifeq ($(TRACE),1) endif ifneq ($(DEBUG),1) - OPTFLAGS += -O0 + OPTFLAGS += -O3 else CXXFLAGS += endif diff --git a/src/cuda-sim/Makefile b/src/cuda-sim/Makefile index d08a696..e8977a6 100644 --- a/src/cuda-sim/Makefile +++ b/src/cuda-sim/Makefile @@ -42,7 +42,7 @@ include ../../version_detection.mk OUTPUT_DIR=$(SIM_OBJ_FILES_DIR)/cuda-sim -OPT := -O0 -g3 -Wall -Wno-unused-function -Wno-sign-compare +OPT := -O3 -g3 -Wall -Wno-unused-function -Wno-sign-compare ifeq ($(DEBUG),1) OPT := -g3 -Wall -Wno-unused-function -Wno-sign-compare endif diff --git a/src/cuda-sim/instructions.cc b/src/cuda-sim/instructions.cc index 59415f4..92c8529 100644 --- a/src/cuda-sim/instructions.cc +++ b/src/cuda-sim/instructions.cc @@ -835,17 +835,6 @@ void add_impl( const ptx_instruction *pI, ptx_thread_info *thread ) unsigned i_type = pI->get_type(); src1_data = thread->get_operand_value(src1, dst, i_type, thread, 1); src2_data = thread->get_operand_value(src2, dst, i_type, thread, 1); - //unsigned warpId_aa,warp_size_aa; - //warpId_aa = pI->warp_id(); - //warp_size_aa=32; - //dim3 t=thread->get_tid(); - //unsigned tid_aa=warp_size_aa*warpId_aa; - - ptx_thread_info *thread2; - thread2=thread; - src1_data = thread2->get_operand_value(src1, dst, i_type, thread2, 1); - src2_data = thread2->get_operand_value(src2, dst, i_type, thread2, 1); - unsigned rounding_mode = pI->rounding_mode(); int orig_rm = fegetround(); diff --git a/src/cuda-sim/ptx.l b/src/cuda-sim/ptx.l index f67dddd..03d6838 100644 --- a/src/cuda-sim/ptx.l +++ b/src/cuda-sim/ptx.l @@ -68,7 +68,7 @@ bra TC; ptx_lval.int_value = BRA_OP; return OPCODE; brx TC; ptx_lval.int_value = BRX_OP; return OPCODE; brev TC; ptx_lval.int_value = BREV_OP; return OPCODE; brkpt TC; ptx_lval.int_value = BRKPT_OP; return OPCODE; -wmma\.mma TC; ptx_lval.int_value = MMA_OP; return OPCODE; +wmma TC; ptx_lval.int_value = MMA_OP; return OPCODE; wmma\.load TC; ptx_lval.int_value = MMA_LD_OP; return OPCODE; wmma\.store TC; ptx_lval.int_value = MMA_ST_OP; return OPCODE; @@ -160,7 +160,7 @@ breakaddr TC; ptx_lval.int_value = BREAKADDR_OP; return OPCODE; \.b\.sync TC; ptx_lval.int_value = LOAD_B; return WMMA_DIRECTIVE; \.c\.sync TC; ptx_lval.int_value = LOAD_C; return WMMA_DIRECTIVE; \.d\.sync TC; ptx_lval.int_value = STORE_D; return WMMA_DIRECTIVE; -\.sync TC;ptx_lval.int_value=MMA; return WMMA_DIRECTIVE; +\.mma\.sync TC;ptx_lval.int_value=MMA; return WMMA_DIRECTIVE; \.row TC; ptx_lval.int_value = ROW; return LAYOUT; \.col TC; ptx_lval.int_value = COL; return LAYOUT; \.m16n16k16 TC; ptx_lval.int_value = M16N16K16; return CONFIGURATION; diff --git a/src/gpgpu-sim/Makefile b/src/gpgpu-sim/Makefile index 4f77699..f10a8a4 100644 --- a/src/gpgpu-sim/Makefile +++ b/src/gpgpu-sim/Makefile @@ -48,7 +48,7 @@ ifeq ($(GNUC_CPP0X), 1) endif ifneq ($(DEBUG),1) - OPTFLAGS += -O0 + OPTFLAGS += -O3 else CXXFLAGS += endif diff --git a/src/intersim2/Makefile b/src/intersim2/Makefile index 4ef21ac..bd42000 100644 --- a/src/intersim2/Makefile +++ b/src/intersim2/Makefile @@ -44,7 +44,7 @@ endif CPPFLAGS += -Wall $(INCPATH) $(DEFINE) ifneq ($(DEBUG),1) -CPPFLAGS += -O0 +CPPFLAGS += -O3 endif CPPFLAGS += -g CPPFLAGS += -fPIC -- cgit v1.3 From 262663ac90d2aa801d6af1eb9bf8a75ee9a5bb18 Mon Sep 17 00:00:00 2001 From: negargoli93 Date: Wed, 20 Jun 2018 18:20:05 -0700 Subject: Tensor core timing model --- cuda-kernels/gpgpusim.config | 7 +++-- src/abstract_hardware_model.h | 2 ++ src/cuda-sim/cuda-sim.cc | 5 ++-- src/gpgpu-sim/gpu-sim.cc | 16 ++++++++-- src/gpgpu-sim/scoreboard.cc | 40 +++++++++++++++++++------ src/gpgpu-sim/shader.cc | 70 +++++++++++++++++++++++++++++++++++++++---- src/gpgpu-sim/shader.h | 47 ++++++++++++++++++++++++++--- 7 files changed, 163 insertions(+), 24 deletions(-) (limited to 'src/gpgpu-sim') diff --git a/cuda-kernels/gpgpusim.config b/cuda-kernels/gpgpusim.config index 306d7f9..69a110f 100755 --- a/cuda-kernels/gpgpusim.config +++ b/cuda-kernels/gpgpusim.config @@ -33,10 +33,10 @@ # ID_OC_SP,ID_OC_SFU,ID_OC_MEM,OC_EX_SP,OC_EX_SFU,OC_EX_MEM,EX_WB ## Pascal GP102 has 4 SP SIMD units and 1 SFU unit ## we need to scale the number of pipeline registers to be equal to the number of SP units --gpgpu_pipeline_widths 4,1,1,4,1,1,6 +-gpgpu_pipeline_widths 4,1,1,1,4,1,1,1,,6 -gpgpu_num_sp_units 4 -gpgpu_num_sfu_units 1 - +-gpgpu_num_tensor_core_units 1 # Instruction latencies and initiation intervals # "ADD,MAX,MUL,MAD,DIV" # SFU is 32-width in pascal, then dp units initiation is 1 cycle @@ -72,11 +72,14 @@ ## larger operand collectors and reg_banks are needed for the 4 warp schedulers and 4 SIMD units -gpgpu_operand_collector_num_units_sp 20 -gpgpu_operand_collector_num_units_sfu 4 +-gpgpu_operand_collector_num_units_tensor_core 24 -gpgpu_operand_collector_num_units_mem 8 -gpgpu_operand_collector_num_in_ports_sp 4 -gpgpu_operand_collector_num_out_ports_sp 4 -gpgpu_operand_collector_num_in_ports_sfu 1 -gpgpu_operand_collector_num_out_ports_sfu 1 +-gpgpu_operand_collector_num_in_ports_tensor_core 1 +-gpgpu_operand_collector_num_out_ports_tensor_core 1 -gpgpu_operand_collector_num_in_ports_mem 1 -gpgpu_operand_collector_num_out_ports_mem 1 # gpgpu_num_reg_banks should be increased to 32, but it gives an error! diff --git a/src/abstract_hardware_model.h b/src/abstract_hardware_model.h index e00c941..e2e116e 100644 --- a/src/abstract_hardware_model.h +++ b/src/abstract_hardware_model.h @@ -77,6 +77,7 @@ enum uarch_op_t { NO_OP=-1, ALU_OP=1, SFU_OP, + TENSOR_CORE_OP, ALU_SFU_OP, LOAD_OP, STORE_OP, @@ -133,6 +134,7 @@ enum operation_pipeline_t { UNKOWN_OP, SP__OP, SFU__OP, + TENSOR_CORE__OP, MEM__OP }; typedef enum operation_pipeline_t operation_pipeline; diff --git a/src/cuda-sim/cuda-sim.cc b/src/cuda-sim/cuda-sim.cc index 7552acf..6da0840 100644 --- a/src/cuda-sim/cuda-sim.cc +++ b/src/cuda-sim/cuda-sim.cc @@ -520,7 +520,7 @@ void ptx_instruction::set_mul_div_or_other_archop(){ sp_op=FP_EXP_OP; break; default: - if(op==ALU_OP) + if((op==ALU_OP)||(op==TENSOR_CORE_OP)) sp_op=FP__OP; break; @@ -542,7 +542,7 @@ void ptx_instruction::set_mul_div_or_other_archop(){ sp_op=INT_DIV_OP; break; default: - if(op==ALU_OP) + if((op==ALU_OP)||(op==TENSOR_CORE_OP)) sp_op=INT__OP; break; } @@ -795,6 +795,7 @@ void ptx_instruction::set_opcode_and_latency() case MMA_OP: latency = 64; initiation_interval = 64; + op=TENSOR_CORE_OP; break; case SHFL_OP: latency = 32; diff --git a/src/gpgpu-sim/gpu-sim.cc b/src/gpgpu-sim/gpu-sim.cc index 3bd1892..cc23051 100644 --- a/src/gpgpu-sim/gpu-sim.cc +++ b/src/gpgpu-sim/gpu-sim.cc @@ -306,6 +306,9 @@ void shader_core_config::reg_options(class OptionParser * opp) option_parser_register(opp, "-gpgpu_operand_collector_num_units_sfu", OPT_INT32, &gpgpu_operand_collector_num_units_sfu, "number of collector units (default = 4)", "4"); + option_parser_register(opp, "-gpgpu_operand_collector_num_units_tensor_core", OPT_INT32, &gpgpu_operand_collector_num_units_tensor_core, + "number of collector units (default = 4)", + "4"); option_parser_register(opp, "-gpgpu_operand_collector_num_units_mem", OPT_INT32, &gpgpu_operand_collector_num_units_mem, "number of collector units (default = 2)", "2"); @@ -318,6 +321,9 @@ void shader_core_config::reg_options(class OptionParser * opp) option_parser_register(opp, "-gpgpu_operand_collector_num_in_ports_sfu", OPT_INT32, &gpgpu_operand_collector_num_in_ports_sfu, "number of collector unit in ports (default = 1)", "1"); + option_parser_register(opp, "-gpgpu_operand_collector_num_in_ports_tensor_core", OPT_INT32, &gpgpu_operand_collector_num_in_ports_tensor_core, + "number of collector unit in ports (default = 1)", + "1"); option_parser_register(opp, "-gpgpu_operand_collector_num_in_ports_mem", OPT_INT32, &gpgpu_operand_collector_num_in_ports_mem, "number of collector unit in ports (default = 1)", "1"); @@ -330,6 +336,9 @@ void shader_core_config::reg_options(class OptionParser * opp) option_parser_register(opp, "-gpgpu_operand_collector_num_out_ports_sfu", OPT_INT32, &gpgpu_operand_collector_num_out_ports_sfu, "number of collector unit in ports (default = 1)", "1"); + option_parser_register(opp, "-gpgpu_operand_collector_num_out_ports_tensor_core", OPT_INT32, &gpgpu_operand_collector_num_out_ports_tensor_core, + "number of collector unit in ports (default = 1)", + "1"); option_parser_register(opp, "-gpgpu_operand_collector_num_out_ports_mem", OPT_INT32, &gpgpu_operand_collector_num_out_ports_mem, "number of collector unit in ports (default = 1)", "1"); @@ -350,14 +359,17 @@ void shader_core_config::reg_options(class OptionParser * opp) "1"); option_parser_register(opp, "-gpgpu_pipeline_widths", OPT_CSTR, &pipeline_widths_string, "Pipeline widths " - "ID_OC_SP,ID_OC_SFU,ID_OC_MEM,OC_EX_SP,OC_EX_SFU,OC_EX_MEM,EX_WB", - "1,1,1,1,1,1,1" ); + "ID_OC_SP,ID_OC_SFU,ID_OC_TENSOR_CORE,ID_OC_MEM,OC_EX_SP,OC_EX_SFU,OC_EX_TENSOR_CORE,OC_EX_MEM,EX_WB", + "1,1,1,1,1,1,1,1,1" ); option_parser_register(opp, "-gpgpu_num_sp_units", OPT_INT32, &gpgpu_num_sp_units, "Number of SP units (default=1)", "1"); option_parser_register(opp, "-gpgpu_num_sfu_units", OPT_INT32, &gpgpu_num_sfu_units, "Number of SF units (default=1)", "1"); + option_parser_register(opp, "-gpgpu_num_tensor_core_units", OPT_INT32, &gpgpu_num_tensor_core_units, + "Number of tensor_core units (default=1)", + "1"); option_parser_register(opp, "-gpgpu_num_mem_units", OPT_INT32, &gpgpu_num_mem_units, "Number if ldst units (default=1) WARNING: not hooked up to anything", "1"); diff --git a/src/gpgpu-sim/scoreboard.cc b/src/gpgpu-sim/scoreboard.cc index b538fdf..4d1b43a 100644 --- a/src/gpgpu-sim/scoreboard.cc +++ b/src/gpgpu-sim/scoreboard.cc @@ -82,7 +82,7 @@ const bool Scoreboard::islongop (unsigned warp_id,unsigned regnum) { void Scoreboard::reserveRegisters(const class warp_inst_t* inst) { - for( unsigned r=0; r < 4; r++) { + for( unsigned r=0; r < 8; r++) { if(inst->out[r] > 0) { reserveRegister(inst->warp_id(), inst->out[r]); SHADER_DPRINTF( SCOREBOARD, @@ -100,7 +100,7 @@ void Scoreboard::reserveRegisters(const class warp_inst_t* inst) inst->space.get_type() == param_space_local || inst->space.get_type() == param_space_unclassified || inst->space.get_type() == tex_space)){ - for ( unsigned r=0; r<4; r++) { + for ( unsigned r=0; r<8; r++) { if(inst->out[r] > 0) { SHADER_DPRINTF( SCOREBOARD, "New longopreg marked - warp:%d, reg: %d\n", @@ -115,7 +115,7 @@ void Scoreboard::reserveRegisters(const class warp_inst_t* inst) // Release registers for an instruction void Scoreboard::releaseRegisters(const class warp_inst_t *inst) { - for( unsigned r=0; r < 4; r++) { + for( unsigned r=0; r < 8; r++) { if(inst->out[r] > 0) { SHADER_DPRINTF( SCOREBOARD, "Register Released - warp:%d, reg: %d\n", @@ -142,15 +142,37 @@ bool Scoreboard::checkCollision( unsigned wid, const class inst_t *inst ) const if(inst->out[1] > 0) inst_regs.insert(inst->out[1]); if(inst->out[2] > 0) inst_regs.insert(inst->out[2]); if(inst->out[3] > 0) inst_regs.insert(inst->out[3]); - if(inst->in[0] > 0) inst_regs.insert(inst->in[0]); + if(inst->out[4] > 0) inst_regs.insert(inst->out[4]); + if(inst->out[5] > 0) inst_regs.insert(inst->out[5]); + if(inst->out[6] > 0) inst_regs.insert(inst->out[6]); + if(inst->out[7] > 0) inst_regs.insert(inst->out[7]); + + if(inst->in[0] > 0) inst_regs.insert(inst->in[0]); if(inst->in[1] > 0) inst_regs.insert(inst->in[1]); if(inst->in[2] > 0) inst_regs.insert(inst->in[2]); if(inst->in[3] > 0) inst_regs.insert(inst->in[3]); - if(inst->in[3] > 0) inst_regs.insert(inst->in[4]); - if(inst->in[3] > 0) inst_regs.insert(inst->in[5]); - if(inst->in[3] > 0) inst_regs.insert(inst->in[6]); - if(inst->in[3] > 0) inst_regs.insert(inst->in[7]); - if(inst->pred > 0) inst_regs.insert(inst->pred); + if(inst->in[4] > 0) inst_regs.insert(inst->in[4]); + if(inst->in[5] > 0) inst_regs.insert(inst->in[5]); + if(inst->in[6] > 0) inst_regs.insert(inst->in[6]); + if(inst->in[7] > 0) inst_regs.insert(inst->in[7]); + if(inst->in[8] > 0) inst_regs.insert(inst->in[8]); + if(inst->in[9] > 0) inst_regs.insert(inst->in[9]); + if(inst->in[10] > 0) inst_regs.insert(inst->in[10]); + if(inst->in[11] > 0) inst_regs.insert(inst->in[11]); + if(inst->in[12] > 0) inst_regs.insert(inst->in[12]); + if(inst->in[13] > 0) inst_regs.insert(inst->in[13]); + if(inst->in[14] > 0) inst_regs.insert(inst->in[14]); + if(inst->in[15] > 0) inst_regs.insert(inst->in[15]); + if(inst->in[16] > 0) inst_regs.insert(inst->in[16]); + if(inst->in[17] > 0) inst_regs.insert(inst->in[17]); + if(inst->in[18] > 0) inst_regs.insert(inst->in[18]); + if(inst->in[19] > 0) inst_regs.insert(inst->in[19]); + if(inst->in[20] > 0) inst_regs.insert(inst->in[20]); + if(inst->in[21] > 0) inst_regs.insert(inst->in[21]); + if(inst->in[22] > 0) inst_regs.insert(inst->in[22]); + if(inst->in[23] > 0) inst_regs.insert(inst->in[23]); + + if(inst->pred > 0) inst_regs.insert(inst->pred); if(inst->ar1 > 0) inst_regs.insert(inst->ar1); if(inst->ar2 > 0) inst_regs.insert(inst->ar2); diff --git a/src/gpgpu-sim/shader.cc b/src/gpgpu-sim/shader.cc index 92cdb5b..fcac755 100644 --- a/src/gpgpu-sim/shader.cc +++ b/src/gpgpu-sim/shader.cc @@ -148,6 +148,7 @@ shader_core_ctx::shader_core_ctx( class gpgpu_sim *gpu, &m_warp, &m_pipeline_reg[ID_OC_SP], &m_pipeline_reg[ID_OC_SFU], + &m_pipeline_reg[ID_OC_TENSOR_CORE], &m_pipeline_reg[ID_OC_MEM], i ) @@ -162,6 +163,7 @@ shader_core_ctx::shader_core_ctx( class gpgpu_sim *gpu, &m_warp, &m_pipeline_reg[ID_OC_SP], &m_pipeline_reg[ID_OC_SFU], + &m_pipeline_reg[ID_OC_TENSOR_CORE], &m_pipeline_reg[ID_OC_MEM], i, config->gpgpu_scheduler_string @@ -177,6 +179,7 @@ shader_core_ctx::shader_core_ctx( class gpgpu_sim *gpu, &m_warp, &m_pipeline_reg[ID_OC_SP], &m_pipeline_reg[ID_OC_SFU], + &m_pipeline_reg[ID_OC_TENSOR_CORE], &m_pipeline_reg[ID_OC_MEM], i ) @@ -191,6 +194,7 @@ shader_core_ctx::shader_core_ctx( class gpgpu_sim *gpu, &m_warp, &m_pipeline_reg[ID_OC_SP], &m_pipeline_reg[ID_OC_SFU], + &m_pipeline_reg[ID_OC_TENSOR_CORE], &m_pipeline_reg[ID_OC_MEM], i, config->gpgpu_scheduler_string @@ -211,9 +215,10 @@ shader_core_ctx::shader_core_ctx( class gpgpu_sim *gpu, } //op collector configuration - enum { SP_CUS, SFU_CUS, MEM_CUS, GEN_CUS }; + enum { SP_CUS, SFU_CUS, TENSOR_CORE_CUS, MEM_CUS, GEN_CUS }; m_operand_collector.add_cu_set(SP_CUS, m_config->gpgpu_operand_collector_num_units_sp, m_config->gpgpu_operand_collector_num_out_ports_sp); m_operand_collector.add_cu_set(SFU_CUS, m_config->gpgpu_operand_collector_num_units_sfu, m_config->gpgpu_operand_collector_num_out_ports_sfu); + m_operand_collector.add_cu_set(TENSOR_CORE_CUS, m_config->gpgpu_operand_collector_num_units_tensor_core, m_config->gpgpu_operand_collector_num_out_ports_tensor_core); m_operand_collector.add_cu_set(MEM_CUS, m_config->gpgpu_operand_collector_num_units_mem, m_config->gpgpu_operand_collector_num_out_ports_mem); m_operand_collector.add_cu_set(GEN_CUS, m_config->gpgpu_operand_collector_num_units_gen, m_config->gpgpu_operand_collector_num_out_ports_gen); @@ -237,7 +242,17 @@ shader_core_ctx::shader_core_ctx( class gpgpu_sim *gpu, m_operand_collector.add_port(in_ports,out_ports,cu_sets); in_ports.clear(),out_ports.clear(),cu_sets.clear(); } + + for (unsigned i = 0; i < m_config->gpgpu_operand_collector_num_in_ports_tensor_core; i++) { + in_ports.push_back(&m_pipeline_reg[ID_OC_TENSOR_CORE]); + out_ports.push_back(&m_pipeline_reg[OC_EX_TENSOR_CORE]); + cu_sets.push_back((unsigned)TENSOR_CORE_CUS); + cu_sets.push_back((unsigned)GEN_CUS); + m_operand_collector.add_port(in_ports,out_ports,cu_sets); + in_ports.clear(),out_ports.clear(),cu_sets.clear(); + } + for (unsigned i = 0; i < m_config->gpgpu_operand_collector_num_in_ports_mem; i++) { in_ports.push_back(&m_pipeline_reg[ID_OC_MEM]); out_ports.push_back(&m_pipeline_reg[OC_EX_MEM]); @@ -251,9 +266,11 @@ shader_core_ctx::shader_core_ctx( class gpgpu_sim *gpu, for (unsigned i = 0; i < m_config->gpgpu_operand_collector_num_in_ports_gen; i++) { in_ports.push_back(&m_pipeline_reg[ID_OC_SP]); in_ports.push_back(&m_pipeline_reg[ID_OC_SFU]); + in_ports.push_back(&m_pipeline_reg[ID_OC_TENSOR_CORE]); in_ports.push_back(&m_pipeline_reg[ID_OC_MEM]); out_ports.push_back(&m_pipeline_reg[OC_EX_SP]); out_ports.push_back(&m_pipeline_reg[OC_EX_SFU]); + out_ports.push_back(&m_pipeline_reg[OC_EX_TENSOR_CORE]); out_ports.push_back(&m_pipeline_reg[OC_EX_MEM]); cu_sets.push_back((unsigned)GEN_CUS); m_operand_collector.add_port(in_ports,out_ports,cu_sets); @@ -263,7 +280,7 @@ shader_core_ctx::shader_core_ctx( class gpgpu_sim *gpu, m_operand_collector.init( m_config->gpgpu_num_reg_banks, this ); // execute - m_num_function_units = m_config->gpgpu_num_sp_units + m_config->gpgpu_num_sfu_units + 1; // sp_unit, sfu, ldst_unit + m_num_function_units = m_config->gpgpu_num_sp_units + m_config->gpgpu_num_sfu_units + m_config->gpgpu_num_tensor_core_units + 1; // sp_unit, sfu, ldst_unit //m_dispatch_port = new enum pipeline_stage_name_t[ m_num_function_units ]; //m_issue_port = new enum pipeline_stage_name_t[ m_num_function_units ]; @@ -280,6 +297,12 @@ shader_core_ctx::shader_core_ctx( class gpgpu_sim *gpu, m_dispatch_port.push_back(ID_OC_SFU); m_issue_port.push_back(OC_EX_SFU); } + + for (int k = 0; k < m_config->gpgpu_num_tensor_core_units; k++) { + m_fu.push_back(new tensor_core( &m_pipeline_reg[EX_WB], m_config, this )); + m_dispatch_port.push_back(ID_OC_TENSOR_CORE); + m_issue_port.push_back(OC_EX_TENSOR_CORE); + } m_ldst_unit = new ldst_unit( m_icnt, m_mem_fetch_allocator, this, &m_operand_collector, m_scoreboard, config, mem_config, stats, shader_id, tpc_id ); m_fu.push_back(m_ldst_unit); @@ -886,7 +909,8 @@ void scheduler_unit::cycle() } else { bool sp_pipe_avail = m_sp_out->has_free(); bool sfu_pipe_avail = m_sfu_out->has_free(); - if( sp_pipe_avail && (pI->op != SFU_OP) ) { + bool tensor_core_pipe_avail = m_tensor_core_out->has_free(); + if( sp_pipe_avail && (pI->op != SFU_OP) && (pI->op != TENSOR_CORE_OP)) { //Jin: special for CDP api if(pI->m_is_cdp && !warp(warp_id).m_cdp_dummy) { @@ -918,6 +942,14 @@ void scheduler_unit::cycle() issued_inst=true; warp_inst_issued = true; } + } + else if ( (pI->op == TENSOR_CORE_OP) ) { + if( tensor_core_pipe_avail ) { + m_shader->issue_warp(*m_tensor_core_out,pI,active_mask,warp_id); + issued++; + issued_inst=true; + warp_inst_issued = true; + } } } } else { SCHED_DPRINTF( "Warp (warp_id %u, dynamic_warp_id %u) fails scoreboard\n", @@ -1083,10 +1115,11 @@ swl_scheduler::swl_scheduler ( shader_core_stats* stats, shader_core_ctx* shader std::vector* warp, register_set* sp_out, register_set* sfu_out, + register_set* tensor_core_out, register_set* mem_out, int id, char* config_string ) - : scheduler_unit ( stats, shader, scoreboard, simt, warp, sp_out, sfu_out, mem_out, id ) + : scheduler_unit ( stats, shader, scoreboard, simt, warp, sp_out, sfu_out, tensor_core_out, mem_out, id ) { unsigned m_prioritization_readin; int ret = sscanf( config_string, @@ -1509,6 +1542,13 @@ sfu:: sfu( register_set* result_port, const shader_core_config *config,shader_c m_name = "SFU"; } +tensor_core:: tensor_core( register_set* result_port, const shader_core_config *config,shader_core_ctx *core ) + : pipelined_simd_unit(result_port,config,config->max_tensor_core_latency,core) +{ + m_name = "TENSOR_CORE"; +} + + void sfu::issue( register_set& source_reg ) { warp_inst_t** ready_reg = source_reg.get_ready(); @@ -1519,6 +1559,17 @@ void sfu::issue( register_set& source_reg ) pipelined_simd_unit::issue(source_reg); } +void tensor_core::issue( register_set& source_reg ) +{ + warp_inst_t** ready_reg = source_reg.get_ready(); + //m_core->incexecstat((*ready_reg)); + + (*ready_reg)->op_pipe= TENSOR_CORE__OP; + m_core->incsfu_stat(m_core->get_config()->warp_size,(*ready_reg)->latency); + pipelined_simd_unit::issue(source_reg); +} + + void ldst_unit::active_lanes_in_pipeline(){ unsigned active_count=pipelined_simd_unit::get_active_lanes_in_pipeline(); assert(active_count<=m_core->get_config()->warp_size); @@ -1540,6 +1591,15 @@ void sfu::active_lanes_in_pipeline(){ m_core->incfumemactivelanes_stat(active_count); } +void tensor_core::active_lanes_in_pipeline(){ + unsigned active_count=pipelined_simd_unit::get_active_lanes_in_pipeline(); + assert(active_count<=m_core->get_config()->warp_size); + m_core->incsfuactivelanes_stat(active_count); + m_core->incfuactivelanes_stat(active_count); + m_core->incfumemactivelanes_stat(active_count); +} + + sp_unit::sp_unit( register_set* result_port, const shader_core_config *config,shader_core_ctx *core) : pipelined_simd_unit(result_port,config,config->max_sp_latency,core) { @@ -2237,7 +2297,7 @@ void shader_core_ctx::incexecstat(warp_inst_t *&inst) switch(inst->sp_op){ case INT__OP: - incialu_stat(inst->active_count(),25); + incialu_stat(inst->active_count(),32); break; case INT_MUL_OP: incimul_stat(inst->active_count(),7.2); diff --git a/src/gpgpu-sim/shader.h b/src/gpgpu-sim/shader.h index db2af01..2c4c43d 100644 --- a/src/gpgpu-sim/shader.h +++ b/src/gpgpu-sim/shader.h @@ -318,11 +318,12 @@ public: std::vector* warp, register_set* sp_out, register_set* sfu_out, + register_set* tensor_core_out, register_set* mem_out, int id) : m_supervised_warps(), m_stats(stats), m_shader(shader), m_scoreboard(scoreboard), m_simt_stack(simt), /*m_pipeline_reg(pipe_regs),*/ m_warp(warp), - m_sp_out(sp_out),m_sfu_out(sfu_out),m_mem_out(mem_out), m_id(id){} + m_sp_out(sp_out),m_sfu_out(sfu_out),m_tensor_core_out(tensor_core_out),m_mem_out(mem_out), m_id(id){} virtual ~scheduler_unit(){} virtual void add_supervised_warp_id(int i) { m_supervised_warps.push_back(&warp(i)); @@ -395,6 +396,7 @@ protected: std::vector* m_warp; register_set* m_sp_out; register_set* m_sfu_out; + register_set* m_tensor_core_out; register_set* m_mem_out; int m_id; @@ -407,9 +409,10 @@ public: std::vector* warp, register_set* sp_out, register_set* sfu_out, + register_set* tensor_core_out, register_set* mem_out, int id ) - : scheduler_unit ( stats, shader, scoreboard, simt, warp, sp_out, sfu_out, mem_out, id ){} + : scheduler_unit ( stats, shader, scoreboard, simt, warp, sp_out, sfu_out, tensor_core_out, mem_out, id ){} virtual ~lrr_scheduler () {} virtual void order_warps (); virtual void done_adding_supervised_warps() { @@ -424,9 +427,10 @@ public: std::vector* warp, register_set* sp_out, register_set* sfu_out, + register_set* tensor_core_out, register_set* mem_out, int id ) - : scheduler_unit ( stats, shader, scoreboard, simt, warp, sp_out, sfu_out, mem_out, id ){} + : scheduler_unit ( stats, shader, scoreboard, simt, warp, sp_out, sfu_out, tensor_core_out, mem_out, id ){} virtual ~gto_scheduler () {} virtual void order_warps (); virtual void done_adding_supervised_warps() { @@ -443,10 +447,11 @@ public: std::vector* warp, register_set* sp_out, register_set* sfu_out, + register_set* tensor_core_out, register_set* mem_out, int id, char* config_str ) - : scheduler_unit ( stats, shader, scoreboard, simt, warp, sp_out, sfu_out, mem_out, id ), + : scheduler_unit ( stats, shader, scoreboard, simt, warp, sp_out, sfu_out, tensor_core_out, mem_out, id ), m_pending_warps() { unsigned inner_level_readin; @@ -493,6 +498,7 @@ public: std::vector* warp, register_set* sp_out, register_set* sfu_out, + register_set* tensor_core_out, register_set* mem_out, int id, char* config_string ); @@ -1062,6 +1068,23 @@ public: virtual void issue( register_set& source_reg ); }; +class tensor_core : public pipelined_simd_unit +{ +public: + tensor_core( register_set* result_port, const shader_core_config *config, shader_core_ctx *core ); + virtual bool can_issue( const warp_inst_t &inst ) const + { + switch(inst.op) { + case TENSOR_CORE_OP: break; + default: return false; + } + return pipelined_simd_unit::can_issue(inst); + } + virtual void active_lanes_in_pipeline(); + virtual void issue( register_set& source_reg ); +}; + + class sp_unit : public pipelined_simd_unit { public: @@ -1201,9 +1224,11 @@ protected: enum pipeline_stage_name_t { ID_OC_SP=0, ID_OC_SFU, + ID_OC_TENSOR_CORE, ID_OC_MEM, OC_EX_SP, OC_EX_SFU, + OC_EX_TENSOR_CORE, OC_EX_MEM, EX_WB, N_PIPELINE_STAGES @@ -1212,9 +1237,11 @@ enum pipeline_stage_name_t { const char* const pipeline_stage_name_decode[] = { "ID_OC_SP", "ID_OC_SFU", + "ID_OC_TENSOR_CORE", "ID_OC_MEM", "OC_EX_SP", "OC_EX_SFU", + "OC_EX_TENSOR_CORE", "OC_EX_MEM", "EX_WB", "N_PIPELINE_STAGES" @@ -1257,6 +1284,7 @@ struct shader_core_config : public core_config max_warps_per_shader = n_thread_per_shader/warp_size; assert( !(n_thread_per_shader % warp_size) ); max_sfu_latency = 512; + max_tensor_core_latency = 512; max_sp_latency = 32; m_L1I_config.init(m_L1I_config.m_config_string,FuncCachePreferNone); m_L1T_config.init(m_L1T_config.m_config_string,FuncCachePreferNone); @@ -1304,21 +1332,25 @@ struct shader_core_config : public core_config //op collector int gpgpu_operand_collector_num_units_sp; int gpgpu_operand_collector_num_units_sfu; + int gpgpu_operand_collector_num_units_tensor_core; int gpgpu_operand_collector_num_units_mem; int gpgpu_operand_collector_num_units_gen; unsigned int gpgpu_operand_collector_num_in_ports_sp; unsigned int gpgpu_operand_collector_num_in_ports_sfu; + unsigned int gpgpu_operand_collector_num_in_ports_tensor_core; unsigned int gpgpu_operand_collector_num_in_ports_mem; unsigned int gpgpu_operand_collector_num_in_ports_gen; unsigned int gpgpu_operand_collector_num_out_ports_sp; unsigned int gpgpu_operand_collector_num_out_ports_sfu; + unsigned int gpgpu_operand_collector_num_out_ports_tensor_core; unsigned int gpgpu_operand_collector_num_out_ports_mem; unsigned int gpgpu_operand_collector_num_out_ports_gen; int gpgpu_num_sp_units; int gpgpu_num_sfu_units; + int gpgpu_num_tensor_core_units; int gpgpu_num_mem_units; //Shader core resources @@ -1331,6 +1363,7 @@ struct shader_core_config : public core_config unsigned max_sp_latency; unsigned max_sfu_latency; + unsigned max_tensor_core_latency; unsigned n_simt_cores_per_cluster; unsigned n_simt_clusters; @@ -1368,12 +1401,14 @@ struct shader_core_stats_pod { unsigned *m_num_fpdiv_acesses; unsigned *m_num_sp_acesses; unsigned *m_num_sfu_acesses; + unsigned *m_num_tensor_core_acesses; unsigned *m_num_trans_acesses; unsigned *m_num_mem_acesses; unsigned *m_num_sp_committed; unsigned *m_num_tlb_hits; unsigned *m_num_tlb_accesses; unsigned *m_num_sfu_committed; + unsigned *m_num_tensor_core_committed; unsigned *m_num_mem_committed; unsigned *m_read_regfile_acesses; unsigned *m_write_regfile_acesses; @@ -1382,6 +1417,7 @@ struct shader_core_stats_pod { unsigned *m_num_imul32_acesses; unsigned *m_active_sp_lanes; unsigned *m_active_sfu_lanes; + unsigned *m_active_tensor_core_lanes; unsigned *m_active_fu_lanes; unsigned *m_active_fu_mem_lanes; unsigned *m_n_diverge; // number of divergence occurring in this shader @@ -1453,6 +1489,7 @@ public: m_num_fpdiv_acesses= (unsigned*) calloc(config->num_shader(),sizeof(unsigned)); m_num_sp_acesses= (unsigned*) calloc(config->num_shader(),sizeof(unsigned)); m_num_sfu_acesses= (unsigned*) calloc(config->num_shader(),sizeof(unsigned)); + m_num_tensor_core_acesses= (unsigned*) calloc(config->num_shader(),sizeof(unsigned)); m_num_trans_acesses= (unsigned*) calloc(config->num_shader(),sizeof(unsigned)); m_num_mem_acesses= (unsigned*) calloc(config->num_shader(),sizeof(unsigned)); m_num_sp_committed= (unsigned*) calloc(config->num_shader(),sizeof(unsigned)); @@ -1460,9 +1497,11 @@ public: m_num_tlb_accesses=(unsigned*) calloc(config->num_shader(),sizeof(unsigned)); m_active_sp_lanes= (unsigned*) calloc(config->num_shader(),sizeof(unsigned)); m_active_sfu_lanes= (unsigned*) calloc(config->num_shader(),sizeof(unsigned)); + m_active_tensor_core_lanes= (unsigned*) calloc(config->num_shader(),sizeof(unsigned)); m_active_fu_lanes= (unsigned*) calloc(config->num_shader(),sizeof(unsigned)); m_active_fu_mem_lanes= (unsigned*) calloc(config->num_shader(),sizeof(unsigned)); m_num_sfu_committed= (unsigned*) calloc(config->num_shader(),sizeof(unsigned)); + m_num_tensor_core_committed= (unsigned*) calloc(config->num_shader(),sizeof(unsigned)); m_num_mem_committed= (unsigned*) calloc(config->num_shader(),sizeof(unsigned)); m_read_regfile_acesses= (unsigned*) calloc(config->num_shader(),sizeof(unsigned)); m_write_regfile_acesses= (unsigned*) calloc(config->num_shader(),sizeof(unsigned)); -- cgit v1.3 From d907c7d848be6ced2b7f2bd2df84b39e57dfbedc Mon Sep 17 00:00:00 2001 From: negargoli Date: Fri, 22 Jun 2018 11:09:18 -0700 Subject: No need to change the config file for adding tensor-core --- cuda-kernels/gpgpusim.config | 9 ++++----- src/gpgpu-sim/shader.cc | 8 ++++---- src/gpgpu-sim/shader.h | 16 +++++++++++----- 3 files changed, 19 insertions(+), 14 deletions(-) (limited to 'src/gpgpu-sim') diff --git a/cuda-kernels/gpgpusim.config b/cuda-kernels/gpgpusim.config index 69a110f..272ad3d 100755 --- a/cuda-kernels/gpgpusim.config +++ b/cuda-kernels/gpgpusim.config @@ -33,10 +33,9 @@ # ID_OC_SP,ID_OC_SFU,ID_OC_MEM,OC_EX_SP,OC_EX_SFU,OC_EX_MEM,EX_WB ## Pascal GP102 has 4 SP SIMD units and 1 SFU unit ## we need to scale the number of pipeline registers to be equal to the number of SP units --gpgpu_pipeline_widths 4,1,1,1,4,1,1,1,,6 +-gpgpu_pipeline_widths 4,1,1,4,1,1,6 -gpgpu_num_sp_units 4 -gpgpu_num_sfu_units 1 --gpgpu_num_tensor_core_units 1 # Instruction latencies and initiation intervals # "ADD,MAX,MUL,MAD,DIV" # SFU is 32-width in pascal, then dp units initiation is 1 cycle @@ -72,14 +71,14 @@ ## larger operand collectors and reg_banks are needed for the 4 warp schedulers and 4 SIMD units -gpgpu_operand_collector_num_units_sp 20 -gpgpu_operand_collector_num_units_sfu 4 --gpgpu_operand_collector_num_units_tensor_core 24 +#-gpgpu_operand_collector_num_units_tensor_core 24 -gpgpu_operand_collector_num_units_mem 8 -gpgpu_operand_collector_num_in_ports_sp 4 -gpgpu_operand_collector_num_out_ports_sp 4 -gpgpu_operand_collector_num_in_ports_sfu 1 -gpgpu_operand_collector_num_out_ports_sfu 1 --gpgpu_operand_collector_num_in_ports_tensor_core 1 --gpgpu_operand_collector_num_out_ports_tensor_core 1 +#-gpgpu_operand_collector_num_in_ports_tensor_core 1 +#-gpgpu_operand_collector_num_out_ports_tensor_core 1 -gpgpu_operand_collector_num_in_ports_mem 1 -gpgpu_operand_collector_num_out_ports_mem 1 # gpgpu_num_reg_banks should be increased to 32, but it gives an error! diff --git a/src/gpgpu-sim/shader.cc b/src/gpgpu-sim/shader.cc index fcac755..c01f867 100644 --- a/src/gpgpu-sim/shader.cc +++ b/src/gpgpu-sim/shader.cc @@ -218,7 +218,7 @@ shader_core_ctx::shader_core_ctx( class gpgpu_sim *gpu, enum { SP_CUS, SFU_CUS, TENSOR_CORE_CUS, MEM_CUS, GEN_CUS }; m_operand_collector.add_cu_set(SP_CUS, m_config->gpgpu_operand_collector_num_units_sp, m_config->gpgpu_operand_collector_num_out_ports_sp); m_operand_collector.add_cu_set(SFU_CUS, m_config->gpgpu_operand_collector_num_units_sfu, m_config->gpgpu_operand_collector_num_out_ports_sfu); - m_operand_collector.add_cu_set(TENSOR_CORE_CUS, m_config->gpgpu_operand_collector_num_units_tensor_core, m_config->gpgpu_operand_collector_num_out_ports_tensor_core); + m_operand_collector.add_cu_set(TENSOR_CORE_CUS, config->gpgpu_operand_collector_num_units_tensor_core, config->gpgpu_operand_collector_num_out_ports_tensor_core); m_operand_collector.add_cu_set(MEM_CUS, m_config->gpgpu_operand_collector_num_units_mem, m_config->gpgpu_operand_collector_num_out_ports_mem); m_operand_collector.add_cu_set(GEN_CUS, m_config->gpgpu_operand_collector_num_units_gen, m_config->gpgpu_operand_collector_num_out_ports_gen); @@ -243,7 +243,7 @@ shader_core_ctx::shader_core_ctx( class gpgpu_sim *gpu, in_ports.clear(),out_ports.clear(),cu_sets.clear(); } - for (unsigned i = 0; i < m_config->gpgpu_operand_collector_num_in_ports_tensor_core; i++) { + for (unsigned i = 0; i < config->gpgpu_operand_collector_num_in_ports_tensor_core; i++) { in_ports.push_back(&m_pipeline_reg[ID_OC_TENSOR_CORE]); out_ports.push_back(&m_pipeline_reg[OC_EX_TENSOR_CORE]); cu_sets.push_back((unsigned)TENSOR_CORE_CUS); @@ -280,7 +280,7 @@ shader_core_ctx::shader_core_ctx( class gpgpu_sim *gpu, m_operand_collector.init( m_config->gpgpu_num_reg_banks, this ); // execute - m_num_function_units = m_config->gpgpu_num_sp_units + m_config->gpgpu_num_sfu_units + m_config->gpgpu_num_tensor_core_units + 1; // sp_unit, sfu, ldst_unit + m_num_function_units = m_config->gpgpu_num_sp_units + m_config->gpgpu_num_sfu_units + config->gpgpu_num_tensor_core_units + 1; // sp_unit, sfu, ldst_unit //m_dispatch_port = new enum pipeline_stage_name_t[ m_num_function_units ]; //m_issue_port = new enum pipeline_stage_name_t[ m_num_function_units ]; @@ -298,7 +298,7 @@ shader_core_ctx::shader_core_ctx( class gpgpu_sim *gpu, m_issue_port.push_back(OC_EX_SFU); } - for (int k = 0; k < m_config->gpgpu_num_tensor_core_units; k++) { + for (int k = 0; k < config->gpgpu_num_tensor_core_units; k++) { m_fu.push_back(new tensor_core( &m_pipeline_reg[EX_WB], m_config, this )); m_dispatch_port.push_back(ID_OC_TENSOR_CORE); m_issue_port.push_back(OC_EX_TENSOR_CORE); diff --git a/src/gpgpu-sim/shader.h b/src/gpgpu-sim/shader.h index 2c4c43d..b7deae6 100644 --- a/src/gpgpu-sim/shader.h +++ b/src/gpgpu-sim/shader.h @@ -1224,15 +1224,15 @@ protected: enum pipeline_stage_name_t { ID_OC_SP=0, ID_OC_SFU, - ID_OC_TENSOR_CORE, ID_OC_MEM, OC_EX_SP, OC_EX_SFU, - OC_EX_TENSOR_CORE, OC_EX_MEM, EX_WB, + ID_OC_TENSOR_CORE, + OC_EX_TENSOR_CORE, N_PIPELINE_STAGES -}; + }; const char* const pipeline_stage_name_decode[] = { "ID_OC_SP", @@ -1266,9 +1266,11 @@ struct shader_core_config : public core_config char* toks = new char[100]; char* tokd = toks; strcpy(toks,pipeline_widths_string); - + toks = strtok(toks,","); - for (unsigned i = 0; i < N_PIPELINE_STAGES; i++) { + pipe_widths[OC_EX_TENSOR_CORE]=1; + pipe_widths[ID_OC_TENSOR_CORE]=1; + for (unsigned i = 0; i < N_PIPELINE_STAGES-2; i++) { assert(toks); ntok = sscanf(toks,"%d", &pipe_widths[i]); assert(ntok == 1); @@ -1286,6 +1288,10 @@ struct shader_core_config : public core_config max_sfu_latency = 512; max_tensor_core_latency = 512; max_sp_latency = 32; + gpgpu_num_tensor_core_units=1; + gpgpu_operand_collector_num_units_tensor_core=24; + gpgpu_operand_collector_num_in_ports_tensor_core=1; + gpgpu_operand_collector_num_out_ports_tensor_core=1; m_L1I_config.init(m_L1I_config.m_config_string,FuncCachePreferNone); m_L1T_config.init(m_L1T_config.m_config_string,FuncCachePreferNone); m_L1C_config.init(m_L1C_config.m_config_string,FuncCachePreferNone); -- cgit v1.3 From e541026cfc0ee4be25e7093cb7ff3acfa3cbb6e7 Mon Sep 17 00:00:00 2001 From: negargoli93 Date: Mon, 16 Jul 2018 13:18:23 -0700 Subject: fix pipline for tensor_core and change config --- cuda-kernels/config_fermi_islip.icnt | 2 +- cuda-kernels/gpgpusim.config | 10 +++++----- src/cuda-sim/cuda-sim.cc | 8 ++++++-- src/gpgpu-sim/gpu-sim.cc | 2 +- src/gpgpu-sim/shader.cc | 10 +++++----- src/gpgpu-sim/shader.h | 18 +++++++++--------- 6 files changed, 27 insertions(+), 23 deletions(-) (limited to 'src/gpgpu-sim') diff --git a/cuda-kernels/config_fermi_islip.icnt b/cuda-kernels/config_fermi_islip.icnt index a788090..3b8b496 100755 --- a/cuda-kernels/config_fermi_islip.icnt +++ b/cuda-kernels/config_fermi_islip.icnt @@ -7,7 +7,7 @@ network_count = 2; // Topology topology = fly; -k = 62; +k = 102; n = 1; // Routing diff --git a/cuda-kernels/gpgpusim.config b/cuda-kernels/gpgpusim.config index 272ad3d..2510d21 100755 --- a/cuda-kernels/gpgpusim.config +++ b/cuda-kernels/gpgpusim.config @@ -10,7 +10,7 @@ -gpgpu_ptx_save_converted_ptxplus 0 # high level architecture configuration --gpgpu_n_clusters 40 +-gpgpu_n_clusters 80 -gpgpu_n_cores_per_cluster 1 -gpgpu_n_mem 11 -gpgpu_n_sub_partition_per_mchannel 2 @@ -33,7 +33,7 @@ # ID_OC_SP,ID_OC_SFU,ID_OC_MEM,OC_EX_SP,OC_EX_SFU,OC_EX_MEM,EX_WB ## Pascal GP102 has 4 SP SIMD units and 1 SFU unit ## we need to scale the number of pipeline registers to be equal to the number of SP units --gpgpu_pipeline_widths 4,1,1,4,1,1,6 +-gpgpu_pipeline_widths 4,1,1,1,4,1,1,1,6 -gpgpu_num_sp_units 4 -gpgpu_num_sfu_units 1 # Instruction latencies and initiation intervals @@ -72,15 +72,15 @@ -gpgpu_operand_collector_num_units_sp 20 -gpgpu_operand_collector_num_units_sfu 4 #-gpgpu_operand_collector_num_units_tensor_core 24 --gpgpu_operand_collector_num_units_mem 8 +-gpgpu_operand_collector_num_units_mem 8 -gpgpu_operand_collector_num_in_ports_sp 4 -gpgpu_operand_collector_num_out_ports_sp 4 -gpgpu_operand_collector_num_in_ports_sfu 1 -gpgpu_operand_collector_num_out_ports_sfu 1 #-gpgpu_operand_collector_num_in_ports_tensor_core 1 #-gpgpu_operand_collector_num_out_ports_tensor_core 1 --gpgpu_operand_collector_num_in_ports_mem 1 --gpgpu_operand_collector_num_out_ports_mem 1 +-gpgpu_operand_collector_num_in_ports_mem 10 +-gpgpu_operand_collector_num_out_ports_mem 10 # gpgpu_num_reg_banks should be increased to 32, but it gives an error! -gpgpu_num_reg_banks 32 diff --git a/src/cuda-sim/cuda-sim.cc b/src/cuda-sim/cuda-sim.cc index 6da0840..506bc95 100644 --- a/src/cuda-sim/cuda-sim.cc +++ b/src/cuda-sim/cuda-sim.cc @@ -648,8 +648,10 @@ void ptx_instruction::set_opcode_and_latency() if ( has_memory_write() ) op = STORE_OP; break; case LD_OP: op = LOAD_OP; break; + case MMA_LD_OP: op = LOAD_OP; break; case LDU_OP: op = LOAD_OP; break; case ST_OP: op = STORE_OP; break; + case MMA_ST_OP: op = STORE_OP; break; case BRA_OP: op = BRANCH_OP; break; case BREAKADDR_OP: op = BRANCH_OP; break; case TEX_OP: op = LOAD_OP; mem_op=TEX; break; @@ -897,9 +899,11 @@ void ptx_instruction::pre_decode() case WB_OPTION: cache_op = CACHE_WRITE_BACK; break; case WT_OPTION: cache_op = CACHE_WRITE_THROUGH; break; default: - if( m_opcode == LD_OP || m_opcode == LDU_OP ) + //if( m_opcode == LD_OP || m_opcode == LDU_OP ) + if( m_opcode == MMA_LD_OP || m_opcode == LD_OP || m_opcode == LDU_OP ) cache_op = CACHE_ALL; - else if( m_opcode == ST_OP ) + //else if( m_opcode == ST_OP ) + else if( m_opcode == ST_OP || m_opcode == ST_OP ) cache_op = CACHE_WRITE_BACK; else if( m_opcode == ATOM_OP ) cache_op = CACHE_GLOBAL; diff --git a/src/gpgpu-sim/gpu-sim.cc b/src/gpgpu-sim/gpu-sim.cc index cc23051..3e064c7 100644 --- a/src/gpgpu-sim/gpu-sim.cc +++ b/src/gpgpu-sim/gpu-sim.cc @@ -1541,7 +1541,7 @@ void gpgpu_sim::cycle() } } - if (!(gpu_sim_cycle % 20000)) { + if (!(gpu_sim_cycle % 50000)) { // deadlock detection if (m_config.gpu_deadlock_detect && gpu_sim_insn == last_gpu_sim_insn) { gpu_deadlock = true; diff --git a/src/gpgpu-sim/shader.cc b/src/gpgpu-sim/shader.cc index c01f867..226e7f0 100644 --- a/src/gpgpu-sim/shader.cc +++ b/src/gpgpu-sim/shader.cc @@ -1387,7 +1387,7 @@ ldst_unit::process_cache_access( cache_t* cache, assert( !read_sent ); inst.accessq_pop_back(); if ( inst.is_load() ) { - for ( unsigned r=0; r < 4; r++) + for ( unsigned r=0; r < 8; r++) if (inst.out[r] > 0) m_pending_writes[inst.warp_id()][inst.out[r]]--; } @@ -1489,7 +1489,7 @@ bool ldst_unit::memory_cycle( warp_inst_t &inst, mem_stage_stall_type &stall_rea inst.accessq_pop_back(); //inst.clear_active( access.get_warp_mask() ); if( inst.is_load() ) { - for( unsigned r=0; r < 4; r++) + for( unsigned r=0; r < 8; r++) if(inst.out[r] > 0) assert( m_pending_writes[inst.warp_id()][inst.out[r]] > 0 ); } else if( inst.is_store() ) @@ -1767,7 +1767,7 @@ void ldst_unit:: issue( register_set ®_set ) if (inst->is_load() and inst->space.get_type() != shared_space) { unsigned warp_id = inst->warp_id(); unsigned n_accesses = inst->accessq_count(); - for (unsigned r = 0; r < 4; r++) { + for (unsigned r = 0; r < 8; r++) { unsigned reg_id = inst->out[r]; if (reg_id > 0) { m_pending_writes[warp_id][reg_id] += n_accesses; @@ -1789,7 +1789,7 @@ void ldst_unit::writeback() if( !m_next_wb.empty() ) { if( m_operand_collector->writeback(m_next_wb) ) { bool insn_completed = false; - for( unsigned r=0; r < 4; r++ ) { + for( unsigned r=0; r < 8; r++ ) { if( m_next_wb.out[r] > 0 ) { if( m_next_wb.space.get_type() != shared_space ) { assert( m_pending_writes[m_next_wb.warp_id()][m_next_wb.out[r]] > 0 ); @@ -1991,7 +1991,7 @@ void ldst_unit::cycle() //} bool pending_requests=false; - for( unsigned r=0; r<4; r++ ) { + for( unsigned r=0; r<8; r++ ) { unsigned reg_id = pipe_reg.out[r]; if( reg_id > 0 ) { if( m_pending_writes[warp_id].find(reg_id) != m_pending_writes[warp_id].end() ) { diff --git a/src/gpgpu-sim/shader.h b/src/gpgpu-sim/shader.h index b7deae6..90a3134 100644 --- a/src/gpgpu-sim/shader.h +++ b/src/gpgpu-sim/shader.h @@ -1224,13 +1224,13 @@ protected: enum pipeline_stage_name_t { ID_OC_SP=0, ID_OC_SFU, + ID_OC_TENSOR_CORE, ID_OC_MEM, OC_EX_SP, OC_EX_SFU, + OC_EX_TENSOR_CORE, OC_EX_MEM, EX_WB, - ID_OC_TENSOR_CORE, - OC_EX_TENSOR_CORE, N_PIPELINE_STAGES }; @@ -1268,9 +1268,9 @@ struct shader_core_config : public core_config strcpy(toks,pipeline_widths_string); toks = strtok(toks,","); - pipe_widths[OC_EX_TENSOR_CORE]=1; - pipe_widths[ID_OC_TENSOR_CORE]=1; - for (unsigned i = 0; i < N_PIPELINE_STAGES-2; i++) { + // pipe_widths[OC_EX_TENSOR_CORE]=1; + // pipe_widths[ID_OC_TENSOR_CORE]=1; + for (unsigned i = 0; i < N_PIPELINE_STAGES; i++) { assert(toks); ntok = sscanf(toks,"%d", &pipe_widths[i]); assert(ntok == 1); @@ -1286,12 +1286,12 @@ struct shader_core_config : public core_config max_warps_per_shader = n_thread_per_shader/warp_size; assert( !(n_thread_per_shader % warp_size) ); max_sfu_latency = 512; - max_tensor_core_latency = 512; + max_tensor_core_latency = 64; max_sp_latency = 32; - gpgpu_num_tensor_core_units=1; + gpgpu_num_tensor_core_units=8; gpgpu_operand_collector_num_units_tensor_core=24; - gpgpu_operand_collector_num_in_ports_tensor_core=1; - gpgpu_operand_collector_num_out_ports_tensor_core=1; + gpgpu_operand_collector_num_in_ports_tensor_core=8; + gpgpu_operand_collector_num_out_ports_tensor_core=8; m_L1I_config.init(m_L1I_config.m_config_string,FuncCachePreferNone); m_L1T_config.init(m_L1T_config.m_config_string,FuncCachePreferNone); m_L1C_config.init(m_L1C_config.m_config_string,FuncCachePreferNone); -- cgit v1.3 From 9e7cd8867d76fb99eadfadfa09947ff057d012d3 Mon Sep 17 00:00:00 2001 From: negargoli93 Date: Thu, 16 Aug 2018 15:20:07 -0700 Subject: Timing model for VCORE --- src/abstract_hardware_model.h | 2 ++ src/cuda-sim/cuda-sim.cc | 28 ++++++++++++++++--- src/cuda-sim/ptx_ir.h | 4 +-- src/gpgpu-sim/gpu-sim.cc | 16 +++++++++-- src/gpgpu-sim/shader.cc | 62 ++++++++++++++++++++++++++++++++++++++++--- src/gpgpu-sim/shader.h | 57 ++++++++++++++++++++++++++++++++++----- 6 files changed, 152 insertions(+), 17 deletions(-) (limited to 'src/gpgpu-sim') diff --git a/src/abstract_hardware_model.h b/src/abstract_hardware_model.h index 7fe5d82..781509f 100644 --- a/src/abstract_hardware_model.h +++ b/src/abstract_hardware_model.h @@ -78,6 +78,7 @@ enum uarch_op_t { ALU_OP=1, SFU_OP, TENSOR_CORE_OP, + VP_CORE_OP, ALU_SFU_OP, LOAD_OP, STORE_OP, @@ -135,6 +136,7 @@ enum operation_pipeline_t { SP__OP, SFU__OP, TENSOR_CORE__OP, + VP_CORE__OP, MEM__OP }; typedef enum operation_pipeline_t operation_pipeline; diff --git a/src/cuda-sim/cuda-sim.cc b/src/cuda-sim/cuda-sim.cc index 2fe5667..1ad12ee 100644 --- a/src/cuda-sim/cuda-sim.cc +++ b/src/cuda-sim/cuda-sim.cc @@ -542,7 +542,7 @@ void ptx_instruction::set_mul_div_or_other_archop(){ sp_op=INT_DIV_OP; break; default: - if((op==ALU_OP)||(op==TENSOR_CORE_OP)) + if((op==ALU_OP)||(op==VP_CORE_OP)) sp_op=INT__OP; break; } @@ -649,9 +649,11 @@ void ptx_instruction::set_opcode_and_latency() break; case LD_OP: op = LOAD_OP; break; case MMA_LD_OP: op = LOAD_OP; break; + case VP_LD_OP: op = LOAD_OP; break; case LDU_OP: op = LOAD_OP; break; case ST_OP: op = STORE_OP; break; case MMA_ST_OP: op = STORE_OP; break; + case VP_ST_OP: op = STORE_OP; break; case BRA_OP: op = BRANCH_OP; break; case BREAKADDR_OP: op = BRANCH_OP; break; case TEX_OP: op = LOAD_OP; mem_op=TEX; break; @@ -799,6 +801,26 @@ void ptx_instruction::set_opcode_and_latency() initiation_interval = 64; op=TENSOR_CORE_OP; break; + case VP_MMA_OP: + if(get_wmma_type()==VP_MMA4) + { + latency = 5; + initiation_interval = 5; + } + if(get_wmma_type()==VP_MMA8) + { + latency = 5; + initiation_interval = 5; + } + if(get_wmma_type()==VP_MMA16) + { + latency = 5; + initiation_interval = 5; + } + op=VP_CORE_OP; + op=VP_CORE_OP; + op=VP_CORE_OP; + break; case SHFL_OP: latency = 32; initiation_interval = 15; @@ -900,10 +922,10 @@ void ptx_instruction::pre_decode() case WT_OPTION: cache_op = CACHE_WRITE_THROUGH; break; default: //if( m_opcode == LD_OP || m_opcode == LDU_OP ) - if( m_opcode == MMA_LD_OP || m_opcode == LD_OP || m_opcode == LDU_OP ) + if( m_opcode ==VP_LD_OP || m_opcode == MMA_LD_OP || m_opcode == LD_OP || m_opcode == LDU_OP ) cache_op = CACHE_ALL; //else if( m_opcode == ST_OP ) - else if( m_opcode == ST_OP || m_opcode == ST_OP ) + else if( m_opcode == VP_ST_OP ||m_opcode == MMA_ST_OP || m_opcode == ST_OP ) cache_op = CACHE_WRITE_BACK; else if( m_opcode == ATOM_OP ) cache_op = CACHE_GLOBAL; diff --git a/src/cuda-sim/ptx_ir.h b/src/cuda-sim/ptx_ir.h index cb4556e..e025013 100644 --- a/src/cuda-sim/ptx_ir.h +++ b/src/cuda-sim/ptx_ir.h @@ -1093,7 +1093,7 @@ public: int membar_level() const { return m_membar_level; } bool has_memory_read() const { - if( m_opcode == LD_OP || m_opcode == LDU_OP || m_opcode == TEX_OP|| m_opcode==MMA_LD_OP) + if( m_opcode == LD_OP || m_opcode == LDU_OP || m_opcode == TEX_OP|| m_opcode==MMA_LD_OP || m_opcode==VP_LD_OP) return true; // Check PTXPlus operand type below // Source operands are memory operands @@ -1105,7 +1105,7 @@ public: return false; } bool has_memory_write() const { - if( m_opcode == ST_OP || m_opcode==MMA_ST_OP ) return true; + if( m_opcode == ST_OP || m_opcode==MMA_ST_OP || m_opcode==VP_ST_OP ) return true; // Check PTXPlus operand type below // Destination operand is a memory operand ptx_instruction::const_iterator op=op_iter_begin(); diff --git a/src/gpgpu-sim/gpu-sim.cc b/src/gpgpu-sim/gpu-sim.cc index 3e064c7..7a797b5 100644 --- a/src/gpgpu-sim/gpu-sim.cc +++ b/src/gpgpu-sim/gpu-sim.cc @@ -309,6 +309,9 @@ void shader_core_config::reg_options(class OptionParser * opp) option_parser_register(opp, "-gpgpu_operand_collector_num_units_tensor_core", OPT_INT32, &gpgpu_operand_collector_num_units_tensor_core, "number of collector units (default = 4)", "4"); + option_parser_register(opp, "-gpgpu_operand_collector_num_units_vp_core", OPT_INT32, &gpgpu_operand_collector_num_units_vp_core, + "number of collector units (default = 4)", + "4"); option_parser_register(opp, "-gpgpu_operand_collector_num_units_mem", OPT_INT32, &gpgpu_operand_collector_num_units_mem, "number of collector units (default = 2)", "2"); @@ -324,6 +327,9 @@ void shader_core_config::reg_options(class OptionParser * opp) option_parser_register(opp, "-gpgpu_operand_collector_num_in_ports_tensor_core", OPT_INT32, &gpgpu_operand_collector_num_in_ports_tensor_core, "number of collector unit in ports (default = 1)", "1"); + option_parser_register(opp, "-gpgpu_operand_collector_num_in_ports_vp_core", OPT_INT32, &gpgpu_operand_collector_num_in_ports_vp_core, + "number of collector unit in ports (default = 1)", + "1"); option_parser_register(opp, "-gpgpu_operand_collector_num_in_ports_mem", OPT_INT32, &gpgpu_operand_collector_num_in_ports_mem, "number of collector unit in ports (default = 1)", "1"); @@ -339,6 +345,9 @@ void shader_core_config::reg_options(class OptionParser * opp) option_parser_register(opp, "-gpgpu_operand_collector_num_out_ports_tensor_core", OPT_INT32, &gpgpu_operand_collector_num_out_ports_tensor_core, "number of collector unit in ports (default = 1)", "1"); + option_parser_register(opp, "-gpgpu_operand_collector_num_out_ports_vp_core", OPT_INT32, &gpgpu_operand_collector_num_out_ports_vp_core, + "number of collector unit in ports (default = 1)", + "1"); option_parser_register(opp, "-gpgpu_operand_collector_num_out_ports_mem", OPT_INT32, &gpgpu_operand_collector_num_out_ports_mem, "number of collector unit in ports (default = 1)", "1"); @@ -359,8 +368,8 @@ void shader_core_config::reg_options(class OptionParser * opp) "1"); option_parser_register(opp, "-gpgpu_pipeline_widths", OPT_CSTR, &pipeline_widths_string, "Pipeline widths " - "ID_OC_SP,ID_OC_SFU,ID_OC_TENSOR_CORE,ID_OC_MEM,OC_EX_SP,OC_EX_SFU,OC_EX_TENSOR_CORE,OC_EX_MEM,EX_WB", - "1,1,1,1,1,1,1,1,1" ); + "ID_OC_SP,ID_OC_SFU,ID_OC_TENSOR_CORE,ID_OC_VP_CORE,ID_OC_MEM,OC_EX_SP,OC_EX_SFU,OC_EX_TENSOR_CORE,OC_EX_TENSOR_CORE,OC_EX_MEM,EX_WB", + "1,1,1,1,1,1,1,1,1,1,1" ); option_parser_register(opp, "-gpgpu_num_sp_units", OPT_INT32, &gpgpu_num_sp_units, "Number of SP units (default=1)", "1"); @@ -370,6 +379,9 @@ void shader_core_config::reg_options(class OptionParser * opp) option_parser_register(opp, "-gpgpu_num_tensor_core_units", OPT_INT32, &gpgpu_num_tensor_core_units, "Number of tensor_core units (default=1)", "1"); + option_parser_register(opp, "-gpgpu_num_vp_core_units", OPT_INT32, &gpgpu_num_vp_core_units, + "Number of vp_core units (default=1)", + "1"); option_parser_register(opp, "-gpgpu_num_mem_units", OPT_INT32, &gpgpu_num_mem_units, "Number if ldst units (default=1) WARNING: not hooked up to anything", "1"); diff --git a/src/gpgpu-sim/shader.cc b/src/gpgpu-sim/shader.cc index 226e7f0..6f11ad9 100644 --- a/src/gpgpu-sim/shader.cc +++ b/src/gpgpu-sim/shader.cc @@ -149,6 +149,7 @@ shader_core_ctx::shader_core_ctx( class gpgpu_sim *gpu, &m_pipeline_reg[ID_OC_SP], &m_pipeline_reg[ID_OC_SFU], &m_pipeline_reg[ID_OC_TENSOR_CORE], + &m_pipeline_reg[ID_OC_VP_CORE], &m_pipeline_reg[ID_OC_MEM], i ) @@ -164,6 +165,7 @@ shader_core_ctx::shader_core_ctx( class gpgpu_sim *gpu, &m_pipeline_reg[ID_OC_SP], &m_pipeline_reg[ID_OC_SFU], &m_pipeline_reg[ID_OC_TENSOR_CORE], + &m_pipeline_reg[ID_OC_VP_CORE], &m_pipeline_reg[ID_OC_MEM], i, config->gpgpu_scheduler_string @@ -180,6 +182,7 @@ shader_core_ctx::shader_core_ctx( class gpgpu_sim *gpu, &m_pipeline_reg[ID_OC_SP], &m_pipeline_reg[ID_OC_SFU], &m_pipeline_reg[ID_OC_TENSOR_CORE], + &m_pipeline_reg[ID_OC_VP_CORE], &m_pipeline_reg[ID_OC_MEM], i ) @@ -195,6 +198,7 @@ shader_core_ctx::shader_core_ctx( class gpgpu_sim *gpu, &m_pipeline_reg[ID_OC_SP], &m_pipeline_reg[ID_OC_SFU], &m_pipeline_reg[ID_OC_TENSOR_CORE], + &m_pipeline_reg[ID_OC_VP_CORE], &m_pipeline_reg[ID_OC_MEM], i, config->gpgpu_scheduler_string @@ -215,10 +219,11 @@ shader_core_ctx::shader_core_ctx( class gpgpu_sim *gpu, } //op collector configuration - enum { SP_CUS, SFU_CUS, TENSOR_CORE_CUS, MEM_CUS, GEN_CUS }; + enum { SP_CUS, SFU_CUS, TENSOR_CORE_CUS, VP_CORE_CUS, MEM_CUS, GEN_CUS }; m_operand_collector.add_cu_set(SP_CUS, m_config->gpgpu_operand_collector_num_units_sp, m_config->gpgpu_operand_collector_num_out_ports_sp); m_operand_collector.add_cu_set(SFU_CUS, m_config->gpgpu_operand_collector_num_units_sfu, m_config->gpgpu_operand_collector_num_out_ports_sfu); m_operand_collector.add_cu_set(TENSOR_CORE_CUS, config->gpgpu_operand_collector_num_units_tensor_core, config->gpgpu_operand_collector_num_out_ports_tensor_core); + m_operand_collector.add_cu_set(VP_CORE_CUS, config->gpgpu_operand_collector_num_units_vp_core, config->gpgpu_operand_collector_num_out_ports_vp_core); m_operand_collector.add_cu_set(MEM_CUS, m_config->gpgpu_operand_collector_num_units_mem, m_config->gpgpu_operand_collector_num_out_ports_mem); m_operand_collector.add_cu_set(GEN_CUS, m_config->gpgpu_operand_collector_num_units_gen, m_config->gpgpu_operand_collector_num_out_ports_gen); @@ -252,6 +257,14 @@ shader_core_ctx::shader_core_ctx( class gpgpu_sim *gpu, in_ports.clear(),out_ports.clear(),cu_sets.clear(); } + for (unsigned i = 0; i < config->gpgpu_operand_collector_num_in_ports_vp_core; i++) { + in_ports.push_back(&m_pipeline_reg[ID_OC_VP_CORE]); + out_ports.push_back(&m_pipeline_reg[OC_EX_VP_CORE]); + cu_sets.push_back((unsigned)VP_CORE_CUS); + cu_sets.push_back((unsigned)GEN_CUS); + m_operand_collector.add_port(in_ports,out_ports,cu_sets); + in_ports.clear(),out_ports.clear(),cu_sets.clear(); + } for (unsigned i = 0; i < m_config->gpgpu_operand_collector_num_in_ports_mem; i++) { in_ports.push_back(&m_pipeline_reg[ID_OC_MEM]); @@ -267,10 +280,12 @@ shader_core_ctx::shader_core_ctx( class gpgpu_sim *gpu, in_ports.push_back(&m_pipeline_reg[ID_OC_SP]); in_ports.push_back(&m_pipeline_reg[ID_OC_SFU]); in_ports.push_back(&m_pipeline_reg[ID_OC_TENSOR_CORE]); + in_ports.push_back(&m_pipeline_reg[ID_OC_VP_CORE]); in_ports.push_back(&m_pipeline_reg[ID_OC_MEM]); out_ports.push_back(&m_pipeline_reg[OC_EX_SP]); out_ports.push_back(&m_pipeline_reg[OC_EX_SFU]); out_ports.push_back(&m_pipeline_reg[OC_EX_TENSOR_CORE]); + out_ports.push_back(&m_pipeline_reg[OC_EX_VP_CORE]); out_ports.push_back(&m_pipeline_reg[OC_EX_MEM]); cu_sets.push_back((unsigned)GEN_CUS); m_operand_collector.add_port(in_ports,out_ports,cu_sets); @@ -280,7 +295,7 @@ shader_core_ctx::shader_core_ctx( class gpgpu_sim *gpu, m_operand_collector.init( m_config->gpgpu_num_reg_banks, this ); // execute - m_num_function_units = m_config->gpgpu_num_sp_units + m_config->gpgpu_num_sfu_units + config->gpgpu_num_tensor_core_units + 1; // sp_unit, sfu, ldst_unit + m_num_function_units = m_config->gpgpu_num_sp_units + m_config->gpgpu_num_sfu_units + config->gpgpu_num_tensor_core_units + config->gpgpu_num_vp_core_units + 1; // sp_unit, sfu, ldst_unit //m_dispatch_port = new enum pipeline_stage_name_t[ m_num_function_units ]; //m_issue_port = new enum pipeline_stage_name_t[ m_num_function_units ]; @@ -304,6 +319,12 @@ shader_core_ctx::shader_core_ctx( class gpgpu_sim *gpu, m_issue_port.push_back(OC_EX_TENSOR_CORE); } + for (int k = 0; k < config->gpgpu_num_vp_core_units; k++) { + m_fu.push_back(new vp_core( &m_pipeline_reg[EX_WB], m_config, this )); + m_dispatch_port.push_back(ID_OC_VP_CORE); + m_issue_port.push_back(OC_EX_VP_CORE); + } + m_ldst_unit = new ldst_unit( m_icnt, m_mem_fetch_allocator, this, &m_operand_collector, m_scoreboard, config, mem_config, stats, shader_id, tpc_id ); m_fu.push_back(m_ldst_unit); m_dispatch_port.push_back(ID_OC_MEM); @@ -910,7 +931,8 @@ void scheduler_unit::cycle() bool sp_pipe_avail = m_sp_out->has_free(); bool sfu_pipe_avail = m_sfu_out->has_free(); bool tensor_core_pipe_avail = m_tensor_core_out->has_free(); - if( sp_pipe_avail && (pI->op != SFU_OP) && (pI->op != TENSOR_CORE_OP)) { + bool vp_core_pipe_avail = m_vp_core_out->has_free(); + if( sp_pipe_avail && (pI->op != SFU_OP) && (pI->op != TENSOR_CORE_OP) && (pI->op !=VP_CORE_OP)) { //Jin: special for CDP api if(pI->m_is_cdp && !warp(warp_id).m_cdp_dummy) { @@ -950,6 +972,14 @@ void scheduler_unit::cycle() issued_inst=true; warp_inst_issued = true; } + } + else if ( (pI->op == VP_CORE_OP) ) { + if( vp_core_pipe_avail ) { + m_shader->issue_warp(*m_vp_core_out,pI,active_mask,warp_id); + issued++; + issued_inst=true; + warp_inst_issued = true; + } } } } else { SCHED_DPRINTF( "Warp (warp_id %u, dynamic_warp_id %u) fails scoreboard\n", @@ -1116,10 +1146,11 @@ swl_scheduler::swl_scheduler ( shader_core_stats* stats, shader_core_ctx* shader register_set* sp_out, register_set* sfu_out, register_set* tensor_core_out, + register_set* vp_core_out, register_set* mem_out, int id, char* config_string ) - : scheduler_unit ( stats, shader, scoreboard, simt, warp, sp_out, sfu_out, tensor_core_out, mem_out, id ) + : scheduler_unit ( stats, shader, scoreboard, simt, warp, sp_out, sfu_out, tensor_core_out, vp_core_out, mem_out, id ) { unsigned m_prioritization_readin; int ret = sscanf( config_string, @@ -1548,6 +1579,12 @@ tensor_core:: tensor_core( register_set* result_port, const shader_core_config m_name = "TENSOR_CORE"; } +vp_core:: vp_core( register_set* result_port, const shader_core_config *config,shader_core_ctx *core ) + : pipelined_simd_unit(result_port,config,config->max_vp_core_latency,core) +{ + m_name = "VP_CORE"; +} + void sfu::issue( register_set& source_reg ) { @@ -1569,6 +1606,16 @@ void tensor_core::issue( register_set& source_reg ) pipelined_simd_unit::issue(source_reg); } +void vp_core::issue( register_set& source_reg ) +{ + warp_inst_t** ready_reg = source_reg.get_ready(); + //m_core->incexecstat((*ready_reg)); + + (*ready_reg)->op_pipe= VP_CORE__OP; + m_core->incsfu_stat(m_core->get_config()->warp_size,(*ready_reg)->latency); + pipelined_simd_unit::issue(source_reg); +} + void ldst_unit::active_lanes_in_pipeline(){ unsigned active_count=pipelined_simd_unit::get_active_lanes_in_pipeline(); @@ -1599,6 +1646,13 @@ void tensor_core::active_lanes_in_pipeline(){ m_core->incfumemactivelanes_stat(active_count); } +void vp_core::active_lanes_in_pipeline(){ + unsigned active_count=pipelined_simd_unit::get_active_lanes_in_pipeline(); + assert(active_count<=m_core->get_config()->warp_size); + m_core->incsfuactivelanes_stat(active_count); + m_core->incfuactivelanes_stat(active_count); + m_core->incfumemactivelanes_stat(active_count); +} sp_unit::sp_unit( register_set* result_port, const shader_core_config *config,shader_core_ctx *core) : pipelined_simd_unit(result_port,config,config->max_sp_latency,core) diff --git a/src/gpgpu-sim/shader.h b/src/gpgpu-sim/shader.h index 90a3134..d292d56 100644 --- a/src/gpgpu-sim/shader.h +++ b/src/gpgpu-sim/shader.h @@ -319,11 +319,12 @@ public: register_set* sp_out, register_set* sfu_out, register_set* tensor_core_out, + register_set* vp_core_out, register_set* mem_out, int id) : m_supervised_warps(), m_stats(stats), m_shader(shader), m_scoreboard(scoreboard), m_simt_stack(simt), /*m_pipeline_reg(pipe_regs),*/ m_warp(warp), - m_sp_out(sp_out),m_sfu_out(sfu_out),m_tensor_core_out(tensor_core_out),m_mem_out(mem_out), m_id(id){} + m_sp_out(sp_out),m_sfu_out(sfu_out),m_tensor_core_out(tensor_core_out),m_vp_core_out(vp_core_out),m_mem_out(mem_out), m_id(id){} virtual ~scheduler_unit(){} virtual void add_supervised_warp_id(int i) { m_supervised_warps.push_back(&warp(i)); @@ -397,6 +398,7 @@ protected: register_set* m_sp_out; register_set* m_sfu_out; register_set* m_tensor_core_out; + register_set* m_vp_core_out; register_set* m_mem_out; int m_id; @@ -410,9 +412,10 @@ public: register_set* sp_out, register_set* sfu_out, register_set* tensor_core_out, + register_set* vp_core_out, register_set* mem_out, int id ) - : scheduler_unit ( stats, shader, scoreboard, simt, warp, sp_out, sfu_out, tensor_core_out, mem_out, id ){} + : scheduler_unit ( stats, shader, scoreboard, simt, warp, sp_out, sfu_out, tensor_core_out, vp_core_out, mem_out, id ){} virtual ~lrr_scheduler () {} virtual void order_warps (); virtual void done_adding_supervised_warps() { @@ -428,9 +431,10 @@ public: register_set* sp_out, register_set* sfu_out, register_set* tensor_core_out, + register_set* vp_core_out, register_set* mem_out, int id ) - : scheduler_unit ( stats, shader, scoreboard, simt, warp, sp_out, sfu_out, tensor_core_out, mem_out, id ){} + : scheduler_unit ( stats, shader, scoreboard, simt, warp, sp_out, sfu_out, tensor_core_out,vp_core_out, mem_out, id ){} virtual ~gto_scheduler () {} virtual void order_warps (); virtual void done_adding_supervised_warps() { @@ -448,10 +452,11 @@ public: register_set* sp_out, register_set* sfu_out, register_set* tensor_core_out, + register_set* vp_core_out, register_set* mem_out, int id, char* config_str ) - : scheduler_unit ( stats, shader, scoreboard, simt, warp, sp_out, sfu_out, tensor_core_out, mem_out, id ), + : scheduler_unit ( stats, shader, scoreboard, simt, warp, sp_out, sfu_out, tensor_core_out, vp_core_out, mem_out, id ), m_pending_warps() { unsigned inner_level_readin; @@ -499,6 +504,7 @@ public: register_set* sp_out, register_set* sfu_out, register_set* tensor_core_out, + register_set* vp_core_out, register_set* mem_out, int id, char* config_string ); @@ -1083,6 +1089,22 @@ public: virtual void active_lanes_in_pipeline(); virtual void issue( register_set& source_reg ); }; +class vp_core : public pipelined_simd_unit +{ +public: + vp_core( register_set* result_port, const shader_core_config *config, shader_core_ctx *core ); + virtual bool can_issue( const warp_inst_t &inst ) const + { + switch(inst.op) { + case VP_CORE_OP: break; + default: return false; + } + return pipelined_simd_unit::can_issue(inst); + } + virtual void active_lanes_in_pipeline(); + virtual void issue( register_set& source_reg ); +}; + class sp_unit : public pipelined_simd_unit @@ -1225,10 +1247,12 @@ enum pipeline_stage_name_t { ID_OC_SP=0, ID_OC_SFU, ID_OC_TENSOR_CORE, + ID_OC_VP_CORE, ID_OC_MEM, OC_EX_SP, OC_EX_SFU, OC_EX_TENSOR_CORE, + OC_EX_VP_CORE, OC_EX_MEM, EX_WB, N_PIPELINE_STAGES @@ -1238,10 +1262,12 @@ const char* const pipeline_stage_name_decode[] = { "ID_OC_SP", "ID_OC_SFU", "ID_OC_TENSOR_CORE", + "ID_OC_VP_CORE", "ID_OC_MEM", "OC_EX_SP", "OC_EX_SFU", "OC_EX_TENSOR_CORE", + "OC_EX_VP_CORE", "OC_EX_MEM", "EX_WB", "N_PIPELINE_STAGES" @@ -1286,13 +1312,21 @@ struct shader_core_config : public core_config max_warps_per_shader = n_thread_per_shader/warp_size; assert( !(n_thread_per_shader % warp_size) ); max_sfu_latency = 512; - max_tensor_core_latency = 64; max_sp_latency = 32; + + max_tensor_core_latency = 64; gpgpu_num_tensor_core_units=8; gpgpu_operand_collector_num_units_tensor_core=24; gpgpu_operand_collector_num_in_ports_tensor_core=8; gpgpu_operand_collector_num_out_ports_tensor_core=8; - m_L1I_config.init(m_L1I_config.m_config_string,FuncCachePreferNone); + + max_vp_core_latency = 64; + gpgpu_num_vp_core_units=8; + gpgpu_operand_collector_num_units_vp_core=24; + gpgpu_operand_collector_num_in_ports_vp_core=8; + gpgpu_operand_collector_num_out_ports_vp_core=8; + + m_L1I_config.init(m_L1I_config.m_config_string,FuncCachePreferNone); m_L1T_config.init(m_L1T_config.m_config_string,FuncCachePreferNone); m_L1C_config.init(m_L1C_config.m_config_string,FuncCachePreferNone); m_L1D_config.init(m_L1D_config.m_config_string,FuncCachePreferNone); @@ -1339,24 +1373,28 @@ struct shader_core_config : public core_config int gpgpu_operand_collector_num_units_sp; int gpgpu_operand_collector_num_units_sfu; int gpgpu_operand_collector_num_units_tensor_core; + int gpgpu_operand_collector_num_units_vp_core; int gpgpu_operand_collector_num_units_mem; int gpgpu_operand_collector_num_units_gen; unsigned int gpgpu_operand_collector_num_in_ports_sp; unsigned int gpgpu_operand_collector_num_in_ports_sfu; unsigned int gpgpu_operand_collector_num_in_ports_tensor_core; + unsigned int gpgpu_operand_collector_num_in_ports_vp_core; unsigned int gpgpu_operand_collector_num_in_ports_mem; unsigned int gpgpu_operand_collector_num_in_ports_gen; unsigned int gpgpu_operand_collector_num_out_ports_sp; unsigned int gpgpu_operand_collector_num_out_ports_sfu; unsigned int gpgpu_operand_collector_num_out_ports_tensor_core; + unsigned int gpgpu_operand_collector_num_out_ports_vp_core; unsigned int gpgpu_operand_collector_num_out_ports_mem; unsigned int gpgpu_operand_collector_num_out_ports_gen; int gpgpu_num_sp_units; int gpgpu_num_sfu_units; int gpgpu_num_tensor_core_units; + int gpgpu_num_vp_core_units; int gpgpu_num_mem_units; //Shader core resources @@ -1370,6 +1408,7 @@ struct shader_core_config : public core_config unsigned max_sp_latency; unsigned max_sfu_latency; unsigned max_tensor_core_latency; + unsigned max_vp_core_latency; unsigned n_simt_cores_per_cluster; unsigned n_simt_clusters; @@ -1408,6 +1447,7 @@ struct shader_core_stats_pod { unsigned *m_num_sp_acesses; unsigned *m_num_sfu_acesses; unsigned *m_num_tensor_core_acesses; + unsigned *m_num_vp_core_acesses; unsigned *m_num_trans_acesses; unsigned *m_num_mem_acesses; unsigned *m_num_sp_committed; @@ -1415,6 +1455,7 @@ struct shader_core_stats_pod { unsigned *m_num_tlb_accesses; unsigned *m_num_sfu_committed; unsigned *m_num_tensor_core_committed; + unsigned *m_num_vp_core_committed; unsigned *m_num_mem_committed; unsigned *m_read_regfile_acesses; unsigned *m_write_regfile_acesses; @@ -1424,6 +1465,7 @@ struct shader_core_stats_pod { unsigned *m_active_sp_lanes; unsigned *m_active_sfu_lanes; unsigned *m_active_tensor_core_lanes; + unsigned *m_active_vp_core_lanes; unsigned *m_active_fu_lanes; unsigned *m_active_fu_mem_lanes; unsigned *m_n_diverge; // number of divergence occurring in this shader @@ -1496,6 +1538,7 @@ public: m_num_sp_acesses= (unsigned*) calloc(config->num_shader(),sizeof(unsigned)); m_num_sfu_acesses= (unsigned*) calloc(config->num_shader(),sizeof(unsigned)); m_num_tensor_core_acesses= (unsigned*) calloc(config->num_shader(),sizeof(unsigned)); + m_num_vp_core_acesses= (unsigned*) calloc(config->num_shader(),sizeof(unsigned)); m_num_trans_acesses= (unsigned*) calloc(config->num_shader(),sizeof(unsigned)); m_num_mem_acesses= (unsigned*) calloc(config->num_shader(),sizeof(unsigned)); m_num_sp_committed= (unsigned*) calloc(config->num_shader(),sizeof(unsigned)); @@ -1504,10 +1547,12 @@ public: m_active_sp_lanes= (unsigned*) calloc(config->num_shader(),sizeof(unsigned)); m_active_sfu_lanes= (unsigned*) calloc(config->num_shader(),sizeof(unsigned)); m_active_tensor_core_lanes= (unsigned*) calloc(config->num_shader(),sizeof(unsigned)); + m_active_vp_core_lanes= (unsigned*) calloc(config->num_shader(),sizeof(unsigned)); m_active_fu_lanes= (unsigned*) calloc(config->num_shader(),sizeof(unsigned)); m_active_fu_mem_lanes= (unsigned*) calloc(config->num_shader(),sizeof(unsigned)); m_num_sfu_committed= (unsigned*) calloc(config->num_shader(),sizeof(unsigned)); m_num_tensor_core_committed= (unsigned*) calloc(config->num_shader(),sizeof(unsigned)); + m_num_vp_core_committed= (unsigned*) calloc(config->num_shader(),sizeof(unsigned)); m_num_mem_committed= (unsigned*) calloc(config->num_shader(),sizeof(unsigned)); m_read_regfile_acesses= (unsigned*) calloc(config->num_shader(),sizeof(unsigned)); m_write_regfile_acesses= (unsigned*) calloc(config->num_shader(),sizeof(unsigned)); -- cgit v1.3 From 7a77d951e6a900d61436df12826bb677aeaee6e6 Mon Sep 17 00:00:00 2001 From: aamir Date: Sun, 9 Sep 2018 15:10:06 -0700 Subject: minor changes for generating mem transaction in timing model. NOTE NOT COMPLETED --- cuda-kernels/gpgpusim.config | 2 +- cutlass-example/gpgpusim.config | 2 +- src/abstract_hardware_model.cc | 2 +- src/abstract_hardware_model.h | 8 +++-- src/cuda-sim/cuda-sim.cc | 65 ++++++++++++++++++++++++----------------- src/cuda-sim/instructions.cc | 64 +++++++++++++++++++++++++++++----------- src/gpgpu-sim/shader.cc | 2 +- src/gpgpu-sim/shader.h | 8 +++++ 8 files changed, 104 insertions(+), 49 deletions(-) (limited to 'src/gpgpu-sim') diff --git a/cuda-kernels/gpgpusim.config b/cuda-kernels/gpgpusim.config index 2510d21..3daa539 100755 --- a/cuda-kernels/gpgpusim.config +++ b/cuda-kernels/gpgpusim.config @@ -33,7 +33,7 @@ # ID_OC_SP,ID_OC_SFU,ID_OC_MEM,OC_EX_SP,OC_EX_SFU,OC_EX_MEM,EX_WB ## Pascal GP102 has 4 SP SIMD units and 1 SFU unit ## we need to scale the number of pipeline registers to be equal to the number of SP units --gpgpu_pipeline_widths 4,1,1,1,4,1,1,1,6 +-gpgpu_pipeline_widths 4,1,1,1,1,4,1,1,1,1,6 -gpgpu_num_sp_units 4 -gpgpu_num_sfu_units 1 # Instruction latencies and initiation intervals diff --git a/cutlass-example/gpgpusim.config b/cutlass-example/gpgpusim.config index 2510d21..3daa539 100644 --- a/cutlass-example/gpgpusim.config +++ b/cutlass-example/gpgpusim.config @@ -33,7 +33,7 @@ # ID_OC_SP,ID_OC_SFU,ID_OC_MEM,OC_EX_SP,OC_EX_SFU,OC_EX_MEM,EX_WB ## Pascal GP102 has 4 SP SIMD units and 1 SFU unit ## we need to scale the number of pipeline registers to be equal to the number of SP units --gpgpu_pipeline_widths 4,1,1,1,4,1,1,1,6 +-gpgpu_pipeline_widths 4,1,1,1,1,4,1,1,1,1,6 -gpgpu_num_sp_units 4 -gpgpu_num_sfu_units 1 # Instruction latencies and initiation intervals diff --git a/src/abstract_hardware_model.cc b/src/abstract_hardware_model.cc index d668de7..b24a77e 100644 --- a/src/abstract_hardware_model.cc +++ b/src/abstract_hardware_model.cc @@ -184,7 +184,7 @@ void warp_inst_t::generate_mem_accesses() { if( empty() || op == MEMORY_BARRIER_OP || m_mem_accesses_created ) return; - if ( !((op == LOAD_OP) || (op == STORE_OP)) ) + if (!((op == LOAD_OP) || (op==TENSOR_CORE_LOAD_OP) || (op==VP_LOAD_OP) || (op == STORE_OP)||(op==TENSOR_CORE_STORE_OP)||(op==VP_STORE_OP))) return; if( m_warp_active_mask.count() == 0 ) return; // predicated off diff --git a/src/abstract_hardware_model.h b/src/abstract_hardware_model.h index 781509f..9c418fa 100644 --- a/src/abstract_hardware_model.h +++ b/src/abstract_hardware_model.h @@ -81,6 +81,10 @@ enum uarch_op_t { VP_CORE_OP, ALU_SFU_OP, LOAD_OP, + VP_LOAD_OP, + TENSOR_CORE_LOAD_OP, + TENSOR_CORE_STORE_OP, + VP_STORE_OP, STORE_OP, BRANCH_OP, BARRIER_OP, @@ -801,8 +805,8 @@ public: { fprintf(fp," [inst @ pc=0x%04x] ", pc ); } - bool is_load() const { return (op == LOAD_OP || memory_op == memory_load); } - bool is_store() const { return (op == STORE_OP || memory_op == memory_store); } + bool is_load() const { return (op == LOAD_OP ||op==TENSOR_CORE_LOAD_OP||op==VP_LOAD_OP || memory_op == memory_load); } + bool is_store() const { return (op == STORE_OP ||op==TENSOR_CORE_STORE_OP||op==VP_STORE_OP || memory_op == memory_store); } unsigned get_num_operands() const {return num_operands;} unsigned get_num_regs() const {return num_regs;} void set_num_regs(unsigned num) {num_regs=num;} diff --git a/src/cuda-sim/cuda-sim.cc b/src/cuda-sim/cuda-sim.cc index bcf64c4..8f684e2 100644 --- a/src/cuda-sim/cuda-sim.cc +++ b/src/cuda-sim/cuda-sim.cc @@ -648,12 +648,12 @@ void ptx_instruction::set_opcode_and_latency() if ( has_memory_write() ) op = STORE_OP; break; case LD_OP: op = LOAD_OP; break; - case MMA_LD_OP: op = LOAD_OP; break; - case VP_LD_OP: op = LOAD_OP; break; + case MMA_LD_OP: op = TENSOR_CORE_LOAD_OP; break; + case VP_LD_OP: op = VP_LOAD_OP; break; case LDU_OP: op = LOAD_OP; break; case ST_OP: op = STORE_OP; break; - case MMA_ST_OP: op = STORE_OP; break; - case VP_ST_OP: op = STORE_OP; break; + case MMA_ST_OP: op = TENSOR_CORE_STORE_OP; break; + case VP_ST_OP: op = VP_STORE_OP; break; case BRA_OP: op = BRANCH_OP; break; case BREAKADDR_OP: op = BRANCH_OP; break; case TEX_OP: op = LOAD_OP; mem_op=TEX; break; @@ -1334,6 +1334,7 @@ void ptx_thread_info::ptx_exec_inst( warp_inst_t &inst, unsigned lane_id) skip = !pred_lookup(pI->get_pred_mod(), pred_value.pred & 0x000F); } } + int inst_opcode=pI->get_opcode(); if( skip ) { inst.set_not_active(lane_id); @@ -1346,17 +1347,21 @@ void ptx_thread_info::ptx_exec_inst( warp_inst_t &inst, unsigned lane_id) pI = pJ; } - int inst_opcode=pI->get_opcode(); + if(((inst_opcode==MMA_OP||inst_opcode==MMA_LD_OP||inst_opcode==MMA_ST_OP||inst_opcode==VP_MMA_OP||inst_opcode==VP_LD_OP||inst_opcode==VP_ST_OP))){ + if(inst.active_count()!=MAX_WARP_SIZE) + while(1); + } if(((inst_opcode!=MMA_OP)&&(inst_opcode!=MMA_LD_OP)&&(inst_opcode!=MMA_ST_OP)&&(inst_opcode!=VP_LD_OP)&&(inst_opcode!=VP_ST_OP)&&(inst_opcode!=VP_MMA_OP))||((inst_opcode==MMA_OP||inst_opcode==MMA_LD_OP||inst_opcode==MMA_ST_OP||inst_opcode==VP_MMA_OP||inst_opcode==VP_LD_OP||inst_opcode==VP_ST_OP)&&(lane_id==0))){ - switch ( inst_opcode ) { -#define OP_DEF(OP,FUNC,STR,DST,CLASSIFICATION) case OP: FUNC(pI,this); op_classification = CLASSIFICATION; break; -#define OP_W_DEF(OP,FUNC,STR,DST,CLASSIFICATION) case OP: FUNC(pI,get_core(),inst); op_classification = CLASSIFICATION; break; -#include "opcodes.def" -#undef OP_DEF -#undef OP_W_DEF - default: printf( "Execution error: Invalid opcode (0x%x)\n", pI->get_opcode() ); break; - }} + switch ( inst_opcode ) { + #define OP_DEF(OP,FUNC,STR,DST,CLASSIFICATION) case OP: FUNC(pI,this); op_classification = CLASSIFICATION; break; + #define OP_W_DEF(OP,FUNC,STR,DST,CLASSIFICATION) case OP: FUNC(pI,get_core(),inst); op_classification = CLASSIFICATION; break; + #include "opcodes.def" + #undef OP_DEF + #undef OP_W_DEF + default: printf( "Execution error: Invalid opcode (0x%x)\n", pI->get_opcode() ); break; + } + } delete pJ; pI = pI_saved; @@ -1398,13 +1403,17 @@ void ptx_thread_info::ptx_exec_inst( warp_inst_t &inst, unsigned lane_id) _memory_op_t insn_memory_op = no_memory_op; unsigned insn_data_size = 0; if ( (pI->has_memory_read() || pI->has_memory_write()) ) { - insn_memaddr = last_eaddr(); - insn_space = last_space(); - unsigned to_type = pI->get_type(); - insn_data_size = datatype2size(to_type); - insn_memory_op = pI->has_memory_read() ? memory_load : memory_store; - } - + //if(!((inst_opcode==MMA_LD_OP||inst_opcode==MMA_ST_OP||inst_opcode==VP_LD_OP||inst_opcode==VP_ST_OP))) + //if(!((inst_opcode==MMA_LD_OP||inst_opcode==VP_LD_OP))) + //{ + insn_memaddr = last_eaddr(); + insn_space = last_space(); + unsigned to_type = pI->get_type(); + insn_data_size = datatype2size(to_type); + insn_memory_op = pI->has_memory_read() ? memory_load : memory_store; + //} + } + if ( pI->get_opcode() == BAR_OP && pI->barrier_op() == RED_OPTION) { inst.add_callback( lane_id, last_callback().function, last_callback().instruction, this,false /*not atomic*/); } @@ -1476,12 +1485,16 @@ void ptx_thread_info::ptx_exec_inst( warp_inst_t &inst, unsigned lane_id) // "Return values" if(!skip) { - inst.space = insn_space; - inst.set_addr(lane_id, insn_memaddr); - inst.data_size = insn_data_size; // simpleAtomicIntrinsics - assert( inst.memory_op == insn_memory_op ); - } - + //if(!((inst_opcode==MMA_LD_OP||inst_opcode==MMA_ST_OP||inst_opcode==VP_LD_OP||inst_opcode==VP_ST_OP))) + //if(!((inst_opcode==MMA_LD_OP||inst_opcode==VP_LD_OP))) + //{ + inst.space = insn_space; + inst.set_addr(lane_id, insn_memaddr); + inst.data_size = insn_data_size; // simpleAtomicIntrinsics + assert( inst.memory_op == insn_memory_op ); + //} + } + } catch ( int x ) { printf("GPGPU-Sim PTX: ERROR (%d) executing intruction (%s:%u)\n", x, pI->source_file(), pI->source_line() ); printf("GPGPU-Sim PTX: '%s'\n", pI->get_source() ); diff --git a/src/cuda-sim/instructions.cc b/src/cuda-sim/instructions.cc index 39b8ba5..77e8e71 100644 --- a/src/cuda-sim/instructions.cc +++ b/src/cuda-sim/instructions.cc @@ -3196,6 +3196,7 @@ void mma_st_impl( const ptx_instruction *pI, core_t *core, warp_inst_t inst ) unsigned wmma_type = pI->get_wmma_type(); unsigned wmma_layout = pI->get_wmma_layout(0); int stride; + _memory_op_t insn_memory_op = pI->has_memory_read() ? memory_load : memory_store; for (thrd=0; thrd < core->get_warp_size(); thrd++) { thread = core->get_thread_info()[tid+thrd]; odd=thrd%2; @@ -3224,6 +3225,7 @@ void mma_st_impl( const ptx_instruction *pI, core_t *core, warp_inst_t inst ) if(g_debug_instruction) printf("mma_st: thrd=%d,addr=%x, fp(size=%d), stride=%d\n",thrd,addr_reg.u32,size,src2_data.u32); addr_t new_addr = addr+thread_group_offset(thrd,wmma_type,wmma_layout,type,stride)*size/8; + addr_t push_addr; ptx_reg_t nw_v[8]; for(k=0;k<8;k++){ @@ -3235,7 +3237,10 @@ void mma_st_impl( const ptx_instruction *pI, core_t *core, warp_inst_t inst ) for(k=0;k<8;k++){ if(type==F32_TYPE){ - mem->write(new_addr+4*acc_float_offset(k,wmma_layout,stride),size/8,&v[k].s64,thread,pI); + //mem->write(new_addr+4*acc_float_offset(k,wmma_layout,stride),size/8,&v[k].s64,thread,pI); + push_addr=new_addr+4*acc_float_offset(k,wmma_layout,stride); + mem->write(push_addr,size/8,&v[k].s64,thread,pI); + if(g_debug_instruction){ printf("wmma:store:thread%d=%x,%x,%x,%x,%x,%x,%x,%x\n",thrd,v[0].s64,v[1].s64,v[2].s64,v[3].s64,v[4].s64,v[5].s64,v[6].s64,v[7].s64); float temp; @@ -3250,10 +3255,16 @@ void mma_st_impl( const ptx_instruction *pI, core_t *core, warp_inst_t inst ) } else if(type==F16_TYPE){ - if(wmma_layout==ROW) - mem->write(new_addr+k*2,size/8,&nw_v[k].s64,thread,pI); - else if(wmma_layout==COL) - mem->write(new_addr+k*2*stride,size/8,&nw_v[k].s64,thread,pI); + if(wmma_layout==ROW){ + //mem->write(new_addr+k*2,size/8,&nw_v[k].s64,thread,pI); + push_addr=new_addr+k*2; + mem->write(push_addr,size/8,&nw_v[k].s64,thread,pI); + } + else if(wmma_layout==COL){ + //mem->write(new_addr+k*2*stride,size/8,&nw_v[k].s64,thread,pI); + push_addr=new_addr+k*2*stride; + mem->write(push_addr,size/8,&nw_v[k].s64,thread,pI); + } if(g_debug_instruction) printf("wmma:store:thread%d=%x,%x,%x,%x,%x,%x,%x,%x\n",thrd,nw_v[0].s64,nw_v[1].s64,nw_v[2].s64,nw_v[3].s64,nw_v[4].s64,nw_v[5].s64,nw_v[6].s64,nw_v[7].s64); } @@ -3427,6 +3438,7 @@ void mma_ld_impl( const ptx_instruction *pI, core_t *core, warp_inst_t inst ) int tid = inst.warp_id_func()*core->get_warp_size(); int thrd,stride; ptx_thread_info *thread; + _memory_op_t insn_memory_op = pI->has_memory_read() ? memory_load : memory_store; for (thrd=0; thrd < core->get_warp_size(); thrd++){ thread = core->get_thread_info()[tid+thrd]; @@ -3452,13 +3464,18 @@ void mma_ld_impl( const ptx_instruction *pI, core_t *core, warp_inst_t inst ) printf("mma_ld: thrd=%d,addr=%x, fpsize=%d, stride=%d\n",thrd,src1_data.u32,size,src2_data.u32); addr_t new_addr = addr+thread_group_offset(thrd,wmma_type,wmma_layout,type,stride)*size/8; - + addr_t fetch_addr; if(wmma_type==LOAD_A){ for(i=0;i<16;i++){ - if(wmma_layout==ROW) - mem->read(new_addr+2*i,size/8,&data[i].s64); + if(wmma_layout==ROW){ + //mem->read(new_addr+2*i,size/8,&data[i].s64); + fetch_addr=new_addr+2*i; + mem->read(fetch_addr,size/8,&data[i].s64); + } else if(wmma_layout==COL){ - mem->read(new_addr+2*(i%4)+2*stride*4*(i/4),size/8,&data[i].s64); + //mem->read(new_addr+2*(i%4)+2*stride*4*(i/4),size/8,&data[i].s64); + fetch_addr=new_addr+2*(i%4)+2*stride*4*(i/4); + mem->read(fetch_addr,size/8,&data[i].s64); } else{ printf("mma_ld:wrong_layout_type\n"); @@ -3468,10 +3485,15 @@ void mma_ld_impl( const ptx_instruction *pI, core_t *core, warp_inst_t inst ) } else if(wmma_type==LOAD_B){ for(i=0;i<16;i++){ - if(wmma_layout==COL) - mem->read(new_addr+2*i,size/8,&data[i].s64); + if(wmma_layout==COL){ + //mem->read(new_addr+2*i,size/8,&data[i].s64); + fetch_addr=new_addr+2*i; + mem->read(fetch_addr,size/8,&data[i].s64); + } else if(wmma_layout==ROW){ - mem->read(new_addr+2*(i%4)+2*stride*4*(i/4),size/8,&data[i].s64); + //mem->read(new_addr+2*(i%4)+2*stride*4*(i/4),size/8,&data[i].s64); + fetch_addr=new_addr+2*(i%4)+2*stride*4*(i/4); + mem->read(fetch_addr,size/8,&data[i].s64); } else{ printf("mma_ld:wrong_layout_type\n"); @@ -3482,17 +3504,25 @@ void mma_ld_impl( const ptx_instruction *pI, core_t *core, warp_inst_t inst ) else if(wmma_type==LOAD_C){ for(i=0;i<8;i++){ if(type==F16_TYPE){ - if(wmma_layout==ROW) - mem->read(new_addr+2*i,size/8,&data[i].s64); - else if(wmma_layout==COL) - mem->read(new_addr+2*stride*i,size/8,&data[i].s64); + if(wmma_layout==ROW){ + //mem->read(new_addr+2*i,size/8,&data[i].s64); + fetch_addr=new_addr+2*i; + mem->read(fetch_addr,size/8,&data[i].s64); + } + else if(wmma_layout==COL){ + //mem->read(new_addr+2*stride*i,size/8,&data[i].s64); + fetch_addr=new_addr+2*stride*i; + mem->read(fetch_addr,size/8,&data[i].s64); + } else{ printf("mma_ld:wrong_type\n"); abort(); } } else if(type==F32_TYPE){ - mem->read(new_addr+4*acc_float_offset(i,wmma_layout,stride),size/8,&data[i].s64); + //mem->read(new_addr+4*acc_float_offset(i,wmma_layout,stride),size/8,&data[i].s64); + fetch_addr=new_addr+4*acc_float_offset(i,wmma_layout,stride); + mem->read(fetch_addr,size/8,&data[i].s64); } else{ printf("wrong type"); diff --git a/src/gpgpu-sim/shader.cc b/src/gpgpu-sim/shader.cc index 6f11ad9..5e80fb1 100644 --- a/src/gpgpu-sim/shader.cc +++ b/src/gpgpu-sim/shader.cc @@ -920,7 +920,7 @@ void scheduler_unit::cycle() ready_inst = true; const active_mask_t &active_mask = m_simt_stack[warp_id]->get_active_mask(); assert( warp(warp_id).inst_in_pipeline() ); - if ( (pI->op == LOAD_OP) || (pI->op == STORE_OP) || (pI->op == MEMORY_BARRIER_OP) ) { + if ( (pI->op == LOAD_OP)||(pI->op ==TENSOR_CORE_LOAD_OP)||(pI->op ==VP_LOAD_OP)|| (pI->op == STORE_OP)|| (pI->op==TENSOR_CORE_STORE_OP) ||(pI->op==VP_STORE_OP) || (pI->op == MEMORY_BARRIER_OP) ) { if( m_mem_out->has_free() ) { m_shader->issue_warp(*m_mem_out,pI,active_mask,warp_id); issued++; diff --git a/src/gpgpu-sim/shader.h b/src/gpgpu-sim/shader.h index d292d56..d9558b0 100644 --- a/src/gpgpu-sim/shader.h +++ b/src/gpgpu-sim/shader.h @@ -1116,7 +1116,11 @@ public: switch(inst.op) { case SFU_OP: return false; case LOAD_OP: return false; + case TENSOR_CORE_LOAD_OP: return false; + case VP_LOAD_OP: return false; case STORE_OP: return false; + case TENSOR_CORE_STORE_OP: return false; + case VP_STORE_OP: return false; case MEMORY_BARRIER_OP: return false; default: break; } @@ -1158,7 +1162,11 @@ public: { switch(inst.op) { case LOAD_OP: break; + case TENSOR_CORE_LOAD_OP: break; + case VP_LOAD_OP: break; case STORE_OP: break; + case TENSOR_CORE_STORE_OP: break; + case VP_STORE_OP: break; case MEMORY_BARRIER_OP: break; default: return false; } -- cgit v1.3 From 3949357047a621a06b2e7fb4fd6099cce1469d27 Mon Sep 17 00:00:00 2001 From: aamir Date: Sun, 16 Sep 2018 20:07:33 -0700 Subject: print for mem transaction --- src/abstract_hardware_model.cc | 1 + src/abstract_hardware_model.h | 13 ++++++++++++- src/cuda-sim/instructions.cc | 15 ++++++++++++++- src/gpgpu-sim/shader.cc | 6 +++++- src/gpgpu-sim/traffic_breakdown.cc | 1 + 5 files changed, 33 insertions(+), 3 deletions(-) (limited to 'src/gpgpu-sim') diff --git a/src/abstract_hardware_model.cc b/src/abstract_hardware_model.cc index 83e76fe..72ece0b 100644 --- a/src/abstract_hardware_model.cc +++ b/src/abstract_hardware_model.cc @@ -352,6 +352,7 @@ void warp_inst_t::generate_mem_accesses() ptx_file_line_stats_add_uncoalesced_gmem( pc, m_accessq.size() - starting_queue_size ); } m_mem_accesses_created=true; + print_m_accessq(); } void warp_inst_t::memory_coalescing_arch_13( bool is_write, mem_access_type access_type ) diff --git a/src/abstract_hardware_model.h b/src/abstract_hardware_model.h index 9c418fa..d628745 100644 --- a/src/abstract_hardware_model.h +++ b/src/abstract_hardware_model.h @@ -933,7 +933,18 @@ public: for(unsigned i=0; i::iterator it; + for (it = m_accessq.begin(); it != m_accessq.end(); ++it){ + printf("MEM_TXN_GEN:%s:%x, Size:%d \n",mem_access_type_str(it->get_type()), it->get_addr(),it->get_size()); + } + } + } struct transaction_info { std::bitset<4> chunks; // bitmask: 32-byte chunks accessed mem_access_byte_mask_t bytes; diff --git a/src/cuda-sim/instructions.cc b/src/cuda-sim/instructions.cc index 42c63ca..52d89f2 100644 --- a/src/cuda-sim/instructions.cc +++ b/src/cuda-sim/instructions.cc @@ -3354,31 +3354,41 @@ void vp_ld_impl(const ptx_instruction *pI, core_t *core, warp_inst_t &inst) printf("vp_ld: thrx=%d,addr=%x, base_addr=%x, size=%d, stride=%d\n",thrd,new_addr,addr,size,src2_data.u32); if(wmma_type==LOAD_A||wmma_type==LOAD_C){ + printf("lda/c:"); for(i=0;i<8;i++){ if(wmma_layout==ROW){ //mem->read(new_addr+4*i,size/8,&data[i].s64); mem->read(new_addr+4*i,size/8,&data[i].s64); + printf("%x ", new_addr+4*i); mem_txn_addr[num_mem_txn++]=new_addr+4*i; } else if(wmma_layout==COL){ //mem->read(new_addr+4*stride*i,size/8,&data[i].s64); + printf("%x ", new_addr+4*stride*i); mem->read(new_addr+4*stride*i,size/8,&data[i].s64); mem_txn_addr[num_mem_txn++]=new_addr+4*stride*i; } } + } + else if(wmma_type==LOAD_B4){ + printf("ldb4:"); if(wmma_layout==ROW){ mem->read(new_addr,size/8,&data[0].s64); + printf("%x ",new_addr); mem_txn_addr[num_mem_txn++]=new_addr; } else if(wmma_layout==COL){ } } else if(wmma_type==LOAD_B8){ + printf("ldb8:"); if(wmma_layout==ROW){ mem->read(new_addr,size/8,&data[0].s64); mem->read(new_addr+4,size/8,&data[1].s64); + printf("%x ",new_addr,new_addr+4); + mem_txn_addr[num_mem_txn++]=new_addr; mem_txn_addr[num_mem_txn++]=new_addr+4; } @@ -3387,12 +3397,13 @@ void vp_ld_impl(const ptx_instruction *pI, core_t *core, warp_inst_t &inst) } } else if(wmma_type==LOAD_B16){ - printf("LOADB16_MODE"); + printf("ldb16:"); if(wmma_layout==ROW){ mem->read(new_addr,size/8,&data[0].s64); mem->read(new_addr+4,size/8,&data[1].s64); mem->read(new_addr+8,size/8,&data[2].s64); mem->read(new_addr+12,size/8,&data[3].s64); + printf("%x ",new_addr,new_addr+4,new_addr+8,new_addr+12); mem_txn_addr[num_mem_txn++]=new_addr; mem_txn_addr[num_mem_txn++]=new_addr+4; mem_txn_addr[num_mem_txn++]=new_addr+8; @@ -3405,6 +3416,8 @@ void vp_ld_impl(const ptx_instruction *pI, core_t *core, warp_inst_t &inst) printf("wrong vp_load type\n");; abort(); } + printf("\n"); + //generate timing memory request inst.space = space; inst.set_addr(thrd, (new_addr_type *)mem_txn_addr , num_mem_txn); diff --git a/src/gpgpu-sim/shader.cc b/src/gpgpu-sim/shader.cc index 5e80fb1..e745f03 100644 --- a/src/gpgpu-sim/shader.cc +++ b/src/gpgpu-sim/shader.cc @@ -741,7 +741,10 @@ void shader_core_ctx::func_exec_inst( warp_inst_t &inst ) { execute_warp_inst_t(inst); if( inst.is_load() || inst.is_store() ) - inst.generate_mem_accesses(); + { + inst.generate_mem_accesses(); + //inst.print_m_accessq(); + } } void shader_core_ctx::issue_warp( register_set& pipe_reg_set, const warp_inst_t* next_inst, const active_mask_t &active_mask, unsigned warp_id ) @@ -1512,6 +1515,7 @@ bool ldst_unit::memory_cycle( warp_inst_t &inst, mem_stage_stall_type &stall_rea // bypass L1 cache unsigned control_size = inst.is_store() ? WRITE_PACKET_SIZE : READ_PACKET_SIZE; unsigned size = access.get_size() + control_size; + //printf("Interconnect:Addr: %x, size=%d\n",access.get_addr(),size); if( m_icnt->full(size, inst.is_store() || inst.isatomic()) ) { stall_cond = ICNT_RC_FAIL; } else { diff --git a/src/gpgpu-sim/traffic_breakdown.cc b/src/gpgpu-sim/traffic_breakdown.cc index 32f0d30..587067f 100644 --- a/src/gpgpu-sim/traffic_breakdown.cc +++ b/src/gpgpu-sim/traffic_breakdown.cc @@ -46,6 +46,7 @@ std::string traffic_breakdown::classify_memfetch(class mem_fetch * mf) break; default: assert(0 && "Unknown traffic type"); } + printf("%s:Icnt:%s:Request: %x,Size%d,DataSize:%d,CntrlSize%d:\n",m_network_name.c_str(),traffic_name.c_str(),mf->get_addr(),mf->size(),mf->get_data_size(),mf->get_ctrl_size()); return traffic_name; } -- cgit v1.3 From 09e6092ace5213a5d5a49bf80b052802c06a4268 Mon Sep 17 00:00:00 2001 From: aamir Date: Sat, 22 Sep 2018 15:28:40 -0700 Subject: debug prints --- cuda-kernels/Makefile | 1 + cuda-kernels/genericMatrixMultiply.cu | 6 +- cuda-kernels/gpgpusim.config | 2 +- cuda-kernels/scripts | 118 +++++++++++++++++++++++++++++----- cuda-kernels/v16p_kernel.cu | 4 +- src/Makefile | 2 +- src/cuda-sim/Makefile | 2 +- src/cuda-sim/cuda-sim.cc | 8 ++- src/gpgpu-sim/Makefile | 2 +- src/gpgpu-sim/gpu-sim.cc | 4 +- src/gpgpu-sim/shader.cc | 16 +++-- src/gpgpu-sim/shader.h | 2 +- 12 files changed, 131 insertions(+), 36 deletions(-) (limited to 'src/gpgpu-sim') diff --git a/cuda-kernels/Makefile b/cuda-kernels/Makefile index 8effd11..b33519f 100755 --- a/cuda-kernels/Makefile +++ b/cuda-kernels/Makefile @@ -7,4 +7,5 @@ clean: rm _cuob* rm gpgpusim_power* rm gpgpu_inst_stats.txt + rm gpgpusim_visualizer* # nvcc -arch=sm_70 --gpu-architecture=compute_50 --gpu-code=compute_50 -lcudart -g -o tensor_core tensor_core.cu diff --git a/cuda-kernels/genericMatrixMultiply.cu b/cuda-kernels/genericMatrixMultiply.cu index 6a5d33f..8b96483 100644 --- a/cuda-kernels/genericMatrixMultiply.cu +++ b/cuda-kernels/genericMatrixMultiply.cu @@ -40,9 +40,9 @@ void cudaErrCheck_(cudaError_t stat, const char *file, int line) { using namespace nvcuda; // Must be multiples of 16 for wmma code to work -#define MATRIX_M (32) -#define MATRIX_N (32) -#define MATRIX_K (32) +#define MATRIX_M (64) +#define MATRIX_N (64) +#define MATRIX_K (64) diff --git a/cuda-kernels/gpgpusim.config b/cuda-kernels/gpgpusim.config index 3daa539..5f1be25 100755 --- a/cuda-kernels/gpgpusim.config +++ b/cuda-kernels/gpgpusim.config @@ -138,7 +138,7 @@ -gpgpu_memlatency_stat 14 -gpgpu_runtime_stat 500 -enable_ptx_file_line_stats 1 --visualizer_enabled 0 +-visualizer_enabled 1 # power model configs -power_simulation_enabled 1 diff --git a/cuda-kernels/scripts b/cuda-kernels/scripts index b9a3b50..a24fb72 100755 --- a/cuda-kernels/scripts +++ b/cuda-kernels/scripts @@ -10,11 +10,16 @@ nvcc --gpu-architecture=compute_70 --gpu-code=compute_70 -lcudart -g -o vp4_256 ./vp4_128>vp4_log128 ./vp4_256>vp4_log256 -grep "kernel_name.*vp_example" -A 589 vp4_log16 >vp4_16_summary -grep "kernel_name.*vp_example" -A 589 vp4_log32 >vp4_32_summary -grep "kernel_name.*vp_example" -A 589 vp4_log64 >vp4_64_summary -grep "kernel_name.*vp_example" -A 589 vp4_log128 >vp4_128_summary -grep "kernel_name.*vp_example" -A 589 vp4_log256 >vp4_256_summary +#grep "kernel_name.*convertInt32To" -A 589 vp4_log16 >vp4_16_summary +#grep "kernel_name.*convertInt32To" -A 589 vp4_log32 >vp4_32_summary +#grep "kernel_name.*convertInt32To" -A 589 vp4_log64 >vp4_64_summary +#grep "kernel_name.*convertInt32To" -A 589 vp4_log128 >vp4_128_summary +#grep "kernel_name.*convertInt32To" -A 589 vp4_log256 >vp4_256_summary +###grep "kernel_name.*vp_example" -A 589 vp4_log16 >>vp4_16_summary +###grep "kernel_name.*vp_example" -A 589 vp4_log32 >>vp4_32_summary +###grep "kernel_name.*vp_example" -A 589 vp4_log64 >>vp4_64_summary +#grep "kernel_name.*vp_example" -A 589 vp4_log128 >>vp4_128_summary +#grep "kernel_name.*vp_example" -A 589 vp4_log256 >>vp4_256_summary nvcc --gpu-architecture=compute_70 --gpu-code=compute_70 -lcudart -g -o vp8_16 v8p_genericMatrixMultiply.cu -DSIZE=16 @@ -28,11 +33,16 @@ nvcc --gpu-architecture=compute_70 --gpu-code=compute_70 -lcudart -g -o vp8_256 ./vp8_128>vp8_log128 ./vp8_256>vp8_log256 -grep "kernel_name.*vp_example" -A 589 vp8_log16 >vp8_16_summary -grep "kernel_name.*vp_example" -A 589 vp8_log32 >vp8_32_summary -grep "kernel_name.*vp_example" -A 589 vp8_log64 >vp8_64_summary -grep "kernel_name.*vp_example" -A 589 vp8_log128 >vp8_128_summary -grep "kernel_name.*vp_example" -A 589 vp8_log256 >vp8_256_summary +#grep "kernel_name.*convertInt32To" -A 589 vp8_log16 >vp8_16_summary +#grep "kernel_name.*convertInt32To" -A 589 vp8_log32 >vp8_32_summary +#grep "kernel_name.*convertInt32To" -A 589 vp8_log64 >vp8_64_summary +#grep "kernel_name.*convertInt32To" -A 589 vp8_log128 >vp8_128_summary +#grep "kernel_name.*convertInt32To" -A 589 vp8_log256 >vp8_256_summary +##grep "kernel_name.*vp_example" -A 589 vp8_log16 >>vp8_16_summary +##grep "kernel_name.*vp_example" -A 589 vp8_log32 >>vp8_32_summary +##grep "kernel_name.*vp_example" -A 589 vp8_log64 >>vp8_64_summary +#grep "kernel_name.*vp_example" -A 589 vp8_log128 >>vp8_128_summary +#grep "kernel_name.*vp_example" -A 589 vp8_log256 >>vp8_256_summary nvcc --gpu-architecture=compute_70 --gpu-code=compute_70 -lcudart -g -o vp16_16 v16p_genericMatrixMultiply.cu -DSIZE=16 nvcc --gpu-architecture=compute_70 --gpu-code=compute_70 -lcudart -g -o vp16_32 v16p_genericMatrixMultiply.cu -DSIZE=32 @@ -45,11 +55,85 @@ nvcc --gpu-architecture=compute_70 --gpu-code=compute_70 -lcudart -g -o vp16_256 ./vp16_128>vp16_log128 ./vp16_256>vp16_log256 -grep "kernel_name.*vp_example" -A 589 vp16_log16 >vp16_16_summary -grep "kernel_name.*vp_example" -A 589 vp16_log32 >vp16_32_summary -grep "kernel_name.*vp_example" -A 589 vp16_log64 >vp16_64_summary -grep "kernel_name.*vp_example" -A 589 vp16_log128 >vp16_128_summary -grep "kernel_name.*vp_example" -A 589 vp16_log256 >vp16_256_summary +#grep "kernel_name.*convertInt32To" -A 589 vp16_log16 >vp16_16_summary +#grep "kernel_name.*convertInt32To" -A 589 vp16_log32 >vp16_32_summary +#grep "kernel_name.*convertInt32To" -A 589 vp16_log64 >vp16_64_summary +#grep "kernel_name.*convertInt32To" -A 589 vp16_log128 >vp16_128_summary +#grep "kernel_name.*convertInt32To" -A 589 vp16_log256 >vp16_256_summary +##grep "kernel_name.*vp_example" -A 589 vp16_log16 >>vp16_16_summary +##grep "kernel_name.*vp_example" -A 589 vp16_log32 >>vp16_32_summary +##grep "kernel_name.*vp_example" -A 589 vp16_log64 >>vp16_64_summary +#grep "kernel_name.*vp_example" -A 589 vp16_log128 >>vp16_128_summary +#grep "kernel_name.*vp_example" -A 589 vp16_log256 >>vp16_256_summary -mv vp* VPlog/ -make clean +###sed -n '/bind.*to.*kernel.*vp_example/,$p' vp4_log16 > kernel_log_vp4_log16 +###sed -n '/bind.*to.*kernel.*vp_example/,$p' vp4_log32 > kernel_log_vp4_log32 +###sed -n '/bind.*to.*kernel.*vp_example/,$p' vp4_log64 > kernel_log_vp4_log64 +####sed -n '/bind.*to.*kernel.*vp_example/,$p' vp4_log128 > kernel_log_vp4_log128 +####sed -n '/bind.*to.*kernel.*vp_example/,$p' vp4_log256 > kernel_log_vp4_log256 +###sed -n '/bind.*to.*kernel.*vp_example/,$p' vp8_log16 > kernel_log_vp8_log16 +###sed -n '/bind.*to.*kernel.*vp_example/,$p' vp8_log32 > kernel_log_vp8_log32 +###sed -n '/bind.*to.*kernel.*vp_example/,$p' vp8_log64 > kernel_log_vp8_log64 +####sed -n '/bind.*to.*kernel.*vp_example/,$p' vp8_log128 > kernel_log_vp8_log128 +####sed -n '/bind.*to.*kernel.*vp_example/,$p' vp8_log256 > kernel_log_vp8_log256 +###sed -n '/bind.*to.*kernel.*vp_example/,$p' vp16_log16 > kernel_log_vp16_log16 +###sed -n '/bind.*to.*kernel.*vp_example/,$p' vp16_log32 > kernel_log_vp16_log32 +###sed -n '/bind.*to.*kernel.*vp_example/,$p' vp16_log64 > kernel_log_vp16_log64 +####sed -n '/bind.*to.*kernel.*vp_example/,$p' vp16_log128 > kernel_log_vp16_log128 +####sed -n '/bind.*to.*kernel.*vp_example/,$p' vp16_log256 > kernel_log_vp16_log256 +### +###grep memtocore:Icnt kernel_log_vp4_log16>memtocore_vp4_log16 +###grep coretomem:Icnt kernel_log_vp4_log16>coretomem_vp4_log16 +###grep MEM_TXN kernel_log_vp4_log16> gen_mem_vp4_log16 +###grep memtocore:Icnt kernel_log_vp4_log32>memtocore_vp4_log32 +###grep coretomem:Icnt kernel_log_vp4_log32>coretomem_vp4_log32 +###grep MEM_TXN kernel_log_vp4_log32> gen_mem_vp4_log32 +###grep memtocore:Icnt kernel_log_vp4_log64>memtocore_vp4_log64 +###grep coretomem:Icnt kernel_log_vp4_log64>coretomem_vp4_log64 +###grep MEM_TXN kernel_log_vp4_log64> gen_mem_vp4_log64 +####grep memtocore:Icnt kernel_log_vp4_log128>memtocore_vp4_log128 +####grep coretomem:Icnt kernel_log_vp4_log128>coretomem_vp4_log128 +####grep MEM_TXN kernel_log_vp4_log128> gen_mem_vp4_log128 +####grep memtocore:Icnt kernel_log_vp4_log256>memtocore_vp4_log256 +####grep coretomem:Icnt kernel_log_vp4_log256>coretomem_vp4_log256 +####grep MEM_TXN kernel_log_vp4_log256> gen_mem_vp4_log256 +### +###grep memtocore:Icnt kernel_log_vp8_log16>memtocore_vp8_log16 +###grep coretomem:Icnt kernel_log_vp8_log16>coretomem_vp8_log16 +###grep MEM_TXN kernel_log_vp8_log16> gen_mem_vp8_log16 +###grep memtocore:Icnt kernel_log_vp8_log32>memtocore_vp8_log32 +###grep coretomem:Icnt kernel_log_vp8_log32>coretomem_vp8_log32 +###grep MEM_TXN kernel_log_vp8_log32> gen_mem_vp8_log32 +###grep memtocore:Icnt kernel_log_vp8_log64>memtocore_vp8_log64 +###grep coretomem:Icnt kernel_log_vp8_log64>coretomem_vp8_log64 +###grep MEM_TXN kernel_log_vp8_log64> gen_mem_vp8_log64 +####grep memtocore:Icnt kernel_log_vp8_log128>memtocore_vp8_log128 +####grep coretomem:Icnt kernel_log_vp8_log128>coretomem_vp8_log128 +####grep MEM_TXN kernel_log_vp8_log128> gen_mem_vp8_log128 +####grep memtocore:Icnt kernel_log_vp8_log256>memtocore_vp8_log256 +####grep coretomem:Icnt kernel_log_vp8_log256>coretomem_vp8_log256 +####grep MEM_TXN kernel_log_vp8_log256> gen_mem_vp8_log256 +### +###grep memtocore:Icnt kernel_log_vp16_log16>memtocore_vp16_log16 +###grep coretomem:Icnt kernel_log_vp16_log16>coretomem_vp16_log16 +###grep MEM_TXN kernel_log_vp16_log16> gen_mem_vp16_log16 +###grep memtocore:Icnt kernel_log_vp16_log32>memtocore_vp16_log32 +###grep coretomem:Icnt kernel_log_vp16_log32>coretomem_vp16_log32 +###grep MEM_TXN kernel_log_vp16_log32> gen_mem_vp16_log32 +###grep memtocore:Icnt kernel_log_vp16_log64>memtocore_vp16_log64 +###grep coretomem:Icnt kernel_log_vp16_log64>coretomem_vp16_log64 +###grep MEM_TXN kernel_log_vp16_log64> gen_mem_vp16_log64 +####grep memtocore:Icnt kernel_log_vp16_log128>memtocore_vp16_log128 +####grep coretomem:Icnt kernel_log_vp16_log128>coretomem_vp16_log128 +####grep MEM_TXN kernel_log_vp16_log128> gen_mem_vp16_log128 +####grep memtocore:Icnt kernel_log_vp16_log256>memtocore_vp16_log256 +####grep coretomem:Icnt kernel_log_vp16_log256>coretomem_vp16_log256 +####grep MEM_TXN kernel_log_vp16_log256> gen_mem_vp16_log256 +### +###mv kernel_* MemTraffic/ +###mv memtocore* MemTraffic/ +###mv coretomem* MemTraffic/ +###mv gen_mem* MemTraffic/ +### +###mv vp* VPlog/ +###make clean diff --git a/cuda-kernels/v16p_kernel.cu b/cuda-kernels/v16p_kernel.cu index 864597d..d84328f 100644 --- a/cuda-kernels/v16p_kernel.cu +++ b/cuda-kernels/v16p_kernel.cu @@ -31,7 +31,7 @@ const int WMMA_M = 16; const int WMMA_N = 16; const int WMMA_K = 16; -__global__ void v4p_example(int *a_int32, int *b_int4, int *c,int *d_int32, int M, int N, int K) { +__global__ void v16p_example(int *a_int32, int *b_int4, int *c,int *d_int32, int M, int N, int K) { int registers_a[8]; int registers_b[8]; @@ -294,7 +294,7 @@ int main(int argc, char* argv[]) { //AAMIR //AAMIR printf("Running with wmma...\n"); cudaErrCheck(cudaEventRecord(startWMMA)); - v4p_example <<< 1, 32>>> (a_int32, b_int16, c_int32, d_int32, MATRIX_M, MATRIX_N, MATRIX_K); + v16p_example <<< 1, 32>>> (a_int32, b_int16, c_int32, d_int32, MATRIX_M, MATRIX_N, MATRIX_K); cudaErrCheck(cudaEventRecord(stopWMMA)); cudaErrCheck(cudaEventSynchronize(stopWMMA)); diff --git a/src/Makefile b/src/Makefile index 6001669..09194f3 100644 --- a/src/Makefile +++ b/src/Makefile @@ -46,7 +46,7 @@ ifeq ($(TRACE),1) endif ifneq ($(DEBUG),1) - OPTFLAGS += -O3 + OPTFLAGS += -O0 else CXXFLAGS += endif diff --git a/src/cuda-sim/Makefile b/src/cuda-sim/Makefile index e8977a6..d08a696 100644 --- a/src/cuda-sim/Makefile +++ b/src/cuda-sim/Makefile @@ -42,7 +42,7 @@ include ../../version_detection.mk OUTPUT_DIR=$(SIM_OBJ_FILES_DIR)/cuda-sim -OPT := -O3 -g3 -Wall -Wno-unused-function -Wno-sign-compare +OPT := -O0 -g3 -Wall -Wno-unused-function -Wno-sign-compare ifeq ($(DEBUG),1) OPT := -g3 -Wall -Wno-unused-function -Wno-sign-compare endif diff --git a/src/cuda-sim/cuda-sim.cc b/src/cuda-sim/cuda-sim.cc index 9735a0e..dd7edd9 100644 --- a/src/cuda-sim/cuda-sim.cc +++ b/src/cuda-sim/cuda-sim.cc @@ -809,7 +809,7 @@ void ptx_instruction::set_opcode_and_latency() } if(get_wmma_type()==VP_MMA8) { - latency = 128; + latency =128; initiation_interval = 128; } if(get_wmma_type()==VP_MMA16) @@ -1295,12 +1295,14 @@ void ptx_thread_info::ptx_exec_inst( warp_inst_t &inst, unsigned lane_id) addr_t pc = next_instr(); assert( pc == inst.pc ); // make sure timing model and functional model are in sync const ptx_instruction *pI = m_func_info->get_instruction(pc); - #if 0 + //#if 0 if(lane_id==0){ + //printf("EXECUTION_FLOW:LINE_NUM:%d\n",pI->source_line()); printf("EXECUTION_FLOW:LINE_NUM:%d\n",pI->source_line()); printf("EXECUTION_FLOW:INST:%s\n",pI->get_source()); + printf("EXECUTION_FLOW:PC%d\n",pc); } - #endif + //#endif set_npc( pc + pI->inst_size() ); diff --git a/src/gpgpu-sim/Makefile b/src/gpgpu-sim/Makefile index f10a8a4..4f77699 100644 --- a/src/gpgpu-sim/Makefile +++ b/src/gpgpu-sim/Makefile @@ -48,7 +48,7 @@ ifeq ($(GNUC_CPP0X), 1) endif ifneq ($(DEBUG),1) - OPTFLAGS += -O3 + OPTFLAGS += -O0 else CXXFLAGS += endif diff --git a/src/gpgpu-sim/gpu-sim.cc b/src/gpgpu-sim/gpu-sim.cc index 7a797b5..eba6f54 100644 --- a/src/gpgpu-sim/gpu-sim.cc +++ b/src/gpgpu-sim/gpu-sim.cc @@ -1469,7 +1469,9 @@ void gpgpu_sim::cycle() if( g_single_step && ((gpu_sim_cycle+gpu_tot_sim_cycle) >= g_single_step) ) { asm("int $03"); } - gpu_sim_cycle++; + printf("gpu_sim_cycle=%d\n",gpu_sim_cycle); + gpu_sim_cycle++; + if( g_interactive_debugger_enabled ) gpgpu_debug(); diff --git a/src/gpgpu-sim/shader.cc b/src/gpgpu-sim/shader.cc index e745f03..eeb59bb 100644 --- a/src/gpgpu-sim/shader.cc +++ b/src/gpgpu-sim/shader.cc @@ -1321,10 +1321,16 @@ void ldst_unit::get_L1T_sub_stats(struct cache_sub_stats &css) const{ void shader_core_ctx::warp_inst_complete(const warp_inst_t &inst) { - #if 0 - printf("[warp_inst_complete] uid=%u core=%u warp=%u pc=%#x @ time=%llu issued@%llu\n", - inst.get_uid(), m_sid, inst.warp_id(), inst.pc, gpu_tot_sim_cycle + gpu_sim_cycle, inst.get_issue_cycle()); - #endif + //#if 0 + //printf("[warp_inst_complete] uid=%u core=%u warp=%u pc=%#x @ time=%llu issued@%llu\n", + //const ptx_instruction *pI = m_func_info->get_instruction(inst.pc); + printf("[warp_inst_complete] uid=%u core=%u warp=%u pc=%d @ time=%llu \n", + inst.get_uid(), m_sid, inst.warp_id(), inst.pc, gpu_sim_cycle); + //printf("^instruction:%s",(pI->m_source).c_str()); + //inst.get_uid(), m_sid, inst.warp_id(), inst.pc, gpu_tot_sim_cycle + gpu_sim_cycle, inst.get_issue_cycle()); + //#endif + + if(inst.op_pipe==SP__OP) m_stats->m_num_sp_committed[m_sid]++; else if(inst.op_pipe==SFU__OP) @@ -1510,7 +1516,6 @@ bool ldst_unit::memory_cycle( warp_inst_t &inst, mem_stage_stall_type &stall_rea if (m_core->get_config()->gmem_skip_L1D) bypassL1D = true; } - if( bypassL1D ) { // bypass L1 cache unsigned control_size = inst.is_store() ? WRITE_PACKET_SIZE : READ_PACKET_SIZE; @@ -3472,6 +3477,7 @@ void simt_core_cluster::icnt_inject_request_packet(class mem_fetch *mf) case CONST_ACC_R: m_stats->gpgpu_n_mem_const++; break; case TEXTURE_ACC_R: m_stats->gpgpu_n_mem_texture++; break; case GLOBAL_ACC_R: m_stats->gpgpu_n_mem_read_global++; break; + //case GLOBAL_ACC_R: m_stats->gpgpu_n_mem_read_global++; printf("read_global%d\n",m_stats->gpgpu_n_mem_read_global); break; case GLOBAL_ACC_W: m_stats->gpgpu_n_mem_write_global++; break; case LOCAL_ACC_R: m_stats->gpgpu_n_mem_read_local++; break; case LOCAL_ACC_W: m_stats->gpgpu_n_mem_write_local++; break; diff --git a/src/gpgpu-sim/shader.h b/src/gpgpu-sim/shader.h index d9558b0..99abb63 100644 --- a/src/gpgpu-sim/shader.h +++ b/src/gpgpu-sim/shader.h @@ -1328,7 +1328,7 @@ struct shader_core_config : public core_config gpgpu_operand_collector_num_in_ports_tensor_core=8; gpgpu_operand_collector_num_out_ports_tensor_core=8; - max_vp_core_latency = 64; + max_vp_core_latency = 512; gpgpu_num_vp_core_units=8; gpgpu_operand_collector_num_units_vp_core=24; gpgpu_operand_collector_num_in_ports_vp_core=8; -- cgit v1.3