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/abstract_hardware_model.cc | 4 +++- 1 file changed, 3 insertions(+), 1 deletion(-) (limited to 'src/abstract_hardware_model.cc') diff --git a/src/abstract_hardware_model.cc b/src/abstract_hardware_model.cc index 84d165c..ab83ba4 100644 --- a/src/abstract_hardware_model.cc +++ b/src/abstract_hardware_model.cc @@ -213,6 +213,7 @@ void warp_inst_t::generate_mem_accesses() access_type = is_write? LOCAL_ACC_W: LOCAL_ACC_R; break; case shared_space: break; + case sstarr_space: break; default: assert(0); break; } @@ -220,7 +221,8 @@ void warp_inst_t::generate_mem_accesses() new_addr_type cache_block_size = 0; // in bytes switch( space.get_type() ) { - case shared_space: { + case shared_space: + case sstarr_space: { unsigned subwarp_size = m_config->warp_size / m_config->mem_warp_parts; unsigned total_accesses=0; for( unsigned subwarp=0; subwarp < m_config->mem_warp_parts; subwarp++ ) { -- 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/abstract_hardware_model.cc') 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 cc14e087c164a25cd9e2374777fdc4e57d9cec39 Mon Sep 17 00:00:00 2001 From: aamir Date: Fri, 14 Sep 2018 01:42:35 -0700 Subject: corrected the bug in vpcore functional model --- cuda-kernels/v16p_genericMatrixMultiply.cu | 24 ++++++++++++++---------- cuda-kernels/v16p_kernel.cu | 9 +++++---- cuda-kernels/v4p_genericMatrixMultiply.cu | 23 +++++++++++++---------- cuda-kernels/v4p_kernel.cu | 9 +++++---- cuda-kernels/v8p_genericMatrixMultiply.cu | 21 +++++++++++---------- cuda-kernels/v8p_kernel.cu | 9 +++++---- src/abstract_hardware_model.cc | 3 ++- src/cuda-sim/instructions.cc | 10 ++++++++-- 8 files changed, 63 insertions(+), 45 deletions(-) (limited to 'src/abstract_hardware_model.cc') diff --git a/cuda-kernels/v16p_genericMatrixMultiply.cu b/cuda-kernels/v16p_genericMatrixMultiply.cu index c36d257..284bd7e 100644 --- a/cuda-kernels/v16p_genericMatrixMultiply.cu +++ b/cuda-kernels/v16p_genericMatrixMultiply.cu @@ -1,5 +1,6 @@ #include #include +#include // Define some error checking macros. #define cudaErrCheck(stat) { cudaErrCheck_((stat), __FILE__, __LINE__); } @@ -20,9 +21,9 @@ void curandErrCheck_(curandStatus_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 (SIZE) +#define MATRIX_N (SIZE) +#define MATRIX_K (SIZE) // The only dimensions currently supported by WMMA @@ -83,7 +84,7 @@ __global__ void vp_example(int *a, int *b, int *c, int M, int N, int K ) { asm("CPTX_BEGIN"); asm("vp.load.b16.sync.row.m16n16k16.s32 {%0,%1,%2,%3},[%4],%5;" : "=r"(b_frag[0]),"=r"(b_frag[1]),"=r"(b_frag[2]),"=r"(b_frag[3]): - "l"(b+bRow*ldb/2+bCol),"r"(ldb/2) + "l"(b+bRow*ldb/2+bCol/2),"r"(ldb)//NOTE pointer arithematic is happenning over here so be extracareful ); asm("CPTX_END"); asm("*/"); @@ -118,7 +119,7 @@ __global__ void vp_example(int *a, int *b, int *c, int M, int N, int K ) { asm("vp.load.c.sync.row.m16n16k16.s32 {%0,%1,%2,%3,%4,%5,%6,%7},[%8],%9;" : "=r"(c_frag[0]), "=r"(c_frag[1]),"=r"(c_frag[2]),"=r"(c_frag[3]), "=r"(c_frag[4]),"=r"(c_frag[5]),"=r"(c_frag[6]),"=r"(c_frag[7]): - "l"(c+cRow*ldc),"r"(ldc) + "l"(c+cRow*ldc+cCol),"r"(ldc) ); asm("CPTX_END"); asm("*/"); @@ -223,7 +224,7 @@ int main(int argc, char* argv[]) { cudaEvent_t startWMMA; cudaEvent_t stopWMMA; - + srand(time(NULL)); cudaErrCheck(cudaEventCreate(&startWMMA)); cudaErrCheck(cudaEventCreate(&stopWMMA)); @@ -249,7 +250,8 @@ int main(int argc, char* argv[]) { printf("a_int32\n"); for(int m=0;m0); printf("D_CALCULATED\n"); diff --git a/cuda-kernels/v16p_kernel.cu b/cuda-kernels/v16p_kernel.cu index 31a1460..864597d 100644 --- a/cuda-kernels/v16p_kernel.cu +++ b/cuda-kernels/v16p_kernel.cu @@ -1,5 +1,6 @@ #include #include +#include // Define some error checking macros. #define cudaErrCheck(stat) { cudaErrCheck_((stat), __FILE__, __LINE__); } @@ -205,7 +206,7 @@ int main(int argc, char* argv[]) { cudaEvent_t startWMMA; cudaEvent_t stopWMMA; - + srand(time(NULL)); cudaErrCheck(cudaEventCreate(&startWMMA)); cudaErrCheck(cudaEventCreate(&stopWMMA)); @@ -231,7 +232,7 @@ int main(int argc, char* argv[]) { printf("a_int32\n"); for(int m=0;m #include +#include // Define some error checking macros. #define cudaErrCheck(stat) { cudaErrCheck_((stat), __FILE__, __LINE__); } @@ -20,9 +21,9 @@ void curandErrCheck_(curandStatus_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 (SIZE) +#define MATRIX_N (SIZE) +#define MATRIX_K (SIZE) // The only dimensions currently supported by WMMA @@ -83,7 +84,7 @@ __global__ void vp_example(int *a, int *b, int *c, int M, int N, int K ) { asm("CPTX_BEGIN"); asm("vp.load.b4.sync.row.m16n16k16.s32 {%0},[%1],%2;" : "=r"(b_frag[0]): - "l"(b+bRow*ldb/8+bCol),"r"(ldb/8) + "l"(b+bRow*ldb/8+bCol/8),"r"(ldb) ); asm("CPTX_END"); asm("*/"); @@ -118,7 +119,7 @@ __global__ void vp_example(int *a, int *b, int *c, int M, int N, int K ) { asm("vp.load.c.sync.row.m16n16k16.s32 {%0,%1,%2,%3,%4,%5,%6,%7},[%8],%9;" : "=r"(c_frag[0]), "=r"(c_frag[1]),"=r"(c_frag[2]),"=r"(c_frag[3]), "=r"(c_frag[4]),"=r"(c_frag[5]),"=r"(c_frag[6]),"=r"(c_frag[7]): - "l"(c+cRow*ldc),"r"(ldc) + "l"(c+cRow*ldc+cCol),"r"(ldc) ); asm("CPTX_END"); asm("*/"); @@ -223,7 +224,8 @@ int main(int argc, char* argv[]) { cudaEvent_t startWMMA; cudaEvent_t stopWMMA; - + + srand (time(NULL)); cudaErrCheck(cudaEventCreate(&startWMMA)); cudaErrCheck(cudaEventCreate(&stopWMMA)); @@ -249,7 +251,8 @@ int main(int argc, char* argv[]) { printf("a_int32\n"); for(int m=0;m0); printf("D_CALCULATED\n"); diff --git a/cuda-kernels/v4p_kernel.cu b/cuda-kernels/v4p_kernel.cu index f1a9dbe..053c07b 100644 --- a/cuda-kernels/v4p_kernel.cu +++ b/cuda-kernels/v4p_kernel.cu @@ -1,4 +1,5 @@ #include +#include #include // Define some error checking macros. @@ -205,7 +206,7 @@ int main(int argc, char* argv[]) { cudaEvent_t startWMMA; cudaEvent_t stopWMMA; - + srand (time(NULL)); cudaErrCheck(cudaEventCreate(&startWMMA)); cudaErrCheck(cudaEventCreate(&stopWMMA)); @@ -231,7 +232,7 @@ int main(int argc, char* argv[]) { printf("a_int32\n"); for(int m=0;m #include +#include // Define some error checking macros. #define cudaErrCheck(stat) { cudaErrCheck_((stat), __FILE__, __LINE__); } @@ -20,9 +21,9 @@ void curandErrCheck_(curandStatus_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 (SIZE) +#define MATRIX_N (SIZE) +#define MATRIX_K (SIZE) // The only dimensions currently supported by WMMA @@ -83,7 +84,7 @@ __global__ void vp_example(int *a, int *b, int *c, int M, int N, int K ) { asm("CPTX_BEGIN"); asm("vp.load.b8.sync.row.m16n16k16.s32 {%0,%1},[%2],%3;" : "=r"(b_frag[0]),"=r"(b_frag[1]): - "l"(b+bRow*ldb/4+bCol),"r"(ldb/4) + "l"(b+bRow*ldb/4+bCol/4),"r"(ldb) ); asm("CPTX_END"); asm("*/"); @@ -118,7 +119,7 @@ __global__ void vp_example(int *a, int *b, int *c, int M, int N, int K ) { asm("vp.load.c.sync.row.m16n16k16.s32 {%0,%1,%2,%3,%4,%5,%6,%7},[%8],%9;" : "=r"(c_frag[0]), "=r"(c_frag[1]),"=r"(c_frag[2]),"=r"(c_frag[3]), "=r"(c_frag[4]),"=r"(c_frag[5]),"=r"(c_frag[6]),"=r"(c_frag[7]): - "l"(c+cRow*ldc),"r"(ldc) + "l"(c+cRow*ldc+cCol),"r"(ldc) ); asm("CPTX_END"); asm("*/"); @@ -223,7 +224,7 @@ int main(int argc, char* argv[]) { cudaEvent_t startWMMA; cudaEvent_t stopWMMA; - + srand(time(NULL)); cudaErrCheck(cudaEventCreate(&startWMMA)); cudaErrCheck(cudaEventCreate(&stopWMMA)); @@ -249,7 +250,7 @@ int main(int argc, char* argv[]) { printf("a_int32\n"); for(int m=0;m0); printf("D_CALCULATED\n"); diff --git a/cuda-kernels/v8p_kernel.cu b/cuda-kernels/v8p_kernel.cu index f824eac..9bfcea7 100644 --- a/cuda-kernels/v8p_kernel.cu +++ b/cuda-kernels/v8p_kernel.cu @@ -1,5 +1,6 @@ #include #include +#include // Define some error checking macros. #define cudaErrCheck(stat) { cudaErrCheck_((stat), __FILE__, __LINE__); } @@ -204,7 +205,7 @@ int main(int argc, char* argv[]) { cudaEvent_t startWMMA; cudaEvent_t stopWMMA; - + srand(time(NULL)); cudaErrCheck(cudaEventCreate(&startWMMA)); cudaErrCheck(cudaEventCreate(&stopWMMA)); @@ -231,7 +232,7 @@ int main(int argc, char* argv[]) { printf("a_int32\n"); for(int m=0;mread(new_addr+4*stride*i,size/8,&data[i].s64); mem->read(new_addr+4*stride*i,size/8,&data[i].s64); - mem_txn_addr[num_mem_txn++]=new_addr+4*i; + mem_txn_addr[num_mem_txn++]=new_addr+4*stride*i; } } } @@ -3410,7 +3410,13 @@ void vp_ld_impl(const ptx_instruction *pI, core_t *core, warp_inst_t &inst) inst.set_addr(thrd, (new_addr_type *)mem_txn_addr , num_mem_txn); inst.data_size = 4; // 4 byte transaction assert( inst.memory_op == insn_memory_op ); - + printf("\nVPLOAD_THREAD%d:",thrd); + for(int lll=0;lll::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