From 4d4d5938d715d2b79a617c32583184426b4a642d Mon Sep 17 00:00:00 2001 From: Mengchi Zhang Date: Tue, 9 Jul 2019 23:16:17 -0400 Subject: Move g_ptx_sim_mode Signed-off-by: Mengchi Zhang --- src/cuda-sim/cuda-sim.cc | 6 ++---- src/cuda-sim/cuda-sim.h | 10 +++++++--- src/cuda-sim/cuda_device_runtime.cc | 7 ++++--- src/cuda-sim/cuda_device_runtime.h | 18 ++++++++++++++++-- 4 files changed, 29 insertions(+), 12 deletions(-) (limited to 'src/cuda-sim') diff --git a/src/cuda-sim/cuda-sim.cc b/src/cuda-sim/cuda-sim.cc index b3e2965..7a7d205 100644 --- a/src/cuda-sim/cuda-sim.cc +++ b/src/cuda-sim/cuda-sim.cc @@ -2033,13 +2033,11 @@ void gpgpu_ptx_sim_memcpy_symbol(const char *hostVar, const void *src, size_t co fflush(stdout); } -int g_ptx_sim_mode; // if non-zero run functional simulation only (i.e., no notion of a clock cycle) - extern int ptx_debug; bool g_cuda_launch_blocking = false; -void read_sim_environment_variables() +void cuda_sim::read_sim_environment_variables() { ptx_debug = 0; g_debug_execution = 0; @@ -2185,7 +2183,7 @@ void cuda_sim::gpgpu_cuda_ptx_sim_main_func( kernel_info_t &kernel, bool openCL cta.execute(cp_count,temp); #if (CUDART_VERSION >= 5000) - launch_all_device_kernels(); + gpgpu_ctx->device_runtime->launch_all_device_kernels(); #endif } else diff --git a/src/cuda-sim/cuda-sim.h b/src/cuda-sim/cuda-sim.h index 25ebf7b..3c4336d 100644 --- a/src/cuda-sim/cuda-sim.h +++ b/src/cuda-sim/cuda-sim.h @@ -36,12 +36,12 @@ #include #include"ptx_sim.h" +class gpgpu_context; class memory_space; class function_info; class symbol_table; extern const char *g_gpgpusim_version_string; -extern int g_ptx_sim_mode; extern int g_debug_execution; extern int g_debug_thread_uid; extern void ** g_inst_classification_stat; @@ -50,7 +50,6 @@ extern void ** g_inst_op_classification_stat; 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 read_sim_environment_variables(); extern void ptxinfo_opencl_addinfo( std::map &kernels ); unsigned ptx_sim_init_thread( kernel_info_t &kernel, class ptx_thread_info** thread_info, @@ -124,9 +123,10 @@ struct gpgpu_ptx_sim_info get_ptxinfo(); class cuda_sim { public: - cuda_sim() { + cuda_sim( gpgpu_context* ctx ) { g_ptx_sim_num_insn = 0; g_ptx_kernel_count = -1; // used for classification stat collection purposes + gpgpu_ctx = ctx; } //global variables char *opcode_latency_int; @@ -147,6 +147,9 @@ class cuda_sim { int g_ptx_kernel_count; // used for classification stat collection purposes std::map g_global_name_lookup; // indexed by hostVar std::map g_const_name_lookup; // indexed by hostVar + int g_ptx_sim_mode; // if non-zero run functional simulation only (i.e., no notion of a clock cycle) + // backward pointer + class gpgpu_context* gpgpu_ctx; //global functions void ptx_opcocde_latency_options (option_parser_t opp); void gpgpu_cuda_ptx_sim_main_func( kernel_info_t &kernel, bool openCL = false ); @@ -159,6 +162,7 @@ class cuda_sim { gpgpu_t *gpu ); void gpgpu_ptx_sim_register_global_variable(void *hostVar, const char *deviceName, size_t size ); void gpgpu_ptx_sim_register_const_variable(void*, const char *deviceName, size_t size ); + void read_sim_environment_variables(); }; #endif diff --git a/src/cuda-sim/cuda_device_runtime.cc b/src/cuda-sim/cuda_device_runtime.cc index be8369f..354fa79 100644 --- a/src/cuda-sim/cuda_device_runtime.cc +++ b/src/cuda-sim/cuda_device_runtime.cc @@ -20,6 +20,7 @@ unsigned long long g_max_total_param_size = 0; #include "../stream_manager.h" #include "../gpgpusim_entrypoint.h" #include "cuda_device_runtime.h" +#include "../../libcuda/gpgpu_context.h" #define DEV_RUNTIME_REPORT(a) \ if( g_debug_execution ) { \ @@ -318,17 +319,17 @@ void gpgpusim_cuda_streamCreateWithFlags(const ptx_instruction * pI, ptx_thread_ } -void launch_one_device_kernel() { +void cuda_device_runtime::launch_one_device_kernel() { if(!g_cuda_device_launch_op.empty()) { device_launch_operation_t &op = g_cuda_device_launch_op.front(); - stream_operation stream_op = stream_operation(op.grid, g_ptx_sim_mode, op.stream); + stream_operation stream_op = stream_operation(op.grid, gpgpu_ctx->func_sim->g_ptx_sim_mode, op.stream); g_stream_manager()->push(stream_op); g_cuda_device_launch_op.pop_front(); } } -void launch_all_device_kernels() { +void cuda_device_runtime::launch_all_device_kernels() { while(!g_cuda_device_launch_op.empty()) { launch_one_device_kernel(); } diff --git a/src/cuda-sim/cuda_device_runtime.h b/src/cuda-sim/cuda_device_runtime.h index 6dbcd71..851fed2 100644 --- a/src/cuda-sim/cuda_device_runtime.h +++ b/src/cuda-sim/cuda_device_runtime.h @@ -6,6 +6,20 @@ void gpgpusim_cuda_getParameterBufferV2(const ptx_instruction * pI, ptx_thread_info * thread, const function_info * target_func); void gpgpusim_cuda_launchDeviceV2(const ptx_instruction * pI, ptx_thread_info * thread, const function_info * target_func); void gpgpusim_cuda_streamCreateWithFlags(const ptx_instruction * pI, ptx_thread_info * thread, const function_info * target_func); -void launch_all_device_kernels(); -void launch_one_device_kernel(); +#endif +#if (CUDART_VERSION >= 5000) + +class gpgpu_context; + +class cuda_device_runtime { + public: + cuda_device_runtime( gpgpu_context* ctx ) { + gpgpu_ctx = ctx; + } + // backward pointer + class gpgpu_context* gpgpu_ctx; + void launch_all_device_kernels(); + void launch_one_device_kernel(); +}; + #endif -- cgit v1.3 From bf3146963f4261c24df76f23b5e21cd62d98cb14 Mon Sep 17 00:00:00 2001 From: Mengchi Zhang Date: Tue, 9 Jul 2019 23:25:11 -0400 Subject: Move gpgpu_param_num_shaders Signed-off-by: Mengchi Zhang --- src/cuda-sim/cuda-sim.cc | 6 ++---- src/cuda-sim/cuda-sim.h | 4 +++- src/gpgpu-sim/gpu-sim.cc | 2 +- 3 files changed, 6 insertions(+), 6 deletions(-) (limited to 'src/cuda-sim') diff --git a/src/cuda-sim/cuda-sim.cc b/src/cuda-sim/cuda-sim.cc index 7a7d205..11ba4ec 100644 --- a/src/cuda-sim/cuda-sim.cc +++ b/src/cuda-sim/cuda-sim.cc @@ -61,8 +61,6 @@ int g_debug_thread_uid = 0; addr_t g_debug_pc = 0xBEEF1518; // Output debug information to file options -unsigned gpgpu_param_num_shaders = 0; - unsigned cdp_latency[5]; void cuda_sim::ptx_opcocde_latency_options (option_parser_t opp) { @@ -1736,7 +1734,7 @@ void ptx_thread_info::ptx_exec_inst( warp_inst_t &inst, unsigned lane_id) } -void set_param_gpgpu_num_shaders(int num_shaders) +void cuda_sim::set_param_gpgpu_num_shaders(int num_shaders) { gpgpu_param_num_shaders = num_shaders; } @@ -1818,7 +1816,7 @@ unsigned ptx_sim_init_thread( kernel_info_t &kernel, assert( max_cta_per_sm > 0 ); //unsigned sm_idx = (tid/cta_size)*gpgpu_param_num_shaders + sid; - unsigned sm_idx = hw_cta_id*gpgpu_param_num_shaders + sid; + unsigned sm_idx = hw_cta_id*gpu->gpgpu_ctx->func_sim->gpgpu_param_num_shaders + sid; if ( shared_memory_lookup.find(sm_idx) == shared_memory_lookup.end() ) { if ( g_debug_execution >= 1 ) { diff --git a/src/cuda-sim/cuda-sim.h b/src/cuda-sim/cuda-sim.h index 3c4336d..aa1fe40 100644 --- a/src/cuda-sim/cuda-sim.h +++ b/src/cuda-sim/cuda-sim.h @@ -66,7 +66,6 @@ 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 ); -void set_param_gpgpu_num_shaders(int num_shaders); /*! @@ -126,6 +125,7 @@ class cuda_sim { cuda_sim( gpgpu_context* ctx ) { g_ptx_sim_num_insn = 0; g_ptx_kernel_count = -1; // used for classification stat collection purposes + gpgpu_param_num_shaders = 0; gpgpu_ctx = ctx; } //global variables @@ -148,6 +148,7 @@ class cuda_sim { std::map g_global_name_lookup; // indexed by hostVar std::map g_const_name_lookup; // indexed by hostVar int g_ptx_sim_mode; // if non-zero run functional simulation only (i.e., no notion of a clock cycle) + unsigned gpgpu_param_num_shaders; // backward pointer class gpgpu_context* gpgpu_ctx; //global functions @@ -163,6 +164,7 @@ class cuda_sim { void gpgpu_ptx_sim_register_global_variable(void *hostVar, const char *deviceName, size_t size ); void gpgpu_ptx_sim_register_const_variable(void*, const char *deviceName, size_t size ); void read_sim_environment_variables(); + void set_param_gpgpu_num_shaders(int num_shaders); }; #endif diff --git a/src/gpgpu-sim/gpu-sim.cc b/src/gpgpu-sim/gpu-sim.cc index bbcc078..30e0aa5 100644 --- a/src/gpgpu-sim/gpu-sim.cc +++ b/src/gpgpu-sim/gpu-sim.cc @@ -882,7 +882,7 @@ void gpgpu_sim::init() gpu_sim_cycle_parition_util = 0; reinit_clock_domains(); - set_param_gpgpu_num_shaders(m_config.num_shader()); + gpgpu_ctx->func_sim->set_param_gpgpu_num_shaders(m_config.num_shader()); for (unsigned i=0;in_simt_clusters;i++) m_cluster[i]->reinit(); m_shader_stats->new_grid(); -- cgit v1.3 From 8b4f4ac167fd57284cebf38746bf571a52cf3a0f Mon Sep 17 00:00:00 2001 From: Mengchi Zhang Date: Tue, 9 Jul 2019 23:51:17 -0400 Subject: Move g_rpts Signed-off-by: Mengchi Zhang --- src/cuda-sim/cuda-sim.cc | 15 +++------------ src/cuda-sim/cuda-sim.h | 10 ++++++++++ 2 files changed, 13 insertions(+), 12 deletions(-) (limited to 'src/cuda-sim') diff --git a/src/cuda-sim/cuda-sim.cc b/src/cuda-sim/cuda-sim.cc index 11ba4ec..832a223 100644 --- a/src/cuda-sim/cuda-sim.cc +++ b/src/cuda-sim/cuda-sim.cc @@ -112,8 +112,6 @@ cudaLaunchDeviceV2_init_perWarp, cudaLaunchDevicV2_perKernel>" "7200,8000,100,12000,1600"); } -static address_type get_converge_point(address_type pc); - void gpgpu_t::gpgpu_ptx_sim_bindNameToTexture(const char* name, const struct textureReference* texref, int dim, int readmode, int ext) { std::string texname(name); @@ -1096,7 +1094,7 @@ void ptx_instruction::pre_decode() } // get reconvergence pc - reconvergence_pc = get_converge_point(pc); + reconvergence_pc = gpgpu_ctx->func_sim->get_converge_point(pc); m_decoded=true; } @@ -2518,14 +2516,7 @@ void ptxinfo_opencl_addinfo( std::map &kernels ) clear_ptxinfo(); } -struct rec_pts { - gpgpu_recon_t *s_kernel_recon_points; - int s_num_recon; -}; - -class std::map g_rpts; - -struct rec_pts find_reconvergence_points( function_info *finfo ) +struct rec_pts cuda_sim::find_reconvergence_points( function_info *finfo ) { rec_pts tmp; std::map::iterator r=g_rpts.find(finfo); @@ -2564,7 +2555,7 @@ address_type get_return_pc( void *thd ) return the_thread->get_return_PC(); } -address_type get_converge_point( address_type pc ) +address_type cuda_sim::get_converge_point( address_type pc ) { // the branch could encode the reconvergence point and/or a bit that indicates the // reconvergence point is the return PC on the call stack in the case the branch has diff --git a/src/cuda-sim/cuda-sim.h b/src/cuda-sim/cuda-sim.h index aa1fe40..0a1401a 100644 --- a/src/cuda-sim/cuda-sim.h +++ b/src/cuda-sim/cuda-sim.h @@ -120,6 +120,13 @@ void print_ptxinfo(); void clear_ptxinfo(); struct gpgpu_ptx_sim_info get_ptxinfo(); +class gpgpu_recon_t; +struct rec_pts { + gpgpu_recon_t *s_kernel_recon_points; + int s_num_recon; +}; + + class cuda_sim { public: cuda_sim( gpgpu_context* ctx ) { @@ -149,6 +156,7 @@ class cuda_sim { std::map g_const_name_lookup; // indexed by hostVar int g_ptx_sim_mode; // if non-zero run functional simulation only (i.e., no notion of a clock cycle) unsigned gpgpu_param_num_shaders; + class std::map g_rpts; // backward pointer class gpgpu_context* gpgpu_ctx; //global functions @@ -165,6 +173,8 @@ class cuda_sim { void gpgpu_ptx_sim_register_const_variable(void*, const char *deviceName, size_t size ); void read_sim_environment_variables(); 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 ); }; #endif -- cgit v1.3 From 717de45b61d12fd4aedea111515433167c7efab7 Mon Sep 17 00:00:00 2001 From: Mengchi Zhang Date: Wed, 10 Jul 2019 00:11:38 -0400 Subject: Move g_cuda_launch_blocking Signed-off-by: Mengchi Zhang --- src/cuda-sim/cuda-sim.cc | 2 -- src/cuda-sim/cuda-sim.h | 2 ++ src/gpgpusim_entrypoint.cc | 4 +--- 3 files changed, 3 insertions(+), 5 deletions(-) (limited to 'src/cuda-sim') diff --git a/src/cuda-sim/cuda-sim.cc b/src/cuda-sim/cuda-sim.cc index 832a223..e9508ee 100644 --- a/src/cuda-sim/cuda-sim.cc +++ b/src/cuda-sim/cuda-sim.cc @@ -2031,8 +2031,6 @@ void gpgpu_ptx_sim_memcpy_symbol(const char *hostVar, const void *src, size_t co extern int ptx_debug; -bool g_cuda_launch_blocking = false; - void cuda_sim::read_sim_environment_variables() { ptx_debug = 0; diff --git a/src/cuda-sim/cuda-sim.h b/src/cuda-sim/cuda-sim.h index 0a1401a..6177986 100644 --- a/src/cuda-sim/cuda-sim.h +++ b/src/cuda-sim/cuda-sim.h @@ -133,6 +133,7 @@ class cuda_sim { g_ptx_sim_num_insn = 0; g_ptx_kernel_count = -1; // used for classification stat collection purposes gpgpu_param_num_shaders = 0; + g_cuda_launch_blocking = false; gpgpu_ctx = ctx; } //global variables @@ -157,6 +158,7 @@ class cuda_sim { int g_ptx_sim_mode; // if non-zero run functional simulation only (i.e., no notion of a clock cycle) unsigned gpgpu_param_num_shaders; class std::map g_rpts; + bool g_cuda_launch_blocking; // backward pointer class gpgpu_context* gpgpu_ctx; //global functions diff --git a/src/gpgpusim_entrypoint.cc b/src/gpgpusim_entrypoint.cc index 683a695..816159f 100644 --- a/src/gpgpusim_entrypoint.cc +++ b/src/gpgpusim_entrypoint.cc @@ -207,8 +207,6 @@ void exit_simulation() fflush(stdout); } -extern bool g_cuda_launch_blocking; - gpgpu_sim *gpgpu_context::gpgpu_ptx_sim_init_perf() { srand(1); @@ -233,7 +231,7 @@ gpgpu_sim *gpgpu_context::gpgpu_ptx_sim_init_perf() GPGPUsim_ctx_ptr()->g_the_gpu_config->init(); GPGPUsim_ctx_ptr()->g_the_gpu = new gpgpu_sim(*(GPGPUsim_ctx_ptr()->g_the_gpu_config), this); - GPGPUsim_ctx_ptr()->g_stream_manager = new stream_manager((GPGPUsim_ctx_ptr()->g_the_gpu),g_cuda_launch_blocking); + GPGPUsim_ctx_ptr()->g_stream_manager = new stream_manager((GPGPUsim_ctx_ptr()->g_the_gpu), func_sim->g_cuda_launch_blocking); GPGPUsim_ctx_ptr()->g_simulation_starttime = time((time_t *)NULL); -- cgit v1.3 From 0a0dbfe33d434d4e3c6988a345b9e8a9779eddc1 Mon Sep 17 00:00:00 2001 From: Mengchi Zhang Date: Wed, 10 Jul 2019 00:25:15 -0400 Subject: Move g_inst_classification_stat and g_inst_op_classification_stat Signed-off-by: Mengchi Zhang --- src/cuda-sim/cuda-sim.cc | 8 +++----- src/cuda-sim/cuda-sim.h | 6 ++++-- src/gpgpu-sim/gpu-sim.cc | 4 ++-- 3 files changed, 9 insertions(+), 9 deletions(-) (limited to 'src/cuda-sim') diff --git a/src/cuda-sim/cuda-sim.cc b/src/cuda-sim/cuda-sim.cc index e9508ee..7e8aab9 100644 --- a/src/cuda-sim/cuda-sim.cc +++ b/src/cuda-sim/cuda-sim.cc @@ -54,8 +54,6 @@ typedef void * yyscan_t; #include "../../libcuda/gpgpu_context.h" int gpgpu_ptx_instruction_classification; -void ** g_inst_classification_stat = NULL; -void ** g_inst_op_classification_stat= NULL; int g_debug_execution = 0; int g_debug_thread_uid = 0; addr_t g_debug_pc = 0xBEEF1518; @@ -1701,9 +1699,9 @@ void ptx_thread_info::ptx_exec_inst( warp_inst_t &inst, unsigned lane_id) space_type = 0 ; break; } - StatAddSample( g_inst_classification_stat[m_gpu->gpgpu_ctx->func_sim->g_ptx_kernel_count], op_classification); - if (space_type) StatAddSample( g_inst_classification_stat[m_gpu->gpgpu_ctx->func_sim->g_ptx_kernel_count], ( int )space_type); - StatAddSample( g_inst_op_classification_stat[m_gpu->gpgpu_ctx->func_sim->g_ptx_kernel_count], (int) pI->get_opcode() ); + StatAddSample( m_gpu->gpgpu_ctx->func_sim->g_inst_classification_stat[m_gpu->gpgpu_ctx->func_sim->g_ptx_kernel_count], op_classification); + if (space_type) StatAddSample( m_gpu->gpgpu_ctx->func_sim->g_inst_classification_stat[m_gpu->gpgpu_ctx->func_sim->g_ptx_kernel_count], ( int )space_type); + StatAddSample( m_gpu->gpgpu_ctx->func_sim->g_inst_op_classification_stat[m_gpu->gpgpu_ctx->func_sim->g_ptx_kernel_count], (int) pI->get_opcode() ); } if ( (m_gpu->gpgpu_ctx->func_sim->g_ptx_sim_num_insn % 100000) == 0 ) { dim3 ctaid = get_ctaid(); diff --git a/src/cuda-sim/cuda-sim.h b/src/cuda-sim/cuda-sim.h index 6177986..628e434 100644 --- a/src/cuda-sim/cuda-sim.h +++ b/src/cuda-sim/cuda-sim.h @@ -44,8 +44,6 @@ class symbol_table; extern const char *g_gpgpusim_version_string; extern int g_debug_execution; extern int g_debug_thread_uid; -extern void ** g_inst_classification_stat; -extern void ** g_inst_op_classification_stat; 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 ); @@ -134,6 +132,8 @@ class cuda_sim { g_ptx_kernel_count = -1; // used for classification stat collection purposes gpgpu_param_num_shaders = 0; g_cuda_launch_blocking = false; + g_inst_classification_stat = NULL; + g_inst_op_classification_stat= NULL; gpgpu_ctx = ctx; } //global variables @@ -159,6 +159,8 @@ class cuda_sim { unsigned gpgpu_param_num_shaders; class std::map g_rpts; bool g_cuda_launch_blocking; + void ** g_inst_classification_stat; + void ** g_inst_op_classification_stat; // backward pointer class gpgpu_context* gpgpu_ctx; //global functions diff --git a/src/gpgpu-sim/gpu-sim.cc b/src/gpgpu-sim/gpu-sim.cc index 30e0aa5..30411f1 100644 --- a/src/gpgpu-sim/gpu-sim.cc +++ b/src/gpgpu-sim/gpu-sim.cc @@ -1201,8 +1201,8 @@ void gpgpu_sim::gpu_print_stat() insn_warp_occ_print(stdout); } if ( gpgpu_ptx_instruction_classification ) { - StatDisp( g_inst_classification_stat[gpgpu_ctx->func_sim->g_ptx_kernel_count]); - StatDisp( g_inst_op_classification_stat[gpgpu_ctx->func_sim->g_ptx_kernel_count]); + StatDisp( gpgpu_ctx->func_sim->g_inst_classification_stat[gpgpu_ctx->func_sim->g_ptx_kernel_count]); + StatDisp( gpgpu_ctx->func_sim->g_inst_op_classification_stat[gpgpu_ctx->func_sim->g_ptx_kernel_count]); } #ifdef GPGPUSIM_POWER_MODEL -- 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') 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') 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 From 57f8e9bd4e73757f4026e8b257fb625d465e0271 Mon Sep 17 00:00:00 2001 From: Mengchi Zhang Date: Wed, 10 Jul 2019 15:09:34 -0400 Subject: Move gpgpu_ptx_instruction_classification Signed-off-by: Mengchi Zhang --- src/cuda-sim/cuda-sim.cc | 3 +-- src/cuda-sim/cuda-sim.h | 1 + src/debug.h | 2 -- src/gpgpu-sim/gpu-sim.cc | 6 +++--- 4 files changed, 5 insertions(+), 7 deletions(-) (limited to 'src/cuda-sim') diff --git a/src/cuda-sim/cuda-sim.cc b/src/cuda-sim/cuda-sim.cc index fb9bc9e..b370400 100644 --- a/src/cuda-sim/cuda-sim.cc +++ b/src/cuda-sim/cuda-sim.cc @@ -53,7 +53,6 @@ typedef void * yyscan_t; #include "cuda_device_runtime.h" #include "../../libcuda/gpgpu_context.h" -int gpgpu_ptx_instruction_classification; int g_debug_execution = 0; int g_debug_thread_uid = 0; addr_t g_debug_pc = 0xBEEF1518; @@ -1681,7 +1680,7 @@ void ptx_thread_info::ptx_exec_inst( warp_inst_t &inst, unsigned lane_id) if(!(this->m_functionalSimulationMode)) ptx_file_line_stats_add_exec_count(pI); - if ( gpgpu_ptx_instruction_classification ) { + if ( m_gpu->gpgpu_ctx->func_sim->gpgpu_ptx_instruction_classification ) { m_gpu->gpgpu_ctx->func_sim->init_inst_classification_stat(); unsigned space_type=0; switch ( pI->get_space().get_type() ) { diff --git a/src/cuda-sim/cuda-sim.h b/src/cuda-sim/cuda-sim.h index c578524..89f67cf 100644 --- a/src/cuda-sim/cuda-sim.h +++ b/src/cuda-sim/cuda-sim.h @@ -161,6 +161,7 @@ class cuda_sim { std::set g_globals; std::set g_constants; std::map g_pc_to_finfo; + int gpgpu_ptx_instruction_classification; // backward pointer class gpgpu_context* gpgpu_ctx; //global functions diff --git a/src/debug.h b/src/debug.h index 1277494..1005bd5 100644 --- a/src/debug.h +++ b/src/debug.h @@ -83,8 +83,6 @@ private: unsigned m_value; }; -extern int gpgpu_ptx_instruction_classification ; - class ptx_thread_info; class ptx_instruction; bool thread_at_brkpt( ptx_thread_info *thd_info, const class brk_pt &b ); diff --git a/src/gpgpu-sim/gpu-sim.cc b/src/gpgpu-sim/gpu-sim.cc index 30411f1..93f041a 100644 --- a/src/gpgpu-sim/gpu-sim.cc +++ b/src/gpgpu-sim/gpu-sim.cc @@ -499,8 +499,8 @@ void gpgpu_sim_config::reg_options(option_parser_t opp) option_parser_register(opp, "-gpgpu_deadlock_detect", OPT_BOOL, &gpu_deadlock_detect, "Stop the simulation at deadlock (1=on (default), 0=off)", "1"); - option_parser_register(opp, "-gpgpu_ptx_instruction_classification", OPT_INT32, - &gpgpu_ptx_instruction_classification, + option_parser_register(opp, "-gpgpu_ptx_instruction_classification", OPT_INT32, + &(gpgpu_ctx->func_sim->gpgpu_ptx_instruction_classification), "if enabled will classify ptx instruction types per kernel (Max 255 kernels now)", "0"); option_parser_register(opp, "-gpgpu_ptx_sim_mode", OPT_INT32, &(gpgpu_ctx->func_sim->g_ptx_sim_mode), @@ -1200,7 +1200,7 @@ void gpgpu_sim::gpu_print_stat() spill_log_to_file (stdout, 1, gpu_sim_cycle); insn_warp_occ_print(stdout); } - if ( gpgpu_ptx_instruction_classification ) { + if ( gpgpu_ctx->func_sim->gpgpu_ptx_instruction_classification ) { StatDisp( gpgpu_ctx->func_sim->g_inst_classification_stat[gpgpu_ctx->func_sim->g_ptx_kernel_count]); StatDisp( gpgpu_ctx->func_sim->g_inst_op_classification_stat[gpgpu_ctx->func_sim->g_ptx_kernel_count]); } -- cgit v1.3 From e83a9c978ebd08eed616eb6fdbdef88160232076 Mon Sep 17 00:00:00 2001 From: Mengchi Zhang Date: Wed, 10 Jul 2019 16:09:54 -0400 Subject: Move cdp_latency Signed-off-by: Mengchi Zhang --- src/cuda-sim/cuda-sim.cc | 8 +++++--- src/cuda-sim/cuda-sim.h | 1 + src/gpgpu-sim/shader.cc | 7 +++---- 3 files changed, 9 insertions(+), 7 deletions(-) (limited to 'src/cuda-sim') diff --git a/src/cuda-sim/cuda-sim.cc b/src/cuda-sim/cuda-sim.cc index b370400..bd96f7e 100644 --- a/src/cuda-sim/cuda-sim.cc +++ b/src/cuda-sim/cuda-sim.cc @@ -58,7 +58,6 @@ int g_debug_thread_uid = 0; addr_t g_debug_pc = 0xBEEF1518; // Output debug information to file options -unsigned cdp_latency[5]; void cuda_sim::ptx_opcocde_latency_options (option_parser_t opp) { option_parser_register(opp, "-ptx_opcode_latency_int", OPT_CSTR, &opcode_latency_int, @@ -680,8 +679,11 @@ void ptx_instruction::set_opcode_and_latency() sscanf(gpgpu_ctx->func_sim->opcode_initiation_tensor, "%u", &tensor_init); sscanf(gpgpu_ctx->func_sim->cdp_latency_str, "%u,%u,%u,%u,%u", - &cdp_latency[0],&cdp_latency[1],&cdp_latency[2], - &cdp_latency[3],&cdp_latency[4]); + &gpgpu_ctx->func_sim->cdp_latency[0], + &gpgpu_ctx->func_sim->cdp_latency[1], + &gpgpu_ctx->func_sim->cdp_latency[2], + &gpgpu_ctx->func_sim->cdp_latency[3], + &gpgpu_ctx->func_sim->cdp_latency[4]); if(!m_operands.empty()){ std::vector::iterator it; diff --git a/src/cuda-sim/cuda-sim.h b/src/cuda-sim/cuda-sim.h index 89f67cf..0b6e84f 100644 --- a/src/cuda-sim/cuda-sim.h +++ b/src/cuda-sim/cuda-sim.h @@ -163,6 +163,7 @@ class cuda_sim { std::map g_pc_to_finfo; int gpgpu_ptx_instruction_classification; // backward pointer + unsigned cdp_latency[5]; class gpgpu_context* gpgpu_ctx; //global functions void ptx_opcocde_latency_options (option_parser_t opp); diff --git a/src/gpgpu-sim/shader.cc b/src/gpgpu-sim/shader.cc index 6cd6d8f..f380560 100644 --- a/src/gpgpu-sim/shader.cc +++ b/src/gpgpu-sim/shader.cc @@ -1084,12 +1084,11 @@ void scheduler_unit::cycle() if(pI->m_is_cdp && !warp(warp_id).m_cdp_dummy) { assert(warp(warp_id).m_cdp_latency == 0); - extern unsigned cdp_latency[5]; if(pI->m_is_cdp == 1) - warp(warp_id).m_cdp_latency = cdp_latency[pI->m_is_cdp - 1]; + warp(warp_id).m_cdp_latency = m_shader->m_config->gpgpu_ctx->func_sim->cdp_latency[pI->m_is_cdp - 1]; else //cudaLaunchDeviceV2 and cudaGetParameterBufferV2 - warp(warp_id).m_cdp_latency = cdp_latency[pI->m_is_cdp - 1] - + cdp_latency[pI->m_is_cdp] * active_mask.count(); + warp(warp_id).m_cdp_latency = m_shader->m_config->gpgpu_ctx->func_sim->cdp_latency[pI->m_is_cdp - 1] + + m_shader->m_config->gpgpu_ctx->func_sim->cdp_latency[pI->m_is_cdp] * active_mask.count(); warp(warp_id).m_cdp_dummy = true; break; } -- cgit v1.3 From 44a2f74ff840bcfced0f21357d79bdad23da9227 Mon Sep 17 00:00:00 2001 From: Mengchi Zhang Date: Wed, 10 Jul 2019 16:13:37 -0400 Subject: Move g_assemble_code_next_pc Signed-off-by: Mengchi Zhang --- src/cuda-sim/cuda-sim.cc | 5 ++--- src/cuda-sim/cuda-sim.h | 4 +++- 2 files changed, 5 insertions(+), 4 deletions(-) (limited to 'src/cuda-sim') diff --git a/src/cuda-sim/cuda-sim.cc b/src/cuda-sim/cuda-sim.cc index bd96f7e..4564919 100644 --- a/src/cuda-sim/cuda-sim.cc +++ b/src/cuda-sim/cuda-sim.cc @@ -217,7 +217,6 @@ void gpgpu_t::gpgpu_ptx_sim_unbindTexture(const struct textureReference* texref) m_NameToTextureInfo.erase(texname); } -unsigned g_assemble_code_next_pc=0; std::vector function_info::s_g_pc_to_insn; #define MAX_INST_SIZE 8 /*bytes*/ @@ -237,7 +236,7 @@ void function_info::ptx_assemble() fflush(stdout); std::list::iterator i; - addr_t PC = g_assemble_code_next_pc; // globally unique address (across functions) + addr_t PC = gpgpu_ctx->func_sim->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); @@ -269,7 +268,7 @@ void function_info::ptx_assemble() PC += pI->inst_size(); } } - g_assemble_code_next_pc=PC; + gpgpu_ctx->func_sim->g_assemble_code_next_pc=PC; 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 || pI->get_opcode() == CALLP_OP) { diff --git a/src/cuda-sim/cuda-sim.h b/src/cuda-sim/cuda-sim.h index 0b6e84f..99fd92c 100644 --- a/src/cuda-sim/cuda-sim.h +++ b/src/cuda-sim/cuda-sim.h @@ -131,6 +131,7 @@ class cuda_sim { g_cuda_launch_blocking = false; g_inst_classification_stat = NULL; g_inst_op_classification_stat= NULL; + g_assemble_code_next_pc=0; gpgpu_ctx = ctx; } //global variables @@ -162,8 +163,9 @@ class cuda_sim { std::set g_constants; std::map g_pc_to_finfo; int gpgpu_ptx_instruction_classification; - // backward pointer unsigned cdp_latency[5]; + unsigned g_assemble_code_next_pc; + // backward pointer class gpgpu_context* gpgpu_ctx; //global functions void ptx_opcocde_latency_options (option_parser_t opp); -- cgit v1.3 From cf13f05351636d178ee30c3a57872c33ef5a4e47 Mon Sep 17 00:00:00 2001 From: Mengchi Zhang Date: Wed, 10 Jul 2019 16:19:26 -0400 Subject: Move g_debug_thread_uid Signed-off-by: Mengchi Zhang --- src/cuda-sim/cuda-sim.cc | 12 ++++++------ src/cuda-sim/cuda-sim.h | 4 +++- 2 files changed, 9 insertions(+), 7 deletions(-) (limited to 'src/cuda-sim') diff --git a/src/cuda-sim/cuda-sim.cc b/src/cuda-sim/cuda-sim.cc index 4564919..7ff869c 100644 --- a/src/cuda-sim/cuda-sim.cc +++ b/src/cuda-sim/cuda-sim.cc @@ -54,7 +54,6 @@ typedef void * yyscan_t; #include "../../libcuda/gpgpu_context.h" int g_debug_execution = 0; -int g_debug_thread_uid = 0; addr_t g_debug_pc = 0xBEEF1518; // Output debug information to file options @@ -1455,7 +1454,7 @@ void function_info::ptx_jit_config(std::map mallocPt } template -bool ptx_debug_exec_dump_cond(int thd_uid, addr_t pc) +bool cuda_sim::ptx_debug_exec_dump_cond(int thd_uid, addr_t pc) { if (g_debug_execution >= activate_level) { // check each type of debug dump constraint to filter out dumps @@ -1541,7 +1540,8 @@ void ptx_thread_info::ptx_exec_inst( warp_inst_t &inst, unsigned lane_id) } if ( g_debug_execution >= 6 || m_gpu->get_config().get_ptx_inst_debug_to_file()) { - if ( (g_debug_thread_uid==0) || (get_uid() == (unsigned)g_debug_thread_uid) ) { + if ( (m_gpu->gpgpu_ctx->func_sim->g_debug_thread_uid==0) + || (get_uid() == (unsigned)(m_gpu->gpgpu_ctx->func_sim->g_debug_thread_uid)) ) { clear_modifiedregs(); enable_debug_trace(); @@ -1613,7 +1613,7 @@ void ptx_thread_info::ptx_exec_inst( warp_inst_t &inst, unsigned lane_id) fflush(m_gpu->get_ptx_inst_debug_file()); } - if ( ptx_debug_exec_dump_cond<5>(get_uid(), pc) ) { + if ( m_gpu->gpgpu_ctx->func_sim->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", @@ -1667,11 +1667,11 @@ void ptx_thread_info::ptx_exec_inst( warp_inst_t &inst, unsigned lane_id) } if ( g_debug_execution >= 6 ) { - if ( ptx_debug_exec_dump_cond<6>(get_uid(), pc) ) + if ( m_gpu->gpgpu_ctx->func_sim->ptx_debug_exec_dump_cond<6>(get_uid(), pc) ) dump_modifiedregs(stdout); } if ( g_debug_execution >= 10 ) { - if ( ptx_debug_exec_dump_cond<10>(get_uid(), pc) ) + if ( m_gpu->gpgpu_ctx->func_sim->ptx_debug_exec_dump_cond<10>(get_uid(), pc) ) dump_regs(stdout); } update_pc(); diff --git a/src/cuda-sim/cuda-sim.h b/src/cuda-sim/cuda-sim.h index 99fd92c..e4d34fe 100644 --- a/src/cuda-sim/cuda-sim.h +++ b/src/cuda-sim/cuda-sim.h @@ -43,7 +43,6 @@ class symbol_table; extern const char *g_gpgpusim_version_string; extern int g_debug_execution; -extern int g_debug_thread_uid; extern void print_splash(); @@ -132,6 +131,7 @@ class cuda_sim { g_inst_classification_stat = NULL; g_inst_op_classification_stat= NULL; g_assemble_code_next_pc=0; + g_debug_thread_uid = 0; gpgpu_ctx = ctx; } //global variables @@ -165,6 +165,7 @@ class cuda_sim { int gpgpu_ptx_instruction_classification; unsigned cdp_latency[5]; unsigned g_assemble_code_next_pc; + int g_debug_thread_uid; // backward pointer class gpgpu_context* gpgpu_ctx; //global functions @@ -186,6 +187,7 @@ class cuda_sim { 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 ); + template bool ptx_debug_exec_dump_cond(int thd_uid, addr_t pc); }; #endif -- cgit v1.3