From 0c16df3c8b108d8720846bb44b9abcc60ddf42f9 Mon Sep 17 00:00:00 2001 From: Mahmoud Date: Wed, 15 May 2019 20:16:54 -0400 Subject: make gpu_tot_cycle local variable not global variable --- src/gpgpu-sim/gpu-sim.cc | 35 ++++++++++++++--------------------- 1 file changed, 14 insertions(+), 21 deletions(-) (limited to 'src/gpgpu-sim/gpu-sim.cc') diff --git a/src/gpgpu-sim/gpu-sim.cc b/src/gpgpu-sim/gpu-sim.cc index c1ba934..72bac92 100644 --- a/src/gpgpu-sim/gpu-sim.cc +++ b/src/gpgpu-sim/gpu-sim.cc @@ -83,23 +83,6 @@ class gpgpu_sim_wrapper {}; bool g_interactive_debugger_enabled=false; -unsigned long long gpu_sim_cycle = 0; -unsigned long long gpu_tot_sim_cycle = 0; - - -// performance counter for stalls due to congestion. -unsigned int gpu_stall_dramfull = 0; -unsigned int gpu_stall_icnt2sh = 0; -unsigned long long partiton_reqs_in_parallel = 0; -unsigned long long partiton_reqs_in_parallel_total = 0; -unsigned long long partiton_reqs_in_parallel_util = 0; -unsigned long long partiton_reqs_in_parallel_util_total = 0; -unsigned long long gpu_sim_cycle_parition_util = 0; -unsigned long long gpu_tot_sim_cycle_parition_util = 0; -unsigned long long partiton_replys_in_parallel = 0; -unsigned long long partiton_replys_in_parallel_total = 0; - -tr1_hash_map address_random_interleaving; /* Clock Domains */ @@ -731,7 +714,7 @@ gpgpu_sim::gpgpu_sim( const gpgpu_sim_config &config ) #endif m_shader_stats = new shader_core_stats(m_shader_config); - m_memory_stats = new memory_stats_t(m_config.num_shader(),m_shader_config,m_memory_config); + m_memory_stats = new memory_stats_t(m_config.num_shader(),m_shader_config,m_memory_config,this); average_pipeline_duty_cycle = (float *)malloc(sizeof(float)); active_sms=(float *)malloc(sizeof(float)); m_power_stats = new power_stat_t(m_shader_config,average_pipeline_duty_cycle,active_sms,m_shader_stats,m_memory_config,m_memory_stats); @@ -742,6 +725,16 @@ gpgpu_sim::gpgpu_sim( const gpgpu_sim_config &config ) m_total_cta_launched = 0; gpu_deadlock = false; + gpu_stall_dramfull = 0; + gpu_stall_icnt2sh = 0; + partiton_reqs_in_parallel = 0; + partiton_reqs_in_parallel_total = 0; + partiton_reqs_in_parallel_util = 0; + partiton_reqs_in_parallel_util_total = 0; + gpu_sim_cycle_parition_util = 0; + gpu_tot_sim_cycle_parition_util = 0; + partiton_replys_in_parallel = 0; + partiton_replys_in_parallel_total = 0; m_cluster = new simt_core_cluster*[m_shader_config->n_simt_clusters]; for (unsigned i=0;in_simt_clusters;i++) @@ -750,7 +743,7 @@ gpgpu_sim::gpgpu_sim( const gpgpu_sim_config &config ) m_memory_partition_unit = new memory_partition_unit*[m_memory_config->m_n_mem]; m_memory_sub_partition = new memory_sub_partition*[m_memory_config->m_n_mem_sub_partition]; for (unsigned i=0;im_n_mem;i++) { - m_memory_partition_unit[i] = new memory_partition_unit(i, m_memory_config, m_memory_stats); + m_memory_partition_unit[i] = new memory_partition_unit(i, m_memory_config, m_memory_stats, this); for (unsigned p = 0; p < m_memory_config->m_n_sub_partition_per_memory_channel; p++) { unsigned submpid = i * m_memory_config->m_n_sub_partition_per_memory_channel + p; m_memory_sub_partition[submpid] = m_memory_partition_unit[i]->get_sub_partition(p); @@ -1504,7 +1497,7 @@ void shader_core_ctx::issue_block2core( kernel_info_t &kernel ) shader_CTA_count_log(m_sid, 1); SHADER_DPRINTF(LIVENESS, "GPGPU-Sim uArch: cta:%2u, start_tid:%4u, end_tid:%4u, initialized @(%lld,%lld)\n", - free_cta_hw_id, start_thread, end_thread, gpu_sim_cycle, gpu_tot_sim_cycle ); + free_cta_hw_id, start_thread, end_thread, m_gpu->gpu_sim_cycle, m_gpu->gpu_tot_sim_cycle ); } @@ -1721,7 +1714,7 @@ void gpgpu_sim::cycle() for (unsigned i=0;in_simt_clusters;i++) { m_cluster[i]->get_current_occupancy(active, total); } - DPRINTF(LIVENESS, "uArch: inst.: %lld (ipc=%4.1f, occ=%0.4f\% [%llu / %llu]) sim_rate=%u (inst/sec) elapsed = %u:%u:%02u:%02u / %s", + DPRINTFG(LIVENESS, "uArch: inst.: %lld (ipc=%4.1f, occ=%0.4f\% [%llu / %llu]) sim_rate=%u (inst/sec) elapsed = %u:%u:%02u:%02u / %s", gpu_tot_sim_insn + gpu_sim_insn, (double)gpu_sim_insn/(double)gpu_sim_cycle, float(active)/float(total) * 100, active, total, -- cgit v1.3 From 8fb484e4120a08e896f53424e9f4f46710966970 Mon Sep 17 00:00:00 2001 From: Mahmoud Date: Thu, 16 May 2019 11:16:13 -0400 Subject: fixing the link symbol error --- src/gpgpu-sim/gpu-sim.cc | 4 +--- src/gpgpu-sim/gpu-sim.h | 3 --- 2 files changed, 1 insertion(+), 6 deletions(-) (limited to 'src/gpgpu-sim/gpu-sim.cc') diff --git a/src/gpgpu-sim/gpu-sim.cc b/src/gpgpu-sim/gpu-sim.cc index 72bac92..6f19640 100644 --- a/src/gpgpu-sim/gpu-sim.cc +++ b/src/gpgpu-sim/gpu-sim.cc @@ -82,7 +82,7 @@ class gpgpu_sim_wrapper {}; bool g_interactive_debugger_enabled=false; - +tr1_hash_map address_random_interleaving; /* Clock Domains */ @@ -95,8 +95,6 @@ bool g_interactive_debugger_enabled=false; #define MEM_LATENCY_STAT_IMPL - - #include "mem_latency_stat.h" void power_config::reg_options(class OptionParser * opp) diff --git a/src/gpgpu-sim/gpu-sim.h b/src/gpgpu-sim/gpu-sim.h index 8e3b6ee..5ea5765 100644 --- a/src/gpgpu-sim/gpu-sim.h +++ b/src/gpgpu-sim/gpu-sim.h @@ -64,7 +64,6 @@ extern tr1_hash_map address_random_interleaving; - enum dram_ctrl_t { DRAM_FIFO=0, DRAM_FRFCFS=1 @@ -575,8 +574,6 @@ public: unsigned long long partiton_replys_in_parallel; unsigned long long partiton_replys_in_parallel_total; - tr1_hash_map address_random_interleaving; - FuncCache get_cache_config(std::string kernel_name); void set_cache_config(std::string kernel_name, FuncCache cacheConfig ); -- cgit v1.3 From 5d55ca32e2d81ee2c9187566050c5d399a142373 Mon Sep 17 00:00:00 2001 From: Mahmoud Date: Fri, 31 May 2019 13:22:29 -0400 Subject: add new params to the global stryct and fixing some bugs --- libcuda/cuda_runtime_api.cc | 6 ++++-- src/gpgpu-sim/gpu-sim.cc | 4 ++-- src/gpgpusim_entrypoint.h | 2 +- 3 files changed, 7 insertions(+), 5 deletions(-) (limited to 'src/gpgpu-sim/gpu-sim.cc') diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc index 4990edf..3d85b62 100644 --- a/libcuda/cuda_runtime_api.cc +++ b/libcuda/cuda_runtime_api.cc @@ -421,7 +421,8 @@ struct _cuda_device_id *GPGPUSim_Init() prop->maxThreadsPerMultiProcessor = the_gpu->threads_per_core(); #endif the_gpu->set_prop(prop); - the_device = new _cuda_device_id(the_gpu); + GPGPUsim_ctx_ptr()->the_cude_device = new _cuda_device_id(the_gpu); + the_device = GPGPUsim_ctx_ptr()->the_cude_device; } start_sim_thread(1); return the_device; @@ -433,7 +434,8 @@ static CUctx_st* GPGPUSim_Context() CUctx_st *the_context = GPGPUsim_ctx_ptr()->the_context; if( the_context == NULL ) { _cuda_device_id *the_gpu = GPGPUSim_Init(); - the_context = new CUctx_st(the_gpu); + GPGPUsim_ctx_ptr()->the_context = new CUctx_st(the_gpu); + the_context = GPGPUsim_ctx_ptr()->the_context; } return the_context; } diff --git a/src/gpgpu-sim/gpu-sim.cc b/src/gpgpu-sim/gpu-sim.cc index 6f19640..a557d6f 100644 --- a/src/gpgpu-sim/gpu-sim.cc +++ b/src/gpgpu-sim/gpu-sim.cc @@ -1131,7 +1131,7 @@ void gpgpu_sim::gpu_print_stat() time_t curr_time; time(&curr_time); - unsigned long long elapsed_time = MAX( curr_time - g_simulation_starttime, 1 ); + unsigned long long elapsed_time = MAX( curr_time - GPGPUsim_ctx_ptr()->g_simulation_starttime, 1 ); printf( "gpu_total_sim_rate=%u\n", (unsigned)( ( gpu_tot_sim_insn + gpu_sim_insn ) / elapsed_time ) ); //shader_print_l1_miss_stat( stdout ); @@ -1701,7 +1701,7 @@ void gpgpu_sim::cycle() time_t days, hrs, minutes, sec; time_t curr_time; time(&curr_time); - unsigned long long elapsed_time = MAX(curr_time - g_simulation_starttime, 1); + unsigned long long elapsed_time = MAX(curr_time - GPGPUsim_ctx_ptr()->g_simulation_starttime, 1); if ( (elapsed_time - last_liveness_message_time) >= m_config.liveness_message_freq && DTRACE(LIVENESS) ) { days = elapsed_time/(3600*24); hrs = elapsed_time/3600 - 24*days; diff --git a/src/gpgpusim_entrypoint.h b/src/gpgpusim_entrypoint.h index 2ad0fdf..406dd00 100644 --- a/src/gpgpusim_entrypoint.h +++ b/src/gpgpusim_entrypoint.h @@ -33,7 +33,7 @@ #include #include -extern time_t g_simulation_starttime; +//extern time_t g_simulation_starttime; struct GPGPUsim_ctx { -- cgit v1.3 From 80067b31e313b40cfd9f88b84593101706567547 Mon Sep 17 00:00:00 2001 From: Mengchi Zhang Date: Wed, 3 Jul 2019 10:43:29 -0400 Subject: Add backward pointer to gpgpu_sim Signed-off-by: Mengchi Zhang --- src/gpgpu-sim/gpu-sim.cc | 5 +++-- src/gpgpu-sim/gpu-sim.h | 4 +++- src/gpgpusim_entrypoint.cc | 2 +- 3 files changed, 7 insertions(+), 4 deletions(-) (limited to 'src/gpgpu-sim/gpu-sim.cc') diff --git a/src/gpgpu-sim/gpu-sim.cc b/src/gpgpu-sim/gpu-sim.cc index a557d6f..c894278 100644 --- a/src/gpgpu-sim/gpu-sim.cc +++ b/src/gpgpu-sim/gpu-sim.cc @@ -699,9 +699,10 @@ void gpgpu_sim::stop_all_running_kernels(){ void set_ptx_warp_size(const struct core_config * warp_size); -gpgpu_sim::gpgpu_sim( const gpgpu_sim_config &config ) +gpgpu_sim::gpgpu_sim( const gpgpu_sim_config &config, gpgpu_context* ctx ) : gpgpu_t(config), m_config(config) -{ +{ + gpgpu_ctx = ctx; m_shader_config = &m_config.m_shader_config; m_memory_config = &m_config.m_memory_config; set_ptx_warp_size(m_shader_config); diff --git a/src/gpgpu-sim/gpu-sim.h b/src/gpgpu-sim/gpu-sim.h index 5ea5765..b1d2e45 100644 --- a/src/gpgpu-sim/gpu-sim.h +++ b/src/gpgpu-sim/gpu-sim.h @@ -416,10 +416,11 @@ struct occupancy_stats { } }; +class gpgpu_context; class gpgpu_sim : public gpgpu_t { public: - gpgpu_sim( const gpgpu_sim_config &config ); + gpgpu_sim( const gpgpu_sim_config &config, gpgpu_context* ctx ); void set_prop( struct cudaDeviceProp *prop ); @@ -505,6 +506,7 @@ private: void gpgpu_debug(); ///// data ///// + class gpgpu_context* gpgpu_ctx; class simt_core_cluster **m_cluster; class memory_partition_unit **m_memory_partition_unit; diff --git a/src/gpgpusim_entrypoint.cc b/src/gpgpusim_entrypoint.cc index c91a9d1..b54f20c 100644 --- a/src/gpgpusim_entrypoint.cc +++ b/src/gpgpusim_entrypoint.cc @@ -232,7 +232,7 @@ gpgpu_sim *gpgpu_context::gpgpu_ptx_sim_init_perf() assert(setlocale(LC_NUMERIC,"C")); GPGPUsim_ctx_ptr()->g_the_gpu_config->init(); - GPGPUsim_ctx_ptr()->g_the_gpu = new gpgpu_sim(*(GPGPUsim_ctx_ptr()->g_the_gpu_config)); + 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_simulation_starttime = time((time_t *)NULL); -- cgit v1.3 From cb60ae4893086175fee49dc6088d46665a7f088b Mon Sep 17 00:00:00 2001 From: Mengchi Zhang Date: Wed, 3 Jul 2019 12:04:22 -0400 Subject: move g_shader_core_config Signed-off-by: Mengchi Zhang --- src/cuda-sim/ptx_parser.cc | 4 +--- src/cuda-sim/ptx_parser.h | 3 +++ src/gpgpu-sim/gpu-sim.cc | 5 ++--- 3 files changed, 6 insertions(+), 6 deletions(-) (limited to 'src/gpgpu-sim/gpu-sim.cc') diff --git a/src/cuda-sim/ptx_parser.cc b/src/cuda-sim/ptx_parser.cc index 9094ec3..09bc15d 100644 --- a/src/cuda-sim/ptx_parser.cc +++ b/src/cuda-sim/ptx_parser.cc @@ -42,8 +42,7 @@ extern FILE *ptx_get_in (yyscan_t yyscanner ); extern int ptx_parse(yyscan_t scanner, ptx_recognizer* recognizer); extern int ptx_lex_destroy(yyscan_t scanner); -static const struct core_config *g_shader_core_config; -void set_ptx_warp_size(const struct core_config * warp_size) +void ptx_recognizer::set_ptx_warp_size(const struct core_config * warp_size) { g_shader_core_config=warp_size; } @@ -154,7 +153,6 @@ symbol_table * gpgpu_context::init_parser( const char *ptx_filename ) return ptx_parser->g_global_symbol_table; } -static int g_entry_point; void ptx_recognizer::start_function( int entry_point ) { diff --git a/src/cuda-sim/ptx_parser.h b/src/cuda-sim/ptx_parser.h index bc9a872..ec300a1 100644 --- a/src/cuda-sim/ptx_parser.h +++ b/src/cuda-sim/ptx_parser.h @@ -105,6 +105,8 @@ class ptx_recognizer { function_info *g_func_info; operand_info g_return_var; bool g_debug_ir_generation; + int g_entry_point; + const struct core_config *g_shader_core_config; // member function list void init_directive_state(); @@ -172,6 +174,7 @@ class ptx_recognizer { void end_inst_group(); bool check_for_duplicates( const char *identifier ); void read_parser_environment_variables(); + void set_ptx_warp_size(const struct core_config * warp_size); }; diff --git a/src/gpgpu-sim/gpu-sim.cc b/src/gpgpu-sim/gpu-sim.cc index c894278..2ff37d1 100644 --- a/src/gpgpu-sim/gpu-sim.cc +++ b/src/gpgpu-sim/gpu-sim.cc @@ -65,6 +65,7 @@ #include "visualizer.h" #include "stats.h" #include "../cuda-sim/cuda_device_runtime.h" +#include "../../libcuda/gpgpu_context.h" #ifdef GPGPUSIM_POWER_MODEL #include "power_interface.h" @@ -697,15 +698,13 @@ void gpgpu_sim::stop_all_running_kernels(){ } } -void set_ptx_warp_size(const struct core_config * warp_size); - gpgpu_sim::gpgpu_sim( const gpgpu_sim_config &config, gpgpu_context* ctx ) : gpgpu_t(config), m_config(config) { gpgpu_ctx = ctx; m_shader_config = &m_config.m_shader_config; m_memory_config = &m_config.m_memory_config; - set_ptx_warp_size(m_shader_config); + ctx->ptx_parser->set_ptx_warp_size(m_shader_config); ptx_file_line_stats_create_exposed_latency_tracker(m_config.num_shader()); #ifdef GPGPUSIM_POWER_MODEL -- cgit v1.3 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 --- libcuda/gpgpu_context.h | 3 +++ src/cuda-sim/cuda-sim.cc | 7 ++++--- src/cuda-sim/cuda-sim.h | 11 +++++++++-- src/cuda-sim/ptx_ir.cc | 4 +++- src/cuda-sim/ptx_ir.h | 7 ++++++- src/cuda-sim/ptx_parser.cc | 3 ++- src/gpgpu-sim/gpu-sim.cc | 2 +- src/gpgpu-sim/gpu-sim.h | 11 ++++++++--- src/gpgpu-sim/mem_latency_stat.cc | 2 +- src/gpgpu-sim/mem_latency_stat.h | 4 ++-- src/gpgpu-sim/power_interface.cc | 2 +- src/gpgpu-sim/power_interface.h | 2 +- src/gpgpu-sim/power_stat.cc | 6 +++--- src/gpgpu-sim/power_stat.h | 6 +++--- src/gpgpu-sim/shader.cc | 7 ++++--- src/gpgpu-sim/shader.h | 18 ++++++++++++------ src/gpgpusim_entrypoint.cc | 4 ++-- 17 files changed, 65 insertions(+), 34 deletions(-) (limited to 'src/gpgpu-sim/gpu-sim.cc') diff --git a/libcuda/gpgpu_context.h b/libcuda/gpgpu_context.h index a8b60f4..2e21009 100644 --- a/libcuda/gpgpu_context.h +++ b/libcuda/gpgpu_context.h @@ -4,6 +4,7 @@ #include "../src/cuda-sim/ptx_loader.h" #include "../src/cuda-sim/ptx_parser.h" #include "../src/gpgpusim_entrypoint.h" +#include "../src/cuda-sim/cuda-sim.h" class gpgpu_context { public: @@ -13,6 +14,7 @@ class gpgpu_context { ptxinfo = new ptxinfo_data(this); ptx_parser = new ptx_recognizer(this); the_gpgpusim = new GPGPUsim_ctx(this); + func_sim = new cuda_sim(); } // global list symbol_table *g_global_allfiles_symbol_table; @@ -22,6 +24,7 @@ class gpgpu_context { ptxinfo_data* ptxinfo; ptx_recognizer* ptx_parser; GPGPUsim_ctx* the_gpgpusim; + cuda_sim* func_sim; // member function list void cuobjdumpParseBinary(unsigned int handle); class symbol_table *gpgpu_ptx_sim_load_ptx_from_string( const char *p, unsigned source_num ); diff --git a/src/cuda-sim/cuda-sim.cc b/src/cuda-sim/cuda-sim.cc index 261d605..df0bbd7 100644 --- a/src/cuda-sim/cuda-sim.cc +++ b/src/cuda-sim/cuda-sim.cc @@ -51,6 +51,7 @@ typedef void * yyscan_t; #include "decuda_pred_table/decuda_pred_table.h" #include "../stream_manager.h" #include "cuda_device_runtime.h" +#include "../../libcuda/gpgpu_context.h" int gpgpu_ptx_instruction_classification; void ** g_inst_classification_stat = NULL; @@ -66,12 +67,12 @@ int cp_cta_resume; unsigned g_ptx_sim_num_insn = 0; unsigned gpgpu_param_num_shaders = 0; -char *opcode_latency_int, *opcode_latency_fp, *opcode_latency_dp,*opcode_latency_sfu,*opcode_latency_tensor; +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]; -void ptx_opcocde_latency_options (option_parser_t opp) { +void cuda_sim::ptx_opcocde_latency_options (option_parser_t opp) { option_parser_register(opp, "-ptx_opcode_latency_int", OPT_CSTR, &opcode_latency_int, "Opcode latencies for integers " "Default 1,1,19,25,145", @@ -667,7 +668,7 @@ void ptx_instruction::set_opcode_and_latency() * [3] MAD * [4] DIV */ - sscanf(opcode_latency_int, "%u,%u,%u,%u,%u", + 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", diff --git a/src/cuda-sim/cuda-sim.h b/src/cuda-sim/cuda-sim.h index e690356..96d34f5 100644 --- a/src/cuda-sim/cuda-sim.h +++ b/src/cuda-sim/cuda-sim.h @@ -47,10 +47,9 @@ 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_int, *opcode_latency_fp, *opcode_latency_dp,*opcode_latency_sfu,*opcode_latency_tensor; +extern char *opcode_latency_fp, *opcode_latency_dp,*opcode_latency_sfu,*opcode_latency_tensor; -void ptx_opcocde_latency_options (option_parser_t opp); 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, @@ -134,4 +133,12 @@ void print_ptxinfo(); void clear_ptxinfo(); struct gpgpu_ptx_sim_info get_ptxinfo(); +class cuda_sim { + public: + //global variables + char *opcode_latency_int; + //global functions + void ptx_opcocde_latency_options (option_parser_t opp); +}; + #endif diff --git a/src/cuda-sim/ptx_ir.cc b/src/cuda-sim/ptx_ir.cc index 8cedf79..1bd409e 100644 --- a/src/cuda-sim/ptx_ir.cc +++ b/src/cuda-sim/ptx_ir.cc @@ -1095,8 +1095,10 @@ ptx_instruction::ptx_instruction( int opcode, const char *file, unsigned line, const char *source, - const core_config *config ) : warp_inst_t(config) + const core_config *config, + gpgpu_context* ctx ) : warp_inst_t(config) { + gpgpu_ctx = ctx; m_uid = ++g_num_ptx_inst_uid; m_PC = 0; m_opcode = opcode; diff --git a/src/cuda-sim/ptx_ir.h b/src/cuda-sim/ptx_ir.h index fd869c6..1604551 100644 --- a/src/cuda-sim/ptx_ir.h +++ b/src/cuda-sim/ptx_ir.h @@ -43,6 +43,8 @@ #include "memory.h" +class gpgpu_context; + class type_info_key { public: type_info_key() @@ -931,7 +933,8 @@ public: const char *file, unsigned line, const char *source, - const core_config *config ); + const core_config *config, + gpgpu_context* ctx); void print_insn() const; virtual void print_insn( FILE *fp ) const; @@ -1187,6 +1190,8 @@ private: virtual void pre_decode(); friend class function_info; static unsigned g_num_ptx_inst_uid; + // backward pointer + class gpgpu_context* gpgpu_ctx; }; class param_info { 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(); diff --git a/src/gpgpu-sim/gpu-sim.cc b/src/gpgpu-sim/gpu-sim.cc index 2ff37d1..39acdd9 100644 --- a/src/gpgpu-sim/gpu-sim.cc +++ b/src/gpgpu-sim/gpu-sim.cc @@ -1826,7 +1826,7 @@ void gpgpu_sim::dump_pipeline( int mask, int s, int m ) const fflush(stdout); } -const struct shader_core_config * gpgpu_sim::getShaderCoreConfig() +const shader_core_config * gpgpu_sim::getShaderCoreConfig() { return m_shader_config; } diff --git a/src/gpgpu-sim/gpu-sim.h b/src/gpgpu-sim/gpu-sim.h index e2c913a..7eeb7dd 100644 --- a/src/gpgpu-sim/gpu-sim.h +++ b/src/gpgpu-sim/gpu-sim.h @@ -295,7 +295,10 @@ extern bool g_interactive_debugger_enabled; class gpgpu_sim_config : public power_config, public gpgpu_functional_sim_config { public: - gpgpu_sim_config() { m_valid = false; } + gpgpu_sim_config(gpgpu_context* ctx): m_shader_config(ctx) { + m_valid = false; + gpgpu_ctx = ctx; + } void reg_options(class OptionParser * opp); void init() { @@ -341,6 +344,8 @@ private: void init_clock_domains(void ); + // backward pointer + class gpgpu_context* gpgpu_ctx; bool m_valid; shader_core_config m_shader_config; memory_config m_memory_config; @@ -473,7 +478,7 @@ public: /*! * Returning the configuration of the shader core, used by the functional simulation only so far */ - const struct shader_core_config * getShaderCoreConfig(); + const shader_core_config * getShaderCoreConfig(); //! Get shader core Memory Configuration @@ -537,7 +542,7 @@ private: const gpgpu_sim_config &m_config; const struct cudaDeviceProp *m_cuda_properties; - const struct shader_core_config *m_shader_config; + const shader_core_config *m_shader_config; const struct memory_config *m_memory_config; // stats diff --git a/src/gpgpu-sim/mem_latency_stat.cc b/src/gpgpu-sim/mem_latency_stat.cc index 04dc75b..d08ba39 100644 --- a/src/gpgpu-sim/mem_latency_stat.cc +++ b/src/gpgpu-sim/mem_latency_stat.cc @@ -42,7 +42,7 @@ #include #include -memory_stats_t::memory_stats_t( unsigned n_shader, const struct shader_core_config *shader_config, const struct memory_config *mem_config, const class gpgpu_sim* gpu ) +memory_stats_t::memory_stats_t( unsigned n_shader, const shader_core_config *shader_config, const struct memory_config *mem_config, const class gpgpu_sim* gpu ) { assert( mem_config->m_valid ); assert( shader_config->m_valid ); diff --git a/src/gpgpu-sim/mem_latency_stat.h b/src/gpgpu-sim/mem_latency_stat.h index b9285c1..6ce568d 100644 --- a/src/gpgpu-sim/mem_latency_stat.h +++ b/src/gpgpu-sim/mem_latency_stat.h @@ -35,7 +35,7 @@ class memory_stats_t { public: memory_stats_t( unsigned n_shader, - const struct shader_core_config *shader_config, + const class shader_core_config *shader_config, const struct memory_config *mem_config, const class gpgpu_sim* gpu); @@ -53,7 +53,7 @@ public: unsigned m_n_shader; - const struct shader_core_config *m_shader_config; + const shader_core_config *m_shader_config; const struct memory_config *m_memory_config; const class gpgpu_sim* m_gpu; diff --git a/src/gpgpu-sim/power_interface.cc b/src/gpgpu-sim/power_interface.cc index 3861b6a..0272aa6 100644 --- a/src/gpgpu-sim/power_interface.cc +++ b/src/gpgpu-sim/power_interface.cc @@ -38,7 +38,7 @@ void init_mcpat(const gpgpu_sim_config &config, class gpgpu_sim_wrapper *wrapper } -void mcpat_cycle(const gpgpu_sim_config &config, const struct shader_core_config *shdr_config, class gpgpu_sim_wrapper *wrapper, class power_stat_t *power_stats, unsigned stat_sample_freq, unsigned tot_cycle, unsigned cycle, unsigned tot_inst, unsigned inst){ +void mcpat_cycle(const gpgpu_sim_config &config, const shader_core_config *shdr_config, class gpgpu_sim_wrapper *wrapper, class power_stat_t *power_stats, unsigned stat_sample_freq, unsigned tot_cycle, unsigned cycle, unsigned tot_inst, unsigned inst){ static bool mcpat_init=true; diff --git a/src/gpgpu-sim/power_interface.h b/src/gpgpu-sim/power_interface.h index afac22b..a388c23 100644 --- a/src/gpgpu-sim/power_interface.h +++ b/src/gpgpu-sim/power_interface.h @@ -36,7 +36,7 @@ #include "gpgpu_sim_wrapper.h" void init_mcpat(const gpgpu_sim_config &config, class gpgpu_sim_wrapper *wrapper, unsigned stat_sample_freq, unsigned tot_inst, unsigned inst); -void mcpat_cycle(const gpgpu_sim_config &config, const struct shader_core_config *shdr_config, class gpgpu_sim_wrapper *wrapper, class power_stat_t *power_stats, +void mcpat_cycle(const gpgpu_sim_config &config, const shader_core_config *shdr_config, class gpgpu_sim_wrapper *wrapper, class power_stat_t *power_stats, unsigned stat_sample_freq, unsigned tot_cycle, unsigned cycle, unsigned tot_inst, unsigned inst); void mcpat_reset_perf_count(class gpgpu_sim_wrapper *wrapper); diff --git a/src/gpgpu-sim/power_stat.cc b/src/gpgpu-sim/power_stat.cc index 4c995e9..007b4c6 100644 --- a/src/gpgpu-sim/power_stat.cc +++ b/src/gpgpu-sim/power_stat.cc @@ -42,7 +42,7 @@ -power_mem_stat_t::power_mem_stat_t(const struct memory_config *mem_config, const struct shader_core_config *shdr_config, memory_stats_t *mem_stats, shader_core_stats *shdr_stats){ +power_mem_stat_t::power_mem_stat_t(const struct memory_config *mem_config, const shader_core_config *shdr_config, memory_stats_t *mem_stats, shader_core_stats *shdr_stats){ assert( mem_config->m_valid ); m_mem_stats = mem_stats; m_config = mem_config; @@ -125,7 +125,7 @@ void power_mem_stat_t::print (FILE *fout) const { } -power_core_stat_t::power_core_stat_t( const struct shader_core_config *shader_config, shader_core_stats *core_stats ) +power_core_stat_t::power_core_stat_t( const shader_core_config *shader_config, shader_core_stats *core_stats ) { assert( shader_config->m_valid ); m_config = shader_config; @@ -266,7 +266,7 @@ for(unsigned i=0; inum_shader(); ++i){ } } -power_stat_t::power_stat_t( const struct shader_core_config *shader_config,float * average_pipeline_duty_cycle,float *active_sms,shader_core_stats * shader_stats, const struct memory_config *mem_config,memory_stats_t * memory_stats) +power_stat_t::power_stat_t( const shader_core_config *shader_config,float * average_pipeline_duty_cycle,float *active_sms,shader_core_stats * shader_stats, const struct memory_config *mem_config,memory_stats_t * memory_stats) { assert( shader_config->m_valid ); assert( mem_config->m_valid ); diff --git a/src/gpgpu-sim/power_stat.h b/src/gpgpu-sim/power_stat.h index 20af2e5..91fade9 100644 --- a/src/gpgpu-sim/power_stat.h +++ b/src/gpgpu-sim/power_stat.h @@ -73,7 +73,7 @@ struct shader_core_power_stats_pod { class power_core_stat_t : public shader_core_power_stats_pod { public: - power_core_stat_t(const struct shader_core_config *shader_config, shader_core_stats *core_stats); + power_core_stat_t(const shader_core_config *shader_config, shader_core_stats *core_stats); void visualizer_print( gzFile visualizer_file ); void print (FILE *fout); void init(); @@ -113,7 +113,7 @@ struct mem_power_stats_pod{ class power_mem_stat_t : public mem_power_stats_pod{ public: - power_mem_stat_t(const struct memory_config *mem_config, const struct shader_core_config *shdr_config, memory_stats_t *mem_stats, shader_core_stats *shdr_stats); + power_mem_stat_t(const struct memory_config *mem_config, const shader_core_config *shdr_config, memory_stats_t *mem_stats, shader_core_stats *shdr_stats); void visualizer_print( gzFile visualizer_file ); void print (FILE *fout) const; void init(); @@ -128,7 +128,7 @@ private: class power_stat_t { public: - power_stat_t( const struct shader_core_config *shader_config,float * average_pipeline_duty_cycle,float * active_sms,shader_core_stats * shader_stats, const struct memory_config *mem_config,memory_stats_t * memory_stats); + power_stat_t( const shader_core_config *shader_config,float * average_pipeline_duty_cycle,float * active_sms,shader_core_stats * shader_stats, const struct memory_config *mem_config,memory_stats_t * memory_stats); void visualizer_print( gzFile visualizer_file ); void print (FILE *fout) const; void save_stats(){ diff --git a/src/gpgpu-sim/shader.cc b/src/gpgpu-sim/shader.cc index 69b619a..4d12068 100644 --- a/src/gpgpu-sim/shader.cc +++ b/src/gpgpu-sim/shader.cc @@ -46,6 +46,7 @@ #include #include "traffic_breakdown.h" #include "shader_trace.h" +#include "../../libcuda/gpgpu_context.h" #define PRIORITIZE_MSHR_OVER_WB 1 #define MAX(a,b) (((a)>(b))?(a):(b)) @@ -69,7 +70,7 @@ shader_core_ctx::shader_core_ctx( class gpgpu_sim *gpu, class simt_core_cluster *cluster, unsigned shader_id, unsigned tpc_id, - const struct shader_core_config *config, + const shader_core_config *config, const struct memory_config *mem_config, shader_core_stats *stats ) : core_t( gpu, NULL, config->warp_size, config->n_thread_per_shader ), @@ -3018,7 +3019,7 @@ void shader_core_config::set_pipeline_latency() { * [3] MAD * [4] DIV */ - sscanf(opcode_latency_int, "%u,%u,%u,%u,%u", + 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", @@ -3786,7 +3787,7 @@ void opndcoll_rfu_t::collector_unit_t::dispatch() simt_core_cluster::simt_core_cluster( class gpgpu_sim *gpu, unsigned cluster_id, - const struct shader_core_config *config, + const shader_core_config *config, const struct memory_config *mem_config, shader_core_stats *stats, class memory_stats_t *mstats ) diff --git a/src/gpgpu-sim/shader.h b/src/gpgpu-sim/shader.h index 25b9607..e0cefac 100644 --- a/src/gpgpu-sim/shader.h +++ b/src/gpgpu-sim/shader.h @@ -69,6 +69,8 @@ #define WRITE_MASK_SIZE 8 +class gpgpu_context; + enum exec_unit_type_t { NONE = 0, @@ -294,7 +296,7 @@ typedef std::bitset warp_set_t; int register_bank(int regnum, int wid, unsigned num_banks, unsigned bank_warp_shift, bool sub_core_model, unsigned banks_per_sched, unsigned sched_id ); class shader_core_ctx; -struct shader_core_config; +class shader_core_config; class shader_core_stats; enum scheduler_prioritization_type @@ -1032,7 +1034,7 @@ struct ifetch_buffer_t { unsigned m_warp_id; }; -struct shader_core_config; +class shader_core_config; class simd_function_unit { public: @@ -1362,10 +1364,12 @@ const char* const pipeline_stage_name_decode[] = { "N_PIPELINE_STAGES" }; -struct shader_core_config : public core_config +class shader_core_config : public core_config { - shader_core_config(){ + public: + shader_core_config(gpgpu_context* ctx){ pipeline_widths_string = NULL; + gpgpu_ctx = ctx; } void init() @@ -1425,6 +1429,8 @@ struct shader_core_config : public core_config unsigned cid_to_sid( unsigned cid, unsigned cluster_id ) const { return cluster_id*n_simt_cores_per_cluster + cid; } void set_pipeline_latency(); + // backward pointer + class gpgpu_context* gpgpu_ctx; // data char *gpgpu_shader_core_pipeline_opt; bool gpgpu_perfect_mem; @@ -1770,7 +1776,7 @@ public: class simt_core_cluster *cluster, unsigned shader_id, unsigned tpc_id, - const struct shader_core_config *config, + const shader_core_config *config, const struct memory_config *mem_config, shader_core_stats *stats ); @@ -2065,7 +2071,7 @@ class simt_core_cluster { public: simt_core_cluster( class gpgpu_sim *gpu, unsigned cluster_id, - const struct shader_core_config *config, + const shader_core_config *config, const struct memory_config *mem_config, shader_core_stats *stats, memory_stats_t *mstats ); diff --git a/src/gpgpusim_entrypoint.cc b/src/gpgpusim_entrypoint.cc index b54f20c..5018305 100644 --- a/src/gpgpusim_entrypoint.cc +++ b/src/gpgpusim_entrypoint.cc @@ -218,10 +218,10 @@ gpgpu_sim *gpgpu_context::gpgpu_ptx_sim_init_perf() option_parser_t opp = option_parser_create(); ptx_reg_options(opp); - ptx_opcocde_latency_options(opp); + func_sim->ptx_opcocde_latency_options(opp); icnt_reg_options(opp); - GPGPUsim_ctx_ptr()->g_the_gpu_config = new gpgpu_sim_config(); + GPGPUsim_ctx_ptr()->g_the_gpu_config = new gpgpu_sim_config(this); GPGPUsim_ctx_ptr()->g_the_gpu_config->reg_options(opp); // register GPU microrachitecture options option_parser_cmdline(opp, sg_argc, sg_argv); // parse configuration options -- 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/gpgpu-sim/gpu-sim.cc') 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 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/gpgpu-sim/gpu-sim.cc') 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 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 --- libcuda/cuda_runtime_api.cc | 6 +++--- libcuda/gpgpu_context.h | 9 ++++++++- libopencl/opencl_runtime_api.cc | 6 +++--- 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 ++++++++++++++++-- src/gpgpu-sim/gpu-sim.cc | 4 ++-- src/gpgpusim_entrypoint.cc | 2 +- 9 files changed, 46 insertions(+), 22 deletions(-) (limited to 'src/gpgpu-sim/gpu-sim.cc') diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc index bbbaf23..59d2a60 100644 --- a/libcuda/cuda_runtime_api.cc +++ b/libcuda/cuda_runtime_api.cc @@ -1076,7 +1076,7 @@ cudaError_t cudaLaunchInternal( const char *hostFun, gpgpu_context* gpgpu_ctx = CUctx_st* context = GPGPUSim_Context(); char *mode = getenv("PTX_SIM_MODE_FUNC"); if( mode ) - sscanf(mode,"%u", &g_ptx_sim_mode); + sscanf(mode,"%u", &(ctx->func_sim->g_ptx_sim_mode)); gpgpusim_ptx_assert( !ctx->api->g_cuda_launch_stack.empty(), "empty launch stack" ); kernel_config config = ctx->api->g_cuda_launch_stack.back(); { @@ -1092,7 +1092,7 @@ cudaError_t cudaLaunchInternal( const char *hostFun, gpgpu_context* gpgpu_ctx = } struct CUstream_st *stream = config.get_stream(); printf("\nGPGPU-Sim PTX: cudaLaunch for 0x%p (mode=%s) on stream %u\n", hostFun, - g_ptx_sim_mode?"functional simulation":"performance simulation", stream?stream->get_uid():0 ); + (ctx->func_sim->g_ptx_sim_mode)?"functional simulation":"performance simulation", stream?stream->get_uid():0 ); kernel_info_t *grid = ctx->api->gpgpu_cuda_ptx_sim_init_grid(hostFun,config.get_args(),config.grid_dim(),config.block_dim(),context); //do dynamic PDOM analysis for performance simulation scenario std::string kname = grid->name(); @@ -1143,7 +1143,7 @@ cudaError_t cudaLaunchInternal( const char *hostFun, gpgpu_context* gpgpu_ctx = } printf("GPGPU-Sim PTX: pushing kernel \'%s\' to stream %u, gridDim= (%u,%u,%u) blockDim = (%u,%u,%u) \n", kname.c_str(), stream?stream->get_uid():0, gridDim.x,gridDim.y,gridDim.z,blockDim.x,blockDim.y,blockDim.z ); - stream_operation op(grid,g_ptx_sim_mode,stream); + stream_operation op(grid,ctx->func_sim->g_ptx_sim_mode,stream); g_stream_manager()->push(op); ctx->api->g_cuda_launch_stack.pop_back(); return g_last_cudaError = cudaSuccess; diff --git a/libcuda/gpgpu_context.h b/libcuda/gpgpu_context.h index a2ae7b6..3c9f87c 100644 --- a/libcuda/gpgpu_context.h +++ b/libcuda/gpgpu_context.h @@ -5,6 +5,7 @@ #include "../src/cuda-sim/ptx_parser.h" #include "../src/gpgpusim_entrypoint.h" #include "../src/cuda-sim/cuda-sim.h" +#include "../src/cuda-sim/cuda_device_runtime.h" class gpgpu_context { public: @@ -14,7 +15,10 @@ class gpgpu_context { ptxinfo = new ptxinfo_data(this); ptx_parser = new ptx_recognizer(this); the_gpgpusim = new GPGPUsim_ctx(this); - func_sim = new cuda_sim(); + func_sim = new cuda_sim(this); +#if (CUDART_VERSION >= 5000) + device_runtime = new cuda_device_runtime(this); +#endif } // global list symbol_table *g_global_allfiles_symbol_table; @@ -25,6 +29,9 @@ class gpgpu_context { ptx_recognizer* ptx_parser; GPGPUsim_ctx* the_gpgpusim; cuda_sim* func_sim; +#if (CUDART_VERSION >= 5000) + cuda_device_runtime* device_runtime; +#endif // member function list void cuobjdumpParseBinary(unsigned int handle); class symbol_table *gpgpu_ptx_sim_load_ptx_from_string( const char *p, unsigned source_num ); diff --git a/libopencl/opencl_runtime_api.cc b/libopencl/opencl_runtime_api.cc index 0a6eb3e..aaaec4f 100644 --- a/libopencl/opencl_runtime_api.cc +++ b/libopencl/opencl_runtime_api.cc @@ -884,9 +884,9 @@ clEnqueueNDRangeKernel(cl_command_queue command_queue, printf("\n\n\n"); char *mode = getenv("PTX_SIM_MODE_FUNC"); if ( mode ) - sscanf(mode,"%u", &g_ptx_sim_mode); + sscanf(mode,"%u", &(ctx->func_sim->g_ptx_sim_mode)); printf("GPGPU-Sim OpenCL API: clEnqueueNDRangeKernel '%s' (mode=%s)\n", kernel->name().c_str(), - g_ptx_sim_mode?"functional simulation":"performance simulation"); + (ctx->func_sim->g_ptx_sim_mode)?"functional simulation":"performance simulation"); if ( !work_dim || work_dim > 3 ) return CL_INVALID_WORK_DIMENSION; size_t _local_size[3]; if( local_work_size != NULL ) { @@ -957,7 +957,7 @@ clEnqueueNDRangeKernel(cl_command_queue command_queue, 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 ( g_ptx_sim_mode ) + if ( ctx->func_sim->g_ptx_sim_mode ) ctx->func_sim->gpgpu_opencl_ptx_sim_main_func( grid ); else gpgpu_opencl_ptx_sim_main_perf( grid ); 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 diff --git a/src/gpgpu-sim/gpu-sim.cc b/src/gpgpu-sim/gpu-sim.cc index 9f47067..bbcc078 100644 --- a/src/gpgpu-sim/gpu-sim.cc +++ b/src/gpgpu-sim/gpu-sim.cc @@ -503,7 +503,7 @@ void gpgpu_sim_config::reg_options(option_parser_t opp) &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, &g_ptx_sim_mode, + option_parser_register(opp, "-gpgpu_ptx_sim_mode", OPT_INT32, &(gpgpu_ctx->func_sim->g_ptx_sim_mode), "Select between Performance (default) or Functional simulation (1)", "0"); option_parser_register(opp, "-gpgpu_clock_domains", OPT_CSTR, &gpgpu_clock_domains, @@ -1753,7 +1753,7 @@ void gpgpu_sim::cycle() #if (CUDART_VERSION >= 5000) //launch device kernel - launch_one_device_kernel(); + gpgpu_ctx->device_runtime->launch_one_device_kernel(); #endif } } diff --git a/src/gpgpusim_entrypoint.cc b/src/gpgpusim_entrypoint.cc index d9d1023..683a695 100644 --- a/src/gpgpusim_entrypoint.cc +++ b/src/gpgpusim_entrypoint.cc @@ -213,7 +213,7 @@ gpgpu_sim *gpgpu_context::gpgpu_ptx_sim_init_perf() { srand(1); print_splash(); - read_sim_environment_variables(); + func_sim->read_sim_environment_variables(); ptx_parser->read_parser_environment_variables(); option_parser_t opp = option_parser_create(); -- 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/gpgpu-sim/gpu-sim.cc') 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 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/gpgpu-sim/gpu-sim.cc') 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 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/gpgpu-sim/gpu-sim.cc') 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 0f5f90ebd896b807270abfae2cd7634733654668 Mon Sep 17 00:00:00 2001 From: Mengchi Zhang Date: Thu, 11 Jul 2019 23:29:21 -0400 Subject: Move g_kernel_launch_latency Signed-off-by: Mengchi Zhang --- src/abstract_hardware_model.cc | 3 +-- src/cuda-sim/cuda_device_runtime.h | 1 + src/gpgpu-sim/gpu-sim.cc | 3 +-- 3 files changed, 3 insertions(+), 4 deletions(-) (limited to 'src/gpgpu-sim/gpu-sim.cc') diff --git a/src/abstract_hardware_model.cc b/src/abstract_hardware_model.cc index 0f4291d..450087b 100644 --- a/src/abstract_hardware_model.cc +++ b/src/abstract_hardware_model.cc @@ -693,7 +693,6 @@ void warp_inst_t::completed( unsigned long long cycle ) const //Jin: CDP support bool g_cdp_enabled; -unsigned g_kernel_launch_latency; unsigned kernel_info_t::m_next_uid = 1; @@ -717,7 +716,7 @@ kernel_info_t::kernel_info_t( dim3 gridDim, dim3 blockDim, class function_info * m_parent_kernel = NULL; //Jin: launch latency management - m_launch_latency = g_kernel_launch_latency; + m_launch_latency = entry->gpgpu_ctx->device_runtime->g_kernel_launch_latency; volta_cache_config_set=false; m_NameToCudaArray = nameToCudaArray; diff --git a/src/cuda-sim/cuda_device_runtime.h b/src/cuda-sim/cuda_device_runtime.h index a6303b9..5927b54 100644 --- a/src/cuda-sim/cuda_device_runtime.h +++ b/src/cuda-sim/cuda_device_runtime.h @@ -48,6 +48,7 @@ class cuda_device_runtime { unsigned long long g_total_param_size; std::map g_cuda_device_launch_param_map; std::list g_cuda_device_launch_op; + unsigned g_kernel_launch_latency; // backward pointer class gpgpu_context* gpgpu_ctx; #if (CUDART_VERSION >= 5000) diff --git a/src/gpgpu-sim/gpu-sim.cc b/src/gpgpu-sim/gpu-sim.cc index 93f041a..9726cf5 100644 --- a/src/gpgpu-sim/gpu-sim.cc +++ b/src/gpgpu-sim/gpu-sim.cc @@ -548,9 +548,8 @@ void gpgpu_sim_config::reg_options(option_parser_t opp) ptx_file_line_stats_options(opp); //Jin: kernel launch latency - extern unsigned g_kernel_launch_latency; option_parser_register(opp, "-gpgpu_kernel_launch_latency", OPT_INT32, - &g_kernel_launch_latency, "Kernel launch latency in cycles. Default: 0", + &(gpgpu_ctx->device_runtime->g_kernel_launch_latency), "Kernel launch latency in cycles. Default: 0", "0"); extern bool g_cdp_enabled; option_parser_register(opp, "-gpgpu_cdp_enabled", OPT_BOOL, -- cgit v1.3 From 524c7e2ff37bba3ec60706605b34c6e354f928f1 Mon Sep 17 00:00:00 2001 From: Mengchi Zhang Date: Fri, 12 Jul 2019 00:02:57 -0400 Subject: Move g_max_total_param_size Signed-off-by: Mengchi Zhang --- src/cuda-sim/cuda_device_runtime.cc | 1 - src/cuda-sim/cuda_device_runtime.h | 2 ++ src/gpgpu-sim/gpu-sim.cc | 3 +-- 3 files changed, 3 insertions(+), 3 deletions(-) (limited to 'src/gpgpu-sim/gpu-sim.cc') diff --git a/src/cuda-sim/cuda_device_runtime.cc b/src/cuda-sim/cuda_device_runtime.cc index 326a1c5..dc3adc3 100644 --- a/src/cuda-sim/cuda_device_runtime.cc +++ b/src/cuda-sim/cuda_device_runtime.cc @@ -5,7 +5,6 @@ #include #include -unsigned long long g_max_total_param_size = 0; #if (CUDART_VERSION >= 5000) diff --git a/src/cuda-sim/cuda_device_runtime.h b/src/cuda-sim/cuda_device_runtime.h index 5927b54..f37849b 100644 --- a/src/cuda-sim/cuda_device_runtime.h +++ b/src/cuda-sim/cuda_device_runtime.h @@ -43,12 +43,14 @@ class cuda_device_runtime { public: cuda_device_runtime( gpgpu_context* ctx ) { g_total_param_size = 0; + g_max_total_param_size = 0; gpgpu_ctx = ctx; } unsigned long long g_total_param_size; std::map g_cuda_device_launch_param_map; std::list g_cuda_device_launch_op; unsigned g_kernel_launch_latency; + unsigned long long g_max_total_param_size; // backward pointer class gpgpu_context* gpgpu_ctx; #if (CUDART_VERSION >= 5000) diff --git a/src/gpgpu-sim/gpu-sim.cc b/src/gpgpu-sim/gpu-sim.cc index 9726cf5..a3d6a8a 100644 --- a/src/gpgpu-sim/gpu-sim.cc +++ b/src/gpgpu-sim/gpu-sim.cc @@ -1106,8 +1106,7 @@ void gpgpu_sim::gpu_print_stat() printf("gpu_tot_occupancy = %.4f\% \n", (gpu_occupancy + gpu_tot_occupancy).get_occ_fraction() * 100); - extern unsigned long long g_max_total_param_size; - fprintf(statfout, "max_total_param_size = %llu\n", g_max_total_param_size); + fprintf(statfout, "max_total_param_size = %llu\n", gpgpu_ctx->device_runtime->g_max_total_param_size); // performance counter for stalls due to congestion. printf("gpu_stall_dramfull = %d\n", gpu_stall_dramfull); -- cgit v1.3 From 4671280cfe7252bf875d071e667f57064250a1b7 Mon Sep 17 00:00:00 2001 From: Mengchi Zhang Date: Fri, 12 Jul 2019 00:09:29 -0400 Subject: Move g_cdp_enabled Signed-off-by: Mengchi Zhang --- libcuda/cuda_api_object.h | 1 + libcuda/cuda_runtime_api.cc | 11 ++++------- src/abstract_hardware_model.cc | 3 --- src/cuda-sim/cuda_device_runtime.h | 2 ++ src/cuda-sim/ptx_loader.cc | 9 +++------ src/gpgpu-sim/gpu-sim.cc | 3 +-- 6 files changed, 11 insertions(+), 18 deletions(-) (limited to 'src/gpgpu-sim/gpu-sim.cc') diff --git a/libcuda/cuda_api_object.h b/libcuda/cuda_api_object.h index db5e6a4..51416f2 100644 --- a/libcuda/cuda_api_object.h +++ b/libcuda/cuda_api_object.h @@ -199,6 +199,7 @@ class cuda_runtime_api { std::list mergeSections(); cuobjdumpELFSection* findELFSection(const std::string identifier); cuobjdumpPTXSection* findPTXSection(const std::string identifier); + cuobjdumpPTXSection* findPTXSectionInList(std::list §ionlist, const std::string identifier); void cuobjdumpRegisterFatBinary(unsigned int handle, const char* filename, CUctx_st *context); kernel_info_t *gpgpu_cuda_ptx_sim_init_grid( const char *kernel_key, gpgpu_ptx_sim_arg_list_t args, diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc index 59d2a60..10a651a 100644 --- a/libcuda/cuda_runtime_api.cc +++ b/libcuda/cuda_runtime_api.cc @@ -2421,7 +2421,6 @@ __host__ cudaError_t CUDARTAPI cudaGetExportTable(const void **ppExportTable, co //extracts all ptx files from binary and dumps into prog_name.unique_no.sm_<>.ptx files void cuda_runtime_api::extract_ptx_files_using_cuobjdump(CUctx_st *context){ - extern bool g_cdp_enabled; char command[1000]; char *pytorch_bin = getenv("PYTORCH_BIN"); std::string app_binary = get_app_binary(); @@ -2442,7 +2441,7 @@ void cuda_runtime_api::extract_ptx_files_using_cuobjdump(CUctx_st *context){ printf("WARNING: Failed to execute cuobjdump to get list of ptx files \n"); exit(0); } - if(!g_cdp_enabled) { + if(!gpgpu_ctx->device_runtime->g_cdp_enabled) { //based on the list above, dump ptx files individually. Format of dumped ptx file is prog_name.unique_no.sm_<>.ptx std::ifstream infile(ptx_list_file_name); @@ -2515,7 +2514,6 @@ void cuda_runtime_api::extract_code_using_cuobjdump(){ } // Running cuobjdump using dynamic link to current process // Needs the option '-all' to extract PTX from CDP-enabled binary - extern bool g_cdp_enabled; //dump ptx for all individial ptx files into sepearte files which is later used by ptxas. int result=0; @@ -2530,7 +2528,7 @@ void cuda_runtime_api::extract_code_using_cuobjdump(){ snprintf(fname,1024,"_cuobjdump_complete_output_XXXXXX"); int fd=mkstemp(fname); close(fd); - if(!g_cdp_enabled) + if(!gpgpu_ctx->device_runtime->g_cdp_enabled) snprintf(command,1000,"$CUDA_INSTALL_PATH/bin/cuobjdump -ptx -elf -sass %s > %s", app_binary.c_str(), fname); else snprintf(command,1000,"$CUDA_INSTALL_PATH/bin/cuobjdump -ptx -elf -sass -all %s > %s", app_binary.c_str(), fname); @@ -2822,7 +2820,7 @@ cuobjdumpELFSection* cuda_runtime_api::findELFSection(const std::string identifi } //! Within the section list, find the PTX section corresponding to a given identifier -cuobjdumpPTXSection* findPTXSectionInList(std::list §ionlist, const std::string identifier){ +cuobjdumpPTXSection* cuda_runtime_api::findPTXSectionInList(std::list §ionlist, const std::string identifier){ std::list::iterator iter; for ( iter = sectionlist.begin(); iter != sectionlist.end(); @@ -2833,8 +2831,7 @@ cuobjdumpPTXSection* findPTXSectionInList(std::list §ionl if(ptxsection->getIdentifier() == identifier) return ptxsection; else { - extern bool g_cdp_enabled; - if(g_cdp_enabled) { + if(gpgpu_ctx->device_runtime->g_cdp_enabled) { printf("Warning: __cudaRegisterFatBinary needs %s, but find PTX section with %s\n", identifier.c_str(), ptxsection->getIdentifier().c_str()); return ptxsection; diff --git a/src/abstract_hardware_model.cc b/src/abstract_hardware_model.cc index 450087b..fde7874 100644 --- a/src/abstract_hardware_model.cc +++ b/src/abstract_hardware_model.cc @@ -691,9 +691,6 @@ void warp_inst_t::completed( unsigned long long cycle ) const ptx_file_line_stats_add_latency(pc, latency * active_count()); } -//Jin: CDP support -bool g_cdp_enabled; - unsigned kernel_info_t::m_next_uid = 1; /*A snapshot of the texture mappings needs to be stored in the kernel's info as diff --git a/src/cuda-sim/cuda_device_runtime.h b/src/cuda-sim/cuda_device_runtime.h index f37849b..7f7a0ca 100644 --- a/src/cuda-sim/cuda_device_runtime.h +++ b/src/cuda-sim/cuda_device_runtime.h @@ -51,6 +51,8 @@ class cuda_device_runtime { std::list g_cuda_device_launch_op; unsigned g_kernel_launch_latency; unsigned long long g_max_total_param_size; + bool g_cdp_enabled; + // backward pointer class gpgpu_context* gpgpu_ctx; #if (CUDART_VERSION >= 5000) diff --git a/src/cuda-sim/ptx_loader.cc b/src/cuda-sim/ptx_loader.cc index 6e36a62..dca3cec 100644 --- a/src/cuda-sim/ptx_loader.cc +++ b/src/cuda-sim/ptx_loader.cc @@ -330,8 +330,7 @@ void gpgpu_context::gpgpu_ptx_info_load_from_filename( const char *filename, uns std::string ptxas_filename(std::string(filename) + "as"); char buff[1024], extra_flags[1024]; extra_flags[0]=0; - extern bool g_cdp_enabled; - if(!g_cdp_enabled) + if(!device_runtime->g_cdp_enabled) snprintf(extra_flags,1024,"--gpu-name=sm_%u",sm_version); else snprintf(extra_flags,1024,"--compile-only --gpu-name=sm_%u",sm_version); @@ -398,8 +397,7 @@ void gpgpu_context::gpgpu_ptxinfo_load_from_string( const char *p_for_info, unsi "A register size/SM mismatch may result in occupancy differences." ); exit(1); } - extern bool g_cdp_enabled; - if(!g_cdp_enabled) + if(!device_runtime->g_cdp_enabled) snprintf(extra_flags,1024,"--gpu-name=sm_%u", g_occupancy_sm_number); else snprintf(extra_flags,1024,"--compile-only --gpu-name=sm_%u",g_occupancy_sm_number); @@ -467,8 +465,7 @@ void gpgpu_context::gpgpu_ptxinfo_load_from_string( const char *p_for_info, unsi #if CUDART_VERSION >= 3000 if (sm_version == 0) sm_version = 20; - extern bool g_cdp_enabled; - if(!g_cdp_enabled) + if(!device_runtime->g_cdp_enabled) snprintf(extra_flags,1024,"--gpu-name=sm_%u",sm_version); else snprintf(extra_flags,1024,"--compile-only --gpu-name=sm_%u",sm_version); diff --git a/src/gpgpu-sim/gpu-sim.cc b/src/gpgpu-sim/gpu-sim.cc index a3d6a8a..0481259 100644 --- a/src/gpgpu-sim/gpu-sim.cc +++ b/src/gpgpu-sim/gpu-sim.cc @@ -551,9 +551,8 @@ void gpgpu_sim_config::reg_options(option_parser_t opp) option_parser_register(opp, "-gpgpu_kernel_launch_latency", OPT_INT32, &(gpgpu_ctx->device_runtime->g_kernel_launch_latency), "Kernel launch latency in cycles. Default: 0", "0"); - extern bool g_cdp_enabled; option_parser_register(opp, "-gpgpu_cdp_enabled", OPT_BOOL, - &g_cdp_enabled, "Turn on CDP", + &(gpgpu_ctx->device_runtime->g_cdp_enabled), "Turn on CDP", "0"); } -- cgit v1.3 From 23c0bb224295dde9651fd915d854e4f7eafdf88f Mon Sep 17 00:00:00 2001 From: Mengchi Zhang Date: Sun, 14 Jul 2019 11:25:42 -0400 Subject: Move sm_next_access_uid Signed-off-by: Mengchi Zhang --- libcuda/gpgpu_context.h | 2 ++ src/abstract_hardware_model.cc | 13 ++++++++++--- src/abstract_hardware_model.h | 22 +++++++++------------- src/gpgpu-sim/dram.cc | 2 +- src/gpgpu-sim/dram.h | 5 +++-- src/gpgpu-sim/gpu-cache.cc | 6 ++++-- src/gpgpu-sim/gpu-sim.cc | 2 +- src/gpgpu-sim/gpu-sim.h | 15 +++++++++------ src/gpgpu-sim/l2cache.cc | 11 ++++++----- src/gpgpu-sim/l2cache.h | 12 ++++++------ src/gpgpu-sim/mem_fetch.cc | 4 ++-- src/gpgpu-sim/mem_fetch.h | 5 +++-- src/gpgpu-sim/mem_latency_stat.cc | 2 +- src/gpgpu-sim/mem_latency_stat.h | 5 +++-- src/gpgpu-sim/power_stat.cc | 4 ++-- src/gpgpu-sim/power_stat.h | 6 +++--- src/gpgpu-sim/shader.cc | 20 ++++++++++++++++---- src/gpgpu-sim/shader.h | 20 ++++---------------- 18 files changed, 85 insertions(+), 71 deletions(-) (limited to 'src/gpgpu-sim/gpu-sim.cc') diff --git a/libcuda/gpgpu_context.h b/libcuda/gpgpu_context.h index c0b250a..07473be 100644 --- a/libcuda/gpgpu_context.h +++ b/libcuda/gpgpu_context.h @@ -11,6 +11,7 @@ class gpgpu_context { public: gpgpu_context() { g_global_allfiles_symbol_table = NULL; + sm_next_access_uid=0; api = new cuda_runtime_api(this); ptxinfo = new ptxinfo_data(this); ptx_parser = new ptx_recognizer(this); @@ -21,6 +22,7 @@ class gpgpu_context { // global list symbol_table *g_global_allfiles_symbol_table; const char *g_filename; + unsigned sm_next_access_uid; // objects pointers for each file cuda_runtime_api* api; ptxinfo_data* ptxinfo; diff --git a/src/abstract_hardware_model.cc b/src/abstract_hardware_model.cc index fde7874..fe10daa 100644 --- a/src/abstract_hardware_model.cc +++ b/src/abstract_hardware_model.cc @@ -41,9 +41,16 @@ #include #include "../libcuda/gpgpu_context.h" -unsigned mem_access_t::sm_next_access_uid = 0; unsigned warp_inst_t::sm_next_uid = 0; +void mem_access_t::init(gpgpu_context* ctx) +{ + gpgpu_ctx = ctx; + m_uid=++(gpgpu_ctx->sm_next_access_uid); + m_addr=0; + m_req_size=0; +} + checkpoint::checkpoint() { @@ -449,7 +456,7 @@ void warp_inst_t::generate_mem_accesses() byte_mask.set(idx+i); } for( a=accesses.begin(); a != accesses.end(); ++a ) - m_accessq.push_back( mem_access_t(access_type,a->first,cache_block_size,is_write,a->second, byte_mask, mem_access_sector_mask_t())); + m_accessq.push_back( mem_access_t(access_type,a->first,cache_block_size,is_write,a->second, byte_mask, mem_access_sector_mask_t(), m_config->gpgpu_ctx)); } if ( space.get_type() == global_space ) { @@ -681,7 +688,7 @@ void warp_inst_t::memory_coalescing_arch_reduce_and_send( bool is_write, mem_acc assert(lower_half_used && upper_half_used); } } - m_accessq.push_back( mem_access_t(access_type,addr,size,is_write,info.active,info.bytes, info.chunks) ); + m_accessq.push_back( mem_access_t(access_type,addr,size,is_write,info.active,info.bytes, info.chunks,m_config->gpgpu_ctx) ); } void warp_inst_t::completed( unsigned long long cycle ) const diff --git a/src/abstract_hardware_model.h b/src/abstract_hardware_model.h index 8ef8376..a22c8c3 100644 --- a/src/abstract_hardware_model.h +++ b/src/abstract_hardware_model.h @@ -732,13 +732,14 @@ enum cache_operator_type { class mem_access_t { public: - mem_access_t() { init(); } + mem_access_t(gpgpu_context* ctx) { init(ctx); } mem_access_t( mem_access_type type, new_addr_type address, unsigned size, - bool wr ) + bool wr, + gpgpu_context* ctx) { - init(); + init(ctx); m_type = type; m_addr = address; m_req_size = size; @@ -750,10 +751,11 @@ public: bool wr, const active_mask_t &active_mask, const mem_access_byte_mask_t &byte_mask, - const mem_access_sector_mask_t §or_mask) + const mem_access_sector_mask_t §or_mask, + gpgpu_context* ctx) : m_warp_mask(active_mask), m_byte_mask(byte_mask), m_sector_mask(sector_mask) { - init(); + init(ctx); m_type = type; m_addr = address; m_req_size = size; @@ -786,13 +788,9 @@ public: } } + gpgpu_context* gpgpu_ctx; private: - void init() - { - m_uid=++sm_next_access_uid; - m_addr=0; - m_req_size=0; - } + void init(gpgpu_context* ctx); unsigned m_uid; new_addr_type m_addr; // request address @@ -802,8 +800,6 @@ private: active_mask_t m_warp_mask; mem_access_byte_mask_t m_byte_mask; mem_access_sector_mask_t m_sector_mask; - - static unsigned sm_next_access_uid; }; class mem_fetch; diff --git a/src/gpgpu-sim/dram.cc b/src/gpgpu-sim/dram.cc index 5e36d4b..d443d79 100644 --- a/src/gpgpu-sim/dram.cc +++ b/src/gpgpu-sim/dram.cc @@ -41,7 +41,7 @@ int PRINT_CYCLE = 0; template class fifo_pipeline; template class fifo_pipeline; -dram_t::dram_t( unsigned int partition_id, const struct memory_config *config, memory_stats_t *stats, +dram_t::dram_t( unsigned int partition_id, const memory_config *config, memory_stats_t *stats, memory_partition_unit *mp, gpgpu_sim* gpu ) { id = partition_id; diff --git a/src/gpgpu-sim/dram.h b/src/gpgpu-sim/dram.h index 7a3a2da..0bd9725 100644 --- a/src/gpgpu-sim/dram.h +++ b/src/gpgpu-sim/dram.h @@ -106,11 +106,12 @@ enum bank_grp_bits_position{ }; class mem_fetch; +class memory_config; class dram_t { public: - dram_t( unsigned int parition_id, const struct memory_config *config, class memory_stats_t *stats, + dram_t( unsigned int parition_id, const memory_config *config, class memory_stats_t *stats, class memory_partition_unit *mp, class gpgpu_sim* gpu ); bool full(bool is_write) const; @@ -145,7 +146,7 @@ public: - const struct memory_config *m_config; + const memory_config *m_config; private: bankgrp_t **bkgrp; diff --git a/src/gpgpu-sim/gpu-cache.cc b/src/gpgpu-sim/gpu-cache.cc index f1f6e19..1705821 100644 --- a/src/gpgpu-sim/gpu-cache.cc +++ b/src/gpgpu-sim/gpu-cache.cc @@ -1251,7 +1251,8 @@ data_cache::wr_miss_wa_naive( new_addr_type addr, false, // Now performing a read mf->get_access_warp_mask(), mf->get_access_byte_mask(), - mf->get_access_sector_mask()); + mf->get_access_sector_mask(), + m_gpu->gpgpu_ctx); mem_fetch *n_mf = new mem_fetch( *ma, NULL, @@ -1365,7 +1366,8 @@ data_cache::wr_miss_wa_fetch_on_write( new_addr_type addr, false, // Now performing a read mf->get_access_warp_mask(), mf->get_access_byte_mask(), - mf->get_access_sector_mask()); + mf->get_access_sector_mask(), + m_gpu->gpgpu_ctx); mem_fetch *n_mf = new mem_fetch( *ma, NULL, diff --git a/src/gpgpu-sim/gpu-sim.cc b/src/gpgpu-sim/gpu-sim.cc index 0481259..0644b44 100644 --- a/src/gpgpu-sim/gpu-sim.cc +++ b/src/gpgpu-sim/gpu-sim.cc @@ -1828,7 +1828,7 @@ const shader_core_config * gpgpu_sim::getShaderCoreConfig() return m_shader_config; } -const struct memory_config * gpgpu_sim::getMemoryConfig() +const memory_config * gpgpu_sim::getMemoryConfig() { return m_memory_config; } diff --git a/src/gpgpu-sim/gpu-sim.h b/src/gpgpu-sim/gpu-sim.h index 78f0505..19e1eb3 100644 --- a/src/gpgpu-sim/gpu-sim.h +++ b/src/gpgpu-sim/gpu-sim.h @@ -33,6 +33,7 @@ #include "../trace.h" #include "addrdec.h" #include "shader.h" +#include "gpu-cache.h" #include #include #include @@ -143,13 +144,14 @@ struct power_config { }; - -struct memory_config { - memory_config() +class memory_config { + public: + memory_config(gpgpu_context* ctx) { m_valid = false; gpgpu_dram_timing_opt=NULL; gpgpu_L2_queue_config=NULL; + gpgpu_ctx = ctx; } void init() { @@ -291,13 +293,14 @@ struct memory_config { unsigned write_high_watermark; unsigned write_low_watermark; bool m_perf_sim_memcpy; + gpgpu_context* gpgpu_ctx; }; extern bool g_interactive_debugger_enabled; class gpgpu_sim_config : public power_config, public gpgpu_functional_sim_config { public: - gpgpu_sim_config(gpgpu_context* ctx): m_shader_config(ctx) { + gpgpu_sim_config(gpgpu_context* ctx): m_shader_config(ctx), m_memory_config(ctx) { m_valid = false; gpgpu_ctx = ctx; } @@ -507,7 +510,7 @@ public: /*! * Returning the memory configuration of the shader core, used by the functional simulation only so far */ - const struct memory_config * getMemoryConfig(); + const memory_config * getMemoryConfig(); //! Get shader core SIMT cluster @@ -567,7 +570,7 @@ private: const struct cudaDeviceProp *m_cuda_properties; const shader_core_config *m_shader_config; - const struct memory_config *m_memory_config; + const memory_config *m_memory_config; // stats class shader_core_stats *m_shader_stats; diff --git a/src/gpgpu-sim/l2cache.cc b/src/gpgpu-sim/l2cache.cc index 62e70a7..6540b52 100644 --- a/src/gpgpu-sim/l2cache.cc +++ b/src/gpgpu-sim/l2cache.cc @@ -49,7 +49,7 @@ mem_fetch * partition_mf_allocator::alloc(new_addr_type addr, mem_access_type type, unsigned size, bool wr, unsigned long long cycle ) const { assert( wr ); - mem_access_t access( type, addr, size, wr ); + mem_access_t access( type, addr, size, wr, m_memory_config->gpgpu_ctx ); mem_fetch *mf = new mem_fetch( access, NULL, WRITE_PACKET_SIZE, @@ -62,7 +62,7 @@ mem_fetch * partition_mf_allocator::alloc(new_addr_type addr, mem_access_type ty } memory_partition_unit::memory_partition_unit( unsigned partition_id, - const struct memory_config *config, + const memory_config *config, class memory_stats_t *stats, class gpgpu_sim* gpu) : m_id(partition_id), m_config(config), m_stats(stats), m_arbitration_metadata(config), m_gpu(gpu) @@ -95,7 +95,7 @@ memory_partition_unit::~memory_partition_unit() delete[] m_sub_partition; } -memory_partition_unit::arbitration_metadata::arbitration_metadata(const struct memory_config *config) +memory_partition_unit::arbitration_metadata::arbitration_metadata(const memory_config *config) : m_last_borrower(config->m_n_sub_partition_per_memory_channel - 1), m_private_credit(config->m_n_sub_partition_per_memory_channel, 0), m_shared_credit(0) @@ -312,7 +312,7 @@ void memory_partition_unit::print( FILE *fp ) const } memory_sub_partition::memory_sub_partition( unsigned sub_partition_id, - const struct memory_config *config, + const memory_config *config, class memory_stats_t *stats, class gpgpu_sim* gpu) { @@ -640,7 +640,8 @@ std::vector memory_sub_partition::breakdown_request_to_sector_reques mf->is_write(), mf->get_access_warp_mask(), mf->get_access_byte_mask() & byte_sector_mask, - std::bitset().set(j)); + std::bitset().set(j), + m_gpu->gpgpu_ctx); mem_fetch *n_mf = new mem_fetch( *ma, NULL, diff --git a/src/gpgpu-sim/l2cache.h b/src/gpgpu-sim/l2cache.h index 8ff2666..1f74c47 100644 --- a/src/gpgpu-sim/l2cache.h +++ b/src/gpgpu-sim/l2cache.h @@ -58,7 +58,7 @@ private: class memory_partition_unit { public: - memory_partition_unit( unsigned partition_id, const struct memory_config *config, class memory_stats_t *stats, class gpgpu_sim* gpu ); + memory_partition_unit( unsigned partition_id, const memory_config *config, class memory_stats_t *stats, class gpgpu_sim* gpu ); ~memory_partition_unit(); bool busy() const; @@ -98,7 +98,7 @@ public: private: unsigned m_id; - const struct memory_config *m_config; + const memory_config *m_config; class memory_stats_t *m_stats; class memory_sub_partition **m_sub_partition; class dram_t *m_dram; @@ -106,7 +106,7 @@ private: class arbitration_metadata { public: - arbitration_metadata(const struct memory_config *config); + arbitration_metadata(const memory_config *config); // check if a subpartition still has credit bool has_credits(int inner_sub_partition_id) const; @@ -130,7 +130,7 @@ private: std::vector m_private_credit; int m_shared_credit; }; - arbitration_metadata m_arbitration_metadata; + arbitration_metadata m_arbitration_metadata; // determine wheither a given subpartition can issue to DRAM bool can_issue_to_dram(int inner_sub_partition_id); @@ -149,7 +149,7 @@ private: class memory_sub_partition { public: - memory_sub_partition( unsigned sub_partition_id, const struct memory_config *config, class memory_stats_t *stats, class gpgpu_sim* gpu ); + memory_sub_partition( unsigned sub_partition_id, const memory_config *config, class memory_stats_t *stats, class gpgpu_sim* gpu ); ~memory_sub_partition(); unsigned get_id() const { return m_id; } @@ -197,7 +197,7 @@ public: private: // data unsigned m_id; //< the global sub partition ID - const struct memory_config *m_config; + const memory_config *m_config; class l2_cache *m_L2cache; class L2interface *m_L2interface; class gpgpu_sim* m_gpu; diff --git a/src/gpgpu-sim/mem_fetch.cc b/src/gpgpu-sim/mem_fetch.cc index c9b0484..6a00889 100644 --- a/src/gpgpu-sim/mem_fetch.cc +++ b/src/gpgpu-sim/mem_fetch.cc @@ -39,10 +39,10 @@ mem_fetch::mem_fetch( const mem_access_t &access, unsigned wid, unsigned sid, unsigned tpc, - const struct memory_config *config, + const memory_config *config, unsigned long long cycle, mem_fetch *m_original_mf, - mem_fetch *m_original_wr_mf) + mem_fetch *m_original_wr_mf):m_access(access) { m_request_uid = sm_next_mf_request_uid++; diff --git a/src/gpgpu-sim/mem_fetch.h b/src/gpgpu-sim/mem_fetch.h index 4eb3a52..1cab9f2 100644 --- a/src/gpgpu-sim/mem_fetch.h +++ b/src/gpgpu-sim/mem_fetch.h @@ -47,6 +47,7 @@ enum mf_type { #undef MF_TUP #undef MF_TUP_END +class memory_config; class mem_fetch { public: mem_fetch( const mem_access_t &access, @@ -55,7 +56,7 @@ public: unsigned wid, unsigned sid, unsigned tpc, - const struct memory_config *config, + const memory_config *config, unsigned long long cycle, mem_fetch *original_mf = NULL, mem_fetch *original_wr_mf = NULL); @@ -149,7 +150,7 @@ private: static unsigned sm_next_mf_request_uid; - const struct memory_config *m_mem_config; + const memory_config *m_mem_config; unsigned icnt_flit_size; mem_fetch* original_mf; //this pointer is set up when a request is divided into sector requests at L2 cache (if the req size > L2 sector size), so the pointer refers to the original request diff --git a/src/gpgpu-sim/mem_latency_stat.cc b/src/gpgpu-sim/mem_latency_stat.cc index d08ba39..4e94991 100644 --- a/src/gpgpu-sim/mem_latency_stat.cc +++ b/src/gpgpu-sim/mem_latency_stat.cc @@ -42,7 +42,7 @@ #include #include -memory_stats_t::memory_stats_t( unsigned n_shader, const shader_core_config *shader_config, const struct memory_config *mem_config, const class gpgpu_sim* gpu ) +memory_stats_t::memory_stats_t( unsigned n_shader, const shader_core_config *shader_config, const memory_config *mem_config, const class gpgpu_sim* gpu ) { assert( mem_config->m_valid ); assert( shader_config->m_valid ); diff --git a/src/gpgpu-sim/mem_latency_stat.h b/src/gpgpu-sim/mem_latency_stat.h index 6ce568d..0c84972 100644 --- a/src/gpgpu-sim/mem_latency_stat.h +++ b/src/gpgpu-sim/mem_latency_stat.h @@ -32,11 +32,12 @@ #include #include +class memory_config; class memory_stats_t { public: memory_stats_t( unsigned n_shader, const class shader_core_config *shader_config, - const struct memory_config *mem_config, + const memory_config *mem_config, const class gpgpu_sim* gpu); unsigned memlatstat_done( class mem_fetch *mf ); @@ -54,7 +55,7 @@ public: unsigned m_n_shader; const shader_core_config *m_shader_config; - const struct memory_config *m_memory_config; + const memory_config *m_memory_config; const class gpgpu_sim* m_gpu; unsigned max_mrq_latency; diff --git a/src/gpgpu-sim/power_stat.cc b/src/gpgpu-sim/power_stat.cc index 007b4c6..2c02082 100644 --- a/src/gpgpu-sim/power_stat.cc +++ b/src/gpgpu-sim/power_stat.cc @@ -42,7 +42,7 @@ -power_mem_stat_t::power_mem_stat_t(const struct memory_config *mem_config, const shader_core_config *shdr_config, memory_stats_t *mem_stats, shader_core_stats *shdr_stats){ +power_mem_stat_t::power_mem_stat_t(const memory_config *mem_config, const shader_core_config *shdr_config, memory_stats_t *mem_stats, shader_core_stats *shdr_stats){ assert( mem_config->m_valid ); m_mem_stats = mem_stats; m_config = mem_config; @@ -266,7 +266,7 @@ for(unsigned i=0; inum_shader(); ++i){ } } -power_stat_t::power_stat_t( const shader_core_config *shader_config,float * average_pipeline_duty_cycle,float *active_sms,shader_core_stats * shader_stats, const struct memory_config *mem_config,memory_stats_t * memory_stats) +power_stat_t::power_stat_t( const shader_core_config *shader_config,float * average_pipeline_duty_cycle,float *active_sms,shader_core_stats * shader_stats, const memory_config *mem_config,memory_stats_t * memory_stats) { assert( shader_config->m_valid ); assert( mem_config->m_valid ); diff --git a/src/gpgpu-sim/power_stat.h b/src/gpgpu-sim/power_stat.h index 91fade9..24ade99 100644 --- a/src/gpgpu-sim/power_stat.h +++ b/src/gpgpu-sim/power_stat.h @@ -113,7 +113,7 @@ struct mem_power_stats_pod{ class power_mem_stat_t : public mem_power_stats_pod{ public: - power_mem_stat_t(const struct memory_config *mem_config, const shader_core_config *shdr_config, memory_stats_t *mem_stats, shader_core_stats *shdr_stats); + power_mem_stat_t(const memory_config *mem_config, const shader_core_config *shdr_config, memory_stats_t *mem_stats, shader_core_stats *shdr_stats); void visualizer_print( gzFile visualizer_file ); void print (FILE *fout) const; void init(); @@ -128,7 +128,7 @@ private: class power_stat_t { public: - power_stat_t( const shader_core_config *shader_config,float * average_pipeline_duty_cycle,float * active_sms,shader_core_stats * shader_stats, const struct memory_config *mem_config,memory_stats_t * memory_stats); + power_stat_t( const shader_core_config *shader_config,float * average_pipeline_duty_cycle,float * active_sms,shader_core_stats * shader_stats, const memory_config *mem_config,memory_stats_t * memory_stats); void visualizer_print( gzFile visualizer_file ); void print (FILE *fout) const; void save_stats(){ @@ -621,7 +621,7 @@ public: float * m_average_pipeline_duty_cycle; float * m_active_sms; const shader_core_config *m_config; - const struct memory_config *m_mem_config; + const memory_config *m_mem_config; }; diff --git a/src/gpgpu-sim/shader.cc b/src/gpgpu-sim/shader.cc index f380560..b7ae95d 100644 --- a/src/gpgpu-sim/shader.cc +++ b/src/gpgpu-sim/shader.cc @@ -28,7 +28,6 @@ #include #include "shader.h" -#include "gpu-sim.h" #include "addrdec.h" #include "dram.h" #include "stat-tool.h" @@ -53,6 +52,19 @@ #define MIN(a,b) (((a)<(b))?(a):(b)) +mem_fetch *shader_core_mem_fetch_allocator::alloc( new_addr_type addr, mem_access_type type, unsigned size, bool wr, unsigned long long cycle ) const +{ + mem_access_t access( type, addr, size, wr, m_memory_config->gpgpu_ctx); + mem_fetch *mf = new mem_fetch( access, + NULL, + wr?WRITE_PACKET_SIZE:READ_PACKET_SIZE, + -1, + m_core_id, + m_cluster_id, + m_memory_config, + cycle); + return mf; +} ///////////////////////////////////////////////////////////////////////////// std::list shader_core_ctx::get_regs_written( const inst_t &fvt ) const @@ -71,7 +83,7 @@ shader_core_ctx::shader_core_ctx( class gpgpu_sim *gpu, unsigned shader_id, unsigned tpc_id, const shader_core_config *config, - const struct memory_config *mem_config, + const memory_config *mem_config, shader_core_stats *stats ) : core_t( gpu, NULL, config->warp_size, config->n_thread_per_shader ), m_barriers( this, config->max_warps_per_shader, config->max_cta_per_core, config->max_barriers_per_cta, config->warp_size ), @@ -809,7 +821,7 @@ void shader_core_ctx::fetch() // TODO: replace with use of allocator // mem_fetch *mf = m_mem_fetch_allocator->alloc() - mem_access_t acc(INST_ACC_R,ppc,nbytes,false); + mem_access_t acc(INST_ACC_R,ppc,nbytes,false, m_gpu->gpgpu_ctx); mem_fetch *mf = new mem_fetch(acc, NULL/*we don't have an instruction yet*/, READ_PACKET_SIZE, @@ -3787,7 +3799,7 @@ void opndcoll_rfu_t::collector_unit_t::dispatch() simt_core_cluster::simt_core_cluster( class gpgpu_sim *gpu, unsigned cluster_id, const shader_core_config *config, - const struct memory_config *mem_config, + const memory_config *mem_config, shader_core_stats *stats, class memory_stats_t *mstats ) { diff --git a/src/gpgpu-sim/shader.h b/src/gpgpu-sim/shader.h index 2837f1b..b0d7f7f 100644 --- a/src/gpgpu-sim/shader.h +++ b/src/gpgpu-sim/shader.h @@ -1727,6 +1727,7 @@ private: friend class LooseRoundRobbinScheduler; }; +class memory_config; class shader_core_mem_fetch_allocator : public mem_fetch_allocator { public: shader_core_mem_fetch_allocator( unsigned core_id, unsigned cluster_id, const memory_config *config ) @@ -1735,20 +1736,7 @@ public: m_cluster_id = cluster_id; m_memory_config = config; } - mem_fetch *alloc( new_addr_type addr, mem_access_type type, unsigned size, bool wr, unsigned long long cycle ) const - { - mem_access_t access( type, addr, size, wr ); - mem_fetch *mf = new mem_fetch( access, - NULL, - wr?WRITE_PACKET_SIZE:READ_PACKET_SIZE, - -1, - m_core_id, - m_cluster_id, - m_memory_config, - cycle); - return mf; - } - + mem_fetch *alloc( new_addr_type addr, mem_access_type type, unsigned size, bool wr, unsigned long long cycle ) const; mem_fetch *alloc( const warp_inst_t &inst, const mem_access_t &access, unsigned long long cycle ) const { warp_inst_t inst_copy = inst; @@ -1777,7 +1765,7 @@ public: unsigned shader_id, unsigned tpc_id, const shader_core_config *config, - const struct memory_config *mem_config, + const memory_config *mem_config, shader_core_stats *stats ); // used by simt_core_cluster: @@ -2072,7 +2060,7 @@ public: simt_core_cluster( class gpgpu_sim *gpu, unsigned cluster_id, const shader_core_config *config, - const struct memory_config *mem_config, + const memory_config *mem_config, shader_core_stats *stats, memory_stats_t *mstats ); -- cgit v1.3 From 36aaec0a30445533f04aaf6b55c5ef135e20484a Mon Sep 17 00:00:00 2001 From: Mengchi Zhang Date: Mon, 15 Jul 2019 13:31:36 -0400 Subject: Move ptx_line_stats_filename and enable_ptx_file_line_stats Signed-off-by: Mengchi Zhang --- libcuda/gpgpu_context.h | 3 +++ src/cuda-sim/ptx-stats.cc | 9 +++------ src/cuda-sim/ptx-stats.h | 23 ++++++++++++++++------- src/gpgpu-sim/gpu-sim.cc | 4 ++-- 4 files changed, 24 insertions(+), 15 deletions(-) (limited to 'src/gpgpu-sim/gpu-sim.cc') diff --git a/libcuda/gpgpu_context.h b/libcuda/gpgpu_context.h index 1e20c62..337ebb2 100644 --- a/libcuda/gpgpu_context.h +++ b/libcuda/gpgpu_context.h @@ -6,6 +6,7 @@ #include "../src/gpgpusim_entrypoint.h" #include "../src/cuda-sim/cuda-sim.h" #include "../src/cuda-sim/cuda_device_runtime.h" +#include "../src/cuda-sim/ptx-stats.h" class gpgpu_context { public: @@ -21,6 +22,7 @@ class gpgpu_context { the_gpgpusim = new GPGPUsim_ctx(this); func_sim = new cuda_sim(this); device_runtime = new cuda_device_runtime(this); + stats = new ptx_stats(this); } // global list symbol_table *g_global_allfiles_symbol_table; @@ -36,6 +38,7 @@ class gpgpu_context { GPGPUsim_ctx* the_gpgpusim; cuda_sim* func_sim; cuda_device_runtime* device_runtime; + ptx_stats* stats; // member function list void cuobjdumpParseBinary(unsigned int handle); class symbol_table *gpgpu_ptx_sim_load_ptx_from_string( const char *p, unsigned source_num ); diff --git a/src/cuda-sim/ptx-stats.cc b/src/cuda-sim/ptx-stats.cc index 0a9f08d..298729f 100644 --- a/src/cuda-sim/ptx-stats.cc +++ b/src/cuda-sim/ptx-stats.cc @@ -32,12 +32,9 @@ #include #include #include "../tr1_hash_map.h" +#include "../../libcuda/gpgpu_context.h" -// options -bool enable_ptx_file_line_stats; -char * ptx_line_stats_filename = NULL; - -void ptx_file_line_stats_options(option_parser_t opp) +void ptx_stats::ptx_file_line_stats_options(option_parser_t opp) { option_parser_register(opp, "-enable_ptx_file_line_stats", OPT_BOOL, &enable_ptx_file_line_stats, @@ -118,7 +115,7 @@ typedef tr1_hash_map ptx static ptx_file_line_stats_map_t ptx_file_line_stats_tracker; // output statistics to a file -void ptx_file_line_stats_write_file() +void ptx_stats::ptx_file_line_stats_write_file() { // check if stat collection is turned on if (enable_ptx_file_line_stats == 0) return; diff --git a/src/cuda-sim/ptx-stats.h b/src/cuda-sim/ptx-stats.h index 4fc6599..c75fc58 100644 --- a/src/cuda-sim/ptx-stats.h +++ b/src/cuda-sim/ptx-stats.h @@ -29,13 +29,6 @@ #include "../option_parser.h" -extern bool enable_ptx_file_line_stats; - -// set options -void ptx_file_line_stats_options(option_parser_t opp); - -// output stats to a file -void ptx_file_line_stats_write_file(); #ifdef __cplusplus // stat collection interface to cuda-sim @@ -56,3 +49,19 @@ void ptx_file_line_stats_commit_exposed_latency(int sc_id, int exposed_latency); void ptx_file_line_stats_add_warp_divergence(unsigned pc, unsigned n_way_divergence); +class gpgpu_context; +class ptx_stats { + public: + ptx_stats(gpgpu_context* ctx) { + ptx_line_stats_filename = NULL; + gpgpu_ctx = ctx; + } + char * ptx_line_stats_filename; + bool enable_ptx_file_line_stats; + gpgpu_context* gpgpu_ctx; + // set options + void ptx_file_line_stats_options(option_parser_t opp); + + // output stats to a file + void ptx_file_line_stats_write_file(); +}; diff --git a/src/gpgpu-sim/gpu-sim.cc b/src/gpgpu-sim/gpu-sim.cc index 0644b44..bdf989a 100644 --- a/src/gpgpu-sim/gpu-sim.cc +++ b/src/gpgpu-sim/gpu-sim.cc @@ -545,7 +545,7 @@ void gpgpu_sim_config::reg_options(option_parser_t opp) option_parser_register(opp, "-trace_sampling_memory_partition", OPT_INT32, &Trace::sampling_memory_partition, "The memory partition which is printed using MEMPART_DPRINTF. Default -1 (i.e. all)", "-1"); - ptx_file_line_stats_options(opp); + gpgpu_ctx->stats->ptx_file_line_stats_options(opp); //Jin: kernel launch latency option_parser_register(opp, "-gpgpu_kernel_launch_latency", OPT_INT32, @@ -932,7 +932,7 @@ void gpgpu_sim::update_stats() { void gpgpu_sim::print_stats() { - ptx_file_line_stats_write_file(); + gpgpu_ctx->stats->ptx_file_line_stats_write_file(); gpu_print_stat(); if (g_network_mode) { -- cgit v1.3 From 352d2a3336b1c8e5258ca9d92d214973e98837c0 Mon Sep 17 00:00:00 2001 From: Mengchi Zhang Date: Mon, 15 Jul 2019 16:17:07 -0400 Subject: Move s_g_pc_to_insn Signed-off-by: Mengchi Zhang --- libcuda/gpgpu_context.h | 5 +++++ src/abstract_hardware_model.cc | 10 +++++----- src/cuda-sim/cuda-sim.cc | 20 +++++++++----------- src/cuda-sim/cuda-sim.h | 2 -- src/cuda-sim/ptx-stats.cc | 28 ++++++++++++++-------------- src/cuda-sim/ptx-stats.h | 15 ++++++++------- src/cuda-sim/ptx_ir.cc | 8 ++++++++ src/cuda-sim/ptx_ir.h | 9 --------- src/gpgpu-sim/gpu-sim.cc | 2 +- src/gpgpu-sim/mem_latency_stat.cc | 3 ++- src/gpgpu-sim/shader.cc | 4 ++-- src/gpgpu-sim/stat-tool.cc | 33 +++++++++++++++++---------------- src/gpgpu-sim/stat-tool.h | 12 +++++++----- 13 files changed, 78 insertions(+), 73 deletions(-) (limited to 'src/gpgpu-sim/gpu-sim.cc') diff --git a/libcuda/gpgpu_context.h b/libcuda/gpgpu_context.h index ed4f746..346a8a4 100644 --- a/libcuda/gpgpu_context.h +++ b/libcuda/gpgpu_context.h @@ -39,6 +39,8 @@ class gpgpu_context { unsigned long long g_ptx_cta_info_uid; unsigned symbol_sm_next_uid; //uid for symbol unsigned function_info_sm_next_uid; + std::vector s_g_pc_to_insn; // a direct mapping from PC to instruction + // objects pointers for each file cuda_runtime_api* api; ptxinfo_data* ptxinfo; @@ -58,6 +60,9 @@ class gpgpu_context { class gpgpu_sim *gpgpu_ptx_sim_init_perf(); struct _cuda_device_id *GPGPUSim_Init(); void ptx_reg_options(option_parser_t opp); + const ptx_instruction* pc_to_instruction(unsigned pc); + const warp_inst_t *ptx_fetch_inst( address_type pc ); + unsigned translate_pc_to_ptxlineno(unsigned pc); }; gpgpu_context* GPGPU_Context(); diff --git a/src/abstract_hardware_model.cc b/src/abstract_hardware_model.cc index 733d602..d8d5fbd 100644 --- a/src/abstract_hardware_model.cc +++ b/src/abstract_hardware_model.cc @@ -426,7 +426,7 @@ void warp_inst_t::generate_mem_accesses() } assert( total_accesses > 0 && total_accesses <= m_config->warp_size ); cycles = total_accesses; // shared memory conflicts modeled as larger initiation interval - ptx_file_line_stats_add_smem_bank_conflict( pc, total_accesses ); + m_config->gpgpu_ctx->stats->ptx_file_line_stats_add_smem_bank_conflict( pc, total_accesses ); break; } @@ -471,7 +471,7 @@ void warp_inst_t::generate_mem_accesses() } if ( space.get_type() == global_space ) { - ptx_file_line_stats_add_uncoalesced_gmem( pc, m_accessq.size() - starting_queue_size ); + m_config->gpgpu_ctx->stats->ptx_file_line_stats_add_uncoalesced_gmem( pc, m_accessq.size() - starting_queue_size ); } m_mem_accesses_created=true; } @@ -706,7 +706,7 @@ void warp_inst_t::completed( unsigned long long cycle ) const { unsigned long long latency = cycle - issue_cycle; assert(latency <= cycle); // underflow detection - ptx_file_line_stats_add_latency(pc, latency * active_count()); + m_config->gpgpu_ctx->stats->ptx_file_line_stats_add_latency(pc, latency * active_count()); } @@ -1110,7 +1110,7 @@ void simt_stack::update( simt_mask_t &thread_done, addr_vector_t &next_pc, addre if (warp_diverged) { - ptx_file_line_stats_add_warp_divergence(top_pc, 1); + m_gpu->gpgpu_ctx->stats->ptx_file_line_stats_add_warp_divergence(top_pc, 1); } } @@ -1157,7 +1157,7 @@ warp_inst_t core_t::getExecuteWarp(unsigned warpId) { unsigned pc,rpc; m_simt_stack[warpId]->get_pdom_stack_top_info(&pc,&rpc); - warp_inst_t wi= *ptx_fetch_inst(pc); + warp_inst_t wi= *(m_gpu->gpgpu_ctx->ptx_fetch_inst(pc)); wi.set_active(m_simt_stack[warpId]->get_active_mask()); return wi; } diff --git a/src/cuda-sim/cuda-sim.cc b/src/cuda-sim/cuda-sim.cc index c06f093..b9e6552 100644 --- a/src/cuda-sim/cuda-sim.cc +++ b/src/cuda-sim/cuda-sim.cc @@ -215,8 +215,6 @@ void gpgpu_t::gpgpu_ptx_sim_unbindTexture(const struct textureReference* texref) m_NameToTextureInfo.erase(texname); } -std::vector function_info::s_g_pc_to_insn; - #define MAX_INST_SIZE 8 /*bytes*/ void function_info::ptx_assemble() @@ -237,14 +235,14 @@ void function_info::ptx_assemble() 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); + gpgpu_ctx->s_g_pc_to_insn.push_back((ptx_instruction*)NULL); PC += PC%MAX_INST_SIZE; m_start_PC = PC; addr_t n=0; // offset in m_instr_mem //Why s_g_pc_to_insn.size() is needed to reserve additional memory for insts? reserve is cumulative. //s_g_pc_to_insn.reserve(s_g_pc_to_insn.size() + MAX_INST_SIZE*m_instructions.size()); - s_g_pc_to_insn.reserve(MAX_INST_SIZE*m_instructions.size()); + gpgpu_ctx->s_g_pc_to_insn.reserve(MAX_INST_SIZE*m_instructions.size()); for ( i=m_instructions.begin(); i != m_instructions.end(); i++ ) { ptx_instruction *pI = *i; if ( pI->is_label() ) { @@ -253,13 +251,13 @@ void function_info::ptx_assemble() } else { 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]); + gpgpu_ctx->s_g_pc_to_insn.push_back(pI); + assert(pI == gpgpu_ctx->s_g_pc_to_insn[PC]); pI->set_m_instr_mem_index(n); pI->set_PC(PC); assert( pI->inst_size() <= MAX_INST_SIZE ); for( unsigned i=1; i < pI->inst_size(); i++ ) { - s_g_pc_to_insn.push_back((ptx_instruction*)NULL); + gpgpu_ctx->s_g_pc_to_insn.push_back((ptx_instruction*)NULL); m_instr_mem[n+i]=NULL; } n += pI->inst_size(); @@ -1738,9 +1736,9 @@ const struct gpgpu_ptx_sim_info* ptx_sim_kernel_info(const function_info *kernel return kernel->get_kernel_info(); } -const warp_inst_t *ptx_fetch_inst( address_type pc ) +const warp_inst_t *gpgpu_context::ptx_fetch_inst( address_type pc ) { - return function_info::pc_to_instruction(pc); + return pc_to_instruction(pc); } unsigned ptx_sim_init_thread( kernel_info_t &kernel, @@ -2366,11 +2364,11 @@ void functionalCoreSim::executeWarp(unsigned i, bool &allAtBarrier, bool & someO if(!m_warpAtBarrier[i]&& m_liveThreadCount[i]>0) allAtBarrier = false; } -unsigned translate_pc_to_ptxlineno(unsigned pc) +unsigned gpgpu_context::translate_pc_to_ptxlineno(unsigned pc) { // this function assumes that the kernel fits inside a single PTX file // function_info *pFunc = g_func_info; // assume that the current kernel is the one in query - const ptx_instruction *pInsn = function_info::pc_to_instruction(pc); + const ptx_instruction *pInsn = pc_to_instruction(pc); unsigned ptx_line_number = pInsn->source_line(); return ptx_line_number; diff --git a/src/cuda-sim/cuda-sim.h b/src/cuda-sim/cuda-sim.h index 5bd4cb2..1be3d19 100644 --- a/src/cuda-sim/cuda-sim.h +++ b/src/cuda-sim/cuda-sim.h @@ -58,10 +58,8 @@ unsigned ptx_sim_init_thread( kernel_info_t &kernel, unsigned hw_warp_id, gpgpu_t *gpu, 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); - /*! * This class functionally executes a kernel. It uses the basic data structures and procedures in core_t */ diff --git a/src/cuda-sim/ptx-stats.cc b/src/cuda-sim/ptx-stats.cc index 298729f..22517df 100644 --- a/src/cuda-sim/ptx-stats.cc +++ b/src/cuda-sim/ptx-stats.cc @@ -151,27 +151,27 @@ void ptx_file_line_stats_add_exec_count(const ptx_instruction *pInsn) // attribute pipeline latency to this ptx instruction (specified by the pc) // pipeline latency is the number of cycles a warp with this instruction spent in the pipeline -void ptx_file_line_stats_add_latency(unsigned pc, unsigned latency) +void ptx_stats::ptx_file_line_stats_add_latency(unsigned pc, unsigned latency) { - const ptx_instruction *pInsn = function_info::pc_to_instruction(pc); + const ptx_instruction *pInsn = gpgpu_ctx->pc_to_instruction(pc); ptx_file_line_stats_tracker[ptx_file_line(pInsn->source_file(), pInsn->source_line())].latency += latency; } // attribute dram traffic to this ptx instruction (specified by the pc) // dram traffic is counted in number of requests -void ptx_file_line_stats_add_dram_traffic(unsigned pc, unsigned dram_traffic) +void ptx_stats::ptx_file_line_stats_add_dram_traffic(unsigned pc, unsigned dram_traffic) { - const ptx_instruction *pInsn = function_info::pc_to_instruction(pc); + const ptx_instruction *pInsn = gpgpu_ctx->pc_to_instruction(pc); ptx_file_line_stats_tracker[ptx_file_line(pInsn->source_file(), pInsn->source_line())].dram_traffic += dram_traffic; } // attribute the number of shared memory access cycles to a ptx instruction // counts both the number of warps doing shared memory access and the number of cycles involved -void ptx_file_line_stats_add_smem_bank_conflict(unsigned pc, unsigned n_way_bkconflict) +void ptx_stats::ptx_file_line_stats_add_smem_bank_conflict(unsigned pc, unsigned n_way_bkconflict) { - const ptx_instruction *pInsn = function_info::pc_to_instruction(pc); + const ptx_instruction *pInsn = gpgpu_ctx->pc_to_instruction(pc); ptx_file_line_stats& line_stats = ptx_file_line_stats_tracker[ptx_file_line(pInsn->source_file(), pInsn->source_line())]; line_stats.smem_n_way_bank_conflict_total += n_way_bkconflict; @@ -180,9 +180,9 @@ void ptx_file_line_stats_add_smem_bank_conflict(unsigned pc, unsigned n_way_bkco // attribute a non-coalesced mem access to a ptx instruction // counts both the number of warps causing this and the number of memory requests generated -void ptx_file_line_stats_add_uncoalesced_gmem(unsigned pc, unsigned n_access) +void ptx_stats::ptx_file_line_stats_add_uncoalesced_gmem(unsigned pc, unsigned n_access) { - const ptx_instruction *pInsn = function_info::pc_to_instruction(pc); + const ptx_instruction *pInsn = gpgpu_ctx->pc_to_instruction(pc); ptx_file_line_stats& line_stats = ptx_file_line_stats_tracker[ptx_file_line(pInsn->source_file(), pInsn->source_line())]; line_stats.gmem_n_access_total += n_access; @@ -239,17 +239,17 @@ void ptx_file_line_stats_create_exposed_latency_tracker(int n_shader_cores) } // add an inflight memory instruction -void ptx_file_line_stats_add_inflight_memory_insn(int sc_id, unsigned pc) +void ptx_stats::ptx_file_line_stats_add_inflight_memory_insn(int sc_id, unsigned pc) { - const ptx_instruction *pInsn = function_info::pc_to_instruction(pc); + const ptx_instruction *pInsn = gpgpu_ctx->pc_to_instruction(pc); inflight_mem_tracker[sc_id].add_count(pInsn); } // remove an inflight memory instruction -void ptx_file_line_stats_sub_inflight_memory_insn(int sc_id, unsigned pc) +void ptx_stats::ptx_file_line_stats_sub_inflight_memory_insn(int sc_id, unsigned pc) { - const ptx_instruction *pInsn = function_info::pc_to_instruction(pc); + const ptx_instruction *pInsn = gpgpu_ctx->pc_to_instruction(pc); inflight_mem_tracker[sc_id].sub_count(pInsn); } @@ -262,9 +262,9 @@ void ptx_file_line_stats_commit_exposed_latency(int sc_id, int exposed_latency) } // attribute the number of warp divergence to a ptx instruction -void ptx_file_line_stats_add_warp_divergence(unsigned pc, unsigned n_way_divergence) +void ptx_stats::ptx_file_line_stats_add_warp_divergence(unsigned pc, unsigned n_way_divergence) { - const ptx_instruction *pInsn = function_info::pc_to_instruction(pc); + const ptx_instruction *pInsn = gpgpu_ctx->pc_to_instruction(pc); ptx_file_line_stats& line_stats = ptx_file_line_stats_tracker[ptx_file_line(pInsn->source_file(), pInsn->source_line())]; line_stats.warp_divergence += n_way_divergence; diff --git a/src/cuda-sim/ptx-stats.h b/src/cuda-sim/ptx-stats.h index c75fc58..246b4ce 100644 --- a/src/cuda-sim/ptx-stats.h +++ b/src/cuda-sim/ptx-stats.h @@ -37,17 +37,10 @@ void ptx_file_line_stats_add_exec_count(const ptx_instruction *pInsn); #endif // stat collection interface to gpgpu-sim -void ptx_file_line_stats_add_latency(unsigned pc, unsigned latency); -void ptx_file_line_stats_add_dram_traffic(unsigned pc, unsigned dram_traffic); -void ptx_file_line_stats_add_smem_bank_conflict(unsigned pc, unsigned n_way_bkconflict); -void ptx_file_line_stats_add_uncoalesced_gmem(unsigned pc, unsigned n_access); void ptx_file_line_stats_create_exposed_latency_tracker(int n_shader_cores); -void ptx_file_line_stats_add_inflight_memory_insn(int sc_id, unsigned pc); -void ptx_file_line_stats_sub_inflight_memory_insn(int sc_id, unsigned pc); void ptx_file_line_stats_commit_exposed_latency(int sc_id, int exposed_latency); -void ptx_file_line_stats_add_warp_divergence(unsigned pc, unsigned n_way_divergence); class gpgpu_context; class ptx_stats { @@ -64,4 +57,12 @@ class ptx_stats { // output stats to a file void ptx_file_line_stats_write_file(); + // stat collection interface to gpgpu-sim + void ptx_file_line_stats_add_latency(unsigned pc, unsigned latency); + void ptx_file_line_stats_add_dram_traffic(unsigned pc, unsigned dram_traffic); + void ptx_file_line_stats_add_smem_bank_conflict(unsigned pc, unsigned n_way_bkconflict); + void ptx_file_line_stats_add_uncoalesced_gmem(unsigned pc, unsigned n_access); + void ptx_file_line_stats_add_inflight_memory_insn(int sc_id, unsigned pc); + void ptx_file_line_stats_sub_inflight_memory_insn(int sc_id, unsigned pc); + void ptx_file_line_stats_add_warp_divergence(unsigned pc, unsigned n_way_divergence); }; diff --git a/src/cuda-sim/ptx_ir.cc b/src/cuda-sim/ptx_ir.cc index 5fa9379..3384d49 100644 --- a/src/cuda-sim/ptx_ir.cc +++ b/src/cuda-sim/ptx_ir.cc @@ -43,6 +43,14 @@ typedef void * yyscan_t; #define STR_SIZE 1024 +const ptx_instruction* gpgpu_context::pc_to_instruction(unsigned pc) +{ + if( pc < s_g_pc_to_insn.size() ) + return s_g_pc_to_insn[pc]; + else + return NULL; +} + unsigned symbol::get_uid() { unsigned result = (gpgpu_ctx->symbol_sm_next_uid)++; diff --git a/src/cuda-sim/ptx_ir.h b/src/cuda-sim/ptx_ir.h index 8fc0a06..f4c5c37 100644 --- a/src/cuda-sim/ptx_ir.h +++ b/src/cuda-sim/ptx_ir.h @@ -1372,13 +1372,6 @@ public: return m_symtab; } - static const ptx_instruction* pc_to_instruction(unsigned pc) - { - if( pc < s_g_pc_to_insn.size() ) - return s_g_pc_to_insn[pc]; - else - return NULL; - } unsigned local_mem_framesize() const { return m_local_mem_framesize; @@ -1436,8 +1429,6 @@ private: symbol_table *m_symtab; - static std::vector s_g_pc_to_insn; // a direct mapping from PC to instruction - //parameter size for device kernels int m_args_aligned_size; diff --git a/src/gpgpu-sim/gpu-sim.cc b/src/gpgpu-sim/gpu-sim.cc index bdf989a..e4ae04f 100644 --- a/src/gpgpu-sim/gpu-sim.cc +++ b/src/gpgpu-sim/gpu-sim.cc @@ -886,7 +886,7 @@ void gpgpu_sim::init() m_shader_stats->new_grid(); // initialize the control-flow, memory access, memory latency logger if (m_config.g_visualizer_enabled) { - create_thread_CFlogger( m_config.num_shader(), m_shader_config->n_thread_per_shader, 0, m_config.gpgpu_cflog_interval ); + create_thread_CFlogger( gpgpu_ctx, m_config.num_shader(), m_shader_config->n_thread_per_shader, 0, m_config.gpgpu_cflog_interval ); } shader_CTA_count_create( m_config.num_shader(), m_config.gpgpu_cflog_interval); if (m_config.gpgpu_cflog_interval != 0) { diff --git a/src/gpgpu-sim/mem_latency_stat.cc b/src/gpgpu-sim/mem_latency_stat.cc index 4e94991..a1b43a8 100644 --- a/src/gpgpu-sim/mem_latency_stat.cc +++ b/src/gpgpu-sim/mem_latency_stat.cc @@ -41,6 +41,7 @@ #include #include #include +#include "../../libcuda/gpgpu_context.h" memory_stats_t::memory_stats_t( unsigned n_shader, const shader_core_config *shader_config, const memory_config *mem_config, const class gpgpu_sim* gpu ) { @@ -195,7 +196,7 @@ void memory_stats_t::memlatstat_dram_access(mem_fetch *mf) mem_access_type_stats[mf->get_access_type()][dram_id][bank]++; } if (mf->get_pc() != (unsigned)-1) - ptx_file_line_stats_add_dram_traffic(mf->get_pc(), mf->get_data_size()); + m_gpu->gpgpu_ctx->stats->ptx_file_line_stats_add_dram_traffic(mf->get_pc(), mf->get_data_size()); } void memory_stats_t::memlatstat_icnt2mem_pop(mem_fetch *mf) diff --git a/src/gpgpu-sim/shader.cc b/src/gpgpu-sim/shader.cc index b7ae95d..c697450 100644 --- a/src/gpgpu-sim/shader.cc +++ b/src/gpgpu-sim/shader.cc @@ -744,7 +744,7 @@ void shader_core_ctx::decode() if( m_inst_fetch_buffer.m_valid ) { // decode 1 or 2 instructions and place them into ibuffer address_type pc = m_inst_fetch_buffer.m_pc; - const warp_inst_t* pI1 = ptx_fetch_inst(pc); + const warp_inst_t* pI1 = m_gpu->gpgpu_ctx->ptx_fetch_inst(pc); m_warp[m_inst_fetch_buffer.m_warp_id].ibuffer_fill(0,pI1); m_warp[m_inst_fetch_buffer.m_warp_id].inc_inst_in_pipeline(); if( pI1 ) { @@ -754,7 +754,7 @@ void shader_core_ctx::decode() }else if(pI1->oprnd_type==FP_OP) { m_stats->m_num_FPdecoded_insn[m_sid]++; } - const warp_inst_t* pI2 = ptx_fetch_inst(pc+pI1->isize); + const warp_inst_t* pI2 = m_gpu->gpgpu_ctx->ptx_fetch_inst(pc+pI1->isize); if( pI2 ) { m_warp[m_inst_fetch_buffer.m_warp_id].ibuffer_fill(1,pI2); m_warp[m_inst_fetch_buffer.m_warp_id].inc_inst_in_pipeline(); diff --git a/src/gpgpu-sim/stat-tool.cc b/src/gpgpu-sim/stat-tool.cc index 6a4c75b..35a4cc3 100644 --- a/src/gpgpu-sim/stat-tool.cc +++ b/src/gpgpu-sim/stat-tool.cc @@ -37,6 +37,7 @@ #include #include #include +#include "../../libcuda/gpgpu_context.h" //////////////////////////////////////////////////////////////////////////////// @@ -110,12 +111,10 @@ void spill_log_to_file (FILE *fout, int final, unsigned long long current_cycle //////////////////////////////////////////////////////////////////////////////// -unsigned translate_pc_to_ptxlineno(unsigned pc); - static int n_thread_CFloggers = 0; static thread_CFlocality** thread_CFlogger = NULL; -void create_thread_CFlogger( int n_loggers, int n_threads, address_type start_pc, unsigned long long logging_interval) +void create_thread_CFlogger(gpgpu_context* ctx, int n_loggers, int n_threads, address_type start_pc, unsigned long long logging_interval) { destroy_thread_CFlogger(); @@ -126,7 +125,7 @@ void create_thread_CFlogger( int n_loggers, int n_threads, address_type start_pc char buffer[32]; for (int i = 0; i < n_thread_CFloggers; i++) { snprintf(buffer, 32, "%02d", i); - thread_CFlogger[i] = new thread_CFlocality( name_tpl + buffer, logging_interval, n_threads, start_pc); + thread_CFlogger[i] = new thread_CFlocality( ctx, name_tpl + buffer, logging_interval, n_threads, start_pc); if (logging_interval != 0) { add_snap_shot_trigger(thread_CFlogger[i]); add_spill_log(thread_CFlogger[i]); @@ -368,10 +367,10 @@ static int s_cache_access_logger_n_types = 0; static std::vector s_cache_access_logger; enum cache_access_logger_types { - NORMAL, TEXTURE, CONSTANT, INSTRUCTION + NORMALS, TEXTURE, CONSTANT, INSTRUCTION }; -int get_shader_normal_cache_id() { return NORMAL; } +int get_shader_normal_cache_id() { return NORMALS; } int get_shader_texture_cache_id() { return TEXTURE; } int get_shader_constant_cache_id() { return CONSTANT; } int get_shader_instruction_cache_id() { return INSTRUCTION; } @@ -394,7 +393,7 @@ void shader_cache_access_log( int logger_id, int type, int miss) { if (s_cache_access_logger_n_types == 0) return; if (logger_id < 0) return; - assert(type == NORMAL || type == TEXTURE || type == CONSTANT || type == INSTRUCTION); + assert(type == NORMALS || type == TEXTURE || type == CONSTANT || type == INSTRUCTION); assert(miss == 0 || miss == 1); s_cache_access_logger[logger_id].log(2 * type + miss); @@ -404,7 +403,7 @@ void shader_cache_access_unlog( int logger_id, int type, int miss) { if (s_cache_access_logger_n_types == 0) return; if (logger_id < 0) return; - assert(type == NORMAL || type == TEXTURE || type == CONSTANT || type == INSTRUCTION); + assert(type == NORMALS || type == TEXTURE || type == CONSTANT || type == INSTRUCTION); assert(miss == 0 || miss == 1); s_cache_access_logger[logger_id].unlog(2 * type + miss); @@ -477,22 +476,24 @@ void shader_CTA_count_visualizer_gzprint( gzFile fout ) //////////////////////////////////////////////////////////////////////////////// //////////////////////////////////////////////////////////////////////////////// -thread_insn_span::thread_insn_span(unsigned long long cycle) +thread_insn_span::thread_insn_span(unsigned long long cycle, gpgpu_context* ctx) : m_cycle(cycle), #if (tr1_hash_map_ismap == 1) m_insn_span_count() #else m_insn_span_count(32*1024) #endif -{ +{ + gpgpu_ctx = ctx; } thread_insn_span::~thread_insn_span() { } -thread_insn_span::thread_insn_span(const thread_insn_span& other) +thread_insn_span::thread_insn_span(const thread_insn_span& other, gpgpu_context* ctx) : m_cycle(other.m_cycle), - m_insn_span_count(other.m_insn_span_count) + m_insn_span_count(other.m_insn_span_count) { + gpgpu_ctx = ctx; } thread_insn_span& thread_insn_span::operator=(const thread_insn_span& other) @@ -551,7 +552,7 @@ void thread_insn_span::print_sparse_histo(FILE *fout) const int n_printed_entries = 0; span_count_map::const_iterator i_sc = m_insn_span_count.begin(); for (; i_sc != m_insn_span_count.end(); ++i_sc) { - unsigned ptx_lineno = translate_pc_to_ptxlineno(i_sc->first); + unsigned ptx_lineno = gpgpu_ctx->translate_pc_to_ptxlineno(i_sc->first); fprintf(fout, "%u %d ", ptx_lineno, i_sc->second); n_printed_entries++; } @@ -566,7 +567,7 @@ void thread_insn_span::print_sparse_histo(gzFile fout) const int n_printed_entries = 0; span_count_map::const_iterator i_sc = m_insn_span_count.begin(); for (; i_sc != m_insn_span_count.end(); ++i_sc) { - unsigned ptx_lineno = translate_pc_to_ptxlineno(i_sc->first); + unsigned ptx_lineno = gpgpu_ctx->translate_pc_to_ptxlineno(i_sc->first); gzprintf(fout, "%u %d ", ptx_lineno, i_sc->second); n_printed_entries++; } @@ -578,14 +579,14 @@ void thread_insn_span::print_sparse_histo(gzFile fout) const //////////////////////////////////////////////////////////////////////////////// -thread_CFlocality::thread_CFlocality(std::string name, +thread_CFlocality::thread_CFlocality( gpgpu_context* ctx, std::string name, unsigned long long snap_shot_interval, int nthreads, address_type start_pc, unsigned long long start_cycle) : snap_shot_trigger(snap_shot_interval), m_name(name), m_nthreads(nthreads), m_thread_pc(nthreads, start_pc), m_cycle(start_cycle), - m_thd_span(start_cycle) + m_thd_span(start_cycle, ctx) { std::fill(m_thread_pc.begin(), m_thread_pc.end(), -1); // so that hw thread with no work assigned will not clobber results } diff --git a/src/gpgpu-sim/stat-tool.h b/src/gpgpu-sim/stat-tool.h index 5646f01..67b3923 100644 --- a/src/gpgpu-sim/stat-tool.h +++ b/src/gpgpu-sim/stat-tool.h @@ -35,6 +35,7 @@ #include #include +class gpgpu_context; ///////////////////////////////////////////////////////////////////////////////////// // logger snapshot trigger: // - automate the snap_shot part of loggers to avoid modifying simulation loop everytime @@ -80,8 +81,8 @@ public: class thread_insn_span { public: - thread_insn_span(unsigned long long cycle); - thread_insn_span(const thread_insn_span& other); + thread_insn_span(unsigned long long cycle, gpgpu_context* ctx); + thread_insn_span(const thread_insn_span& other, gpgpu_context* ctx); ~thread_insn_span(); thread_insn_span& operator=(const thread_insn_span& other); @@ -94,7 +95,8 @@ public: void print_sparse_histo(FILE *fout) const; void print_sparse_histo(gzFile fout) const; -private: +private: + gpgpu_context* gpgpu_ctx; typedef tr1_hash_map span_count_map; unsigned long long m_cycle; span_count_map m_insn_span_count; @@ -102,7 +104,7 @@ private: class thread_CFlocality : public snap_shot_trigger, public spill_log_interface { public: - thread_CFlocality(std::string name, unsigned long long snap_shot_interval, + thread_CFlocality(gpgpu_context* ctx, std::string name, unsigned long long snap_shot_interval, int nthreads, address_type start_pc, unsigned long long start_cycle = 0); ~thread_CFlocality(); @@ -270,7 +272,7 @@ void try_snap_shot (unsigned long long current_cycle); void set_spill_interval (unsigned long long interval); void spill_log_to_file (FILE *fout, int final, unsigned long long current_cycle); -void create_thread_CFlogger( int n_loggers, int n_threads, address_type start_pc, unsigned long long logging_interval); +void create_thread_CFlogger(gpgpu_context* ctx, int n_loggers, int n_threads, address_type start_pc, unsigned long long logging_interval); void destroy_thread_CFlogger( ); void cflog_update_thread_pc( int logger_id, int thread_id, address_type pc ); void cflog_snapshot( int logger_id, unsigned long long cycle ); -- cgit v1.3 From da0c8dff3b4e89acaf2f2dd31bf8940ab4a1e71c Mon Sep 17 00:00:00 2001 From: Mahmoud Date: Mon, 26 Aug 2019 12:23:43 -0400 Subject: missing some files to add --- src/gpgpu-sim/gpu-sim.cc | 9 ++++++--- src/gpgpu-sim/local_interconnect.h | 30 +++++++++++++++++++++--------- 2 files changed, 27 insertions(+), 12 deletions(-) (limited to 'src/gpgpu-sim/gpu-sim.cc') diff --git a/src/gpgpu-sim/gpu-sim.cc b/src/gpgpu-sim/gpu-sim.cc index 92d5366..4f071c7 100644 --- a/src/gpgpu-sim/gpu-sim.cc +++ b/src/gpgpu-sim/gpu-sim.cc @@ -257,9 +257,12 @@ void shader_core_config::reg_options(class OptionParser * opp) "per-shader L1 data cache config " " {::,:::,::, | none}", "none" ); + option_parser_register(opp, "-l1_banks", OPT_UINT32, &m_L1D_config.l1_banks, + "The number of L1 cache banks", + "1"); option_parser_register(opp, "-l1_latency", OPT_UINT32, &m_L1D_config.l1_latency, "L1 Hit Latency", - "0"); + "1"); option_parser_register(opp, "-smem_latency", OPT_UINT32, &smem_latency, "smem Latency", "3"); @@ -320,8 +323,8 @@ void shader_core_config::reg_options(class OptionParser * opp) option_parser_register(opp, "-gpgpu_shmem_size", OPT_UINT32, &gpgpu_shmem_size, "Size of shared memory per shader core (default 16kB)", "16384"); - option_parser_register(opp, "-adaptive_volta_cache_config", OPT_BOOL, &adaptive_volta_cache_config, - "adaptive_volta_cache_config", + option_parser_register(opp, "-adaptive_cache_config", OPT_BOOL, &adaptive_volta_cache_config, + "adaptive_cache_config", "0"); option_parser_register(opp, "-gpgpu_shmem_sizeDefault", OPT_UINT32, &gpgpu_shmem_sizeDefault, "Size of shared memory per shader core (default 16kB)", diff --git a/src/gpgpu-sim/local_interconnect.h b/src/gpgpu-sim/local_interconnect.h index f4a2af1..a784da8 100644 --- a/src/gpgpu-sim/local_interconnect.h +++ b/src/gpgpu-sim/local_interconnect.h @@ -35,27 +35,35 @@ using namespace std; +enum Interconnect_type { + REQ_NET=0, + REPLY_NET=1 +}; + +enum Arbiteration_type { + NAIVE_RR=0, + iSLIP=1 +}; + struct inct_config { - //config for local interconnect unsigned in_buffer_limit; unsigned out_buffer_limit; unsigned subnets; + Arbiteration_type arbiter_algo; }; -enum Interconnect_type { - REQ_NET=0, - REPLY_NET=1 -}; class xbar_router { public: - xbar_router(unsigned router_id, enum Interconnect_type m_type, unsigned n_shader, unsigned n_mem, unsigned m_in_buffer_limit, unsigned m_out_buffer_limit); + xbar_router(unsigned router_id, enum Interconnect_type m_type, unsigned n_shader, unsigned n_mem, unsigned m_in_buffer_limit, unsigned m_out_buffer_limit, enum Arbiteration_type m_arbit_type); ~xbar_router(); void Push(unsigned input_deviceID, unsigned output_deviceID, void* data, unsigned int size); void* Pop(unsigned ouput_deviceID); - void Advance(); + void Advance( ); + + bool Busy() const; bool Has_Buffer_In(unsigned input_deviceID, unsigned size, bool update_counter=false); bool Has_Buffer_Out(unsigned output_deviceID, unsigned size); @@ -70,6 +78,9 @@ public: unsigned long long packets_num; private: + void iSLIP_Advance(); + void RR_Advance(); + struct Packet{ Packet(void* m_data, unsigned m_output_deviceID) { data = m_data; @@ -82,11 +93,12 @@ private: vector > out_buffers; unsigned _n_shader, _n_mem, total_nodes; unsigned in_buffer_limit, out_buffer_limit; - vector next_node; -// unsigned next_node; + vector next_node; //used for iSLIP arbit + unsigned next_node_id; //used for RR arbit unsigned m_id; enum Interconnect_type router_type; unsigned active_in_buffers,active_out_buffers; + Arbiteration_type arbit_type; friend class LocalInterconnect; -- cgit v1.3 From 93597743b24b7d7726096d337e6c9c4516cb2273 Mon Sep 17 00:00:00 2001 From: Mahmoud Date: Fri, 30 Aug 2019 18:43:21 -0400 Subject: update the adaptive cache behaviour and make L1 fully assoc in Volta --- configs/tested-cfgs/SM7_QV100/gpgpusim.config | 2 +- configs/tested-cfgs/SM7_TITANV/gpgpusim.config | 2 +- src/abstract_hardware_model.cc | 4 +- src/abstract_hardware_model.h | 9 ++++- src/gpgpu-sim/addrdec.cc | 4 +- src/gpgpu-sim/gpu-cache.cc | 7 +++- src/gpgpu-sim/gpu-cache.h | 5 +++ src/gpgpu-sim/gpu-sim.cc | 2 +- src/gpgpu-sim/shader.cc | 54 ++++++++++++++++---------- 9 files changed, 59 insertions(+), 30 deletions(-) (limited to 'src/gpgpu-sim/gpu-sim.cc') diff --git a/configs/tested-cfgs/SM7_QV100/gpgpusim.config b/configs/tested-cfgs/SM7_QV100/gpgpusim.config index f807e11..c0d22ee 100644 --- a/configs/tested-cfgs/SM7_QV100/gpgpusim.config +++ b/configs/tested-cfgs/SM7_QV100/gpgpusim.config @@ -91,7 +91,7 @@ # Volta unified cache has four banks -l1_banks 4 #-mem_unit_ports 4 --gpgpu_cache:dl1 S:4:128:64,L:L:s:N:L,A:256:8,16:0,32 +-gpgpu_cache:dl1 S:1:128:256,L:L:s:N:L,A:256:8,16:0,32 -gpgpu_shmem_size 98304 -gpgpu_shmem_sizeDefault 98304 -gpgpu_shmem_per_block 65536 diff --git a/configs/tested-cfgs/SM7_TITANV/gpgpusim.config b/configs/tested-cfgs/SM7_TITANV/gpgpusim.config index 888ce71..0339b0d 100644 --- a/configs/tested-cfgs/SM7_TITANV/gpgpusim.config +++ b/configs/tested-cfgs/SM7_TITANV/gpgpusim.config @@ -91,7 +91,7 @@ # Volta unified cache has four banks -l1_banks 4 #-mem_unit_ports 4 --gpgpu_cache:dl1 S:4:128:64,L:L:s:N:L,A:256:8,16:0,32 +-gpgpu_cache:dl1 S:1:128:256,L:L:s:N:L,A:256:8,16:0,32 -gpgpu_shmem_size 98304 -gpgpu_shmem_sizeDefault 98304 -gpgpu_shmem_per_block 65536 diff --git a/src/abstract_hardware_model.cc b/src/abstract_hardware_model.cc index 023f51b..35a3984 100644 --- a/src/abstract_hardware_model.cc +++ b/src/abstract_hardware_model.cc @@ -710,7 +710,7 @@ kernel_info_t::kernel_info_t( dim3 gridDim, dim3 blockDim, class function_info * //Jin: launch latency management m_launch_latency = g_kernel_launch_latency; - volta_cache_config_set=false; + cache_config_set=false; } /*A snapshot of the texture mappings needs to be stored in the kernel's info as @@ -735,7 +735,7 @@ kernel_info_t::kernel_info_t( dim3 gridDim, dim3 blockDim, class function_info * //Jin: launch latency management m_launch_latency = g_kernel_launch_latency; - volta_cache_config_set=false; + cache_config_set=false; m_NameToCudaArray = nameToCudaArray; m_NameToTextureInfo = nameToTextureInfo; } diff --git a/src/abstract_hardware_model.h b/src/abstract_hardware_model.h index 27a1ba6..231b6a2 100644 --- a/src/abstract_hardware_model.h +++ b/src/abstract_hardware_model.h @@ -65,6 +65,11 @@ enum FuncCache FuncCachePreferL1 = 2 }; +enum AdaptiveCache +{ + FIXED = 0, + VOLTA = 1 +}; #ifdef __cplusplus @@ -345,7 +350,7 @@ public: unsigned long long end_cycle; unsigned m_launch_latency; - mutable bool volta_cache_config_set; + mutable bool cache_config_set; }; struct core_config { @@ -388,7 +393,7 @@ struct core_config { unsigned gpgpu_max_insn_issue_per_warp; bool gmem_skip_L1D; // on = global memory access always skip the L1 cache - bool adaptive_volta_cache_config; + unsigned adaptive_cache_config; }; // bounded stack that implements simt reconvergence using pdom mechanism from MICRO'07 paper diff --git a/src/gpgpu-sim/addrdec.cc b/src/gpgpu-sim/addrdec.cc index b0db034..3262456 100644 --- a/src/gpgpu-sim/addrdec.cc +++ b/src/gpgpu-sim/addrdec.cc @@ -124,8 +124,8 @@ void linear_to_raw_address_translation::addrdec_tlx(new_addr_type addr, addrdec_ * Rau, B. R et al. * ISCA 1991 * - * equations are adopted from: - * "Sacat: streaming-aware conflict-avoiding thrashing-resistant gpgpu cache management scheme." + * equations are corresponding to IPOLY(37) and are adopted from: + * "SACAT: streaming-aware conflict-avoiding thrashing-resistant gpgpu cache management scheme." * Khairy et al. * IEEE TPDS 2017. */ diff --git a/src/gpgpu-sim/gpu-cache.cc b/src/gpgpu-sim/gpu-cache.cc index db9701d..8d00ea9 100644 --- a/src/gpgpu-sim/gpu-cache.cc +++ b/src/gpgpu-sim/gpu-cache.cc @@ -65,8 +65,13 @@ const char * cache_fail_status_str(enum cache_reservation_fail_reason status) unsigned l1d_cache_config::set_bank(new_addr_type addr) const{ - if(m_cache_type == SECTOR) + //For sector cache, we select one sector per bank (sector interleaving) + //This is what was found in Volta (one sector per bank, sector interleaving) + //otherwise, line interleaving + if(m_cache_type == SECTOR) { + //assert(l1_banks == SECTOR_CHUNCK_SIZE); return (addr >> m_sector_sz_log2) & (l1_banks-1); + } else return (addr >> m_line_sz_log2) & (l1_banks-1); } diff --git a/src/gpgpu-sim/gpu-cache.h b/src/gpgpu-sim/gpu-cache.h index 90adbb5..dd22886 100644 --- a/src/gpgpu-sim/gpu-cache.h +++ b/src/gpgpu-sim/gpu-cache.h @@ -656,6 +656,11 @@ public: assert( m_valid ); return MAX_DEFAULT_CACHE_SIZE_MULTIBLIER * m_nset * original_m_assoc; } + unsigned get_max_assoc() const + { + assert( m_valid ); + return MAX_DEFAULT_CACHE_SIZE_MULTIBLIER * original_m_assoc; + } void print( FILE *fp ) const { fprintf( fp, "Size = %d B (%d Set x %d-way x %d byte line)\n", diff --git a/src/gpgpu-sim/gpu-sim.cc b/src/gpgpu-sim/gpu-sim.cc index 4f071c7..ed94865 100644 --- a/src/gpgpu-sim/gpu-sim.cc +++ b/src/gpgpu-sim/gpu-sim.cc @@ -323,7 +323,7 @@ void shader_core_config::reg_options(class OptionParser * opp) option_parser_register(opp, "-gpgpu_shmem_size", OPT_UINT32, &gpgpu_shmem_size, "Size of shared memory per shader core (default 16kB)", "16384"); - option_parser_register(opp, "-adaptive_cache_config", OPT_BOOL, &adaptive_volta_cache_config, + option_parser_register(opp, "-adaptive_cache_config", OPT_UINT32, &adaptive_cache_config, "adaptive_cache_config", "0"); option_parser_register(opp, "-gpgpu_shmem_sizeDefault", OPT_UINT32, &gpgpu_shmem_sizeDefault, diff --git a/src/gpgpu-sim/shader.cc b/src/gpgpu-sim/shader.cc index ffd3035..ff2bf3f 100644 --- a/src/gpgpu-sim/shader.cc +++ b/src/gpgpu-sim/shader.cc @@ -2980,33 +2980,47 @@ unsigned int shader_core_config::max_cta( const kernel_info_t &k ) const abort(); } - if(adaptive_volta_cache_config && !k.volta_cache_config_set) { - //For Volta, we assign the remaining shared memory to L1 cache - //For more info, see https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#shared-memory-7-x + if(adaptive_cache_config && !k.cache_config_set) { + //For more info about adaptive cache, see https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#shared-memory-7-x unsigned total_shmed = kernel_info->smem * result; assert(total_shmed >=0 && total_shmed <= gpgpu_shmem_size); - assert(gpgpu_shmem_size == 98304); //Volta has 96 KB shared - assert(m_L1D_config.get_nset() == 4); //Volta L1 has four sets + //assert(gpgpu_shmem_size == 98304); //Volta has 96 KB shared + //assert(m_L1D_config.get_nset() == 4); //Volta L1 has four sets if(total_shmed < gpgpu_shmem_size){ - if(total_shmed == 0) - m_L1D_config.set_assoc(256); //L1 is 128KB ans shd=0 - else if(total_shmed > 0 && total_shmed <= 8192) - m_L1D_config.set_assoc(240); //L1 is 120KB ans shd=8KB - else if(total_shmed > 8192 && total_shmed <= 16384) - m_L1D_config.set_assoc(224); //L1 is 112KB ans shd=16KB - else if(total_shmed > 16384 && total_shmed <= 32768) - m_L1D_config.set_assoc(192); //L1 is 96KB ans shd=32KB - else if(total_shmed > 32768 && total_shmed <= 65536) - m_L1D_config.set_assoc(128); //L1 is 64KB ans shd=64KB - else if(total_shmed > 65536 && total_shmed <= gpgpu_shmem_size) - m_L1D_config.set_assoc(64); //L1 is 32KB and shd=96KB - else + switch (adaptive_cache_config) { + case FIXED: + break; + case VOLTA: { + //For Volta, we assign the remaining shared memory to L1 cache + //For more info about adaptive cache, see https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#shared-memory-7-x + assert(gpgpu_shmem_size == 98304); //Volta has 96 KB shared + + //To Do: make it flexible and not tuned to 9KB share memory + unsigned max_assoc = m_L1D_config.get_max_assoc(); + if(total_shmed == 0) + m_L1D_config.set_assoc(max_assoc); //L1 is 128KB and shd=0 + else if(total_shmed > 0 && total_shmed <= 8192) + m_L1D_config.set_assoc(0.9375 * max_assoc); //L1 is 120KB and shd=8KB + else if(total_shmed > 8192 && total_shmed <= 16384) + m_L1D_config.set_assoc(0.875 * max_assoc); //L1 is 112KB and shd=16KB + else if(total_shmed > 16384 && total_shmed <= 32768) + m_L1D_config.set_assoc(0.75 * max_assoc); //L1 is 96KB and shd=32KB + else if(total_shmed > 32768 && total_shmed <= 65536) + m_L1D_config.set_assoc(0.5 * max_assoc); //L1 is 64KB and shd=64KB + else if(total_shmed > 65536 && total_shmed <= gpgpu_shmem_size) + m_L1D_config.set_assoc(0.25 * max_assoc); //L1 is 32KB and shd=96KB + else + assert(0); + break; + } + default: assert(0); + } - printf ("GPGPU-Sim: Reconfigure L1 cache in Volta Archi to %uKB\n", m_L1D_config.get_total_size_inKB()); + printf ("GPGPU-Sim: Reconfigure L1 cache to %uKB\n", m_L1D_config.get_total_size_inKB()); } - k.volta_cache_config_set = true; + k.cache_config_set = true; } return result; -- cgit v1.3