From 11b308e7363e937966b035b4891db32b4eece3bf Mon Sep 17 00:00:00 2001 From: Tor Aamodt Date: Fri, 1 Oct 2010 08:55:28 -0800 Subject: integrating recent changes from fermi-test into fermi (i'll use "fermi" for more disruptive changes to the pipeline model such as updating the MSHRs and getting rid of the warp tracker, ripping out DWF, etc...) [git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 7805] --- src/cuda-sim/cuda-math.h | 234 ++++++++++++- src/cuda-sim/cuda-sim.cc | 789 ++++++++++++++++++------------------------- src/cuda-sim/cuda-sim.h | 49 +-- src/cuda-sim/dram_callback.h | 6 - src/cuda-sim/instructions.cc | 149 ++++++-- src/cuda-sim/memory.cc | 78 ++++- src/cuda-sim/memory.h | 1 + src/cuda-sim/opcodes.def | 4 +- src/cuda-sim/ptx-stats.cc | 9 +- src/cuda-sim/ptx-stats.h | 6 +- src/cuda-sim/ptx.l | 11 +- src/cuda-sim/ptx.y | 11 +- src/cuda-sim/ptx_ir.cc | 72 ++-- src/cuda-sim/ptx_ir.h | 50 ++- src/cuda-sim/ptx_loader.cc | 137 +++++++- src/cuda-sim/ptx_loader.h | 5 +- src/cuda-sim/ptx_parser.cc | 61 +++- src/cuda-sim/ptx_parser.h | 3 +- src/cuda-sim/ptx_sim.cc | 80 ++++- src/cuda-sim/ptx_sim.h | 144 ++++---- 20 files changed, 1168 insertions(+), 731 deletions(-) (limited to 'src/cuda-sim') diff --git a/src/cuda-sim/cuda-math.h b/src/cuda-sim/cuda-math.h index 5abf4f3..367563f 100644 --- a/src/cuda-sim/cuda-math.h +++ b/src/cuda-sim/cuda-math.h @@ -47,6 +47,7 @@ namespace cuda_math { #define __attribute__(a) // to remove warnings inside math_functions.h #undef INT_MAX +#if CUDART_VERSION < 3000 // DEVICE_BUILTIN struct int4 { int x, y, z, w; @@ -70,18 +71,245 @@ namespace cuda_math { extern float rsqrtf(float); // CUDA 2.3 beta -#if CUDART_VERSION < 3000 #define CUDA_FLOAT_MATH_FUNCTIONS #include #define __CUDA_INTERNAL_COMPILATION__ #include #undef __CUDA_INTERNAL_COMPILATION__ #undef __attribute__ + +// float to integer conversion +int float2int(float a, enum cudaRoundMode mode) +{ + return __internal_float2uint(a, mode); +} + +// float to unsigned integer conversion +unsigned int float2uint(float a, enum cudaRoundMode mode) +{ + return __internal_float2uint(a, mode); +} + +float __ll2float_rz(long long int a) { + int orig_rnd_mode = fegetround(); + fesetround(FE_TOWARDZERO); + float b = a; + fesetround(orig_rnd_mode); + return b; +} +float __ll2float_ru(long long int a) { + int orig_rnd_mode = fegetround(); + fesetround(FE_UPWARD); + float b = a; + fesetround(orig_rnd_mode); + return b; +} +float __ll2float_rd(long long int a) { + int orig_rnd_mode = fegetround(); + fesetround(FE_DOWNWARD); + float b = a; + fesetround(orig_rnd_mode); + return b; +} + #else + #define CUDA_FLOAT_MATH_FUNCTIONS -#include "/home/taamodt/nvcuda/2.3/cuda/include/device_types.h" +#define __CUDACC__ + +// implementing int to float intrinsics with different rounding modes +#include +#include + +// 32-bit integer to float +float __int2float_rn(int a) { + int orig_rnd_mode = fegetround(); + fesetround(FE_TONEAREST); + float b = a; + fesetround(orig_rnd_mode); + return b; +} +float __int2float_rz(int a) { + int orig_rnd_mode = fegetround(); + fesetround(FE_TOWARDZERO); + float b = a; + fesetround(orig_rnd_mode); + return b; +} +float __int2float_ru(int a) { + int orig_rnd_mode = fegetround(); + fesetround(FE_UPWARD); + float b = a; + fesetround(orig_rnd_mode); + return b; +} +float __int2float_rd(int a) { + int orig_rnd_mode = fegetround(); + fesetround(FE_DOWNWARD); + float b = a; + fesetround(orig_rnd_mode); + return b; +} + +// 32-bit unsigned integer to float +float __uint2float_rn(unsigned int a) { + int orig_rnd_mode = fegetround(); + fesetround(FE_TONEAREST); + float b = a; + fesetround(orig_rnd_mode); + return b; +} +float __uint2float_rz(unsigned int a) { + int orig_rnd_mode = fegetround(); + fesetround(FE_TOWARDZERO); + float b = a; + fesetround(orig_rnd_mode); + return b; +} +float __uint2float_ru(unsigned int a) { + int orig_rnd_mode = fegetround(); + fesetround(FE_UPWARD); + float b = a; + fesetround(orig_rnd_mode); + return b; +} +float __uint2float_rd(unsigned int a) { + int orig_rnd_mode = fegetround(); + fesetround(FE_DOWNWARD); + float b = a; + fesetround(orig_rnd_mode); + return b; +} + +// 64-bit integer to float +float __ll2float_rn(long long int a) { + int orig_rnd_mode = fegetround(); + fesetround(FE_TONEAREST); + float b = a; + fesetround(orig_rnd_mode); + return b; +} +float __ll2float_rz(long long int a) { + int orig_rnd_mode = fegetround(); + fesetround(FE_TOWARDZERO); + float b = a; + fesetround(orig_rnd_mode); + return b; +} +float __ll2float_ru(long long int a) { + int orig_rnd_mode = fegetround(); + fesetround(FE_UPWARD); + float b = a; + fesetround(orig_rnd_mode); + return b; +} +float __ll2float_rd(long long int a) { + int orig_rnd_mode = fegetround(); + fesetround(FE_DOWNWARD); + float b = a; + fesetround(orig_rnd_mode); + return b; +} + +// 64-bit unsigned integer to float +float __ull2float_rn(unsigned long long int a) { + int orig_rnd_mode = fegetround(); + fesetround(FE_TONEAREST); + float b = a; + fesetround(orig_rnd_mode); + return b; +} +float __ull2float_rz(unsigned long long int a) { + int orig_rnd_mode = fegetround(); + fesetround(FE_TOWARDZERO); + float b = a; + fesetround(orig_rnd_mode); + return b; +} +float __ull2float_ru(unsigned long long int a) { + int orig_rnd_mode = fegetround(); + fesetround(FE_UPWARD); + float b = a; + fesetround(orig_rnd_mode); + return b; +} +float __ull2float_rd(unsigned long long int a) { + int orig_rnd_mode = fegetround(); + fesetround(FE_DOWNWARD); + float b = a; + fesetround(orig_rnd_mode); + return b; +} + +// float to integer conversion +int float2int(float a, enum cudaRoundMode mode) +{ + int tmp; + switch (mode) { + case cuda_math::cudaRoundZero: tmp = truncf(a); break; + case cuda_math::cudaRoundNearest: tmp = nearbyintf(a); break; + case cuda_math::cudaRoundMinInf: tmp = floorf(a); break; + case cuda_math::cudaRoundPosInf: tmp = ceilf(a); break; + default: abort(); + } + return tmp; +} + +int __internal_float2int(float a, enum cudaRoundMode mode) +{ + return float2int(a, mode); +} + +// float to unsigned integer conversion +unsigned int float2uint(float a, enum cudaRoundMode mode) +{ + unsigned int tmp; + switch (mode) { + case cuda_math::cudaRoundZero: tmp = truncf(a); break; + case cuda_math::cudaRoundNearest: tmp = nearbyintf(a); break; + case cuda_math::cudaRoundMinInf: tmp = floorf(a); break; + case cuda_math::cudaRoundPosInf: tmp = ceilf(a); break; + default: abort(); + } + return tmp; +} + +unsigned int __internal_float2uint(float a, enum cudaRoundMode mode) +{ + return float2uint(a, mode); +} + +// intrinsic for division +float fdividef(float a, float b) +{ + return (a / b); +} + +float __internal_accurate_fdividef(float a, float b) +{ + return fdividef(a, b); +} + +// intrinsic for saturate (clamp values beyond 0 and 1) +float __saturatef(float a) +{ + float b; + if (isnan(a)) b = 0.0f; + else if (a >= 1.0f) b = 1.0f; + else if (a <= 0.0f) b = 0.0f; + else b = a; + return b; +} + +// intrinsic for power +float __powf(float a, float b) +{ + return powf(a, b); +} + +#undef __CUDACC__ #define __CUDA_INTERNAL_COMPILATION__ -#include "/home/taamodt/nvcuda/2.3/cuda/include/math_functions.h" +#include #undef __CUDA_INTERNAL_COMPILATION__ #undef __attribute__ #endif diff --git a/src/cuda-sim/cuda-sim.cc b/src/cuda-sim/cuda-sim.cc index 0a6ccd6..de4445a 100644 --- a/src/cuda-sim/cuda-sim.cc +++ b/src/cuda-sim/cuda-sim.cc @@ -99,7 +99,6 @@ std::map TextureToArrayM std::map TextureToInfoMap; std::map NameToTextureMap; unsigned int g_texcache_linesize; -int gpgpu_option_spread_blocks_across_cores = 0; unsigned gpgpu_param_num_shaders = 0; void gpgpu_ptx_sim_bindNameToTexture(const char* name, const struct textureReference* texref) @@ -208,10 +207,12 @@ int gpgpu_ptx_sim_sizeofTexture(const char* name) return array->size; } -unsigned g_assemble_code_next_pc=1; +unsigned g_assemble_code_next_pc=0; std::map g_pc_to_finfo; std::vector function_info::s_g_pc_to_insn; +#define MAX_INST_SIZE 8 /*bytes*/ + void function_info::ptx_assemble() { if( m_assembled ) { @@ -220,16 +221,22 @@ void function_info::ptx_assemble() // get the instructions into instruction memory... unsigned num_inst = m_instructions.size(); - m_instr_mem = new ptx_instruction*[ num_inst ]; - m_instr_mem_size = num_inst; + m_instr_mem_size = MAX_INST_SIZE*(num_inst+1); + m_instr_mem = new ptx_instruction*[ m_instr_mem_size ]; printf("GPGPU-Sim PTX: instruction assembly for function \'%s\'... ", m_name.c_str() ); fflush(stdout); std::list::iterator i; - addr_t n=0; // offset in m_instr_mem + addr_t PC = g_assemble_code_next_pc; // globally unique address (across functions) + // start function on an aligned address + for( unsigned i=0; i < (PC%MAX_INST_SIZE); i++ ) + s_g_pc_to_insn.push_back((ptx_instruction*)NULL); + PC += PC%MAX_INST_SIZE; m_start_PC = PC; - s_g_pc_to_insn.reserve(s_g_pc_to_insn.size() + m_instructions.size()); + + addr_t n=0; // offset in m_instr_mem + s_g_pc_to_insn.reserve(s_g_pc_to_insn.size() + MAX_INST_SIZE*m_instructions.size()); for ( i=m_instructions.begin(); i != m_instructions.end(); i++ ) { ptx_instruction *pI = *i; if ( pI->is_label() ) { @@ -239,17 +246,22 @@ void function_info::ptx_assemble() g_pc_to_finfo[PC] = this; m_instr_mem[n] = pI; s_g_pc_to_insn.push_back(pI); - assert(pI == s_g_pc_to_insn[PC - 1]); + assert(pI == s_g_pc_to_insn[PC]); pI->set_m_instr_mem_index(n); pI->set_PC(PC); - n++; - PC++; + assert( pI->inst_size() <= MAX_INST_SIZE ); + for( unsigned i=1; i < pI->inst_size(); i++ ) { + s_g_pc_to_insn.push_back((ptx_instruction*)NULL); + m_instr_mem[n+i]=NULL; + } + n += pI->inst_size(); + PC += pI->inst_size(); } } g_assemble_code_next_pc=PC; - for ( unsigned ii=0; ii < n; ii++ ) { // handle branch instructions + for ( unsigned ii=0; ii < n; ii += m_instr_mem[ii]->inst_size() ) { // handle branch instructions ptx_instruction *pI = m_instr_mem[ii]; - if ( pI->get_opcode() == BRA_OP || pI->get_opcode() == BREAKADDR_OP ) { + if ( pI->get_opcode() == BRA_OP || pI->get_opcode() == BREAKADDR_OP || pI->get_opcode() == CALLP_OP) { operand_info &target = pI->dst(); //get operand, e.g. target name if ( labels.find(target.name()) == labels.end() ) { printf("GPGPU-Sim PTX: Loader error (%s:%u): Branch label \"%s\" does not appear in assembly code.", @@ -262,6 +274,11 @@ void function_info::ptx_assemble() target.set_type(label_t); } } + for ( unsigned ii=0; ii < n; ii += m_instr_mem[ii]->inst_size() ) { // handle branch instructions + ptx_instruction *pI = m_instr_mem[ii]; + pI->pre_decode(); + } + printf(" done.\n"); fflush(stdout); @@ -328,12 +345,12 @@ bool isspace_shared( unsigned smid, addr_t addr ) bool isspace_global( addr_t addr ) { - return (addr > GLOBAL_HEAP_START) || (addr < STATIC_ALLOC_LIMIT); + return (addr >= GLOBAL_HEAP_START) || (addr < STATIC_ALLOC_LIMIT); } memory_space_t whichspace( addr_t addr ) { - if( (addr > GLOBAL_HEAP_START) || (addr < STATIC_ALLOC_LIMIT) ) { + if( (addr >= GLOBAL_HEAP_START) || (addr < STATIC_ALLOC_LIMIT) ) { return global_space; } else if( addr > SHARED_GENERIC_START ) { return shared_space; @@ -446,14 +463,6 @@ void gpgpu_ptx_sim_memset( size_t dst_start_addr, int c, size_t count ) fflush(stdout); } -int ptx_thread_done( void *thd ) -{ - ptx_thread_info *the_thread = (ptx_thread_info *) thd; - int result = 0; - result = (the_thread==NULL) || the_thread->is_done(); - return result; -} - const char * ptx_get_fname( unsigned PC ) { static const char *null_ptr = ""; @@ -471,14 +480,6 @@ unsigned ptx_thread_donecycle( void *thr ) return the_thread->donecycle(); } -int ptx_thread_get_next_pc( void *thd ) -{ - ptx_thread_info *the_thread = (ptx_thread_info *) thd; - if ( the_thread == NULL ) - return -1; - return the_thread->get_pc(); // PC should already be updatd to next PC at this point (was set in shader_decode() last time thread ran) -} - void* ptx_thread_get_next_finfo( void *thd ) { ptx_thread_info *the_thread = (ptx_thread_info *) thd; @@ -531,7 +532,7 @@ void ptx_print_insn( address_type pc, FILE *fp ) { std::map::iterator f = g_pc_to_finfo.find(pc); if( f == g_pc_to_finfo.end() ) { - fprintf(fp,"", pc, pc ); + fprintf(fp,"", pc ); return; } function_info *finfo = f->second; @@ -539,108 +540,118 @@ void ptx_print_insn( address_type pc, FILE *fp ) finfo->print_insn(pc,fp); } -static void get_opcode_info( const ptx_instruction *pI, unsigned opcode, unsigned *cycles, unsigned *op_type ) +static void get_opcode_info( const ptx_instruction *pI, unsigned opcode, unsigned *cycles, op_type *op ) { - *op_type = ALU_OP; + *op = ALU_OP; *cycles = 1; if ( opcode == LD_OP ) { - *op_type = LOAD_OP; + *op = LOAD_OP; } else if ( opcode == ST_OP ) { - *op_type = STORE_OP; + *op = STORE_OP; } else if ( opcode == BRA_OP ) { - *op_type = BRANCH_OP; + *op = BRANCH_OP; } else if ( opcode == BREAKADDR_OP ) { - *op_type = BRANCH_OP; + *op = BRANCH_OP; } else if ( opcode == TEX_OP ) { - *op_type = LOAD_OP; + *op = LOAD_OP; } else if ( opcode == ATOM_OP ) { - *op_type = LOAD_OP; // make atomics behave more like a load. + *op = LOAD_OP; // timing model treats this like load for now } else if ( opcode == BAR_OP ) { - *op_type = BARRIER_OP; - } + *op = BARRIER_OP; + } else if ( opcode == MEMBAR_OP ) + *op = MEMORY_BARRIER_OP; // Floating point instructions - if( opcode == RCP_OP || opcode == LG2_OP || opcode == RSQRT_OP ) { - *cycles = 4; - *op_type = SFU_OP; + if( opcode == RCP_OP ) { + *cycles = 2; + *op = SFU_OP; + } else if ( opcode == LG2_OP || opcode == RSQRT_OP ) { + *cycles = 4; + *op = SFU_OP; } else if( opcode == SQRT_OP || opcode == SIN_OP || opcode == COS_OP || opcode == EX2_OP ) { *cycles = 4; - *op_type = SFU_OP; + *op = SFU_OP; } else if( opcode == DIV_OP ) { // Floating point only if( pI->get_type() == F32_TYPE || pI->get_type() == F64_TYPE ) { - *cycles = 8; - *op_type = SFU_OP; + *cycles = 4; + *op = SFU_OP; } } // Integer instructions if( opcode == MUL_OP ) { if( pI->get_type() == B32_TYPE || pI->get_type() == U32_TYPE || pI->get_type() == S32_TYPE ) { // 32-bit integer instruction - *cycles = 4; - *op_type = SFU_OP; + *cycles = 5; + *op = SFU_OP; } - } + if( pI->get_type() == F32_TYPE || pI->get_type() == F64_TYPE ) + *op = ALU_SFU_OP; + } + if( opcode == MAD_OP ) { + if( pI->get_type() == B32_TYPE || pI->get_type() == U32_TYPE || pI->get_type() == S32_TYPE ) { + // 32-bit integer instruction + *cycles = 6; + *op = SFU_OP; + } + } +} + +void ptx_thread_info::ptx_fetch_inst( inst_t &inst ) const +{ + addr_t pc = get_pc(); + const ptx_instruction *pI = m_func_info->get_instruction(pc); + inst = (const inst_t&)*pI; + assert( inst.valid() ); } -void function_info::ptx_decode_inst( ptx_thread_info *thread, - unsigned *op_type, - int *i1, int *i2, int *i3, int *i4, - int *o1, int *o2, int *o3, int *o4, - int *vectorin, - int *vectorout, - int *arch_reg, - int *pred, - int *ar1, int *ar2 ) +void ptx_instruction::pre_decode() { - addr_t pc = thread->get_pc(); - unsigned index = pc - m_start_PC; - assert( index < m_instr_mem_size ); - ptx_instruction *pI = m_instr_mem[index]; //get instruction from m_instr_mem[PC] + pc = m_PC; + isize = m_inst_size; + for( unsigned i=0; i<4; i++) { + out[i] = 0; + in[i] = 0; + } + is_vectorin = 0; + is_vectorout = 0; + std::fill_n(arch_reg, MAX_REG_OPERANDS, -1); + pred = 0; + ar1 = 0; + ar2 = 0; bool has_dst = false ; - int opcode = pI->get_opcode(); //determine the opcode + int opcode = get_opcode(); //determine the opcode - switch ( pI->get_opcode() ) { + switch ( get_opcode() ) { #define OP_DEF(OP,FUNC,STR,DST,CLASSIFICATION) case OP: has_dst = (DST!=0); break; #include "opcodes.def" #undef OP_DEF default: - printf( "Execution error: Invalid opcode (0x%x)\n", pI->get_opcode() ); + printf( "Execution error: Invalid opcode (0x%x)\n", get_opcode() ); break; } - unsigned cycles; - get_opcode_info(pI,opcode,&cycles,op_type); - - // Quick fix for memory operands in ALU instructions - if( pI->has_memory_read() ) - *op_type = LOAD_OP; - else if( pI->has_memory_write() ) - *op_type = STORE_OP; - - if( pI->has_memory_read() && pI->has_memory_write() ) { - printf("Instruction has both a memory read and memory write - not supported by timing simulator."); - assert(0); - } + get_opcode_info(this,opcode,&cycles,&op); // Get register operands int n=0,m=0; - ptx_instruction::const_iterator op=pI->op_iter_begin(); - for ( ; op != pI->op_iter_end(); op++, n++ ) { //process operands + ptx_instruction::const_iterator opr=op_iter_begin(); + for ( ; opr != op_iter_end(); opr++, n++ ) { //process operands - const operand_info &o = *op; + const operand_info &o = *opr; if ( has_dst && n==0 ) { if ( o.is_reg() ) { //but is destination an actual register? (seems like it fails if it's a vector) - *o1 = o.reg_num(); + out[0] = o.reg_num(); arch_reg[0] = o.arch_reg_num(); } else if ( o.is_vector() ) { //but is destination an actual register? (seems like it fails if it's a vector) - *vectorin = 1; - *o1 = o.reg1_num(); - *o2 = o.reg2_num(); - *o3 = o.reg3_num(); - *o4 = o.reg4_num(); - for (int i = 0; i < 4; i++) + is_vectorin = 1; + unsigned num_elem = o.get_vect_nelem(); + if( num_elem >= 1 ) out[0] = o.reg1_num(); + if( num_elem >= 2 ) out[1] = o.reg2_num(); + if( num_elem >= 3 ) out[2] = o.reg3_num(); + if( num_elem >= 4 ) out[3] = o.reg4_num(); + for (int i = 0; i < num_elem; i++) arch_reg[i] = o.arch_reg_num(i); } } else { @@ -648,21 +659,21 @@ void function_info::ptx_decode_inst( ptx_thread_info *thread, int reg_num = o.reg_num(); arch_reg[m + 4] = o.arch_reg_num(); switch ( m ) { - case 0: *i1 = reg_num; break; - case 1: *i2 = reg_num; break; - case 2: *i3 = reg_num; break; - default: - break; + case 0: in[0] = reg_num; break; + case 1: in[1] = reg_num; break; + case 2: in[2] = reg_num; break; + default: break; } m++; } else if ( o.is_vector() ) { assert(m == 0); //only support 1 vector operand (for textures) right now - *vectorout = 1; - *i1 = o.reg1_num(); - *i2 = o.reg2_num(); - *i3 = o.reg3_num(); - *i4 = o.reg4_num(); - for (int i = 0; i < 4; i++) + is_vectorout = 1; + unsigned num_elem = o.get_vect_nelem(); + if( num_elem >= 1 ) in[0] = o.reg1_num(); + if( num_elem >= 2 ) in[1] = o.reg2_num(); + if( num_elem >= 3 ) in[2] = o.reg3_num(); + if( num_elem >= 4 ) in[3] = o.reg4_num(); + for (int i = 0; i < num_elem; i++) arch_reg[i + 4] = o.arch_reg_num(i); m+=4; } @@ -670,64 +681,34 @@ void function_info::ptx_decode_inst( ptx_thread_info *thread, } // Get predicate - if(pI->has_pred()) { - const operand_info &p = pI->get_pred(); - *pred = p.reg_num(); + if(has_pred()) { + const operand_info &p = get_pred(); + pred = p.reg_num(); } // Get address registers inside memory operands. // Assuming only one memory operand per instruction, // and maximum of two address registers for one memory operand. - if( pI->has_memory_read() || pI->has_memory_write() ) { - ptx_instruction::const_iterator op=pI->op_iter_begin(); - for ( ; op != pI->op_iter_end(); op++, n++ ) { //process operands + if( has_memory_read() || has_memory_write() ) { + ptx_instruction::const_iterator op=op_iter_begin(); + for ( ; op != op_iter_end(); op++, n++ ) { //process operands const operand_info &o = *op; // memory operand with addressing (ex. s[0x4] or g[$r1]) if(o.is_memory_operand2()) { // memory operand with one address register (ex. g[$r1] or s[$r2+0x4]) if(o.get_double_operand_type() == 0 && o.is_memory_operand()){ - *ar1 = o.reg_num(); + ar1 = o.reg_num(); } // memory operand with two address register (ex. s[$r1+$r1] or g[$r1+=$r2]) else if(o.get_double_operand_type() == 1 || o.get_double_operand_type() == 2) { - *ar1 = o.reg1_num(); - *ar2 = o.reg2_num(); - } - } - } - } - - // Get predicate - if(pI->has_pred()) { - const operand_info &p = pI->get_pred(); - *pred = p.reg_num(); - } - - // Get address registers inside memory operands. - // Assuming only one memory operand per instruction, - // and maximum of two address registers for one memory operand. - if( pI->has_memory_read() || pI->has_memory_write() ) { - ptx_instruction::const_iterator op=pI->op_iter_begin(); - for ( ; op != pI->op_iter_end(); op++, n++ ) { //process operands - const operand_info &o = *op; - - // memory operand with addressing (ex. s[0x4] or g[$r1]) - if(o.is_memory_operand2()) { - // memory operand with one address register (ex. g[$r1] or s[$r2+0x4]) - if(o.get_double_operand_type() == 0 && o.is_memory_operand()){ - const symbol *base_addr = o.get_symbol(); - if( base_addr->is_reg() ) - *ar1 = base_addr->reg_num(); - } - // memory operand with two address register (ex. s[$r1+$r1] or g[$r1+=$r2]) - else if(o.get_double_operand_type() == 1 || o.get_double_operand_type() == 2) { - *ar1 = o.reg1_num(); - *ar2 = o.reg2_num(); + ar1 = o.reg1_num(); + ar2 = o.reg2_num(); } } } } + m_decoded=true; } void function_info::add_param_name_type_size( unsigned index, std::string name, int type, size_t size ) @@ -855,7 +836,7 @@ void function_info::param_to_shared( memory_space *shared_mem, symbol_table *sym type_info_key::type_decode(xtype,size,tmp); // Write to shared memory - offset + 0x10 - shared_mem->write(offset+0x10,size/8,&value,NULL,NULL); + shared_mem->write(offset+0x10,size/8,value.pdata,NULL,NULL); } } @@ -923,49 +904,45 @@ unsigned datatype2size( unsigned data_type ) unsigned g_warp_active_mask; -void function_info::ptx_exec_inst( ptx_thread_info *thread, - addr_t *addr, - memory_space_t *space, - unsigned *data_size, - unsigned *cycles, - dram_callback_t* callback, - unsigned warp_active_mask ) +void ptx_thread_info::ptx_exec_inst( inst_t &inst ) { + inst.memory_op = no_memory_op; bool skip = false; int op_classification = 0; - addr_t pc = thread->next_instr(); - unsigned index = pc - m_start_PC; - assert( index < m_instr_mem_size ); - ptx_instruction *pI = m_instr_mem[index]; + 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); + set_npc( pc + pI->inst_size() ); + try { - thread->clearRPC(); - thread->m_last_set_operand_value.u64 = 0; + clearRPC(); + m_last_set_operand_value.u64 = 0; - if(thread->is_done()) + if(is_done()) { printf("attempted to execute instruction on a thread that is already done.\n"); assert(0); } - if ( g_debug_execution >= 6 ) { - if ( (g_debug_thread_uid==0) || (thread->get_uid() == (unsigned)g_debug_thread_uid) ) { - thread->clear_modifiedregs(); - thread->enable_debug_trace(); + if ( g_debug_execution >= 6 || g_ptx_inst_debug_to_file) { + if ( (g_debug_thread_uid==0) || (get_uid() == (unsigned)g_debug_thread_uid) ) { + clear_modifiedregs(); + enable_debug_trace(); } } if( pI->has_pred() ) { const operand_info &pred = pI->get_pred(); - ptx_reg_t pred_value = thread->get_operand_value(pred, pred, PRED_TYPE, thread, 0); + ptx_reg_t pred_value = get_operand_value(pred, pred, PRED_TYPE, this, 0); if(pI->get_pred_mod() == -1) { skip = (pred_value.pred & 0x0001) ^ pI->get_pred_neg(); //ptxplus inverts the zero flag } else { skip = !pred_lookup(pI->get_pred_mod(), pred_value.pred & 0x000F); } } - g_warp_active_mask = warp_active_mask; + g_warp_active_mask = inst.warp_active_mask; if( !skip ) { switch ( pI->get_opcode() ) { -#define OP_DEF(OP,FUNC,STR,DST,CLASSIFICATION) case OP: FUNC(pI,thread); op_classification = CLASSIFICATION; break; +#define OP_DEF(OP,FUNC,STR,DST,CLASSIFICATION) case OP: FUNC(pI,this); op_classification = CLASSIFICATION; break; #include "opcodes.def" #undef OP_DEF @@ -976,110 +953,108 @@ void function_info::ptx_exec_inst( ptx_thread_info *thread, // Run exit instruction if exit option included if(pI->is_exit()) - exit_impl(pI,thread); + exit_impl(pI,this); } // Output instruction information to file and stdout if( g_ptx_inst_debug_to_file != 0 && - (g_ptx_inst_debug_thread_uid == 0 || g_ptx_inst_debug_thread_uid == thread->get_uid()) ) { - dim3 ctaid = thread->get_ctaid(); - dim3 tid = thread->get_tid(); + (g_ptx_inst_debug_thread_uid == 0 || g_ptx_inst_debug_thread_uid == get_uid()) ) { + dim3 ctaid = get_ctaid(); + dim3 tid = get_tid(); fprintf(ptx_inst_debug_file, "[thd=%u] : (%s:%u - %s)\n", - thread->get_uid(), + get_uid(), pI->source_file(), pI->source_line(), pI->get_source() ); //fprintf(ptx_inst_debug_file, "has memory read=%d, has memory write=%d\n", pI->has_memory_read(), pI->has_memory_write()); fflush(ptx_inst_debug_file); } - if ( ptx_debug_exec_dump_cond<5>(thread->get_uid(), pc) ) { - dim3 ctaid = thread->get_ctaid(); - dim3 tid = thread->get_tid(); - printf("%u [cyc=%u][thd=%u][i=%u] : ctaid=(%u,%u,%u) tid=(%u,%u,%u) icount=%u [pc=%u] (%s:%u - %s) [0x%llx]\n", + if ( ptx_debug_exec_dump_cond<5>(get_uid(), pc) ) { + dim3 ctaid = get_ctaid(); + dim3 tid = get_tid(); + printf("%u [thd=%u][i=%u] : ctaid=(%u,%u,%u) tid=(%u,%u,%u) icount=%u [pc=%u] (%s:%u - %s) [0x%llx]\n", g_ptx_sim_num_insn, - (unsigned)gpu_sim_cycle, - thread->get_uid(), + get_uid(), pI->uid(), ctaid.x,ctaid.y,ctaid.z,tid.x,tid.y,tid.z, - thread->get_icount(), + get_icount(), pc, pI->source_file(), pI->source_line(), pI->get_source(), - thread->m_last_set_operand_value.u64 ); + m_last_set_operand_value.u64 ); fflush(stdout); } addr_t insn_memaddr = 0xFEEBDAED; memory_space_t insn_space = undefined_space; + _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 = thread->last_eaddr(); - insn_space = thread->last_space(); + 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() == ATOM_OP ) { - insn_memaddr = thread->last_eaddr(); - insn_space = thread->last_space(); - callback->function = thread->last_callback().function; - callback->instruction = thread->last_callback().instruction; - callback->thread = thread; + insn_memaddr = last_eaddr(); + insn_space = last_space(); + inst.callback.function = last_callback().function; + inst.callback.instruction = last_callback().instruction; + inst.callback.thread = this; unsigned to_type = pI->get_type(); insn_data_size = datatype2size(to_type); } else { // make sure that the callback isn't set - callback->function = NULL; - callback->instruction = NULL; + inst.callback.function = NULL; + inst.callback.instruction = NULL; + inst.callback.thread = NULL; } - // Set number of cycles for this instruction - int opcode = pI->get_opcode(); //determine the opcode - unsigned op_type; - get_opcode_info(pI,opcode,cycles,&op_type); - if (pI->get_opcode() == TEX_OP) { - *addr = thread->last_eaddr(); - *space = thread->last_space(); + inst.memreqaddr = last_eaddr(); + inst.space = last_space(); unsigned to_type = pI->get_type(); switch ( to_type ) { case B8_TYPE: case S8_TYPE: case U8_TYPE: - *data_size = 1; break; + inst.data_size = 1; break; case B16_TYPE: case S16_TYPE: case U16_TYPE: case F16_TYPE: - *data_size = 2; break; + inst.data_size = 2; break; case B32_TYPE: case S32_TYPE: case U32_TYPE: case F32_TYPE: - *data_size = 4; break; + inst.data_size = 4; break; case B64_TYPE: case S64_TYPE: case U64_TYPE: case F64_TYPE: - *data_size = 8; break; + inst.data_size = 8; break; default: assert(0); break; } } // Output register information to file and stdout if( g_ptx_inst_debug_to_file != 0 && - (g_ptx_inst_debug_thread_uid == 0 || g_ptx_inst_debug_thread_uid == thread->get_uid()) ) { - thread->dump_modifiedregs(ptx_inst_debug_file); + (g_ptx_inst_debug_thread_uid == 0 || g_ptx_inst_debug_thread_uid == get_uid()) ) { + dump_modifiedregs(ptx_inst_debug_file); + dump_regs(ptx_inst_debug_file); } if ( g_debug_execution >= 6 ) { - if ( ptx_debug_exec_dump_cond<6>(thread->get_uid(), pc) ) - thread->dump_modifiedregs(stdout); + if ( ptx_debug_exec_dump_cond<6>(get_uid(), pc) ) + dump_modifiedregs(stdout); } if ( g_debug_execution >= 10 ) { - if ( ptx_debug_exec_dump_cond<10>(thread->get_uid(), pc) ) - thread->dump_regs(stdout); + if ( ptx_debug_exec_dump_cond<10>(get_uid(), pc) ) + dump_regs(stdout); } - thread->update_pc(); + update_pc( pI->inst_size() ); g_ptx_sim_num_insn++; ptx_file_line_stats_add_exec_count(pI); if ( gpgpu_ptx_instruction_classification ) { @@ -1103,17 +1078,24 @@ void function_info::ptx_exec_inst( ptx_thread_info *thread, StatAddSample( g_inst_op_classification_stat[g_ptx_kernel_count], (int) pI->get_opcode() ); } if ( (g_ptx_sim_num_insn % 100000) == 0 ) { - dim3 ctaid = thread->get_ctaid(); - dim3 tid = thread->get_tid(); + dim3 ctaid = get_ctaid(); + dim3 tid = get_tid(); printf("GPGPU-Sim PTX: %u instructions simulated : ctaid=(%u,%u,%u) tid=(%u,%u,%u)\n", g_ptx_sim_num_insn, ctaid.x,ctaid.y,ctaid.z,tid.x,tid.y,tid.z ); fflush(stdout); } // "Return values" - *space = insn_space; - *addr = insn_memaddr; - *data_size = insn_data_size; + if(!skip) { + inst.space = insn_space; + inst.memreqaddr = insn_memaddr; + inst.data_size = insn_data_size; + inst.memory_op = insn_memory_op; + } else { + inst.memreqaddr = 0xFEEBDAED; + inst.space = undefined_space; + inst.memory_op = no_memory_op; + } } catch ( int x ) { printf("GPGPU-Sim PTX: ERROR (%d) executing intruction (%s:%u)\n", x, pI->source_file(), pI->source_line() ); @@ -1122,10 +1104,6 @@ void function_info::ptx_exec_inst( ptx_thread_info *thread, } } -unsigned g_gx, g_gy, g_gz; - -dim3 g_cudaGridDim, g_cudaBlockDim; - unsigned g_cta_launch_sid; std::list g_active_threads; std::map g_sm_idx_offset_next; @@ -1134,45 +1112,39 @@ std::map g_shared_memory_lookup; std::map g_ptx_cta_lookup; std::map > g_local_memory_lookup; -// return number of blocks in grid -unsigned ptx_sim_grid_size() -{ - return g_cudaGridDim.x * g_cudaGridDim.y * g_cudaGridDim.z; -} - -void set_option_gpgpu_spread_blocks_across_cores(int option) -{ - gpgpu_option_spread_blocks_across_cores = option; -} - void set_param_gpgpu_num_shaders(int num_shaders) { gpgpu_param_num_shaders = num_shaders; } -unsigned ptx_sim_cta_size() +const struct gpgpu_ptx_sim_kernel_info* ptx_sim_kernel_info(function_info *kernel) { - return g_cudaBlockDim.x * g_cudaBlockDim.y * g_cudaBlockDim.z; -} - -const struct gpgpu_ptx_sim_kernel_info* ptx_sim_kernel_info() { - return g_entrypoint_func_info->get_kernel_info(); + return kernel->get_kernel_info(); } -void ptx_sim_free_sm( ptx_thread_info** thread_info ) +const inst_t *ptx_fetch_inst( address_type pc ) { + return function_info::pc_to_instruction(pc); } -unsigned ptx_sim_init_thread( ptx_thread_info** thread_info,int sid,unsigned tid,unsigned threads_left,unsigned num_threads, core_t *core, unsigned hw_cta_id, unsigned hw_warp_id ) +unsigned ptx_sim_init_thread( kernel_info_t &kernel, + ptx_thread_info** thread_info, + int sid, + unsigned tid, + unsigned threads_left, + unsigned num_threads, + core_t *core, + unsigned hw_cta_id, + unsigned hw_warp_id ) { if ( *thread_info != NULL ) { ptx_thread_info *thd = *thread_info; assert( thd->is_done() ); if ( g_debug_execution==-1 ) { dim3 ctaid = thd->get_ctaid(); - dim3 tid = thd->get_tid(); + dim3 t = thd->get_tid(); printf("GPGPU-Sim PTX simulator: thread exiting ctaid=(%u,%u,%u) tid=(%u,%u,%u) uid=%u\n", - ctaid.x,ctaid.y,ctaid.z,tid.x,tid.y,tid.z, thd->get_uid() ); + ctaid.x,ctaid.y,ctaid.z,t.x,t.y,t.z, thd->get_uid() ); fflush(stdout); } thd->m_cta_info->assert_barrier_empty(); @@ -1197,11 +1169,11 @@ unsigned ptx_sim_init_thread( ptx_thread_info** thread_info,int sid,unsigned tid return 1; } - if ( g_gx >= g_cudaGridDim.x || g_gy >= g_cudaGridDim.y || g_gz >= g_cudaGridDim.z ) { + if ( kernel.no_more_ctas_to_run() ) { return 0; //finished! } - if ( threads_left < ptx_sim_cta_size() ) { + if ( threads_left < kernel.threads_per_cta() ) { return 0; } @@ -1214,25 +1186,12 @@ unsigned ptx_sim_init_thread( ptx_thread_info** thread_info,int sid,unsigned tid ptx_cta_info *cta_info = NULL; memory_space *shared_mem = NULL; - unsigned cta_size = ptx_sim_cta_size(); //blocksize + unsigned cta_size = kernel.threads_per_cta(); unsigned sm_offset = g_sm_idx_offset_next[sid]; unsigned max_cta_per_sm = num_threads/cta_size; // e.g., 256 / 48 = 5 assert( max_cta_per_sm > 0 ); - unsigned sm_idx = sid*max_cta_per_sm + sm_offset; - sm_idx = max_cta_per_sm*sid + tid/cta_size; - - if (!gpgpu_option_spread_blocks_across_cores) { - // update offset... - if ( (sm_offset + 1) >= max_cta_per_sm ) { - sm_offset = 0; - } else { - sm_offset++; - } - g_sm_idx_offset_next[sid] = sm_offset; - } else { - sm_idx = (tid/cta_size)*gpgpu_param_num_shaders + sid; - } + unsigned sm_idx = (tid/cta_size)*gpgpu_param_num_shaders + sid; if ( g_shared_memory_lookup.find(sm_idx) == g_shared_memory_lookup.end() ) { if ( g_debug_execution >= 1 ) { @@ -1256,67 +1215,57 @@ unsigned ptx_sim_init_thread( ptx_thread_info** thread_info,int sid,unsigned tid } std::map &local_mem_lookup = g_local_memory_lookup[sid]; - unsigned new_tid; - for ( unsigned tz=0; tz < g_cudaBlockDim.z; tz++ ) { - for ( unsigned ty=0; ty < g_cudaBlockDim.y; ty++ ) { - for ( unsigned tx=0; tx < g_cudaBlockDim.x; tx++ ) { - new_tid = tx + g_cudaBlockDim.x*ty + g_cudaBlockDim.x*g_cudaBlockDim.y*tz; - new_tid += tid; - ptx_thread_info *thd = new ptx_thread_info(); - - memory_space *local_mem = NULL; - std::map::iterator l = local_mem_lookup.find(new_tid); - if ( l != local_mem_lookup.end() ) { - local_mem = l->second; - } else { - char buf[512]; - snprintf(buf,512,"local_%u_%u", sid, new_tid); - local_mem = new memory_space_impl<32>(buf,32); - local_mem_lookup[new_tid] = local_mem; - } - thd->set_info(g_entrypoint_func_info); - thd->set_nctaid(g_cudaGridDim.x,g_cudaGridDim.y,g_cudaGridDim.z); - thd->set_ntid(g_cudaBlockDim.x,g_cudaBlockDim.y,g_cudaBlockDim.z); - thd->set_ctaid(g_gx,g_gy,g_gz); - thd->set_tid(tx,ty,tz); - if( g_entrypoint_func_info->get_ptx_version().extensions() ) - thd->cpy_tid_to_reg(tx,ty,tz); - thd->set_hw_tid((unsigned)-1); - thd->set_hw_wid((unsigned)-1); - thd->set_hw_ctaid((unsigned)-1); - thd->set_core(NULL); - thd->set_hw_sid((unsigned)-1); - thd->set_valid(); - thd->m_shared_mem = shared_mem; - function_info *finfo = thd->func_info(); - symbol_table *st = finfo->get_symtab(); - thd->func_info()->param_to_shared(thd->m_shared_mem,st); - thd->m_cta_info = cta_info; - cta_info->add_thread(thd); - thd->m_local_mem = local_mem; - if ( g_debug_execution==-1 ) { - printf("GPGPU-Sim PTX simulator: allocating thread ctaid=(%u,%u,%u) tid=(%u,%u,%u) @ 0x%Lx\n", - g_gx,g_gy,g_gz,tx,ty,tz, (unsigned long long)thd ); - fflush(stdout); - } - g_active_threads.push_back(thd); - } + while( kernel.more_threads_in_cta() ) { + dim3 ctaid3d = kernel.get_next_cta_id(); + unsigned new_tid = kernel.get_next_thread_id(); + dim3 tid3d = kernel.get_next_thread_id_3d(); + kernel.increment_thread_id(); + new_tid += tid; + ptx_thread_info *thd = new ptx_thread_info(); + + memory_space *local_mem = NULL; + std::map::iterator l = local_mem_lookup.find(new_tid); + if ( l != local_mem_lookup.end() ) { + local_mem = l->second; + } else { + char buf[512]; + snprintf(buf,512,"local_%u_%u", sid, new_tid); + local_mem = new memory_space_impl<32>(buf,32); + local_mem_lookup[new_tid] = local_mem; } + thd->set_info(kernel.entry()); + thd->set_nctaid(kernel.get_grid_dim()); + thd->set_ntid(kernel.get_cta_dim()); + thd->set_ctaid(ctaid3d); + thd->set_tid(tid3d); + if( kernel.entry()->get_ptx_version().extensions() ) + thd->cpy_tid_to_reg(tid3d); + thd->set_hw_tid((unsigned)-1); + thd->set_hw_wid((unsigned)-1); + thd->set_hw_ctaid((unsigned)-1); + thd->set_core(NULL); + thd->set_hw_sid((unsigned)-1); + thd->set_valid(); + thd->m_shared_mem = shared_mem; + function_info *finfo = thd->func_info(); + symbol_table *st = finfo->get_symtab(); + thd->func_info()->param_to_shared(thd->m_shared_mem,st); + thd->m_cta_info = cta_info; + cta_info->add_thread(thd); + thd->m_local_mem = local_mem; + if ( g_debug_execution==-1 ) { + printf("GPGPU-Sim PTX simulator: allocating thread ctaid=(%u,%u,%u) tid=(%u,%u,%u) @ 0x%Lx\n", + ctaid3d.x,ctaid3d.y,ctaid3d.z,tid3d.x,tid3d.y,tid3d.z, (unsigned long long)thd ); + fflush(stdout); + } + g_active_threads.push_back(thd); } if ( g_debug_execution==-1 ) { printf("GPGPU-Sim PTX simulator: <-- FINISHING THREAD ALLOCATION\n"); fflush(stdout); } - g_gx++; - if ( g_gx >= g_cudaGridDim.x ) { - g_gx = 0; - g_gy++; - if ( g_gy >= g_cudaGridDim.y ) { - g_gy = 0; - g_gz++; - } - } + kernel.increment_cta_id(); g_cta_launch_sid = -1; @@ -1336,7 +1285,7 @@ unsigned ptx_sim_init_thread( ptx_thread_info** thread_info,int sid,unsigned tid void init_inst_classification_stat() { char kernelname[256] =""; -#define MAX_CLASS_KER 256 +#define MAX_CLASS_KER 1024 if (!g_inst_classification_stat) g_inst_classification_stat = (void**)calloc(MAX_CLASS_KER, sizeof(void*)); snprintf(kernelname, MAX_CLASS_KER, "Kernel %d Classification\n",g_ptx_kernel_count ); assert( g_ptx_kernel_count < MAX_CLASS_KER ) ; // a static limit on number of kernels increase it if it fails! @@ -1350,21 +1299,12 @@ void init_inst_classification_stat() { std::map *g_kernel_name_to_function_lookup=NULL; std::map *g_host_to_kernel_entrypoint_name_lookup=NULL; -void gpgpu_cuda_ptx_sim_init_grid( const char *kernel_key, struct gpgpu_ptx_sim_arg* args, - struct dim3 gridDim, struct dim3 blockDim ) +function_info *get_kernel(const char *kernel_key, std::string &kernel_func_name_mangled ) { - g_gx=0; - g_gy=0; - g_gz=0; - g_cudaGridDim = gridDim; - g_cudaBlockDim = blockDim; - g_sm_idx_offset_next.clear(); - g_sm_next_index = 0; - if ( g_host_to_kernel_entrypoint_name_lookup->find(kernel_key) == g_host_to_kernel_entrypoint_name_lookup->end() ) { - printf("GPGPU-Sim PTX: ERROR ** cannot locate PTX entry point\n" ); - printf("GPGPU-Sim PTX: existing entry points: \n"); + printf("GPGPU-Sim PTX: ERROR ** cannot locate __global__ function from hostPtr\n" ); + printf("GPGPU-Sim PTX: registered PTX kernels: \n"); std::map::iterator i_eptr = g_host_to_kernel_entrypoint_name_lookup->begin(); for (; i_eptr != g_host_to_kernel_entrypoint_name_lookup->end(); ++i_eptr) { printf("GPGPU-Sim PTX: (%p,%s)\n", i_eptr->first, i_eptr->second.c_str()); @@ -1372,71 +1312,77 @@ void gpgpu_cuda_ptx_sim_init_grid( const char *kernel_key, struct gpgpu_ptx_sim_ printf("\n"); abort(); } - - std::string kname = (*g_host_to_kernel_entrypoint_name_lookup)[kernel_key]; - printf("GPGPU-Sim PTX: Launching kernel \'%s\' gridDim= (%u,%u,%u) blockDim = (%u,%u,%u); ntuid=%u\n", - kname.c_str(), g_cudaGridDim.x,g_cudaGridDim.y,g_cudaGridDim.z,g_cudaBlockDim.x,g_cudaBlockDim.y,g_cudaBlockDim.z, - g_ptx_thread_info_uid_next ); - - if ( g_kernel_name_to_function_lookup->find(kname) == + kernel_func_name_mangled = (*g_host_to_kernel_entrypoint_name_lookup)[kernel_key]; + if ( g_kernel_name_to_function_lookup->find(kernel_func_name_mangled) == g_kernel_name_to_function_lookup->end() ) { - printf("GPGPU-Sim PTX: ERROR ** function \'%s\' not found in ptx file\n", kname.c_str() ); + printf("GPGPU-Sim PTX: ERROR ** function \'%s\' not found in ptx file\n", + kernel_func_name_mangled.c_str() ); abort(); } - g_entrypoint_func_info = g_func_info = (*g_kernel_name_to_function_lookup)[kname]; + return (*g_kernel_name_to_function_lookup)[kernel_func_name_mangled]; +} + +const struct gpgpu_ptx_sim_kernel_info * get_kernel_info(const char *kernel_key) +{ + std::string kname; + function_info *finfo = get_kernel(kernel_key,kname); + return finfo->get_kernel_info(); +} + +size_t get_kernel_code_size( class function_info *entry ) +{ + return entry->get_function_size(); +} + +kernel_info_t gpgpu_cuda_ptx_sim_init_grid( const char *kernel_key, gpgpu_ptx_sim_arg_list_t args, + struct dim3 gridDim, struct dim3 blockDim ) +{ + g_sm_idx_offset_next.clear(); + g_sm_next_index = 0; + std::string kname; + function_info *entry = get_kernel(kernel_key,kname); + + printf("GPGPU-Sim PTX: Launching kernel \'%s\' gridDim= (%u,%u,%u) blockDim = (%u,%u,%u); ntuid=%u\n", + kname.c_str(), gridDim.x,gridDim.y,gridDim.z,blockDim.x,blockDim.y,blockDim.z, + g_ptx_thread_info_uid_next ); - unsigned argcount=0; - struct gpgpu_ptx_sim_arg *tmparg = args; - while (tmparg) { - tmparg = tmparg->m_next; - argcount++; - } + unsigned argcount=args.size(); unsigned argn=1; - while (args) { - g_func_info->add_param_data(argcount-argn,args); - args = args->m_next; + for( gpgpu_ptx_sim_arg_list_t::iterator a = args.begin(); a != args.end(); a++ ) { + entry->add_param_data(argcount-argn,&(*a)); argn++; } - g_func_info->finalize(g_param_mem); + + entry->finalize(g_param_mem); g_ptx_kernel_count++; if ( gpgpu_ptx_instruction_classification ) { init_inst_classification_stat(); } fflush(stdout); + + return kernel_info_t(gridDim,blockDim,entry); } -void gpgpu_opencl_ptx_sim_init_grid(class function_info *entry,struct gpgpu_ptx_sim_arg *args, struct dim3 gridDim, struct dim3 blockDim ) +kernel_info_t gpgpu_opencl_ptx_sim_init_grid(class function_info *entry,gpgpu_ptx_sim_arg_list_t args, struct dim3 gridDim, struct dim3 blockDim ) { - g_gx=0; - g_gy=0; - g_gz=0; - g_cudaGridDim = gridDim; - g_cudaBlockDim = blockDim; g_sm_idx_offset_next.clear(); g_sm_next_index = 0; - g_entrypoint_func_info = g_func_info = entry; - - unsigned argcount=0; - struct gpgpu_ptx_sim_arg *tmparg = args; - while (tmparg) { - tmparg = tmparg->m_next; - argcount++; - } - + unsigned argcount=args.size(); unsigned argn=1; - while (args) { - g_func_info->add_param_data(argcount-argn,args); - args = args->m_next; + for( gpgpu_ptx_sim_arg_list_t::iterator a = args.begin(); a != args.end(); a++ ) { + entry->add_param_data(argcount-argn,&(*a)); argn++; } - g_func_info->finalize(g_param_mem); + entry->finalize(g_param_mem); g_ptx_kernel_count++; if ( gpgpu_ptx_instruction_classification ) { init_inst_classification_stat(); } fflush(stdout); + + return kernel_info_t(gridDim,blockDim,entry); } const char *g_gpgpusim_version_string = "2.1.1b (beta)"; @@ -1625,7 +1571,7 @@ ptx_cta_info *g_func_cta_info = NULL; #define MAX(a,b) (((a)>(b))?(a):(b)) -void gpgpu_cuda_ptx_sim_main_func( const char *kernel_key, dim3 gridDim, dim3 blockDim, struct gpgpu_ptx_sim_arg *args) +void gpgpu_cuda_ptx_sim_main_func( const char *kernel_key, dim3 gridDim, dim3 blockDim, gpgpu_ptx_sim_arg_list_t args) { printf("GPGPU-Sim: Performing Functional Simulation...\n"); @@ -1634,12 +1580,11 @@ void gpgpu_cuda_ptx_sim_main_func( const char *kernel_key, dim3 gridDim, dim3 bl exit(1); time_t end_time, elapsed_time, days, hrs, minutes, sec; - int i1, i2, i3, i4, o1, o2, o3, o4; - int vectorin, vectorout; - int pred; - int ar1, ar2; - gpgpu_cuda_ptx_sim_init_grid(kernel_key, args,gridDim,blockDim); + kernel_info_t kernel = gpgpu_cuda_ptx_sim_init_grid(kernel_key, args,gridDim,blockDim); + + std::string kname; + function_info *finfo = get_kernel(kernel_key,kname); memory_space *shared_mem = new memory_space_impl<16*1024>("shared",4); @@ -1647,46 +1592,40 @@ void gpgpu_cuda_ptx_sim_main_func( const char *kernel_key, dim3 gridDim, dim3 bl if ( g_func_cta_info == NULL ) g_func_cta_info = new ptx_cta_info(0); - - for ( unsigned gx=0; gx < gridDim.x; gx++ ) { - for ( unsigned gy=0; gy < gridDim.y; gy++ ) { - for ( unsigned gz=0; gz < gridDim.z; gz++ ) { + while( !kernel.no_more_ctas_to_run() ) { std::list active_threads; std::list blocked_threads; + dim3 ctaid3d = kernel.get_next_cta_id(); + kernel.increment_cta_id(); g_func_cta_info->check_cta_thread_status_and_reset(); - - for ( unsigned tx=0; tx < blockDim.x; tx++ ) { - for ( unsigned ty=0; ty < blockDim.y; ty++ ) { - for ( unsigned tz=0; tz < blockDim.z; tz++ ) { + while( kernel.more_threads_in_cta() ) { memory_space *local_mem = NULL; ptx_thread_info *thd = new ptx_thread_info(); - - unsigned lm_idx = blockDim.x*blockDim.y*tz + blockDim.x * ty + tx; + dim3 tid3d = kernel.get_next_thread_id_3d(); + unsigned lm_idx = kernel.get_next_thread_id(); + kernel.increment_thread_id(); std::map::iterator lm=lm_lookup.find(lm_idx); if ( lm == lm_lookup.end() ) { char buf[1024]; - snprintf(buf,1024,"local_(%u,%u,%u)", tx, ty, tz ); + snprintf(buf,1024,"local_(%u,%u,%u)", tid3d.x, tid3d.y, tid3d.z ); local_mem = new memory_space_impl<32>(buf,32); lm_lookup[lm_idx] = local_mem; } else { local_mem = lm->second; } - - thd->set_info(g_func_info); - thd->set_nctaid(gridDim.x,gridDim.y,gridDim.z); - thd->set_ntid(blockDim.x, blockDim.y, blockDim.z); - thd->set_ctaid(gx,gy,gz); - thd->set_tid(tx,ty,tz); + thd->set_info(finfo); + thd->set_nctaid(ctaid3d); + thd->set_ntid(kernel.get_cta_dim()); + thd->set_ctaid(ctaid3d); + thd->set_tid(tid3d); thd->set_valid(); thd->m_shared_mem = shared_mem; thd->m_local_mem = local_mem; thd->m_cta_info = g_func_cta_info; g_func_cta_info->add_thread(thd); active_threads.push_back(thd); - } - } } while ( !(active_threads.empty() && blocked_threads.empty()) ) { @@ -1720,22 +1659,12 @@ void gpgpu_cuda_ptx_sim_main_func( const char *kernel_key, dim3 gridDim, dim3 bl break; } - unsigned op_type; - addr_t addr; - unsigned cycles; - memory_space_t space; - int arch_reg[MAX_REG_OPERANDS] = { -1 }; - unsigned data_size; - dram_callback_t callback; - unsigned warp_active_mask = (unsigned)-1; // vote instruction with diverged warps won't execute correctly - // in functional simulation mode - - g_func_info->ptx_decode_inst( thread, &op_type, &i1, &i2, &i3, &i4, &o1, &o2, &o3, &o4, &vectorin, &vectorout, arch_reg, &pred, &ar1, &ar2 ); - g_func_info->ptx_exec_inst( thread, &addr, &space, &data_size, &cycles, &callback, warp_active_mask ); + inst_t inst; + inst.warp_active_mask = (unsigned)-1; // vote instruction with diverged warps won't execute correctly + // in functional simulation mode + thread->ptx_exec_inst( inst ); } } - } - } } printf( "GPGPU-Sim: Done functional simulation (%u instructions simulated).\n", g_ptx_sim_num_insn ); if ( gpgpu_ptx_instruction_classification ) { @@ -1757,70 +1686,12 @@ void gpgpu_cuda_ptx_sim_main_func( const char *kernel_key, dim3 gridDim, dim3 bl fflush(stdout); } -void ptx_decode_inst( void *thd, unsigned *op, int *i1, int *i2, int *i3, int *i4, int *o1, int *o2, int *o3, int *o4, int *vectorin, int *vectorout, int *arch_reg, int *pred, int *ar1, int *ar2 ) -{ - *op = NO_OP; - *o1 = 0; - *o2 = 0; - *o3 = 0; - *o4 = 0; - *i1 = 0; - *i2 = 0; - *i3 = 0; - *i4 = 0; - *vectorin = 0; - *vectorout = 0; - std::fill_n(arch_reg, MAX_REG_OPERANDS, -1); - *pred = 0; - *ar1 = 0; - *ar2 = 0; - - if ( thd == NULL ) - return; - - ptx_thread_info *thread = (ptx_thread_info *) thd; - g_func_info = thread->func_info(); - g_func_info->ptx_decode_inst(thread,op,i1,i2,i3,i4,o1,o2,o3,o4,vectorin,vectorout,arch_reg,pred,ar1,ar2); -} - -extern "C" unsigned ptx_get_inst_op( void *thd) -{ - if ( thd == NULL ) - return NO_OP; - - ptx_thread_info *thread = (ptx_thread_info *) thd; - return(thread->func_info())->ptx_get_inst_op(thread); -} - -void ptx_exec_inst( void *thd, address_type *addr, memory_space_t *space, unsigned *data_size, unsigned *cycles, dram_callback_t* callback, unsigned warp_active_mask ) -{ - if ( thd == NULL ) - return; - *cycles = 1; - ptx_thread_info *thread = (ptx_thread_info *) thd; - g_func_info = thread->func_info(); - g_func_info->ptx_exec_inst( thread, addr, space, data_size, cycles, callback, warp_active_mask ); -} - -void ptx_dump_regs( void *thd ) -{ - if ( thd == NULL ) - return; - ptx_thread_info *t = (ptx_thread_info *) thd; - t->dump_regs(stdout); -} - unsigned ptx_set_tex_cache_linesize(unsigned linesize) { g_texcache_linesize = linesize; return 0; } -unsigned ptx_kernel_program_size() -{ - return g_func_info->get_function_size(); -} - unsigned translate_pc_to_ptxlineno(unsigned pc) { // this function assumes that the kernel fits inside a single PTX file @@ -1837,12 +1708,11 @@ int g_ptxinfo_error_detected; static char *g_ptxinfo_kname = NULL; static struct gpgpu_ptx_sim_kernel_info g_ptxinfo_kinfo; +static void clear_ptxinfo(); + extern "C" void ptxinfo_function(const char *fname ) { - g_ptxinfo_kinfo.regs=0; - g_ptxinfo_kinfo.lmem=0; - g_ptxinfo_kinfo.smem=0; - g_ptxinfo_kinfo.cmem=0; + clear_ptxinfo(); g_ptxinfo_kname = strdup(fname); } @@ -1874,6 +1744,8 @@ void clear_ptxinfo() g_ptxinfo_kinfo.lmem=0; g_ptxinfo_kinfo.smem=0; g_ptxinfo_kinfo.cmem=0; + g_ptxinfo_kinfo.ptx_version=0; + g_ptxinfo_kinfo.sm_target=0; } void ptxinfo_opencl_addinfo( std::map &kernels ) @@ -1991,14 +1863,9 @@ unsigned int get_converge_point( unsigned int pc, void *thd ) abort(); // returning garbage! } -void find_reconvergence_points() -{ - find_reconvergence_points(g_func_info); -} - -void dwf_process_reconv_pts() +void dwf_process_reconv_pts(function_info *entry) { - rec_pts tmp = find_reconvergence_points(g_func_info); + rec_pts tmp = find_reconvergence_points(entry); for (int i = 0; i < tmp.s_num_recon; ++i) { dwf_insert_reconv_pt(tmp.s_kernel_recon_points[i].target_pc); } diff --git a/src/cuda-sim/cuda-sim.h b/src/cuda-sim/cuda-sim.h index d7d3f8a..57f2837 100644 --- a/src/cuda-sim/cuda-sim.h +++ b/src/cuda-sim/cuda-sim.h @@ -23,15 +23,15 @@ extern int g_ptx_kernel_count; // used for classification stat collection purpos extern FILE* ptx_inst_debug_file; -extern void gpgpu_cuda_ptx_sim_init_grid( const char *kernel_key, - struct gpgpu_ptx_sim_arg *args, +extern class kernel_info_t gpgpu_cuda_ptx_sim_init_grid( const char *kernel_key, + gpgpu_ptx_sim_arg_list_t args, struct dim3 gridDim, struct dim3 blockDim ); -extern void gpgpu_opencl_ptx_sim_init_grid(class function_info *entry, - struct gpgpu_ptx_sim_arg *args, +extern class kernel_info_t gpgpu_opencl_ptx_sim_init_grid(class function_info *entry, + gpgpu_ptx_sim_arg_list_t args, struct dim3 gridDim, struct dim3 blockDim ); -extern void gpgpu_cuda_ptx_sim_main_func( const char *kernel_key, dim3 gridDim, dim3 blockDim, struct gpgpu_ptx_sim_arg *); +extern void gpgpu_cuda_ptx_sim_main_func( const char *kernel_key, dim3 gridDim, dim3 blockDim, gpgpu_ptx_sim_arg_list_t ); extern void print_splash(); extern void* gpgpu_ptx_sim_malloc( size_t count ); extern void* gpgpu_ptx_sim_mallocarray( size_t count ); @@ -55,8 +55,8 @@ extern void read_sim_environment_variables(); extern void register_function_implementation( const char *name, function_info *impl ); extern void ptxinfo_cuda_addinfo(); extern void ptxinfo_opencl_addinfo( std::map &kernels ); -extern void ptx_sim_free_sm( class ptx_thread_info** thread_info ); -unsigned ptx_sim_init_thread( class ptx_thread_info** thread_info, +unsigned ptx_sim_init_thread( kernel_info_t &kernel, + class ptx_thread_info** thread_info, int sid, unsigned tid, unsigned threads_left, @@ -64,12 +64,9 @@ unsigned ptx_sim_init_thread( class ptx_thread_info** thread_info, class core_t *core, unsigned hw_cta_id, unsigned hw_warp_id ); -unsigned ptx_sim_cta_size(); -const struct gpgpu_ptx_sim_kernel_info* ptx_sim_kernel_info(); -void set_option_gpgpu_spread_blocks_across_cores(int option); -int ptx_thread_done( void *thd ); +const inst_t *ptx_fetch_inst( address_type pc ); +const struct gpgpu_ptx_sim_kernel_info* ptx_sim_kernel_info(class function_info *kernel); unsigned ptx_thread_donecycle( void *thr ); -int ptx_thread_get_next_pc( void *thd ); void* ptx_thread_get_next_finfo( void *thd ); int ptx_thread_at_barrier( void *thd ); int ptx_thread_all_at_barrier( void *thd ); @@ -78,28 +75,10 @@ void ptx_thread_reset_barrier( void *thd ); void ptx_thread_release_barrier( void *thd ); void ptx_print_insn( address_type pc, FILE *fp ); unsigned int ptx_set_tex_cache_linesize( unsigned linesize); -void ptx_decode_inst( void *thd, - unsigned *op, - int *i1, - int *i2, - int *i3, - int *i4, - int *o1, - int *o2, - int *o3, - int *o4, - int *vectorin, - int *vectorout, - int *arch_reg, - int *pred, - int *ar1, - int *ar2 ); -void ptx_exec_inst( void *thd, - address_type *addr, - memory_space_t *space, - unsigned *data_size, - unsigned *cycles, - dram_callback_t* callback, - unsigned warp_active_mask ); + +function_info *get_kernel(const char *kernel_key, std::string &kernel_func_name_mangled ); +void dwf_process_reconv_pts(function_info *entry); +void set_param_gpgpu_num_shaders(int num_shaders); +unsigned int get_converge_point(unsigned int pc, void *thd); #endif diff --git a/src/cuda-sim/dram_callback.h b/src/cuda-sim/dram_callback.h index 9336ee0..fee97eb 100644 --- a/src/cuda-sim/dram_callback.h +++ b/src/cuda-sim/dram_callback.h @@ -75,10 +75,4 @@ #ifndef __DRAM_CALLBACK_H__ #define __DRAM_CALLBACK_H__ -typedef struct { - void (*function)(void* pI, void* gOldGThread); - void* instruction; - void* thread;// callback has to abuse g_thread when executed -}dram_callback_t; - #endif // #ifndef __DRAM_CALLBACK_H__ diff --git a/src/cuda-sim/instructions.cc b/src/cuda-sim/instructions.cc index d9a9630..78996bb 100644 --- a/src/cuda-sim/instructions.cc +++ b/src/cuda-sim/instructions.cc @@ -651,16 +651,20 @@ void ptx_thread_info::set_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, - unsigned num_elements ) + const ptx_reg_t &data4 ) { - set_reg(dst.vec_symbol(0), data1); - set_reg(dst.vec_symbol(1), data2); - if (num_elements > 2) { - set_reg(dst.vec_symbol(2), data3); - if (num_elements > 3) { - set_reg(dst.vec_symbol(3), data4); - } + unsigned num_elements = dst.get_vect_nelem(); + if (num_elements > 0) { + set_reg(dst.vec_symbol(0), data1); + if (num_elements > 1) { + set_reg(dst.vec_symbol(1), data2); + if (num_elements > 2) { + set_reg(dst.vec_symbol(2), data3); + if (num_elements > 3) { + set_reg(dst.vec_symbol(3), data4); + } + } + } } m_last_set_operand_value = data1; @@ -1084,6 +1088,8 @@ void atom_callback( void* ptx_inst, void* thd ) // Write operation result into global memory // (i.e. copy src1_data to dst) g_global_mem->write(src1_data.u32,size/8,&op_result.s64,thread,pI); + gpgpu_sim *gpu = thread->get_gpu(); + gpu->decrement_atomic_count(thread->get_hw_sid(),thread->get_hw_wid()); } // atom_impl will now result in a callback being called in mem_ctrl_pop (gpu-sim.c) @@ -1113,7 +1119,7 @@ void bar_sync_impl( const ptx_instruction *pI, ptx_thread_info *thread ) { const operand_info &dst = pI->dst(); ptx_reg_t b = thread->get_operand_value(dst, dst, U32_TYPE, thread, 1); - assert( b.u32 == 0 ); // not clear what should happen if this is not zero + assert( b.u32 == 0 ); // support for bar.sync a{,b}; where a != 0 not yet implemented } void bfe_impl( const ptx_instruction *pI, ptx_thread_info *thread ) { inst_not_implemented(pI); } @@ -1194,20 +1200,80 @@ void call_impl( const ptx_instruction *pI, ptx_thread_info *thread ) unsigned sid = thread->get_hw_sid(); unsigned tid = thread->get_hw_tid(); + gpgpu_sim *gpu = thread->get_gpu(); unsigned callee_pc=0, callee_rpc=0; - if( gpgpu_simd_model == POST_DOMINATOR ) { - get_pdom_stack_top_info(sid,tid,&callee_pc,&callee_rpc); + if( gpu->simd_model() == POST_DOMINATOR ) { + gpu->get_pdom_stack_top_info(sid,tid,&callee_pc,&callee_rpc); assert( callee_pc == thread->get_pc() ); } - thread->callstack_push(callee_pc+1,callee_rpc,return_var_src,return_var_dst,call_uid_next++); + thread->callstack_push(callee_pc + pI->inst_size(), callee_rpc, return_var_src, return_var_dst, call_uid_next++); copy_buffer_list_into_frame(thread, arg_values); thread->set_npc(target_func); } -void clz_impl( const ptx_instruction *pI, ptx_thread_info *thread ) { inst_not_implemented(pI); } +//Ptxplus version of call instruction. Jumps to a label not a different Kernel. +void callp_impl( const ptx_instruction *pI, ptx_thread_info *thread ) +{ + + static unsigned call_uid_next = 1; + + const operand_info &target = pI->dst(); + ptx_reg_t target_pc = thread->get_operand_value(target, target, U32_TYPE, thread, 1); + + const symbol *return_var_src = NULL; + const symbol *return_var_dst = NULL; + + unsigned sid = thread->get_hw_sid(); + unsigned tid = thread->get_hw_tid(); + gpgpu_sim *gpu = thread->get_gpu(); + unsigned callee_pc=0, callee_rpc=0; + if( gpu->simd_model() == POST_DOMINATOR ) { + gpu->get_pdom_stack_top_info(sid,tid,&callee_pc,&callee_rpc); + assert( callee_pc == thread->get_pc() ); + } + + thread->callstack_push_plus(callee_pc + pI->inst_size(), callee_rpc, return_var_src, return_var_dst, call_uid_next++); + thread->set_npc(target_pc); +} + +void clz_impl( const ptx_instruction *pI, ptx_thread_info *thread ) +{ + ptx_reg_t a, d; + const operand_info &dst = pI->dst(); + const operand_info &src1 = pI->src1(); + + unsigned i_type = pI->get_type(); + a = thread->get_operand_value(src1, dst, i_type, thread, 1); + + int max; + unsigned long long mask; + d.u64 = 0; + + switch ( i_type ) { + case B32_TYPE: + max = 32; + mask = 0x80000000; + break; + case B64_TYPE: + max = 64; + mask = 0x8000000000000000; + break; + default: + printf("Execution error: type mismatch with instruction\n"); + assert(0); + break; + } + + while ((d.u32 < max) && ((a.u64&mask) == 0) ) { + d.u32++; + a.u64 = a.u64 << 1; + } + + thread->set_operand_value(dst,d, B32_TYPE, thread, pI); +} void cnot_impl( const ptx_instruction *pI, ptx_thread_info *thread ) { @@ -1313,7 +1379,7 @@ ptx_reg_t f2x( ptx_reg_t x, unsigned from_width, unsigned to_width, int to_sign, ptx_reg_t y; if ( to_sign == 1 ) { // convert to 64-bit number first? - int tmp = cuda_math::__internal_float2int(x.f32, mode); + int tmp = cuda_math::float2int(x.f32, mode); if ((x.u32 & 0x7f800000) == 0) tmp = 0; // round denorm. FP to 0 if (saturation_mode && to_width < 32) { @@ -1327,7 +1393,7 @@ ptx_reg_t f2x( ptx_reg_t x, unsigned from_width, unsigned to_width, int to_sign, default: assert(0); break; } } else if ( to_sign == 0 ) { - unsigned int tmp = cuda_math::__internal_float2uint(x.f32, mode); + unsigned int tmp = cuda_math::float2uint(x.f32, mode); if ((x.u32 & 0x7f800000) == 0) tmp = 0; // round denorm. FP to 0 if (saturation_mode && to_width < 32) { @@ -1432,10 +1498,10 @@ ptx_reg_t s2f( ptx_reg_t x, unsigned from_width, unsigned to_width, int to_sign, case 16: assert(0); break; case 32: switch (rounding_mode) { - case RZ_OPTION: y.f32 = cuda_math::__ll2float_rn(y.s64); break; + case RZ_OPTION: y.f32 = cuda_math::__ll2float_rz(y.s64); break; case RN_OPTION: y.f32 = cuda_math::__ll2float_rn(y.s64); break; - case RM_OPTION: y.f32 = cuda_math::__ll2float_rn(y.s64); break; - case RP_OPTION: y.f32 = cuda_math::__ll2float_rn(y.s64); break; + case RM_OPTION: y.f32 = cuda_math::__ll2float_rd(y.s64); break; + case RP_OPTION: y.f32 = cuda_math::__ll2float_ru(y.s64); break; default: break; } break; @@ -1549,9 +1615,9 @@ ptx_reg_t d2d( ptx_reg_t x, unsigned from_width, unsigned to_width, int to_sign, break; case RNI_OPTION: #if CUDART_VERSION >= 3000 - y.f64 = nearbyint(x.f32); + y.f64 = nearbyint(x.f64); #else - y.f64 = cuda_math::__internal_nearbyint(x.f64); + y.f64 = cuda_math::__internal_nearbyintf(x.f64); #endif break; case RMI_OPTION: @@ -1816,7 +1882,8 @@ void cvta_impl( const ptx_instruction *pI, ptx_thread_info *thread ) } else { switch( space.get_type() ) { case shared_space: to_addr_hw = shared_to_generic( smid, from_addr_hw ); break; - case local_space: to_addr_hw = local_to_generic( smid, hwtid, from_addr_hw ); break; + case local_space: to_addr_hw = local_to_generic( smid, hwtid, from_addr_hw ) + + thread->get_local_mem_stack_pointer(); break; // add stack ptr here so that it can be passed as a pointer at function call case global_space: to_addr_hw = global_to_generic( from_addr_hw ); break; default: abort(); } @@ -2018,11 +2085,11 @@ void ld_exec( const ptx_instruction *pI, ptx_thread_info *thread ) mem->read(addr+2*size/8,size/8,&data3.s64); if (vector_spec != V3_TYPE) { //v4 mem->read(addr+3*size/8,size/8,&data4.s64); - thread->set_vector_operand_values(dst,data1,data2,data3,data4, 4); + thread->set_vector_operand_values(dst,data1,data2,data3,data4); } else //v3 - thread->set_vector_operand_values(dst,data1,data2,data3,data3,3); + thread->set_vector_operand_values(dst,data1,data2,data3,data3); } else //v2 - thread->set_vector_operand_values(dst,data1,data2,data2,data2,2); + thread->set_vector_operand_values(dst,data1,data2,data2,data2); } thread->m_last_effective_address = addr; thread->m_last_memory_space = space; @@ -2249,7 +2316,10 @@ void max_impl( const ptx_instruction *pI, ptx_thread_info *thread ) thread->set_operand_value(dst,d, i_type, thread, pI); } -void membar_impl( const ptx_instruction *pI, ptx_thread_info *thread ) { inst_not_implemented(pI); } +void membar_impl( const ptx_instruction *pI, ptx_thread_info *thread ) +{ + // handled by timing simulator +} void min_impl( const ptx_instruction *pI, ptx_thread_info *thread ) { @@ -2332,16 +2402,16 @@ void mov_impl( const ptx_instruction *pI, ptx_thread_info *thread ) unsigned bits_per_dst_elem = nbits_to_move / nelem; for( unsigned i=0; i < nelem; i++ ) { switch(bits_per_dst_elem) { - case 8: v[i].u8 = tmp_bits.u64 & (((unsigned long long) 0xFF) << (8*i)); break; - case 16: v[i].u16 = tmp_bits.u64 & (((unsigned long long) 0xFFFF) << (16*i)); break; - case 32: v[i].u32 = tmp_bits.u64 & (((unsigned long long) 0xFFFFFFFF) << (32*i)); break; + case 8: v[i].u8 = (tmp_bits.u64 >> (8*i)) & ((unsigned long long) 0xFF); break; + case 16: v[i].u16 = (tmp_bits.u64 >> (16*i)) & ((unsigned long long) 0xFFFF); break; + case 32: v[i].u32 = (tmp_bits.u64 >> (32*i)) & ((unsigned long long) 0xFFFFFFFF); break; default: printf("Execution error: mov pack/unpack with unsupported source/dst size ratio (dst)\n"); assert(0); break; } } - thread->set_vector_operand_values(dst,v[0],v[1],v[2],v[3],nelem); + thread->set_vector_operand_values(dst,v[0],v[1],v[2],v[3]); } else { thread->set_operand_value(dst,tmp_bits, i_type, thread, pI); } @@ -2707,6 +2777,19 @@ void ret_impl( const ptx_instruction *pI, ptx_thread_info *thread ) } } +//Ptxplus version of ret instruction. +void retp_impl( const ptx_instruction *pI, ptx_thread_info *thread ) +{ + bool empty = thread->callstack_pop_plus(); + if( empty ) { + core_t *sc = thread->get_core(); + unsigned warp_id = thread->get_hw_wid(); + sc->warp_exit(warp_id); + thread->m_cta_info->register_thread_exit(thread); + thread->set_done(); + } +} + void rsqrt_impl( const ptx_instruction *pI, ptx_thread_info *thread ) { ptx_reg_t a, d; @@ -3689,7 +3772,7 @@ void tex_impl( const ptx_instruction *pI, ptx_thread_info *thread ) assert(0); } thread->m_last_memory_space = tex_space; - thread->set_vector_operand_values(dst,data1,data2,data3,data4,4); + thread->set_vector_operand_values(dst,data1,data2,data3,data4); } void txq_impl( const ptx_instruction *pI, ptx_thread_info *thread ) { inst_not_implemented(pI); } @@ -3792,10 +3875,6 @@ void inst_not_implemented( const ptx_instruction * pI ) abort(); } -void print_instruction(const ptx_instruction *instruction) -{ -} - ptx_reg_t srcOperandModifiers(ptx_reg_t opData, operand_info opInfo, operand_info dstInfo, unsigned type, ptx_thread_info *thread) { ptx_reg_t result; diff --git a/src/cuda-sim/memory.cc b/src/cuda-sim/memory.cc index b825f12..aecfd31 100644 --- a/src/cuda-sim/memory.cc +++ b/src/cuda-sim/memory.cc @@ -83,10 +83,35 @@ template memory_space_impl::memory_space_impl( std::strin template void memory_space_impl::write( mem_addr_t addr, size_t length, const void *data, class ptx_thread_info *thd, const ptx_instruction *pI) { mem_addr_t index = addr >> m_log2_block_size; - unsigned offset = addr & (BSIZE-1); - unsigned nbytes = length; - assert( (addr+length) <= (index+1)*BSIZE ); - m_data[index].write(offset,nbytes,(const unsigned char*)data); + if ( (addr+length) <= (index+1)*BSIZE ) { + // fast route for intra-block access + unsigned offset = addr & (BSIZE-1); + unsigned nbytes = length; + m_data[index].write(offset,nbytes,(const unsigned char*)data); + } else { + // slow route for inter-block access + unsigned nbytes_remain = length; + unsigned src_offset = 0; + mem_addr_t current_addr = addr; + + while (nbytes_remain > 0) { + unsigned offset = current_addr & (BSIZE-1); + mem_addr_t page = current_addr >> m_log2_block_size; + mem_addr_t access_limit = offset + nbytes_remain; + if (access_limit > BSIZE) { + access_limit = BSIZE; + } + + size_t tx_bytes = access_limit - offset; + m_data[page].write(offset, tx_bytes, &((const unsigned char*)data)[src_offset]); + + // advance pointers + src_offset += tx_bytes; + current_addr += tx_bytes; + nbytes_remain -= tx_bytes; + } + assert(nbytes_remain == 0); + } if( !m_watchpoints.empty() ) { std::map::iterator i; for( i=m_watchpoints.begin(); i!=m_watchpoints.end(); i++ ) { @@ -97,28 +122,59 @@ template void memory_space_impl::write( mem_addr_t addr, } } -template void memory_space_impl::read( mem_addr_t addr, size_t length, void *data ) const +template void memory_space_impl::read_single_block( mem_addr_t blk_idx, mem_addr_t addr, size_t length, void *data) const { - mem_addr_t index = addr >> m_log2_block_size; - unsigned offset = addr & (BSIZE-1); - unsigned nbytes = length; - typename map_t::const_iterator i = m_data.find(index); - if( (addr+length) > (index+1)*BSIZE ) { + if ((addr + length) > (blk_idx + 1) * BSIZE) { printf("GPGPU-Sim PTX: ERROR * access to memory \'%s\' is unaligned : addr=0x%x, length=%zu\n", m_name.c_str(), addr, length); printf("GPGPU-Sim PTX: (addr+length)=0x%lx > 0x%x=(index+1)*BSIZE, index=0x%x, BSIZE=0x%x\n", - (addr+length),(index+1)*BSIZE, index, BSIZE); + (addr+length),(blk_idx+1)*BSIZE, blk_idx, BSIZE); throw 1; } + typename map_t::const_iterator i = m_data.find(blk_idx); if( i == m_data.end() ) { for( size_t n=0; n < length; n++ ) ((unsigned char*)data)[n] = (unsigned char) 0; //printf("GPGPU-Sim PTX: WARNING reading %zu bytes from unititialized memory at address 0x%x in space %s\n", length, addr, m_name.c_str() ); } else { + unsigned offset = addr & (BSIZE-1); + unsigned nbytes = length; i->second.read(offset,nbytes,(unsigned char*)data); } } +template void memory_space_impl::read( mem_addr_t addr, size_t length, void *data ) const +{ + mem_addr_t index = addr >> m_log2_block_size; + if ((addr+length) <= (index+1)*BSIZE ) { + // fast route for intra-block access + read_single_block(index, addr, length, data); + } else { + // slow route for inter-block access + unsigned nbytes_remain = length; + unsigned dst_offset = 0; + mem_addr_t current_addr = addr; + + while (nbytes_remain > 0) { + unsigned offset = current_addr & (BSIZE-1); + mem_addr_t page = current_addr >> m_log2_block_size; + mem_addr_t access_limit = offset + nbytes_remain; + if (access_limit > BSIZE) { + access_limit = BSIZE; + } + + size_t tx_bytes = access_limit - offset; + read_single_block(page, current_addr, tx_bytes, &((unsigned char*)data)[dst_offset]); + + // advance pointers + dst_offset += tx_bytes; + current_addr += tx_bytes; + nbytes_remain -= tx_bytes; + } + assert(nbytes_remain == 0); + } +} + template void memory_space_impl::print( const char *format, FILE *fout ) const { typename map_t::const_iterator i_page; diff --git a/src/cuda-sim/memory.h b/src/cuda-sim/memory.h index bac4d24..4487b36 100644 --- a/src/cuda-sim/memory.h +++ b/src/cuda-sim/memory.h @@ -164,6 +164,7 @@ public: virtual void set_watch( addr_t addr, unsigned watchpoint ); private: + void read_single_block( mem_addr_t blk_idx, mem_addr_t addr, size_t length, void *data) const; std::string m_name; unsigned m_log2_block_size; typedef mem_map > map_t; diff --git a/src/cuda-sim/opcodes.def b/src/cuda-sim/opcodes.def index c1217ee..5f9caaa 100644 --- a/src/cuda-sim/opcodes.def +++ b/src/cuda-sim/opcodes.def @@ -77,7 +77,7 @@ OP_DEF(ADD_OP,add_impl,"add",1,1) OP_DEF(ADDC_OP,addc_impl,"addc",1,1) OP_DEF(AND_OP,and_impl,"and",1,1) OP_DEF(ANDN_OP,andn_impl,"andn",1,1) -OP_DEF(ATOM_OP,atom_impl,"atom",0,3) +OP_DEF(ATOM_OP,atom_impl,"atom",1,3) OP_DEF(BAR_OP,bar_sync_impl,"bar.sync",1,3) OP_DEF(BFE_OP,bfe_impl,"bfe",1,1) OP_DEF(BFI_OP,bfi_impl,"bfi",1,1) @@ -86,6 +86,7 @@ OP_DEF(BRA_OP,bra_impl,"bra",0,3) OP_DEF(BREV_OP,brev_impl,"brev",1,1) OP_DEF(BRKPT_OP,brkpt_impl,"brkpt",1,9) 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) OP_DEF(CNOT_OP,cnot_impl,"cnot",1,1) OP_DEF(COS_OP,cos_impl,"cos",1,4) @@ -122,6 +123,7 @@ OP_DEF(RCP_OP,rcp_impl,"rcp",1,4) OP_DEF(RED_OP,red_impl,"red",1,7) OP_DEF(REM_OP,rem_impl,"rem",1,1) OP_DEF(RET_OP,ret_impl,"ret",0,3) +OP_DEF(RETP_OP,retp_impl,"retp",0,3) OP_DEF(RSQRT_OP,rsqrt_impl,"rsqrt",1,4) OP_DEF(SAD_OP,sad_impl,"sad",1,1) OP_DEF(SELP_OP,selp_impl,"selp",1,1) diff --git a/src/cuda-sim/ptx-stats.cc b/src/cuda-sim/ptx-stats.cc index 567c62b..40e574a 100644 --- a/src/cuda-sim/ptx-stats.cc +++ b/src/cuda-sim/ptx-stats.cc @@ -67,9 +67,6 @@ #include #include "../tr1_hash_map.h" -// external dependencies -extern function_info *g_func_info; - // options bool enable_ptx_file_line_stats; char * ptx_line_stats_filename = NULL; @@ -192,7 +189,7 @@ void ptx_file_line_stats_add_exec_count(const ptx_instruction *pInsn) // attribute pipeline latency to this ptx instruction (specified by the pc) // pipeline latency is the number of cycles a warp with this instruction spent in the pipeline -void ptx_file_line_stats_add_latency(void * ptx_thd, unsigned pc, unsigned latency) +void ptx_file_line_stats_add_latency(unsigned pc, unsigned latency) { const ptx_instruction *pInsn = function_info::pc_to_instruction(pc); @@ -210,7 +207,7 @@ void ptx_file_line_stats_add_dram_traffic(unsigned pc, unsigned dram_traffic) // attribute the number of shared memory access cycles to a ptx instruction // counts both the number of warps doing shared memory access and the number of cycles involved -void ptx_file_line_stats_add_smem_bank_conflict(void * ptx_thd, unsigned pc, unsigned n_way_bkconflict) +void ptx_file_line_stats_add_smem_bank_conflict(unsigned pc, unsigned n_way_bkconflict) { const ptx_instruction *pInsn = function_info::pc_to_instruction(pc); @@ -221,7 +218,7 @@ void ptx_file_line_stats_add_smem_bank_conflict(void * ptx_thd, unsigned pc, uns // attribute a non-coalesced mem access to a ptx instruction // counts both the number of warps causing this and the number of memory requests generated -void ptx_file_line_stats_add_uncoalesced_gmem(void * ptx_thd, unsigned pc, unsigned n_access) +void ptx_file_line_stats_add_uncoalesced_gmem(unsigned pc, unsigned n_access) { const ptx_instruction *pInsn = function_info::pc_to_instruction(pc); diff --git a/src/cuda-sim/ptx-stats.h b/src/cuda-sim/ptx-stats.h index 004433b..019586c 100644 --- a/src/cuda-sim/ptx-stats.h +++ b/src/cuda-sim/ptx-stats.h @@ -78,10 +78,10 @@ void ptx_file_line_stats_add_exec_count(const ptx_instruction *pInsn); #endif // stat collection interface to gpgpu-sim -void ptx_file_line_stats_add_latency(void * ptx_thd, unsigned pc, unsigned latency); +void ptx_file_line_stats_add_latency(unsigned pc, unsigned latency); void ptx_file_line_stats_add_dram_traffic(unsigned pc, unsigned dram_traffic); -void ptx_file_line_stats_add_smem_bank_conflict(void * ptx_thd, unsigned pc, unsigned n_way_bkconflict); -void ptx_file_line_stats_add_uncoalesced_gmem(void * ptx_thd, unsigned pc, unsigned n_access); +void ptx_file_line_stats_add_smem_bank_conflict(unsigned pc, unsigned n_way_bkconflict); +void ptx_file_line_stats_add_uncoalesced_gmem(unsigned pc, unsigned n_access); void ptx_file_line_stats_create_exposed_latency_tracker(int n_shader_cores); void ptx_file_line_stats_add_inflight_memory_insn(int sc_id, unsigned pc); diff --git a/src/cuda-sim/ptx.l b/src/cuda-sim/ptx.l index 82daef9..1b87c7d 100644 --- a/src/cuda-sim/ptx.l +++ b/src/cuda-sim/ptx.l @@ -98,6 +98,7 @@ bra TC; ptx_lval.int_value = BRA_OP; return OPCODE; brev TC; ptx_lval.int_value = BREV_OP; return OPCODE; brkpt TC; ptx_lval.int_value = BRKPT_OP; return OPCODE; call TC; ptx_lval.int_value = CALL_OP; return OPCODE; +callp TC; ptx_lval.int_value = CALLP_OP; return OPCODE; clz TC; ptx_lval.int_value = CLZ_OP; return OPCODE; cnot TC; ptx_lval.int_value = CNOT_OP; return OPCODE; cos TC; ptx_lval.int_value = COS_OP; return OPCODE; @@ -135,6 +136,7 @@ rcp TC; ptx_lval.int_value = RCP_OP; return OPCODE; red TC; ptx_lval.int_value = RED_OP; return OPCODE; rem TC; ptx_lval.int_value = REM_OP; return OPCODE; ret TC; ptx_lval.int_value = RET_OP; return OPCODE; +retp TC; ptx_lval.int_value = RETP_OP; return OPCODE; rsqrt TC; ptx_lval.int_value = RSQRT_OP; return OPCODE; sad TC; ptx_lval.int_value = SAD_OP; return OPCODE; selp TC; ptx_lval.int_value = SELP_OP; return OPCODE; @@ -263,7 +265,7 @@ breakaddr TC; ptx_lval.int_value = BREAKADDR_OP; return OPCODE; \.bb128 TC; return BB128_TYPE; \.pred TC; return PRED_TYPE; -\.texref TC; return TEXREF_TYPE; +\.texref TC; BEGIN(NOT_OPCODE); return TEXREF_TYPE; \.samplerref TC; return SAMPLERREF_TYPE; \.surfref TC; return SURFREF_TYPE; @@ -271,6 +273,8 @@ breakaddr TC; ptx_lval.int_value = BREAKADDR_OP; return OPCODE; \.v3 TC; return V3_TYPE; \.v4 TC; return V4_TYPE; +\.half TC; return HALF_OPTION; /* ptxplus */ + \.equ TC; return EQU_OPTION; \.neu TC; return NEU_OPTION; \.ltu TC; return LTU_OPTION; @@ -322,6 +326,7 @@ breakaddr TC; ptx_lval.int_value = BREAKADDR_OP; return OPCODE; \.all TC; return ALL_OPTION; \.gl TC; return GLOBAL_OPTION; \.cta TC; return CTA_OPTION; +\.sys TC; return SYS_OPTION; \.exit TC; return EXIT_OPTION; @@ -368,11 +373,11 @@ breakaddr TC; ptx_lval.int_value = BREAKADDR_OP; return OPCODE; ">" TC; return RIGHT_ANGLE_BRACKET; "(" TC; return LEFT_PAREN; ")" TC; return RIGHT_PAREN; -":" TC; return COLON; +":" TC; BEGIN(INITIAL); return COLON; ";" TC; BEGIN(INITIAL); return SEMI_COLON; "!" TC; return EXCLAMATION; "=" TC; return EQUALS; -"{" TC; BEGIN(INITIAL); return LEFT_BRACE; +"{" TC; return LEFT_BRACE; "}" TC; return RIGHT_BRACE; \. TC; return PERIOD; "/" TC; return BACKSLASH; diff --git a/src/cuda-sim/ptx.y b/src/cuda-sim/ptx.y index 1551f22..833e082 100644 --- a/src/cuda-sim/ptx.y +++ b/src/cuda-sim/ptx.y @@ -133,6 +133,7 @@ %token V4_TYPE %token COMMA %token PRED +%token HALF_OPTION %token EQ_OPTION %token NE_OPTION %token LT_OPTION @@ -205,6 +206,7 @@ %token ALL_OPTION %token GLOBAL_OPTION %token CTA_OPTION +%token SYS_OPTION %token EXIT_OPTION %token ABS_OPTION %token TO_OPTION @@ -260,7 +262,8 @@ function_decl_header: ENTRY_DIRECTIVE { $$ = 1; g_func_decl=1; func_header(".ent | EXTERN_DIRECTIVE FUNC_DIRECTIVE { $$ = 2; g_func_decl=1; func_header(".func"); } ; -param_list: param_entry { add_directive(); } +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 identifier_spec { add_function_arg(); } @@ -277,8 +280,10 @@ statement_list: directive_statement { add_directive(); } ; directive_statement: variable_declaration SEMI_COLON - | VERSION_DIRECTIVE DOUBLE_OPERAND { add_version_info($2); } + | VERSION_DIRECTIVE DOUBLE_OPERAND { add_version_info($2, 0); } + | VERSION_DIRECTIVE DOUBLE_OPERAND PLUS { add_version_info($2,1); } | 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); } | LOC_DIRECTIVE INT_OPERAND INT_OPERAND INT_OPERAND @@ -416,6 +421,7 @@ option: type_spec | ALL_OPTION { add_option(ALL_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); } @@ -428,6 +434,7 @@ option: type_spec | ABS_OPTION { add_option(ABS_OPTION); } | atomic_operation_spec ; | TO_OPTION { add_option(TO_OPTION); } + | HALF_OPTION { add_option(HALF_OPTION); } ; atomic_operation_spec: ATOMIC_AND { add_option(ATOMIC_AND); } diff --git a/src/cuda-sim/ptx_ir.cc b/src/cuda-sim/ptx_ir.cc index cd925a5..db915fe 100644 --- a/src/cuda-sim/ptx_ir.cc +++ b/src/cuda-sim/ptx_ir.cc @@ -118,10 +118,10 @@ symbol_table::symbol_table( const char *scope_name, unsigned entry_point, symbol { m_scope_name = std::string(scope_name); m_reg_allocator=0; - m_shared_next = 0x100; // for debug with valgrind: make zero imply undefined address - m_const_next = 0x100; // for debug with valgrind: make zero imply undefined address - m_global_next = 0x100; // for debug with valgrind: make zero imply undefined address - m_local_next = 0x100; // for debug with valgrind: make zero imply undefined address + m_shared_next = 0x100; + m_const_next = 0x100; + m_global_next = 0x100; + m_local_next = 0; m_parent = parent; if ( m_parent ) { m_shared_next = m_parent->m_shared_next; @@ -140,12 +140,24 @@ const ptx_version &symbol_table::get_ptx_version() const else return m_parent->get_ptx_version(); } +unsigned symbol_table::get_sm_target() const +{ + if( m_parent == NULL ) + return m_ptx_version.target(); + else return m_parent->get_sm_target(); +} + void symbol_table::set_ptx_version( float ver, unsigned ext ) { m_ptx_version = ptx_version(ver,ext); assert( m_ptx_version.ver() < 3 ); } +void symbol_table::set_sm_target( const char *target, const char *ext, const char *ext2 ) +{ + m_ptx_version.set_target(target,ext,ext2); +} + symbol *symbol_table::lookup( const char *identifier ) { std::string key(identifier); @@ -306,13 +318,13 @@ void function_info::create_basic_blocks() i = find_next_real_instruction(++i); } else { switch( pI->get_opcode() ) { - case BRA_OP: case RET_OP: case EXIT_OP: + case BRA_OP: case RET_OP: case EXIT_OP: case RETP_OP: i++; if( i != m_instructions.end() ) leaders.push_back(*i); i = find_next_real_instruction(i); break; - case CALL_OP: + case CALL_OP: case CALLP_OP: if( pI->has_pred() ) { printf("GPGPU-Sim PTX: Warning found predicated call\n"); i++; @@ -424,7 +436,7 @@ void function_info::connect_basic_blocks( ) //iterate across m_basic_blocks of f ptx_instruction *pI = (*bb_itr)->ptx_end; if ((*bb_itr)->is_exit) //reached last basic block, no successors to link continue; - if (pI->get_opcode() == RET_OP || pI->get_opcode() == EXIT_OP ) { + if (pI->get_opcode() == RETP_OP || pI->get_opcode() == RET_OP || pI->get_opcode() == EXIT_OP ) { (*bb_itr)->successor_ids.insert(exit_bb->bb_id); exit_bb->predecessor_ids.insert((*bb_itr)->bb_id); if( pI->has_pred() ) { @@ -451,7 +463,7 @@ void function_info::connect_basic_blocks( ) //iterate across m_basic_blocks of f // if basic block does not end in an unpredicated branch, // then next basic block is also successor // (this is better than testing for .uni) - unsigned next_addr = pI->get_m_instr_mem_index() + 1; + unsigned next_addr = pI->get_m_instr_mem_index() + pI->inst_size(); basic_block_t *next_bb = m_instr_mem[next_addr]->get_bb(); (*bb_itr)->successor_ids.insert(next_bb->bb_id); next_bb->predecessor_ids.insert((*bb_itr)->bb_id); @@ -653,12 +665,12 @@ void function_info::find_ipostdominators( ) void function_info::find_idominators( ) { - // find immediate postdominator blocks, using algorithm of + // find immediate dominator blocks, using algorithm of // Muchnick's Adv. Compiler Design & Implemmntation Fig 7.15 - printf("GPGPU-Sim PTX: Finding immediate postdominators for \'%s\'...\n", m_name.c_str() ); + printf("GPGPU-Sim PTX: Finding immediate dominators for \'%s\'...\n", m_name.c_str() ); fflush(stdout); - assert( m_basic_blocks.size() >= 2 ); // must have a distinquished exit block - for (unsigned i=0; i= 2 ); // must have a distinquished entry block + for (unsigned i=0; iTmp_ids = m_basic_blocks[i]->dominator_ids; assert( m_basic_blocks[i]->bb_id == i ); m_basic_blocks[i]->Tmp_ids.erase(i); @@ -683,9 +695,9 @@ void function_info::find_idominators( ) unsigned num_idoms=0; unsigned num_nopred = 0; for ( int n = 0; n < m_basic_blocks.size(); ++n) { - assert( m_basic_blocks[n]->Tmp_ids.size() <= 1 ); - // if the above assert fails we have an error in either postdominator - // computation, the flow graph does not have a unique exit, or some other error + //assert( m_basic_blocks[n]->Tmp_ids.size() <= 1 ); + // if the above assert fails we have an error in either dominator + // computation, the flow graph does not have a unique entry, or some other error if( !m_basic_blocks[n]->Tmp_ids.empty() ) { m_basic_blocks[n]->immediatedominator_id = *m_basic_blocks[n]->Tmp_ids.begin(); num_idoms++; @@ -843,6 +855,8 @@ unsigned type_info_key::type_decode( int type, size_t &size, int &basic_type ) case B16_TYPE: size=16; basic_type=0; return 13; case B32_TYPE: size=32; basic_type=0; return 14; case B64_TYPE: size=64; basic_type=0; return 15; + case BB64_TYPE: size=64; basic_type=0; return 15; + case BB128_TYPE: size=128; basic_type=0; return 16; case TEXREF_TYPE: case SAMPLERREF_TYPE: case SURFREF_TYPE: size=32; basic_type=3; return 16; default: @@ -968,7 +982,8 @@ ptx_instruction::ptx_instruction( int opcode, memory_space_t space_spec, const char *file, unsigned line, - const char *source ) + const char *source, + unsigned warp_size ) : inst_t() { m_uid = ++g_num_ptx_inst_uid; m_PC = 0; @@ -996,8 +1011,9 @@ ptx_instruction::ptx_instruction( int opcode, m_geom_spec = 0; m_vector_spec = 0; m_atomic_spec = 0; - m_warp_size = ::warp_size; + m_warp_size = warp_size; m_membar_level = 0; + m_inst_size = 8; // bytes std::list::const_iterator i; unsigned n=1; @@ -1096,6 +1112,9 @@ ptx_instruction::ptx_instruction( int opcode, case CTA_OPTION: m_membar_level = CTA_OPTION; break; + case SYS_OPTION: + m_membar_level = SYS_OPTION; + break; case FTZ_OPTION: break; case EXIT_OPTION: @@ -1113,7 +1132,9 @@ ptx_instruction::ptx_instruction( int opcode, case CA_OPTION: case CG_OPTION: case CS_OPTION: case LU_OPTION: case CV_OPTION: m_cache_option = last_ptx_inst_option; break; - + case HALF_OPTION: + m_inst_size = 4; // bytes + break; default: assert(0); break; @@ -1141,7 +1162,7 @@ void ptx_instruction::print_insn( FILE *fp ) const snprintf(buf,1024,"%s", m_source.c_str()); p = strtok(buf,";"); if( !is_label() ) - fprintf(fp,"PC=%3u [idx=%3u] ", m_PC, m_instr_mem_index ); + fprintf(fp," PC=%3u [%3u] ", m_PC, m_instr_mem_index ); else fprintf(fp," " ); fprintf(fp,"(%s:%u) %s", m_source_file.c_str(), m_source_line, p ); @@ -1168,7 +1189,13 @@ function_info::function_info(int entry_point ) void function_info::print_insn( unsigned pc, FILE * fp ) const { unsigned index = pc - m_start_PC; - fprintf(fp,"FUNC[%s]",m_name.c_str() ); + char command[1024]; + char buffer[1024]; + snprintf(command,1024,"c++filt -p %s",m_name.c_str()); + FILE *p = popen(command,"r"); + buffer[0]=0; + fscanf(p,"%1023s",buffer); + fprintf(fp,"%s",buffer); if ( index >= m_instr_mem_size ) { fprintf(fp, "", m_start_PC + m_instr_mem_size - 1 ); } else { @@ -1191,10 +1218,13 @@ extern "C" FILE *ptxinfo_in; void gpgpu_ptx_assemble( std::string kname, void *kinfo ) { function_info *func_info = (function_info *)kinfo; + if((function_info *)kinfo == NULL) { + printf("GPGPU-Sim PTX: Error - missing function definition \'%s\'\n", kname.c_str()); + abort(); + } if( func_info->is_extern() ) { printf("GPGPU-Sim PTX: skipping assembly for extern declared function \'%s\'\n", func_info->get_name().c_str() ); return; } - g_func_info = func_info; func_info->ptx_assemble(); } diff --git a/src/cuda-sim/ptx_ir.h b/src/cuda-sim/ptx_ir.h index 43ac92a..afb2476 100644 --- a/src/cuda-sim/ptx_ir.h +++ b/src/cuda-sim/ptx_ir.h @@ -319,7 +319,9 @@ public: symbol_table( const char *scope_name, unsigned entry_point, symbol_table *parent ); void set_name( const char *name ); const ptx_version &get_ptx_version() const; + unsigned get_sm_target() const; void set_ptx_version( float ver, unsigned ext ); + void set_sm_target( const char *target, const char *ext, const char *ext2 ); symbol* lookup( const char *identifier ); std::string get_scope_name() const { return m_scope_name; } symbol *add_variable( const char *identifier, const type_info *type, unsigned size, const char *filename, unsigned line ); @@ -784,9 +786,9 @@ struct gpgpu_recon_t { address_type target_pc; }; -class ptx_instruction { +class ptx_instruction : public inst_t { public: - ptx_instruction( int opcode, + ptx_instruction( int opcode, const symbol *pred, int neg_pred, int pred_mod, @@ -798,9 +800,13 @@ public: memory_space_t space_spec, const char *file, unsigned line, - const char *source ); + const char *source, + unsigned warp_size ); + + void print_insn() const; - void print_insn( FILE *fp ) const; + virtual void print_insn( FILE *fp ) const; + unsigned inst_size() const { return m_inst_size; } unsigned uid() const { return m_uid;} int get_opcode() const { return m_opcode;} const char *get_opcode_cstr() const @@ -961,6 +967,7 @@ public: } private: + basic_block_t *m_basic_block; unsigned m_uid; addr_t m_PC; @@ -999,6 +1006,10 @@ private: enum vote_mode_t m_vote_mode; int m_membar_level; int m_instr_mem_index; //index into m_instr_mem array + unsigned m_inst_size; // bytes + + virtual void pre_decode(); + friend class function_info; }; class param_info { @@ -1036,6 +1047,7 @@ class function_info { public: function_info(int entry_point ); const ptx_version &get_ptx_version() const { return m_symtab->get_ptx_version(); } + unsigned get_sm_target() const { return m_symtab->get_sm_target(); } bool is_extern() const { return m_extern; } void set_name(const char *name) { @@ -1094,22 +1106,7 @@ public: unsigned get_function_size() { return m_instructions.size();} void ptx_assemble(); - void ptx_decode_inst( ptx_thread_info *thd, - unsigned *op_type, - int *i1, - int *i2, - int *i3, - int *i4, - int *o1, - int *o2, - int *o3, - int *o4, - int *vectorin, - int *vectorout, - int *arch_reg, - int *pred, - int *ar1, int *ar2 ); - void ptx_exec_inst( ptx_thread_info *thd, addr_t *addr, memory_space_t *space, unsigned *data_size, unsigned *cycles, dram_callback_t* callback, unsigned warp_active_mask ); + unsigned ptx_get_inst_op( ptx_thread_info *thread ); void add_param( const char *name, struct param_t value ) { @@ -1169,6 +1166,8 @@ public: const void set_kernel_info (const struct gpgpu_ptx_sim_kernel_info &info) { m_kernel_info = info; + m_kernel_info.ptx_version = 10*get_ptx_version().ver(); + m_kernel_info.sm_target = get_ptx_version().target(); } symbol_table *get_symtab() { @@ -1177,9 +1176,8 @@ public: static const ptx_instruction* pc_to_instruction(unsigned pc) { - assert(pc > 0); assert(pc <= s_g_pc_to_insn.size()); - return s_g_pc_to_insn[pc - 1]; + return s_g_pc_to_insn[pc]; } unsigned local_mem_framesize() const { @@ -1404,16 +1402,12 @@ struct textureInfo { unsigned int texel_size_numbits; //log2(texel_size) }; - -extern function_info *g_func_info; - -extern function_info *g_entrypoint_func_info; extern std::map g_sym_name_to_symbol_table; -#define GLOBAL_HEAP_START 0x10000000 +#define GLOBAL_HEAP_START 0x80000000 // start allocating from this address (lower values used for allocating globals in .ptx file) #define SHARED_MEM_SIZE_MAX (64*1024) -#define LOCAL_MEM_SIZE_MAX 1024 +#define LOCAL_MEM_SIZE_MAX (16*1024) #define MAX_STREAMING_MULTIPROCESSORS 64 #define MAX_THREAD_PER_SM 1024 #define TOTAL_LOCAL_MEM_PER_SM (MAX_THREAD_PER_SM*LOCAL_MEM_SIZE_MAX) diff --git a/src/cuda-sim/ptx_loader.cc b/src/cuda-sim/ptx_loader.cc index 783ce91..1aa5101 100644 --- a/src/cuda-sim/ptx_loader.cc +++ b/src/cuda-sim/ptx_loader.cc @@ -67,6 +67,7 @@ #include "cuda-sim.h" #include "ptx_parser.h" #include +#include /// globals @@ -78,6 +79,7 @@ bool g_override_embedded_ptx = false; struct ptx_info_t { char *str; + char *cubin_str; char *fname; ptx_info_t *next; }; @@ -93,6 +95,9 @@ extern "C" int ptxinfo_parse(); extern "C" int ptxinfo_debug; extern "C" FILE *ptxinfo_in; +extern int g_ptx_convert_to_ptxplus; +extern int g_ptx_save_converted_ptxplus; + /// static functions static int load_static_globals( symbol_table *symtab, unsigned min_gaddr, unsigned max_gaddr) @@ -223,11 +228,23 @@ void gpgpu_ptx_sim_load_gpu_kernels() printf("GPGPU-Sim PTX: USING EMBEDDED .ptx files...\n"); ptx_info_t *s; for ( s=g_ptx_source_array; s!=NULL; s=s->next ) { - symbol_table *symtab=gpgpu_ptx_sim_load_ptx_from_string(s->str, ++source_num); + symbol_table *symtab; + source_num++; + if(g_ptx_convert_to_ptxplus) { + char *ptxplus_str = gpgpu_ptx_sim_convert_ptx_to_ptxplus(s->str, s->cubin_str, source_num); + symtab=gpgpu_ptx_sim_load_ptx_from_string(ptxplus_str, s->str, source_num); + delete[] ptxplus_str; + } else { + symtab=gpgpu_ptx_sim_load_ptx_from_string(s->str, s->str, source_num); + } load_static_globals(symtab,STATIC_ALLOC_LIMIT,0xFFFFFFFF); load_constants(symtab,STATIC_ALLOC_LIMIT); } } else { + if(g_ptx_convert_to_ptxplus) { + perror("GPGPU-Sim PTX: convert_to_ptxplus option enabled. Cannot use this option with external ptx files.\n"); + assert(0); + } const char *filename = NULL; struct dirent **namelist; int n = scandir(".", &namelist, ptx_file_filter, alphasort); @@ -235,20 +252,17 @@ void gpgpu_ptx_sim_load_gpu_kernels() perror("GPGPU-Sim PTX: no PTX files returned by scandir"); else { while (n--) { - if ( filename != NULL ) { - printf("Loader error: support for multiple .ptx files not yet enabled\n"); - abort(); - } filename = namelist[n]->d_name; printf("Parsing %s..\n", filename); ptx_in = fopen( filename, "r" ); - free(namelist[n]); symbol_table *symtab=init_parser(filename); ptx_parse (); ptxinfo_in = open_ptxinfo(filename); ptxinfo_parse(); load_static_globals(symtab,STATIC_ALLOC_LIMIT,0xFFFFFFFF); load_constants(symtab,STATIC_ALLOC_LIMIT); + + free(namelist[n]); } free(namelist); } @@ -278,11 +292,17 @@ void gpgpu_ptx_sim_load_gpu_kernels() } } -void gpgpu_ptx_sim_add_ptxstring( const char *ptx_string, const char *sourcefname ) +void gpgpu_ptx_sim_add_ptxstring( const char *ptx_string, const char *cubin_string, const char *sourcefname ) { ptx_info_t *t = new ptx_info_t; t->next = NULL; t->str = strdup(ptx_string); + if (cubin_string != NULL) { + t->cubin_str = strdup(cubin_string); + } else { + assert(g_ptx_convert_to_ptxplus == 0); + t->cubin_str = NULL; + } t->fname = strdup(sourcefname); // put ptx source into a fifo @@ -338,7 +358,106 @@ void print_ptx_file( const char *p, unsigned source_num, const char *filename ) fflush(stdout); } -symbol_table *gpgpu_ptx_sim_load_ptx_from_string( const char *p, unsigned source_num ) +char* gpgpu_ptx_sim_convert_ptx_to_ptxplus(const char *ptx_str, const char *cubin_str, unsigned source_num) +{ + printf("GPGPU-Sim PTX: converting EMBEDDED .ptx file to ptxplus \n"); + + // Extract ptx to a file + char fname_ptx[1024]; + snprintf(fname_ptx,1024,"_ptx_XXXXXX"); + int fd=mkstemp(fname_ptx); + close(fd); + + printf("GPGPU-Sim PTX: extracting embedded .ptx to temporary file \"%s\"\n", fname_ptx); + FILE *ptxfile = fopen(fname_ptx,"w"); + fprintf(ptxfile,"%s",ptx_str); + fclose(ptxfile); + + // Extract cubin to a file + char fname_cubin[1024]; + snprintf(fname_cubin,1024,"_cubin_XXXXXX"); + int fd2=mkstemp(fname_cubin); + close(fd2); + + printf("GPGPU-Sim PTX: extracting embedded cubin to temporary file \"%s\"\n", fname_cubin); + FILE *cubinfile = fopen(fname_cubin,"w"); + fprintf(cubinfile,"%s",cubin_str); + fclose(cubinfile); + + // Run decuda + char fname_decuda[1024]; + snprintf(fname_decuda,1024,"_decuda_XXXXXX"); + int fd3=mkstemp(fname_decuda); + close(fd3); + + char decuda_commandline[1024]; + snprintf(decuda_commandline,1024,"$DECUDA_INSTALL_PATH/decuda.py -o %s %s", fname_decuda, fname_cubin); + + printf("GPGPU-Sim PTX: calling decuda on cubin file, decuda output file = \"%s\"\n", fname_decuda); + int decuda_result = system(decuda_commandline); + if( decuda_result != 0 ) { + printf("GPGPU-Sim PTX: ERROR ** while calling decuda (b) %d\n", decuda_result); + printf(" Ensure env variable DECUDA_INSTALL_PATH is set and points to decuda base directory.\n"); + exit(1); + } + + // Run decuda_to_ptxplus + char fname_ptxplus[1024]; + snprintf(fname_ptxplus,1024,"_ptxplus_XXXXXX"); + int fd4=mkstemp(fname_ptxplus); + close(fd4); + + char d2pp_commandline[1024]; + snprintf(d2pp_commandline,1024,"$D2PP_INSTALL_PATH/decuda_to_ptxplus %s %s %s %s > /dev/null", fname_decuda, fname_ptx, fname_cubin, fname_ptxplus); + + printf("GPGPU-Sim PTX: calling decuda_to_ptxplus, ptxplus output file = \"%s\"\n", fname_ptxplus); + int d2pp_result = system(d2pp_commandline); + if( d2pp_result != 0 ) { + printf("GPGPU-Sim PTX: ERROR ** while calling decuda_to_ptxplus %d\n", d2pp_result); + printf(" Ensure env variable D2PP_INSTALL_PATH is set and points to decuda_to_ptxplus base directory.\n"); + exit(1); + } + + // Get ptxplus from file + std::ifstream fileStream(fname_ptxplus, std::ios::in); + std::string text, line; + while(getline(fileStream,line)) { + text += (line + "\n"); + } + fileStream.close(); + + char* ptxplus_str = new char [strlen(text.c_str())+1]; + strcpy(ptxplus_str, text.c_str()); + + // Save ptxplus to file if specified + if(g_ptx_save_converted_ptxplus) { + char fname_ptxplus_save[1024]; + snprintf(fname_ptxplus_save,1024,"_%u.ptxplus", source_num ); + printf("GPGPU-Sim PTX: saving converted ptxplus to file \"%s\"\n", fname_ptxplus_save); + + FILE *file_ptxplus_save = fopen(fname_ptxplus_save,"w"); + fprintf(file_ptxplus_save,"%s",ptxplus_str); + fclose(file_ptxplus_save); + } + + // Remove temporary files + char rm_commandline[1024]; + snprintf(rm_commandline,1024,"rm -f %s %s %s %s", fname_ptx, fname_cubin, fname_decuda, fname_ptxplus); + printf("GPGPU-Sim PTX: removing temporary files using \"%s\"\n", rm_commandline); + int rm_result = system(rm_commandline); + if( rm_result != 0 ) { + printf("GPGPU-Sim PTX: ERROR ** while removing temporary files %d\n", rm_result); + exit(1); + } + + printf("GPGPU-Sim PTX: DONE converting EMBEDDED .ptx file to ptxplus \n"); + + return ptxplus_str; + +} + + +symbol_table *gpgpu_ptx_sim_load_ptx_from_string( const char *p, const char *p_for_info, unsigned source_num ) { char buf[1024]; snprintf(buf,1024,"_%u.ptx", source_num ); @@ -375,7 +494,7 @@ symbol_table *gpgpu_ptx_sim_load_ptx_from_string( const char *p, unsigned source printf("GPGPU-Sim PTX: extracting embedded .ptx to temporary file \"%s\"\n", fname); FILE *ptxfile = fopen(fname,"w"); - fprintf(ptxfile,"%s",p); + fprintf(ptxfile,"%s", p_for_info); fclose(ptxfile); char fname2[1024]; diff --git a/src/cuda-sim/ptx_loader.h b/src/cuda-sim/ptx_loader.h index 95130f1..01dc0b7 100644 --- a/src/cuda-sim/ptx_loader.h +++ b/src/cuda-sim/ptx_loader.h @@ -75,7 +75,8 @@ extern memory_space *g_param_mem; extern bool g_override_embedded_ptx; void gpgpu_ptx_sim_load_gpu_kernels(); -void gpgpu_ptx_sim_add_ptxstring( const char *ptx_string, const char *sourcefname ); -class symbol_table *gpgpu_ptx_sim_load_ptx_from_string( const char *p, unsigned source_num ); +void gpgpu_ptx_sim_add_ptxstring( const char *ptx_string, const char *cubin_string, const char *sourcefname ); +class symbol_table *gpgpu_ptx_sim_load_ptx_from_string( const char *p, const char *p_for_info, unsigned source_num ); +char* gpgpu_ptx_sim_convert_ptx_to_ptxplus(const char *ptx_str, const char *cubin_str, unsigned source_num); #endif diff --git a/src/cuda-sim/ptx_parser.cc b/src/cuda-sim/ptx_parser.cc index d4a892f..5a65481 100644 --- a/src/cuda-sim/ptx_parser.cc +++ b/src/cuda-sim/ptx_parser.cc @@ -68,6 +68,12 @@ extern "C" int ptx_error( const char *s ); extern int ptx_lineno; +static unsigned g_warp_size; +void set_ptx_warp_size(unsigned warp_size) +{ + g_warp_size=warp_size; +} + static bool g_debug_ir_generation=false; const char *g_filename; unsigned g_max_regs_per_thread = 0; @@ -110,12 +116,10 @@ std::list g_scalar_type; fflush(stdout); \ } -unsigned g_entry_func_param_index=0; -function_info *g_func_info = NULL; -function_info *g_entrypoint_func_info = NULL; -symbol_table *g_entrypoint_symbol_table = NULL; -std::map g_ptx_token_decode; -operand_info g_return_var; +static unsigned g_entry_func_param_index=0; +static function_info *g_func_info = NULL; +static std::map g_ptx_token_decode; +static operand_info g_return_var; const char *decode_token( int type ) { @@ -194,9 +198,6 @@ void add_function_name( const char *name ) { DPRINTF("add_function_name %s %s", name, ((g_entry_point==1)?"(entrypoint)":((g_entry_point==2)?"(extern)":""))); bool prior_decl = g_global_symbol_table->add_function_decl( name, g_entry_point, &g_func_info, &g_current_symbol_table ); - if( g_entry_point ) { - g_entrypoint_func_info = g_func_info; - } if( g_add_identifier_cached__identifier ) { add_identifier( g_add_identifier_cached__identifier, g_add_identifier_cached__array_dim, @@ -270,7 +271,7 @@ extern "C" char linebuf[1024]; void set_return() { - parse_assert( (g_opcode == CALL_OP), "only call can have return value"); + parse_assert( (g_opcode == CALL_OP || g_opcode == CALLP_OP), "only call can have return value"); g_operands.front().set_return(); g_return_var = g_operands.front(); } @@ -291,6 +292,7 @@ const ptx_instruction *ptx_instruction_lookup( const char *filename, unsigned li void add_instruction() { DPRINTF("add_instruction: %s", ((g_opcode>0)?g_opcode_string[g_opcode]:"