From b3e786e3d8d720217f36a214e9b5be9a19ab9dd2 Mon Sep 17 00:00:00 2001 From: Mengchi Zhang Date: Mon, 8 Jul 2019 12:37:52 -0400 Subject: Move opcode_latency_int thus pass gpgpu_context into many classes Signed-off-by: Mengchi Zhang --- src/cuda-sim/ptx_parser.cc | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) (limited to 'src/cuda-sim/ptx_parser.cc') diff --git a/src/cuda-sim/ptx_parser.cc b/src/cuda-sim/ptx_parser.cc index 05fc618..0139534 100644 --- a/src/cuda-sim/ptx_parser.cc +++ b/src/cuda-sim/ptx_parser.cc @@ -285,7 +285,8 @@ void ptx_recognizer::add_instruction() gpgpu_ctx->g_filename, ptx_get_lineno(scanner), linebuf, - g_shader_core_config ); + g_shader_core_config, + gpgpu_ctx ); g_instructions.push_back(i); g_inst_lookup[gpgpu_ctx->g_filename][ptx_get_lineno(scanner)] = i; init_instruction_state(); -- cgit v1.3 From c58d52ccde1eea6c45d0f8d2d5411bcf1ebab4f9 Mon Sep 17 00:00:00 2001 From: Mengchi Zhang Date: Wed, 10 Jul 2019 00:45:39 -0400 Subject: Move g_globals and g_constants Signed-off-by: Mengchi Zhang --- libopencl/opencl_runtime_api.cc | 10 +++++----- src/cuda-sim/cuda-sim.cc | 5 +---- src/cuda-sim/cuda-sim.h | 4 +++- src/cuda-sim/ptx_parser.cc | 7 ++----- src/stream_manager.cc | 5 +++-- 5 files changed, 14 insertions(+), 17 deletions(-) (limited to 'src/cuda-sim/ptx_parser.cc') diff --git a/libopencl/opencl_runtime_api.cc b/libopencl/opencl_runtime_api.cc index aaaec4f..7f029c7 100644 --- a/libopencl/opencl_runtime_api.cc +++ b/libopencl/opencl_runtime_api.cc @@ -950,11 +950,11 @@ clEnqueueNDRangeKernel(cl_command_queue command_queue, gpgpu_t *gpu = command_queue->get_device()->the_device(); if (kernel->get_implementation()->get_ptx_version().ver() <3.0){ - gpgpu_ptx_sim_memcpy_symbol( "%_global_size", _global_size, 3 * sizeof(int), 0, 1, gpu ); - gpgpu_ptx_sim_memcpy_symbol( "%_work_dim", &work_dim, 1 * sizeof(int), 0, 1, gpu ); - gpgpu_ptx_sim_memcpy_symbol( "%_global_num_groups", &GridDim, 3 * sizeof(int), 0, 1, gpu ); - gpgpu_ptx_sim_memcpy_symbol( "%_global_launch_offset", zeros, 3 * sizeof(int), 0, 1, gpu ); - gpgpu_ptx_sim_memcpy_symbol( "%_global_block_offset", zeros, 3 * sizeof(int), 0, 1, gpu ); + ctx->func_sim->gpgpu_ptx_sim_memcpy_symbol( "%_global_size", _global_size, 3 * sizeof(int), 0, 1, gpu ); + ctx->func_sim->gpgpu_ptx_sim_memcpy_symbol( "%_work_dim", &work_dim, 1 * sizeof(int), 0, 1, gpu ); + ctx->func_sim->gpgpu_ptx_sim_memcpy_symbol( "%_global_num_groups", &GridDim, 3 * sizeof(int), 0, 1, gpu ); + ctx->func_sim->gpgpu_ptx_sim_memcpy_symbol( "%_global_launch_offset", zeros, 3 * sizeof(int), 0, 1, gpu ); + ctx->func_sim->gpgpu_ptx_sim_memcpy_symbol( "%_global_block_offset", zeros, 3 * sizeof(int), 0, 1, gpu ); } kernel_info_t *grid = ctx->func_sim->gpgpu_opencl_ptx_sim_init_grid(kernel->get_implementation(),params,GridDim,BlockDim,gpu); if ( ctx->func_sim->g_ptx_sim_mode ) diff --git a/src/cuda-sim/cuda-sim.cc b/src/cuda-sim/cuda-sim.cc index 7e8aab9..3d1da62 100644 --- a/src/cuda-sim/cuda-sim.cc +++ b/src/cuda-sim/cuda-sim.cc @@ -1942,9 +1942,6 @@ void print_splash() } } -std::set g_globals; -std::set g_constants; - void cuda_sim::gpgpu_ptx_sim_register_const_variable(void *hostVar, const char *deviceName, size_t size ) { printf("GPGPU-Sim PTX registering constant %s (%zu bytes) to name mapping\n", deviceName, size ); @@ -1957,7 +1954,7 @@ void cuda_sim::gpgpu_ptx_sim_register_global_variable(void *hostVar, const char g_global_name_lookup[hostVar] = deviceName; } -void gpgpu_ptx_sim_memcpy_symbol(const char *hostVar, const void *src, size_t count, size_t offset, int to, gpgpu_t *gpu ) +void cuda_sim::gpgpu_ptx_sim_memcpy_symbol(const char *hostVar, const void *src, size_t count, size_t offset, int to, gpgpu_t *gpu ) { printf("GPGPU-Sim PTX: starting gpgpu_ptx_sim_memcpy_symbol with hostVar 0x%p\n", hostVar); bool found_sym = false; diff --git a/src/cuda-sim/cuda-sim.h b/src/cuda-sim/cuda-sim.h index 628e434..5c95100 100644 --- a/src/cuda-sim/cuda-sim.h +++ b/src/cuda-sim/cuda-sim.h @@ -46,7 +46,6 @@ extern int g_debug_execution; extern int g_debug_thread_uid; extern void print_splash(); -extern void gpgpu_ptx_sim_memcpy_symbol(const char *hostVar, const void *src, size_t count, size_t offset, int to, gpgpu_t *gpu ); extern void ptxinfo_opencl_addinfo( std::map &kernels ); unsigned ptx_sim_init_thread( kernel_info_t &kernel, @@ -161,6 +160,8 @@ class cuda_sim { bool g_cuda_launch_blocking; void ** g_inst_classification_stat; void ** g_inst_op_classification_stat; + std::set g_globals; + std::set g_constants; // backward pointer class gpgpu_context* gpgpu_ctx; //global functions @@ -179,6 +180,7 @@ class cuda_sim { void set_param_gpgpu_num_shaders(int num_shaders); struct rec_pts find_reconvergence_points( function_info *finfo ); address_type get_converge_point( address_type pc ); + void gpgpu_ptx_sim_memcpy_symbol(const char *hostVar, const void *src, size_t count, size_t offset, int to, gpgpu_t *gpu ); }; #endif diff --git a/src/cuda-sim/ptx_parser.cc b/src/cuda-sim/ptx_parser.cc index 0139534..269ec4d 100644 --- a/src/cuda-sim/ptx_parser.cc +++ b/src/cuda-sim/ptx_parser.cc @@ -322,9 +322,6 @@ bool ptx_recognizer::check_for_duplicates( const char *identifier ) return ( s != NULL ); } -extern std::set g_globals; -extern std::set g_constants; - // Returns padding that needs to be inserted ahead of address to make it aligned to min(size, maxalign) /* * @param address the address in bytes @@ -453,7 +450,7 @@ void ptx_recognizer::add_identifier( const char *identifier, int array_dim, unsi g_current_symbol_table->alloc_global( num_bits/8 + addr_pad ); } if( g_current_symbol_table == g_global_symbol_table ) { - g_constants.insert( identifier ); + gpgpu_ctx->func_sim->g_constants.insert( identifier ); } assert( g_current_symbol_table != NULL ); g_sym_name_to_symbol_table[ identifier ] = g_current_symbol_table; @@ -471,7 +468,7 @@ void ptx_recognizer::add_identifier( const char *identifier, int array_dim, unsi fflush(stdout); g_last_symbol->set_address( addr+addr_pad ); g_current_symbol_table->alloc_global( num_bits/8 + addr_pad ); - g_globals.insert( identifier ); + gpgpu_ctx->func_sim->g_globals.insert( identifier ); assert( g_current_symbol_table != NULL ); g_sym_name_to_symbol_table[ identifier ] = g_current_symbol_table; break; diff --git a/src/stream_manager.cc b/src/stream_manager.cc index e07f4e4..be3dd71 100644 --- a/src/stream_manager.cc +++ b/src/stream_manager.cc @@ -29,6 +29,7 @@ #include "gpgpusim_entrypoint.h" #include "cuda-sim/cuda-sim.h" #include "gpgpu-sim/gpu-sim.h" +#include "../libcuda/gpgpu_context.h" unsigned CUstream_st::sm_next_stream_uid = 0; @@ -150,13 +151,13 @@ bool stream_operation::do_operation( gpgpu_sim *gpu ) case stream_memcpy_to_symbol: if(g_debug_execution >= 3) printf("memcpy to symbol\n"); - gpgpu_ptx_sim_memcpy_symbol(m_symbol,m_host_address_src,m_cnt,m_offset,1,gpu); + gpu->gpgpu_ctx->func_sim->gpgpu_ptx_sim_memcpy_symbol(m_symbol,m_host_address_src,m_cnt,m_offset,1,gpu); m_stream->record_next_done(); break; case stream_memcpy_from_symbol: if(g_debug_execution >= 3) printf("memcpy from symbol\n"); - gpgpu_ptx_sim_memcpy_symbol(m_symbol,m_host_address_dst,m_cnt,m_offset,0,gpu); + gpu->gpgpu_ctx->func_sim->gpgpu_ptx_sim_memcpy_symbol(m_symbol,m_host_address_dst,m_cnt,m_offset,0,gpu); m_stream->record_next_done(); break; case stream_kernel_launch: -- cgit v1.3 From 87726d32ada00fcd93f2cf24ccae4ba593c4f9ec Mon Sep 17 00:00:00 2001 From: Mengchi Zhang Date: Wed, 10 Jul 2019 15:03:16 -0400 Subject: Move g_pc_to_finfo Signed-off-by: Mengchi Zhang --- src/abstract_hardware_model.cc | 3 ++- src/abstract_hardware_model.h | 13 ++++++++----- src/cuda-sim/cuda-sim.cc | 7 +++---- src/cuda-sim/cuda-sim.h | 5 +++-- src/cuda-sim/ptx_ir.cc | 12 +++++++----- src/cuda-sim/ptx_ir.h | 9 +++++++-- src/cuda-sim/ptx_parser.cc | 2 +- src/gpgpu-sim/shader.cc | 4 ++-- src/gpgpu-sim/shader.h | 2 +- 9 files changed, 34 insertions(+), 23 deletions(-) (limited to 'src/cuda-sim/ptx_parser.cc') diff --git a/src/abstract_hardware_model.cc b/src/abstract_hardware_model.cc index 248e7a5..ef09051 100644 --- a/src/abstract_hardware_model.cc +++ b/src/abstract_hardware_model.cc @@ -39,6 +39,7 @@ #include #include #include +#include "../libcuda/gpgpu_context.h" unsigned mem_access_t::sm_next_access_uid = 0; unsigned warp_inst_t::sm_next_uid = 0; @@ -945,7 +946,7 @@ void simt_stack::print (FILE *fout) const } else { fprintf(fout," " ); } - ptx_print_insn( stack_entry.m_pc, fout ); + m_gpu->gpgpu_ctx->func_sim->ptx_print_insn( stack_entry.m_pc, fout ); fprintf(fout,"\n"); } diff --git a/src/abstract_hardware_model.h b/src/abstract_hardware_model.h index da29a11..8ef8376 100644 --- a/src/abstract_hardware_model.h +++ b/src/abstract_hardware_model.h @@ -347,9 +347,11 @@ public: mutable bool volta_cache_config_set; }; -struct core_config { - core_config() - { +class core_config { + public: + core_config(gpgpu_context* ctx) + { + gpgpu_ctx = ctx; m_valid = false; num_shmem_bank=16; shmem_limited_broadcast = false; @@ -361,6 +363,8 @@ struct core_config { bool m_valid; unsigned warp_size; + // backward pointer + class gpgpu_context* gpgpu_ctx; // off-chip memory request architecture parameters int gpgpu_coalesce_arch; @@ -934,7 +938,7 @@ public: m_empty=true; m_config=NULL; } - warp_inst_t( const core_config *config ) + warp_inst_t( const core_config *config ) { m_uid=0; assert(config->warp_size<=MAX_WARP_SIZE); @@ -1105,7 +1109,6 @@ public: unsigned get_uid() const { return m_uid; } unsigned get_schd_id() const { return m_scheduler_id; } - protected: unsigned m_uid; diff --git a/src/cuda-sim/cuda-sim.cc b/src/cuda-sim/cuda-sim.cc index 3d1da62..fb9bc9e 100644 --- a/src/cuda-sim/cuda-sim.cc +++ b/src/cuda-sim/cuda-sim.cc @@ -220,7 +220,6 @@ void gpgpu_t::gpgpu_ptx_sim_unbindTexture(const struct textureReference* texref) } 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*/ @@ -257,7 +256,7 @@ void function_info::ptx_assemble() const symbol *l = pI->get_label(); labels[l->name()] = n; } else { - g_pc_to_finfo[PC] = this; + gpgpu_ctx->func_sim->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]); @@ -497,7 +496,7 @@ void gpgpu_t::gpu_memset( size_t dst_start_addr, int c, size_t count ) } } -void ptx_print_insn( address_type pc, FILE *fp ) +void cuda_sim::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() ) { @@ -509,7 +508,7 @@ void ptx_print_insn( address_type pc, FILE *fp ) finfo->print_insn(pc,fp); } -std::string ptx_get_insn_str( address_type pc ) +std::string cuda_sim::ptx_get_insn_str( address_type pc ) { std::map::iterator f = g_pc_to_finfo.find(pc); if( f == g_pc_to_finfo.end() ) { diff --git a/src/cuda-sim/cuda-sim.h b/src/cuda-sim/cuda-sim.h index 5c95100..c578524 100644 --- a/src/cuda-sim/cuda-sim.h +++ b/src/cuda-sim/cuda-sim.h @@ -61,8 +61,6 @@ unsigned ptx_sim_init_thread( kernel_info_t &kernel, bool functionalSimulationMode = false); const warp_inst_t *ptx_fetch_inst( address_type pc ); const struct gpgpu_ptx_sim_info* ptx_sim_kernel_info(const class function_info *kernel); -void ptx_print_insn( address_type pc, FILE *fp ); -std::string ptx_get_insn_str( address_type pc ); /*! @@ -162,6 +160,7 @@ class cuda_sim { void ** g_inst_op_classification_stat; std::set g_globals; std::set g_constants; + std::map g_pc_to_finfo; // backward pointer class gpgpu_context* gpgpu_ctx; //global functions @@ -181,6 +180,8 @@ class cuda_sim { struct rec_pts find_reconvergence_points( function_info *finfo ); address_type get_converge_point( address_type pc ); void gpgpu_ptx_sim_memcpy_symbol(const char *hostVar, const void *src, size_t count, size_t offset, int to, gpgpu_t *gpu ); + void ptx_print_insn( address_type pc, FILE *fp ); + std::string ptx_get_insn_str( address_type pc ); }; #endif diff --git a/src/cuda-sim/ptx_ir.cc b/src/cuda-sim/ptx_ir.cc index 1bd409e..c537091 100644 --- a/src/cuda-sim/ptx_ir.cc +++ b/src/cuda-sim/ptx_ir.cc @@ -83,8 +83,9 @@ symbol_table::symbol_table() assert(0); } -symbol_table::symbol_table( const char *scope_name, unsigned entry_point, symbol_table *parent ) +symbol_table::symbol_table( const char *scope_name, unsigned entry_point, symbol_table *parent, gpgpu_context* ctx ) { + gpgpu_ctx = ctx; m_scope_name = std::string(scope_name); m_reg_allocator=0; m_shared_next = 0; @@ -183,7 +184,7 @@ symbol_table* symbol_table::start_inst_group() { //previous added assert(m_inst_group_symtab.find(std::string(inst_group_name)) == m_inst_group_symtab.end()); - symbol_table *sym_table = new symbol_table(inst_group_name, 3/*inst group*/, this ); + symbol_table *sym_table = new symbol_table(inst_group_name, 3/*inst group*/, this, gpgpu_ctx ); sym_table->m_global_next = m_global_next; sym_table->m_shared_next = m_shared_next; @@ -221,7 +222,7 @@ bool symbol_table::add_function_decl( const char *name, int entry_point, functio *func_info = m_function_info_lookup[key]; prior_decl = true; } else { - *func_info = new function_info(entry_point); + *func_info = new function_info(entry_point, gpgpu_ctx); (*func_info)->set_name(name); (*func_info)->set_maxnt_id(0); m_function_info_lookup[key] = *func_info; @@ -232,7 +233,7 @@ bool symbol_table::add_function_decl( const char *name, int entry_point, functio *sym_table = m_function_symtab_lookup[key]; } else { assert( !prior_decl ); - *sym_table = new symbol_table( "", entry_point, this ); + *sym_table = new symbol_table( "", entry_point, this, gpgpu_ctx ); // Initial setup code to support a register represented as "_". // This register is used when an instruction operand is @@ -1373,8 +1374,9 @@ std::string ptx_instruction::to_string() const unsigned function_info::sm_next_uid = 1; -function_info::function_info(int entry_point ) +function_info::function_info(int entry_point, gpgpu_context* ctx ) { + gpgpu_ctx = ctx; m_uid = sm_next_uid++; m_entry_point = (entry_point==1)?true:false; m_extern = (entry_point==2)?true:false; diff --git a/src/cuda-sim/ptx_ir.h b/src/cuda-sim/ptx_ir.h index 1604551..babd54b 100644 --- a/src/cuda-sim/ptx_ir.h +++ b/src/cuda-sim/ptx_ir.h @@ -306,7 +306,7 @@ private: class symbol_table { public: symbol_table(); - symbol_table( const char *scope_name, unsigned entry_point, symbol_table *parent ); + symbol_table( const char *scope_name, unsigned entry_point, symbol_table *parent, gpgpu_context* ctx); void set_name( const char *name ); const ptx_version &get_ptx_version() const; unsigned get_sm_target() const; @@ -348,6 +348,9 @@ public: symbol_table* start_inst_group(); symbol_table* end_inst_group(); + // backward pointer + class gpgpu_context* gpgpu_ctx; + private: unsigned m_reg_allocator; unsigned m_shared_next; @@ -1233,7 +1236,7 @@ private: class function_info { public: - function_info(int entry_point ); + function_info(int entry_point, gpgpu_context* ctx ); 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; } @@ -1403,6 +1406,8 @@ public: void set_maxnt_id(unsigned maxthreads) { maxnt_id = maxthreads;} unsigned get_maxnt_id() { return maxnt_id;} + // backward pointer + class gpgpu_context* gpgpu_ctx; private: unsigned maxnt_id; diff --git a/src/cuda-sim/ptx_parser.cc b/src/cuda-sim/ptx_parser.cc index 269ec4d..5a94679 100644 --- a/src/cuda-sim/ptx_parser.cc +++ b/src/cuda-sim/ptx_parser.cc @@ -112,7 +112,7 @@ symbol_table * gpgpu_context::init_parser( const char *ptx_filename ) { g_filename = strdup(ptx_filename); if (g_global_allfiles_symbol_table == NULL) { - g_global_allfiles_symbol_table = new symbol_table("global_allfiles", 0, NULL); + g_global_allfiles_symbol_table = new symbol_table("global_allfiles", 0, NULL, this); ptx_parser->g_global_symbol_table = ptx_parser->g_current_symbol_table = g_global_allfiles_symbol_table; } /*else { diff --git a/src/gpgpu-sim/shader.cc b/src/gpgpu-sim/shader.cc index 69790fc..6cd6d8f 100644 --- a/src/gpgpu-sim/shader.cc +++ b/src/gpgpu-sim/shader.cc @@ -1025,7 +1025,7 @@ void scheduler_unit::cycle() m_simt_stack[warp_id]->get_pdom_stack_top_info(&pc,&rpc); SCHED_DPRINTF( "Warp (warp_id %u, dynamic_warp_id %u) has valid instruction (%s)\n", (*iter)->get_warp_id(), (*iter)->get_dynamic_warp_id(), - ptx_get_insn_str( pc).c_str() ); + m_shader->m_config->gpgpu_ctx->func_sim->ptx_get_insn_str( pc).c_str() ); if( pI ) { assert(valid); if( pc != pI->pc ) { @@ -2690,7 +2690,7 @@ void warp_inst_t::print( FILE *fout ) const for (unsigned j=0; jwarp_size; j++) fprintf(fout, "%c", (active(j)?'1':'0') ); fprintf(fout, "]: "); - ptx_print_insn( pc, fout ); + m_config->gpgpu_ctx->func_sim->ptx_print_insn( pc, fout ); fprintf(fout, "\n"); } void shader_core_ctx::incexecstat(warp_inst_t *&inst) diff --git a/src/gpgpu-sim/shader.h b/src/gpgpu-sim/shader.h index e0cefac..2837f1b 100644 --- a/src/gpgpu-sim/shader.h +++ b/src/gpgpu-sim/shader.h @@ -1367,7 +1367,7 @@ const char* const pipeline_stage_name_decode[] = { class shader_core_config : public core_config { public: - shader_core_config(gpgpu_context* ctx){ + shader_core_config(gpgpu_context* ctx):core_config(ctx){ pipeline_widths_string = NULL; gpgpu_ctx = ctx; } -- cgit v1.3