From 98bb2046a70715c8599561ca78425a7ff99ce53b Mon Sep 17 00:00:00 2001 From: Mengchi Zhang Date: Mon, 8 Jul 2019 14:21:08 -0400 Subject: Move cp_count Signed-off-by: Mengchi Zhang --- src/cuda-sim/cuda-sim.h | 4 +++- 1 file changed, 3 insertions(+), 1 deletion(-) (limited to 'src/cuda-sim/cuda-sim.h') diff --git a/src/cuda-sim/cuda-sim.h b/src/cuda-sim/cuda-sim.h index 96d34f5..4c8d953 100644 --- a/src/cuda-sim/cuda-sim.h +++ b/src/cuda-sim/cuda-sim.h @@ -55,7 +55,6 @@ extern class kernel_info_t *gpgpu_opencl_ptx_sim_init_grid(class function_info * struct dim3 gridDim, struct dim3 blockDim, class gpgpu_t *gpu ); -extern void gpgpu_cuda_ptx_sim_main_func( kernel_info_t &kernel, bool openCL = false ); extern void print_splash(); extern void gpgpu_ptx_sim_register_const_variable(void*, const char *deviceName, size_t size ); extern void gpgpu_ptx_sim_register_global_variable(void *hostVar, const char *deviceName, size_t size ); @@ -137,8 +136,11 @@ class cuda_sim { public: //global variables char *opcode_latency_int; + int cp_count; //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 ); + int gpgpu_opencl_ptx_sim_main_func( kernel_info_t *grid ); }; #endif -- cgit v1.3 From c643efb70a02fcbf4a1326cf96d625cf2e51a0b5 Mon Sep 17 00:00:00 2001 From: Mengchi Zhang Date: Mon, 8 Jul 2019 14:34:47 -0400 Subject: Move cp_cta_resume Signed-off-by: Mengchi Zhang --- src/cuda-sim/cuda-sim.cc | 7 +++---- src/cuda-sim/cuda-sim.h | 1 + 2 files changed, 4 insertions(+), 4 deletions(-) (limited to 'src/cuda-sim/cuda-sim.h') diff --git a/src/cuda-sim/cuda-sim.cc b/src/cuda-sim/cuda-sim.cc index 12f38af..cba256b 100644 --- a/src/cuda-sim/cuda-sim.cc +++ b/src/cuda-sim/cuda-sim.cc @@ -61,7 +61,6 @@ int g_debug_execution = 0; int g_debug_thread_uid = 0; addr_t g_debug_pc = 0xBEEF1518; // Output debug information to file options -int cp_cta_resume; unsigned g_ptx_sim_num_insn = 0; unsigned gpgpu_param_num_shaders = 0; @@ -2271,7 +2270,7 @@ void functionalCoreSim::initializeCTA(unsigned ctaid_cp) assert(m_thread[i]!=NULL && !m_thread[i]->is_done()); char fname[2048]; snprintf(fname,2048,"checkpoint_files/thread_%d_0_reg.txt",i ); - if(cp_cta_resume==1) + if(m_gpu->gpgpu_ctx->func_sim->cp_cta_resume==1) m_thread[i]->resume_reg_thread(fname,symtab); ctaLiveThreads++; } @@ -2295,7 +2294,7 @@ void functionalCoreSim::createWarp(unsigned warpId) char fname[2048]; snprintf(fname,2048,"checkpoint_files/warp_%d_0_simt.txt",warpId ); - if(cp_cta_resume==1) + if(m_gpu->gpgpu_ctx->func_sim->cp_cta_resume==1) { unsigned pc,rpc; m_simt_stack[warpId]->resume(fname); @@ -2312,7 +2311,7 @@ void functionalCoreSim::createWarp(unsigned warpId) void functionalCoreSim::execute(int inst_count, unsigned ctaid_cp) { m_gpu->gpgpu_ctx->func_sim->cp_count= m_gpu->checkpoint_insn_Y; - cp_cta_resume= m_gpu->checkpoint_CTA_t; + m_gpu->gpgpu_ctx->func_sim->cp_cta_resume= m_gpu->checkpoint_CTA_t; initializeCTA(ctaid_cp); int count=0; diff --git a/src/cuda-sim/cuda-sim.h b/src/cuda-sim/cuda-sim.h index 4c8d953..6b40300 100644 --- a/src/cuda-sim/cuda-sim.h +++ b/src/cuda-sim/cuda-sim.h @@ -137,6 +137,7 @@ class cuda_sim { //global variables char *opcode_latency_int; int cp_count; + int cp_cta_resume; //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 ); -- cgit v1.3 From 7c13f6c7cc7fd598b268810c903983b79606f3ca Mon Sep 17 00:00:00 2001 From: Mengchi Zhang Date: Mon, 8 Jul 2019 14:53:27 -0400 Subject: Move g_ptxinfo_error_detected Signed-off-by: Mengchi Zhang --- src/cuda-sim/cuda-sim.cc | 2 -- src/cuda-sim/cuda-sim.h | 1 + src/cuda-sim/ptxinfo.l | 4 +--- 3 files changed, 2 insertions(+), 5 deletions(-) (limited to 'src/cuda-sim/cuda-sim.h') diff --git a/src/cuda-sim/cuda-sim.cc b/src/cuda-sim/cuda-sim.cc index cba256b..a143aa5 100644 --- a/src/cuda-sim/cuda-sim.cc +++ b/src/cuda-sim/cuda-sim.cc @@ -2404,8 +2404,6 @@ unsigned translate_pc_to_ptxlineno(unsigned pc) extern std::map get_duplicate(); -int g_ptxinfo_error_detected; - static char *g_ptxinfo_kname = NULL; static struct gpgpu_ptx_sim_info g_ptxinfo; static std::map g_duplicate; diff --git a/src/cuda-sim/cuda-sim.h b/src/cuda-sim/cuda-sim.h index 6b40300..76450dc 100644 --- a/src/cuda-sim/cuda-sim.h +++ b/src/cuda-sim/cuda-sim.h @@ -138,6 +138,7 @@ class cuda_sim { char *opcode_latency_int; int cp_count; int cp_cta_resume; + int g_ptxinfo_error_detected; //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 ); diff --git a/src/cuda-sim/ptxinfo.l b/src/cuda-sim/ptxinfo.l index 3a152b0..51371e3 100644 --- a/src/cuda-sim/ptxinfo.l +++ b/src/cuda-sim/ptxinfo.l @@ -88,13 +88,11 @@ OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. %% -extern int g_ptxinfo_error_detected; - int ptxinfo_error(yyscan_t yyscanner, ptxinfo_data* ptxinfo, const char* msg) { struct yyguts_t * yyg = (struct yyguts_t*)yyscanner; int i; - g_ptxinfo_error_detected = 1; + ptxinfo->gpgpu_ctx->func_sim->g_ptxinfo_error_detected = 1; fflush(stdout); printf("GPGPU-Sim: ERROR while parsing output of ptxas (used to capture resource usage information)\n"); if( msg != NULL ) -- cgit v1.3 From 8f6668941cf2728dba9700e45f11f61401a1fcf4 Mon Sep 17 00:00:00 2001 From: Mengchi Zhang Date: Mon, 8 Jul 2019 15:45:05 -0400 Subject: Move g_ptx_sim_num_insn Signed-off-by: Mengchi Zhang --- src/abstract_hardware_model.cc | 3 ++- src/abstract_hardware_model.h | 6 +++++- src/cuda-sim/cuda-sim.cc | 9 ++++----- src/cuda-sim/cuda-sim.h | 4 ++++ src/gpgpu-sim/gpu-sim.cc | 2 +- src/gpgpu-sim/gpu-sim.h | 2 ++ 6 files changed, 18 insertions(+), 8 deletions(-) (limited to 'src/cuda-sim/cuda-sim.h') diff --git a/src/abstract_hardware_model.cc b/src/abstract_hardware_model.cc index 7755477..248e7a5 100644 --- a/src/abstract_hardware_model.cc +++ b/src/abstract_hardware_model.cc @@ -173,9 +173,10 @@ void gpgpu_functional_sim_config::ptx_set_tex_cache_linesize(unsigned linesize) m_texcache_linesize = linesize; } -gpgpu_t::gpgpu_t( const gpgpu_functional_sim_config &config ) +gpgpu_t::gpgpu_t( const gpgpu_functional_sim_config &config, gpgpu_context* ctx ) : m_function_model_config(config) { + gpgpu_ctx = ctx; m_global_mem = new memory_space_impl<8192>("global",64*1024); m_tex_mem = new memory_space_impl<8192>("tex",64*1024); diff --git a/src/abstract_hardware_model.h b/src/abstract_hardware_model.h index 68cb693..da29a11 100644 --- a/src/abstract_hardware_model.h +++ b/src/abstract_hardware_model.h @@ -31,6 +31,8 @@ // Forward declarations class gpgpu_sim; class kernel_info_t; +class gpgpu_context; + //Set a hard limit of 32 CTAs per shader [cuda only has 8] #define MAX_CTA_PER_SHADER 32 @@ -529,7 +531,9 @@ private: class gpgpu_t { public: - gpgpu_t( const gpgpu_functional_sim_config &config ); + gpgpu_t( const gpgpu_functional_sim_config &config, gpgpu_context* ctx ); + // backward pointer + class gpgpu_context* gpgpu_ctx; int checkpoint_option; int checkpoint_kernel; int checkpoint_CTA; diff --git a/src/cuda-sim/cuda-sim.cc b/src/cuda-sim/cuda-sim.cc index a143aa5..939358b 100644 --- a/src/cuda-sim/cuda-sim.cc +++ b/src/cuda-sim/cuda-sim.cc @@ -62,7 +62,6 @@ int g_debug_thread_uid = 0; addr_t g_debug_pc = 0xBEEF1518; // Output debug information to file options -unsigned g_ptx_sim_num_insn = 0; unsigned gpgpu_param_num_shaders = 0; char *opcode_latency_fp, *opcode_latency_dp,*opcode_latency_sfu,*opcode_latency_tensor; @@ -1629,7 +1628,7 @@ void ptx_thread_info::ptx_exec_inst( warp_inst_t &inst, unsigned lane_id) dim3 ctaid = get_ctaid(); dim3 tid = get_tid(); printf("%u [thd=%u][i=%u] : ctaid=(%u,%u,%u) tid=(%u,%u,%u) icount=%u [pc=%u] (%s:%u - %s) [0x%llx]\n", - g_ptx_sim_num_insn, + m_gpu->gpgpu_ctx->func_sim->g_ptx_sim_num_insn, get_uid(), pI->uid(), ctaid.x,ctaid.y,ctaid.z,tid.x,tid.y,tid.z, get_icount(), @@ -1687,7 +1686,7 @@ void ptx_thread_info::ptx_exec_inst( warp_inst_t &inst, unsigned lane_id) dump_regs(stdout); } update_pc(); - g_ptx_sim_num_insn++; + m_gpu->gpgpu_ctx->func_sim->g_ptx_sim_num_insn++; //not using it with functional simulation mode if(!(this->m_functionalSimulationMode)) @@ -1714,11 +1713,11 @@ void ptx_thread_info::ptx_exec_inst( warp_inst_t &inst, unsigned lane_id) if (space_type) StatAddSample( g_inst_classification_stat[g_ptx_kernel_count], ( int )space_type); StatAddSample( g_inst_op_classification_stat[g_ptx_kernel_count], (int) pI->get_opcode() ); } - if ( (g_ptx_sim_num_insn % 100000) == 0 ) { + if ( (m_gpu->gpgpu_ctx->func_sim->g_ptx_sim_num_insn % 100000) == 0 ) { dim3 ctaid = get_ctaid(); dim3 tid = get_tid(); DPRINTF(LIVENESS, "GPGPU-Sim PTX: %u instructions simulated : ctaid=(%u,%u,%u) tid=(%u,%u,%u)\n", - g_ptx_sim_num_insn, ctaid.x,ctaid.y,ctaid.z,tid.x,tid.y,tid.z ); + m_gpu->gpgpu_ctx->func_sim->g_ptx_sim_num_insn, ctaid.x,ctaid.y,ctaid.z,tid.x,tid.y,tid.z ); fflush(stdout); } diff --git a/src/cuda-sim/cuda-sim.h b/src/cuda-sim/cuda-sim.h index 76450dc..4566dc2 100644 --- a/src/cuda-sim/cuda-sim.h +++ b/src/cuda-sim/cuda-sim.h @@ -134,11 +134,15 @@ struct gpgpu_ptx_sim_info get_ptxinfo(); class cuda_sim { public: + cuda_sim() { + g_ptx_sim_num_insn = 0; + } //global variables char *opcode_latency_int; int cp_count; int cp_cta_resume; int g_ptxinfo_error_detected; + unsigned g_ptx_sim_num_insn; //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 ); diff --git a/src/gpgpu-sim/gpu-sim.cc b/src/gpgpu-sim/gpu-sim.cc index 39acdd9..4f9ccbf 100644 --- a/src/gpgpu-sim/gpu-sim.cc +++ b/src/gpgpu-sim/gpu-sim.cc @@ -699,7 +699,7 @@ void gpgpu_sim::stop_all_running_kernels(){ } gpgpu_sim::gpgpu_sim( const gpgpu_sim_config &config, gpgpu_context* ctx ) - : gpgpu_t(config), m_config(config) + : gpgpu_t(config, ctx), m_config(config) { gpgpu_ctx = ctx; m_shader_config = &m_config.m_shader_config; diff --git a/src/gpgpu-sim/gpu-sim.h b/src/gpgpu-sim/gpu-sim.h index b47ab16..119b934 100644 --- a/src/gpgpu-sim/gpu-sim.h +++ b/src/gpgpu-sim/gpu-sim.h @@ -62,6 +62,8 @@ #define SAMPLELOG 222 #define DUMPLOG 333 +class gpgpu_context; + extern tr1_hash_map address_random_interleaving; enum dram_ctrl_t { -- cgit v1.3 From 39910217c8137f5ff964d2f991a688020e2d5c8c Mon Sep 17 00:00:00 2001 From: Mengchi Zhang Date: Tue, 9 Jul 2019 09:53:55 -0400 Subject: Move opcode_latency_* Signed-off-by: Mengchi Zhang --- src/cuda-sim/cuda-sim.cc | 9 ++++----- src/cuda-sim/cuda-sim.h | 6 ++++-- src/gpgpu-sim/shader.cc | 8 ++++---- 3 files changed, 12 insertions(+), 11 deletions(-) (limited to 'src/cuda-sim/cuda-sim.h') diff --git a/src/cuda-sim/cuda-sim.cc b/src/cuda-sim/cuda-sim.cc index 939358b..f8c21bc 100644 --- a/src/cuda-sim/cuda-sim.cc +++ b/src/cuda-sim/cuda-sim.cc @@ -64,7 +64,6 @@ addr_t g_debug_pc = 0xBEEF1518; unsigned gpgpu_param_num_shaders = 0; -char *opcode_latency_fp, *opcode_latency_dp,*opcode_latency_sfu,*opcode_latency_tensor; char *opcode_initiation_int, *opcode_initiation_fp, *opcode_initiation_dp,*opcode_initiation_sfu,*opcode_initiation_tensor; char *cdp_latency_str; unsigned cdp_latency[5]; @@ -668,15 +667,15 @@ void ptx_instruction::set_opcode_and_latency() sscanf(gpgpu_ctx->func_sim->opcode_latency_int, "%u,%u,%u,%u,%u", &int_latency[0],&int_latency[1],&int_latency[2], &int_latency[3],&int_latency[4]); - sscanf(opcode_latency_fp, "%u,%u,%u,%u,%u", + sscanf(gpgpu_ctx->func_sim->opcode_latency_fp, "%u,%u,%u,%u,%u", &fp_latency[0],&fp_latency[1],&fp_latency[2], &fp_latency[3],&fp_latency[4]); - sscanf(opcode_latency_dp, "%u,%u,%u,%u,%u", + sscanf(gpgpu_ctx->func_sim->opcode_latency_dp, "%u,%u,%u,%u,%u", &dp_latency[0],&dp_latency[1],&dp_latency[2], &dp_latency[3],&dp_latency[4]); - sscanf(opcode_latency_sfu, "%u", + sscanf(gpgpu_ctx->func_sim->opcode_latency_sfu, "%u", &sfu_latency); - sscanf(opcode_latency_tensor, "%u", + sscanf(gpgpu_ctx->func_sim->opcode_latency_tensor, "%u", &tensor_latency); sscanf(opcode_initiation_int, "%u,%u,%u,%u,%u", &int_init[0],&int_init[1],&int_init[2], diff --git a/src/cuda-sim/cuda-sim.h b/src/cuda-sim/cuda-sim.h index 4566dc2..977e3fc 100644 --- a/src/cuda-sim/cuda-sim.h +++ b/src/cuda-sim/cuda-sim.h @@ -47,8 +47,6 @@ extern int g_debug_thread_uid; extern void ** g_inst_classification_stat; extern void ** g_inst_op_classification_stat; extern int g_ptx_kernel_count; // used for classification stat collection purposes -extern char *opcode_latency_fp, *opcode_latency_dp,*opcode_latency_sfu,*opcode_latency_tensor; - extern class kernel_info_t *gpgpu_opencl_ptx_sim_init_grid(class function_info *entry, gpgpu_ptx_sim_arg_list_t args, @@ -139,6 +137,10 @@ class cuda_sim { } //global variables char *opcode_latency_int; + char *opcode_latency_fp; + char *opcode_latency_dp; + char *opcode_latency_sfu; + char *opcode_latency_tensor; int cp_count; int cp_cta_resume; int g_ptxinfo_error_detected; diff --git a/src/gpgpu-sim/shader.cc b/src/gpgpu-sim/shader.cc index 4d12068..69790fc 100644 --- a/src/gpgpu-sim/shader.cc +++ b/src/gpgpu-sim/shader.cc @@ -3022,15 +3022,15 @@ void shader_core_config::set_pipeline_latency() { sscanf(gpgpu_ctx->func_sim->opcode_latency_int, "%u,%u,%u,%u,%u", &int_latency[0],&int_latency[1],&int_latency[2], &int_latency[3],&int_latency[4]); - sscanf(opcode_latency_fp, "%u,%u,%u,%u,%u", + sscanf(gpgpu_ctx->func_sim->opcode_latency_fp, "%u,%u,%u,%u,%u", &fp_latency[0],&fp_latency[1],&fp_latency[2], &fp_latency[3],&fp_latency[4]); - sscanf(opcode_latency_dp, "%u,%u,%u,%u,%u", + sscanf(gpgpu_ctx->func_sim->opcode_latency_dp, "%u,%u,%u,%u,%u", &dp_latency[0],&dp_latency[1],&dp_latency[2], &dp_latency[3],&dp_latency[4]); - sscanf(opcode_latency_sfu, "%u", + sscanf(gpgpu_ctx->func_sim->opcode_latency_sfu, "%u", &sfu_latency); - sscanf(opcode_latency_tensor, "%u", + sscanf(gpgpu_ctx->func_sim->opcode_latency_tensor, "%u", &tensor_latency); //all div operation are executed on sfu -- cgit v1.3 From 61ce0709ba3c6b5563d75293cdb2ce489b5be585 Mon Sep 17 00:00:00 2001 From: Mengchi Zhang Date: Tue, 9 Jul 2019 09:57:35 -0400 Subject: Move opcode_initiation_* Signed-off-by: Mengchi Zhang --- src/cuda-sim/cuda-sim.cc | 11 +++++------ src/cuda-sim/cuda-sim.h | 5 +++++ 2 files changed, 10 insertions(+), 6 deletions(-) (limited to 'src/cuda-sim/cuda-sim.h') diff --git a/src/cuda-sim/cuda-sim.cc b/src/cuda-sim/cuda-sim.cc index f8c21bc..c299b4f 100644 --- a/src/cuda-sim/cuda-sim.cc +++ b/src/cuda-sim/cuda-sim.cc @@ -64,7 +64,6 @@ addr_t g_debug_pc = 0xBEEF1518; unsigned gpgpu_param_num_shaders = 0; -char *opcode_initiation_int, *opcode_initiation_fp, *opcode_initiation_dp,*opcode_initiation_sfu,*opcode_initiation_tensor; char *cdp_latency_str; unsigned cdp_latency[5]; @@ -677,18 +676,18 @@ void ptx_instruction::set_opcode_and_latency() &sfu_latency); sscanf(gpgpu_ctx->func_sim->opcode_latency_tensor, "%u", &tensor_latency); - sscanf(opcode_initiation_int, "%u,%u,%u,%u,%u", + sscanf(gpgpu_ctx->func_sim->opcode_initiation_int, "%u,%u,%u,%u,%u", &int_init[0],&int_init[1],&int_init[2], &int_init[3],&int_init[4]); - sscanf(opcode_initiation_fp, "%u,%u,%u,%u,%u", + sscanf(gpgpu_ctx->func_sim->opcode_initiation_fp, "%u,%u,%u,%u,%u", &fp_init[0],&fp_init[1],&fp_init[2], &fp_init[3],&fp_init[4]); - sscanf(opcode_initiation_dp, "%u,%u,%u,%u,%u", + sscanf(gpgpu_ctx->func_sim->opcode_initiation_dp, "%u,%u,%u,%u,%u", &dp_init[0],&dp_init[1],&dp_init[2], &dp_init[3],&dp_init[4]); - sscanf(opcode_initiation_sfu, "%u", + sscanf(gpgpu_ctx->func_sim->opcode_initiation_sfu, "%u", &sfu_init); - sscanf(opcode_initiation_tensor, "%u", + sscanf(gpgpu_ctx->func_sim->opcode_initiation_tensor, "%u", &tensor_init); sscanf(cdp_latency_str, "%u,%u,%u,%u,%u", &cdp_latency[0],&cdp_latency[1],&cdp_latency[2], diff --git a/src/cuda-sim/cuda-sim.h b/src/cuda-sim/cuda-sim.h index 977e3fc..adb4448 100644 --- a/src/cuda-sim/cuda-sim.h +++ b/src/cuda-sim/cuda-sim.h @@ -141,6 +141,11 @@ class cuda_sim { char *opcode_latency_dp; char *opcode_latency_sfu; char *opcode_latency_tensor; + char *opcode_initiation_int; + char *opcode_initiation_fp; + char *opcode_initiation_dp; + char *opcode_initiation_sfu; + char *opcode_initiation_tensor; int cp_count; int cp_cta_resume; int g_ptxinfo_error_detected; -- cgit v1.3 From 5b015d4347c4945de5a6a7b0559e67e3824d76d6 Mon Sep 17 00:00:00 2001 From: Mengchi Zhang Date: Tue, 9 Jul 2019 10:27:35 -0400 Subject: Move some more vars Signed-off-by: Mengchi Zhang --- src/cuda-sim/cuda-sim.cc | 5 +---- src/cuda-sim/cuda-sim.h | 1 + src/cuda-sim/ptx_parser.h | 1 - 3 files changed, 2 insertions(+), 5 deletions(-) (limited to 'src/cuda-sim/cuda-sim.h') diff --git a/src/cuda-sim/cuda-sim.cc b/src/cuda-sim/cuda-sim.cc index c299b4f..e86395d 100644 --- a/src/cuda-sim/cuda-sim.cc +++ b/src/cuda-sim/cuda-sim.cc @@ -64,7 +64,6 @@ addr_t g_debug_pc = 0xBEEF1518; unsigned gpgpu_param_num_shaders = 0; -char *cdp_latency_str; unsigned cdp_latency[5]; void cuda_sim::ptx_opcocde_latency_options (option_parser_t opp) { @@ -689,7 +688,7 @@ void ptx_instruction::set_opcode_and_latency() &sfu_init); sscanf(gpgpu_ctx->func_sim->opcode_initiation_tensor, "%u", &tensor_init); - sscanf(cdp_latency_str, "%u,%u,%u,%u,%u", + 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]); @@ -2101,8 +2100,6 @@ void read_sim_environment_variables() } } -ptx_cta_info *g_func_cta_info = NULL; - #define MAX(a,b) (((a)>(b))?(a):(b)) unsigned max_cta (const struct gpgpu_ptx_sim_info *kernel_info, unsigned threads_per_cta, unsigned int warp_size, unsigned int n_thread_per_shader, unsigned int gpgpu_shmem_size, unsigned int gpgpu_shader_registers, unsigned int max_cta_per_core) diff --git a/src/cuda-sim/cuda-sim.h b/src/cuda-sim/cuda-sim.h index adb4448..16ee46e 100644 --- a/src/cuda-sim/cuda-sim.h +++ b/src/cuda-sim/cuda-sim.h @@ -150,6 +150,7 @@ class cuda_sim { int cp_cta_resume; int g_ptxinfo_error_detected; unsigned g_ptx_sim_num_insn; + char *cdp_latency_str; //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 ); diff --git a/src/cuda-sim/ptx_parser.h b/src/cuda-sim/ptx_parser.h index e09aab1..ef422e5 100644 --- a/src/cuda-sim/ptx_parser.h +++ b/src/cuda-sim/ptx_parser.h @@ -30,7 +30,6 @@ #include "../abstract_hardware_model.h" #include "ptx_ir.h" -extern int g_error_detected; #ifdef __cplusplus const class ptx_instruction *ptx_instruction_lookup( const char *filename, unsigned linenumber ); -- cgit v1.3 From cda7a145b9e28eff0f3e9ac8197c2b6215755fc8 Mon Sep 17 00:00:00 2001 From: Mengchi Zhang Date: Tue, 9 Jul 2019 14:32:24 -0400 Subject: Move g_ptx_kernel_count Signed-off-by: Mengchi Zhang --- libcuda/cuda_api_object.h | 5 ++++- libcuda/cuda_runtime_api.cc | 2 +- libcuda/gpgpu_context.h | 2 +- libopencl/opencl_runtime_api.cc | 2 +- src/cuda-sim/cuda-sim.cc | 13 ++++++------- src/cuda-sim/cuda-sim.h | 14 ++++++++------ src/gpgpu-sim/gpu-sim.cc | 4 ++-- 7 files changed, 23 insertions(+), 19 deletions(-) (limited to 'src/cuda-sim/cuda-sim.h') diff --git a/libcuda/cuda_api_object.h b/libcuda/cuda_api_object.h index 0054697..db5e6a4 100644 --- a/libcuda/cuda_api_object.h +++ b/libcuda/cuda_api_object.h @@ -169,9 +169,10 @@ private: class cuda_runtime_api { public: - cuda_runtime_api() { + cuda_runtime_api( gpgpu_context* ctx ) { g_glbmap = NULL; g_active_device = 0; //active gpu that runs the code + gpgpu_ctx = ctx; } // global list std::list cuobjdumpSectionList; @@ -187,6 +188,8 @@ class cuda_runtime_api { std::map pinned_memory_size; glbmap_entry_t* g_glbmap; int g_active_device; //active gpu that runs the code + // backward pointer + class gpgpu_context* gpgpu_ctx; // member function list void cuobjdumpInit(); void extract_code_using_cuobjdump(); diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc index 43c5e1f..43c8bae 100644 --- a/libcuda/cuda_runtime_api.cc +++ b/libcuda/cuda_runtime_api.cc @@ -3545,7 +3545,7 @@ kernel_info_t * cuda_runtime_api::gpgpu_cuda_ptx_sim_init_grid( const char *host } entry->finalize(result->get_param_memory()); - g_ptx_kernel_count++; + gpgpu_ctx->func_sim->g_ptx_kernel_count++; fflush(stdout); if(g_debug_execution >= 4){ diff --git a/libcuda/gpgpu_context.h b/libcuda/gpgpu_context.h index 2e21009..a2ae7b6 100644 --- a/libcuda/gpgpu_context.h +++ b/libcuda/gpgpu_context.h @@ -10,7 +10,7 @@ class gpgpu_context { public: gpgpu_context() { g_global_allfiles_symbol_table = NULL; - api = new cuda_runtime_api(); + api = new cuda_runtime_api(this); ptxinfo = new ptxinfo_data(this); ptx_parser = new ptx_recognizer(this); the_gpgpusim = new GPGPUsim_ctx(this); diff --git a/libopencl/opencl_runtime_api.cc b/libopencl/opencl_runtime_api.cc index d302ff8..0a6eb3e 100644 --- a/libopencl/opencl_runtime_api.cc +++ b/libopencl/opencl_runtime_api.cc @@ -956,7 +956,7 @@ clEnqueueNDRangeKernel(cl_command_queue command_queue, 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 ); } - kernel_info_t *grid = gpgpu_opencl_ptx_sim_init_grid(kernel->get_implementation(),params,GridDim,BlockDim,gpu); + kernel_info_t *grid = ctx->func_sim->gpgpu_opencl_ptx_sim_init_grid(kernel->get_implementation(),params,GridDim,BlockDim,gpu); if ( g_ptx_sim_mode ) ctx->func_sim->gpgpu_opencl_ptx_sim_main_func( grid ); else diff --git a/src/cuda-sim/cuda-sim.cc b/src/cuda-sim/cuda-sim.cc index e86395d..0ed125a 100644 --- a/src/cuda-sim/cuda-sim.cc +++ b/src/cuda-sim/cuda-sim.cc @@ -56,7 +56,6 @@ typedef void * yyscan_t; int gpgpu_ptx_instruction_classification; void ** g_inst_classification_stat = NULL; void ** g_inst_op_classification_stat= NULL; -int g_ptx_kernel_count = -1; // used for classification stat collection purposes int g_debug_execution = 0; int g_debug_thread_uid = 0; addr_t g_debug_pc = 0xBEEF1518; @@ -1480,7 +1479,7 @@ bool ptx_debug_exec_dump_cond(int thd_uid, addr_t pc) return false; } -void init_inst_classification_stat() +void cuda_sim::init_inst_classification_stat() { static std::set init; if( init.find(g_ptx_kernel_count) != init.end() ) @@ -1690,7 +1689,7 @@ void ptx_thread_info::ptx_exec_inst( warp_inst_t &inst, unsigned lane_id) ptx_file_line_stats_add_exec_count(pI); if ( gpgpu_ptx_instruction_classification ) { - init_inst_classification_stat(); + m_gpu->gpgpu_ctx->func_sim->init_inst_classification_stat(); unsigned space_type=0; switch ( pI->get_space().get_type() ) { case global_space: space_type = 10; break; @@ -1706,9 +1705,9 @@ void ptx_thread_info::ptx_exec_inst( warp_inst_t &inst, unsigned lane_id) space_type = 0 ; break; } - StatAddSample( g_inst_classification_stat[g_ptx_kernel_count], op_classification); - if (space_type) StatAddSample( g_inst_classification_stat[g_ptx_kernel_count], ( int )space_type); - StatAddSample( g_inst_op_classification_stat[g_ptx_kernel_count], (int) pI->get_opcode() ); + 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() ); } if ( (m_gpu->gpgpu_ctx->func_sim->g_ptx_sim_num_insn % 100000) == 0 ) { dim3 ctaid = get_ctaid(); @@ -1917,7 +1916,7 @@ size_t get_kernel_code_size( class function_info *entry ) } -kernel_info_t *gpgpu_opencl_ptx_sim_init_grid(class function_info *entry, +kernel_info_t *cuda_sim::gpgpu_opencl_ptx_sim_init_grid(class function_info *entry, gpgpu_ptx_sim_arg_list_t args, struct dim3 gridDim, struct dim3 blockDim, diff --git a/src/cuda-sim/cuda-sim.h b/src/cuda-sim/cuda-sim.h index 16ee46e..e259f1f 100644 --- a/src/cuda-sim/cuda-sim.h +++ b/src/cuda-sim/cuda-sim.h @@ -46,13 +46,7 @@ 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 int g_ptx_kernel_count; // used for classification stat collection purposes -extern class kernel_info_t *gpgpu_opencl_ptx_sim_init_grid(class function_info *entry, - gpgpu_ptx_sim_arg_list_t args, - struct dim3 gridDim, - struct dim3 blockDim, - class gpgpu_t *gpu ); extern void print_splash(); extern void gpgpu_ptx_sim_register_const_variable(void*, const char *deviceName, size_t size ); extern void gpgpu_ptx_sim_register_global_variable(void *hostVar, const char *deviceName, size_t size ); @@ -134,6 +128,7 @@ class cuda_sim { public: cuda_sim() { g_ptx_sim_num_insn = 0; + g_ptx_kernel_count = -1; // used for classification stat collection purposes } //global variables char *opcode_latency_int; @@ -151,10 +146,17 @@ class cuda_sim { int g_ptxinfo_error_detected; unsigned g_ptx_sim_num_insn; char *cdp_latency_str; + int g_ptx_kernel_count; // used for classification stat collection purposes //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 ); int gpgpu_opencl_ptx_sim_main_func( kernel_info_t *grid ); + void init_inst_classification_stat(); + kernel_info_t *gpgpu_opencl_ptx_sim_init_grid(class function_info *entry, + gpgpu_ptx_sim_arg_list_t args, + struct dim3 gridDim, + struct dim3 blockDim, + gpgpu_t *gpu ); }; #endif diff --git a/src/gpgpu-sim/gpu-sim.cc b/src/gpgpu-sim/gpu-sim.cc index 4f9ccbf..9f47067 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[g_ptx_kernel_count]); - StatDisp( g_inst_op_classification_stat[g_ptx_kernel_count]); + 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]); } #ifdef GPGPUSIM_POWER_MODEL -- cgit v1.3 From b3655bf28a7402db347f9d7f87049806b9315a05 Mon Sep 17 00:00:00 2001 From: Mengchi Zhang Date: Tue, 9 Jul 2019 15:03:18 -0400 Subject: Move g_global_name_lookup and g_constant_name_lookup Signed-off-by: Mengchi Zhang --- libcuda/cuda_runtime_api.cc | 4 ++-- src/cuda-sim/cuda-sim.cc | 14 ++++++-------- src/cuda-sim/cuda-sim.h | 6 ++++-- 3 files changed, 12 insertions(+), 12 deletions(-) (limited to 'src/cuda-sim/cuda-sim.h') diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc index 43c8bae..bbbaf23 100644 --- a/libcuda/cuda_runtime_api.cc +++ b/libcuda/cuda_runtime_api.cc @@ -767,9 +767,9 @@ void cudaRegisterVarInternal( ctx->cuobjdumpParseBinary((unsigned)(unsigned long long)fatCubinHandle); fflush(stdout); if ( constant && !global && !ext ) { - gpgpu_ptx_sim_register_const_variable(hostVar,deviceName,size); + ctx->func_sim->gpgpu_ptx_sim_register_const_variable(hostVar,deviceName,size); } else if ( !constant && !global && !ext ) { - gpgpu_ptx_sim_register_global_variable(hostVar,deviceName,size); + ctx->func_sim->gpgpu_ptx_sim_register_global_variable(hostVar,deviceName,size); } else cuda_not_implemented(__my_func__,__LINE__); } diff --git a/src/cuda-sim/cuda-sim.cc b/src/cuda-sim/cuda-sim.cc index 0ed125a..b3e2965 100644 --- a/src/cuda-sim/cuda-sim.cc +++ b/src/cuda-sim/cuda-sim.cc @@ -1948,18 +1948,16 @@ void print_splash() } } -std::map g_const_name_lookup; // indexed by hostVar -std::map g_global_name_lookup; // indexed by hostVar std::set g_globals; std::set g_constants; -void gpgpu_ptx_sim_register_const_variable(void *hostVar, const char *deviceName, size_t size ) +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 ); g_const_name_lookup[hostVar] = deviceName; } -void gpgpu_ptx_sim_register_global_variable(void *hostVar, const char *deviceName, size_t size ) +void cuda_sim::gpgpu_ptx_sim_register_global_variable(void *hostVar, const char *deviceName, size_t size ) { printf("GPGPU-Sim PTX registering global %s hostVar to name mapping\n", deviceName ); g_global_name_lookup[hostVar] = deviceName; @@ -1972,14 +1970,14 @@ void gpgpu_ptx_sim_memcpy_symbol(const char *hostVar, const void *src, size_t co memory_space_t mem_region = undefined_space; std::string sym_name; - std::map::iterator c=g_const_name_lookup.find(hostVar); - if ( c!=g_const_name_lookup.end() ) { + std::map::iterator c=gpu->gpgpu_ctx->func_sim->g_const_name_lookup.find(hostVar); + if ( c!=gpu->gpgpu_ctx->func_sim->g_const_name_lookup.end() ) { found_sym = true; sym_name = c->second; mem_region = const_space; } - std::map::iterator g=g_global_name_lookup.find(hostVar); - if ( g!=g_global_name_lookup.end() ) { + std::map::iterator g=gpu->gpgpu_ctx->func_sim->g_global_name_lookup.find(hostVar); + if ( g!=gpu->gpgpu_ctx->func_sim->g_global_name_lookup.end() ) { if ( found_sym ) { printf("Execution error: PTX symbol \"%s\" w/ hostVar=0x%Lx is declared both const and global?\n", sym_name.c_str(), (unsigned long long)hostVar ); diff --git a/src/cuda-sim/cuda-sim.h b/src/cuda-sim/cuda-sim.h index e259f1f..25ebf7b 100644 --- a/src/cuda-sim/cuda-sim.h +++ b/src/cuda-sim/cuda-sim.h @@ -48,8 +48,6 @@ extern void ** g_inst_classification_stat; extern void ** g_inst_op_classification_stat; extern void print_splash(); -extern void gpgpu_ptx_sim_register_const_variable(void*, const char *deviceName, size_t size ); -extern void gpgpu_ptx_sim_register_global_variable(void *hostVar, const char *deviceName, size_t size ); 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(); @@ -147,6 +145,8 @@ class cuda_sim { unsigned g_ptx_sim_num_insn; char *cdp_latency_str; 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 //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 ); @@ -157,6 +157,8 @@ class cuda_sim { struct dim3 gridDim, struct dim3 blockDim, 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 ); }; #endif -- cgit v1.3