summaryrefslogtreecommitdiff
path: root/src
diff options
context:
space:
mode:
Diffstat (limited to 'src')
-rw-r--r--src/cuda-sim/cuda-sim.cc311
-rw-r--r--src/cuda-sim/ptx_ir.cc258
-rw-r--r--src/cuda-sim/ptx_ir.h3
-rw-r--r--src/cuda-sim/ptx_sim.cc13
-rw-r--r--src/cuda-sim/ptx_sim.h2
5 files changed, 289 insertions, 298 deletions
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<operand_info> check_operands( int opcode,
- const std::list<int> &scalar_type,
- const std::list<operand_info> &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<operand_info>::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<operand_info> result;
- int inst_type = scalar_type.front();
- std::list<operand_info>::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<operand_info> &operands,
- const operand_info &return_var,
- const std::list<int> &options,
- const std::list<int> &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<operand_info> 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<int>::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:"<unknown>";
- 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<unsigned,function_info*> g_pc_to_finfo;
std::vector<ptx_instruction*> 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, "<past last instruction (max pc=%u)>", 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, "<no instruction at pc = %u>", 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<unsigned,param_info>::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<unsigned,param_info>::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<std::string,symbol_table*>::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);
diff --git a/src/cuda-sim/ptx_ir.cc b/src/cuda-sim/ptx_ir.cc
index b2c8552..d049a6a 100644
--- a/src/cuda-sim/ptx_ir.cc
+++ b/src/cuda-sim/ptx_ir.cc
@@ -78,7 +78,7 @@ extern "C" int ptx_error( const char *s );
// the program intermediate representation...
symbol_table *g_global_symbol_table = NULL;
-symbol_table *g_current_symbol_table = NULL;
+static symbol_table *g_current_symbol_table = NULL;
std::list<ptx_instruction*> g_instructions;
symbol *g_last_symbol = NULL;
std::map<std::string,symbol_table*> g_sym_name_to_symbol_table;
@@ -1424,3 +1424,259 @@ void copy_buffer_list_into_frame(ptx_thread_info * thread, arg_buffer_list_t &ar
}
}
+
+
+static std::list<operand_info> check_operands( int opcode,
+ const std::list<int> &scalar_type,
+ const std::list<operand_info> &operands )
+{
+ static int g_warn_literal_operands_two_type_inst;
+ 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<operand_info>::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<operand_info> result;
+ int inst_type = scalar_type.front();
+ std::list<operand_info>::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<operand_info> &operands,
+ const operand_info &return_var,
+ const std::list<int> &options,
+ const std::list<int> &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<operand_info> 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<int>::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:"<unknown>";
+ 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 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, "<past last instruction (max pc=%u)>", 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, "<no instruction at pc = %u>", pc );
+ }
+}
+
diff --git a/src/cuda-sim/ptx_ir.h b/src/cuda-sim/ptx_ir.h
index 1daf9f1..839cceb 100644
--- a/src/cuda-sim/ptx_ir.h
+++ b/src/cuda-sim/ptx_ir.h
@@ -1010,7 +1010,7 @@ public:
return m_start_PC;
}
- void finalize( memory_space *param_mem, symbol_table *symtab );
+ void finalize( memory_space *param_mem );
void list_param( FILE *fout ) const;
const struct gpgpu_ptx_sim_kernel_info* get_kernel_info () {
@@ -1259,7 +1259,6 @@ extern function_info *g_func_info;
extern int g_error_detected;
extern bool g_debug_ir_generation;
extern std::list<ptx_instruction*> g_instructions;
-extern symbol_table *g_current_symbol_table;
extern symbol_table *g_entrypoint_symbol_table;
extern function_info *g_entrypoint_func_info;
extern symbol_table *g_global_symbol_table;
diff --git a/src/cuda-sim/ptx_sim.cc b/src/cuda-sim/ptx_sim.cc
index 9a2acf6..31c8c81 100644
--- a/src/cuda-sim/ptx_sim.cc
+++ b/src/cuda-sim/ptx_sim.cc
@@ -301,9 +301,9 @@ unsigned ptx_thread_info::get_builtin( int builtin_id, unsigned dim_mod )
return 0;
}
-void ptx_thread_info::set_info( symbol_table *symtab, function_info *func )
+void ptx_thread_info::set_info( function_info *func )
{
- m_symbol_table = symtab;
+ m_symbol_table = func->get_symtab();
m_func_info = func;
m_PC = func->get_start_PC();
}
@@ -313,9 +313,9 @@ void ptx_thread_info::print_insn( unsigned pc, FILE * fp ) const
m_func_info->print_insn(pc,fp);
}
-static void print_reg( std::string name, ptx_reg_t value )
+static void print_reg( std::string name, ptx_reg_t value, symbol_table *symtab )
{
- const symbol *sym = g_current_symbol_table->lookup(name.c_str());
+ const symbol *sym = symtab->lookup(name.c_str());
printf(" %8s ", name.c_str() );
if( sym == NULL ) {
printf("<unknown type> 0x%llx\n", (unsigned long long ) value.u64 );
@@ -451,8 +451,7 @@ void ptx_thread_info::dump_regs()
for ( r=m_regs.back().begin(); r != m_regs.back().end(); ++r ) {
std::string name = r->first->name();
ptx_reg_t value = r->second;
- print_reg(name,value);
-
+ print_reg(name,value,m_symbol_table);
}
}
@@ -465,7 +464,7 @@ void ptx_thread_info::dump_modifiedregs()
for ( r=m_debug_trace_regs_modified.begin(); r != m_debug_trace_regs_modified.end(); ++r ) {
std::string name = r->first->name();
ptx_reg_t value = r->second;
- print_reg(name,value);
+ print_reg(name,value,m_symbol_table);
}
}
diff --git a/src/cuda-sim/ptx_sim.h b/src/cuda-sim/ptx_sim.h
index 12075c8..0379686 100644
--- a/src/cuda-sim/ptx_sim.h
+++ b/src/cuda-sim/ptx_sim.h
@@ -270,7 +270,7 @@ public:
return m_func_info;
}
void print_insn( unsigned pc, FILE * fp ) const;
- void set_info( symbol_table *symtab, function_info *func );
+ void set_info( function_info *func );
unsigned get_uid() const
{
return m_uid;
generated by cgit v1.3 (git 2.53.0) at 2026-07-11 02:20:23 +0000