From 43eb0e822424b39b52091f45725376cc903b08ee Mon Sep 17 00:00:00 2001 From: Tor Aamodt Date: Sun, 8 Aug 2010 08:34:45 -0800 Subject: refactoring: making g_current_symbol_table local to ptx_ir.cc [git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 7165] --- src/cuda-sim/cuda-sim.cc | 313 ++++------------------------------------------- 1 file changed, 25 insertions(+), 288 deletions(-) (limited to 'src/cuda-sim/cuda-sim.cc') diff --git a/src/cuda-sim/cuda-sim.cc b/src/cuda-sim/cuda-sim.cc index 4acefe5..74b18bc 100644 --- a/src/cuda-sim/cuda-sim.cc +++ b/src/cuda-sim/cuda-sim.cc @@ -222,266 +222,10 @@ int gpgpu_ptx_sim_sizeofTexture(const char* name) return array->size; } -extern unsigned int warp_size; - -int g_warn_literal_operands_two_type_inst; - -std::list check_operands( int opcode, - const std::list &scalar_type, - const std::list &operands ) -{ - if( (opcode == CVT_OP) || (opcode == SET_OP) || (opcode == SLCT_OP) || (opcode == TEX_OP) ) { - // just make sure these do not have have const operands... - if( !g_warn_literal_operands_two_type_inst ) { - std::list::const_iterator o; - for( o = operands.begin(); o != operands.end(); o++ ) { - const operand_info &op = *o; - if( op.is_literal() ) { - printf("GPGPU-Sim PTX: PTX uses two scalar type intruction with literal operand.\n"); - g_warn_literal_operands_two_type_inst = 1; - } - } - } - } else { - assert( scalar_type.size() < 2 ); - if( scalar_type.size() == 1 ) { - std::list result; - int inst_type = scalar_type.front(); - std::list::const_iterator o; - for( o = operands.begin(); o != operands.end(); o++ ) { - const operand_info &op = *o; - if( op.is_literal() ) { - if( (op.get_type() == double_op_t) && (inst_type == F32_TYPE) ) { - ptx_reg_t v = op.get_literal_value(); - float u = (float)v.f64; - operand_info n(u); - result.push_back(n); - } else { - result.push_back(op); - } - } else { - result.push_back(op); - } - } - return result; - } - } - return operands; -} - -ptx_instruction::ptx_instruction( int opcode, - const symbol *pred, - int neg_pred, - symbol *label, - const std::list &operands, - const operand_info &return_var, - const std::list &options, - const std::list &scalar_type, - memory_space_t space_spec, - const char *file, - unsigned line, - const char *source ) -{ - m_uid = ++g_num_ptx_inst_uid; - m_PC = 0; - m_opcode = opcode; - m_pred = pred; - m_neg_pred = neg_pred; - m_label = label; - const std::list checked_operands = check_operands(opcode,scalar_type,operands); - m_operands.insert(m_operands.begin(), checked_operands.begin(), checked_operands.end() ); - m_return_var = return_var; - m_options = options; - m_wide = false; - m_hi = false; - m_lo = false; - m_uni = false; - m_to_option = false; - m_cache_option = 0; - m_rounding_mode = RN_OPTION; - m_compare_op = -1; - m_saturation_mode = 0; - m_geom_spec = 0; - m_vector_spec = 0; - m_atomic_spec = 0; - m_warp_size = ::warp_size; - m_membar_level = 0; - - std::list::const_iterator i; - unsigned n=1; - for ( i=options.begin(); i!= options.end(); i++, n++ ) { - int last_ptx_inst_option = *i; - switch ( last_ptx_inst_option ) { - case EQU_OPTION: - case NEU_OPTION: - case LTU_OPTION: - case LEU_OPTION: - case GTU_OPTION: - case GEU_OPTION: - case EQ_OPTION: - case NE_OPTION: - case LT_OPTION: - case LE_OPTION: - case GT_OPTION: - case GE_OPTION: - case LS_OPTION: - case HS_OPTION: - m_compare_op = last_ptx_inst_option; - break; - case NUM_OPTION: - case NAN_OPTION: - m_compare_op = last_ptx_inst_option; - // assert(0); // finish this - break; - case SAT_OPTION: - m_saturation_mode = 1; - break; - case RNI_OPTION: - case RZI_OPTION: - case RMI_OPTION: - case RPI_OPTION: - case RN_OPTION: - case RZ_OPTION: - case RM_OPTION: - case RP_OPTION: - m_rounding_mode = last_ptx_inst_option; - break; - case HI_OPTION: - m_compare_op = last_ptx_inst_option; - m_hi = true; - assert( !m_lo ); - assert( !m_wide ); - break; - case LO_OPTION: - m_compare_op = last_ptx_inst_option; - m_lo = true; - assert( !m_hi ); - assert( !m_wide ); - break; - case WIDE_OPTION: - m_wide = true; - assert( !m_lo ); - assert( !m_hi ); - break; - case UNI_OPTION: - m_uni = true; // don't care... < now we DO care when constructing flowgraph> - break; - case GEOM_MODIFIER_1D: - case GEOM_MODIFIER_2D: - case GEOM_MODIFIER_3D: - m_geom_spec = last_ptx_inst_option; - break; - case V2_TYPE: - case V3_TYPE: - case V4_TYPE: - m_vector_spec = last_ptx_inst_option; - break; - case ATOMIC_AND: - case ATOMIC_OR: - case ATOMIC_XOR: - case ATOMIC_CAS: - case ATOMIC_EXCH: - case ATOMIC_ADD: - case ATOMIC_INC: - case ATOMIC_DEC: - case ATOMIC_MIN: - case ATOMIC_MAX: - m_atomic_spec = last_ptx_inst_option; - break; - case APPROX_OPTION: - break; - case FULL_OPTION: - break; - case ANY_OPTION: - m_vote_mode = vote_any; - break; - case ALL_OPTION: - m_vote_mode = vote_all; - break; - case GLOBAL_OPTION: - m_membar_level = GLOBAL_OPTION; - break; - case CTA_OPTION: - m_membar_level = CTA_OPTION; - break; - case FTZ_OPTION: - break; - case TO_OPTION: - m_to_option = true; - break; - case CA_OPTION: case CG_OPTION: case CS_OPTION: case LU_OPTION: case CV_OPTION: - m_cache_option = last_ptx_inst_option; - break; - - default: - assert(0); - break; - } - } - m_scalar_type = scalar_type; - m_space_spec = space_spec; - if( ( opcode == ST_OP || opcode == LD_OP ) && (space_spec == undefined_space) ) { - m_space_spec = generic_space; - } - m_source_file = file?file:""; - m_source_line = line; - m_source = source; -} - -void ptx_instruction::print_insn() const -{ - print_insn(stdout); - fflush(stdout); -} - -void ptx_instruction::print_insn( FILE *fp ) const -{ - char buf[1024], *p; - 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 ); - else - fprintf(fp," " ); - fprintf(fp,"(%s:%u) %s", m_source_file.c_str(), m_source_line, p ); -} - unsigned g_assemble_code_next_pc=1; std::map g_pc_to_finfo; std::vector function_info::s_g_pc_to_insn; -unsigned function_info::sm_next_uid = 1; - -function_info::function_info(int entry_point ) -{ - m_uid = sm_next_uid++; - m_entry_point = (entry_point==1)?true:false; - m_extern = (entry_point==2)?true:false; - num_reconvergence_pairs = 0; - m_symtab = NULL; - m_assembled = false; - m_return_var_sym = NULL; - m_kernel_info.cmem = 0; - m_kernel_info.lmem = 0; - m_kernel_info.regs = 0; - m_kernel_info.smem = 0; - m_local_mem_framesize = 0; -} - -void function_info::print_insn( unsigned pc, FILE * fp ) const -{ - unsigned index = pc - m_start_PC; - fprintf(fp,"FUNC[%s]",m_name.c_str() ); - if ( index >= m_instr_mem_size ) { - fprintf(fp, "", m_start_PC + m_instr_mem_size - 1 ); - } else { - if ( m_instr_mem[index] != NULL ) - m_instr_mem[index]->print_insn(fp); - else - fprintf(fp, "", pc ); - } -} - void function_info::ptx_assemble() { if( m_assembled ) { @@ -528,7 +272,7 @@ void function_info::ptx_assemble() } unsigned index = labels[ target.name() ]; //determine address from name unsigned PC = m_instr_mem[index]->get_PC(); - g_current_symbol_table->set_label_address( target.get_symbol(), PC ); + m_symtab->set_label_address( target.get_symbol(), PC ); target.set_type(label_t); } } @@ -1036,7 +780,7 @@ void function_info::add_param_data( unsigned argn, struct gpgpu_ptx_sim_arg *arg } } -void function_info::finalize( memory_space *param_mem, symbol_table *symtab ) +void function_info::finalize( memory_space *param_mem ) { unsigned param_address = 0; for( std::map::iterator i=m_ptx_kernel_param_info.begin(); i!=m_ptx_kernel_param_info.end(); i++ ) { @@ -1045,7 +789,7 @@ void function_info::finalize( memory_space *param_mem, symbol_table *symtab ) int type = p.get_type(); param_t param_value = p.get_value(); param_value.type = type; - symbol *param = symtab->lookup(name.c_str()); + symbol *param = m_symtab->lookup(name.c_str()); unsigned xtype = param->type()->get_key().scalar_type(); assert(xtype==(unsigned)type); size_t size; @@ -1069,14 +813,11 @@ void function_info::finalize( memory_space *param_mem, symbol_table *symtab ) void function_info::list_param( FILE *fout ) const { - symbol_table *symtab = g_current_symbol_table; for( std::map::const_iterator i=m_ptx_kernel_param_info.begin(); i!=m_ptx_kernel_param_info.end(); i++ ) { const param_info &p = i->second; std::string name = p.get_name(); - symbol *param = symtab->lookup(name.c_str()); - + symbol *param = m_symtab->lookup(name.c_str()); addr_t param_addr = param->get_address(); - fprintf(fout, "%s: %#08x\n", name.c_str(), param_addr); } fflush(fout); @@ -1146,7 +887,6 @@ void function_info::ptx_exec_inst( ptx_thread_info *thread, ptx_instruction *pI = m_instr_mem[index]; try { - g_current_symbol_table = thread->get_finfo()->get_symtab(); thread->clearRPC(); thread->m_last_set_operand_value.u64 = 0; @@ -1415,7 +1155,7 @@ unsigned ptx_sim_init_thread( ptx_thread_info** thread_info,int sid,unsigned tid local_mem = new memory_space_impl<32>(buf,32); local_mem_lookup[new_tid] = local_mem; } - thd->set_info(g_entrypoint_symbol_table,g_entrypoint_func_info); + 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); @@ -1510,20 +1250,20 @@ void gpgpu_ptx_sim_init_grid( const char *kernel_key, struct gpgpu_ptx_sim_arg* } printf("\n"); abort(); - } else { - 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) == - g_kernel_name_to_function_lookup->end() ) { - printf("GPGPU-Sim PTX: ERROR ** function \'%s\' not found in ptx file\n", kname.c_str() ); - abort(); - } - g_entrypoint_func_info = g_func_info = (*g_kernel_name_to_function_lookup)[kname]; - g_entrypoint_symbol_table = g_current_symbol_table = g_kernel_name_to_symtab_lookup[kname]; + } + + 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) == + g_kernel_name_to_function_lookup->end() ) { + printf("GPGPU-Sim PTX: ERROR ** function \'%s\' not found in ptx file\n", kname.c_str() ); + abort(); } + g_entrypoint_func_info = g_func_info = (*g_kernel_name_to_function_lookup)[kname]; + g_entrypoint_symbol_table = g_kernel_name_to_symtab_lookup[kname]; unsigned argcount=0; struct gpgpu_ptx_sim_arg *tmparg = args; @@ -1538,7 +1278,7 @@ void gpgpu_ptx_sim_init_grid( const char *kernel_key, struct gpgpu_ptx_sim_arg* args = args->m_next; argn++; } - g_func_info->finalize(g_param_mem,g_current_symbol_table); + g_func_info->finalize(g_param_mem); g_ptx_kernel_count++; if ( gpgpu_ptx_instruction_classification ) { init_inst_classification_stat(); @@ -1661,11 +1401,10 @@ void gpgpu_ptx_sim_memcpy_symbol(const char *hostVar, const void *src, size_t co memory_space *mem = NULL; std::map::iterator st = g_sym_name_to_symbol_table.find(sym_name.c_str()); - if( st != g_sym_name_to_symbol_table.end() ) { - g_current_symbol_table = st->second; - } + assert( st != g_sym_name_to_symbol_table.end() ); + symbol_table *symtab = st->second; - symbol *sym = g_current_symbol_table->lookup(sym_name.c_str()); + symbol *sym = symtab->lookup(sym_name.c_str()); assert(sym); unsigned dst = sym->get_address() + offset; switch (mem_region.get_type()) { @@ -1931,15 +1670,13 @@ void gpgpu_ptx_assemble( std::string kname, void *kinfo ) printf("GPGPU-Sim PTX: skipping assembly for extern declared function \'%s\'\n", func_info->get_name().c_str() ); return; } - g_func_info = func_info; - g_current_symbol_table = g_kernel_name_to_symtab_lookup[ kname ]; - if( g_current_symbol_table == NULL ) { + if( g_kernel_name_to_symtab_lookup[ kname ] == NULL ) { printf("\nGPGPU-Sim PTX: ERROR no information for kernel \'%s\'\n" " this can happen for kernels contained in CUDA\n" " libraries (such as CUBLAS, CUFFT, or CUDPP)\n", kname.c_str() ); exit(1); } - + g_func_info = func_info; func_info->ptx_assemble(); } @@ -2075,7 +1812,7 @@ void gpgpu_ptx_sim_main_func( const char *kernel_key, dim3 gridDim, dim3 blockDi } - thd->set_info(g_current_symbol_table,g_func_info); + 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); -- cgit v1.3