From b3e786e3d8d720217f36a214e9b5be9a19ab9dd2 Mon Sep 17 00:00:00 2001 From: Mengchi Zhang Date: Mon, 8 Jul 2019 12:37:52 -0400 Subject: Move opcode_latency_int thus pass gpgpu_context into many classes Signed-off-by: Mengchi Zhang --- src/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 ++++++++++++------ 10 files changed, 36 insertions(+), 24 deletions(-) (limited to 'src/gpgpu-sim') 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 ); -- cgit v1.3 From 98bb2046a70715c8599561ca78425a7ff99ce53b Mon Sep 17 00:00:00 2001 From: Mengchi Zhang Date: Mon, 8 Jul 2019 14:21:08 -0400 Subject: Move cp_count Signed-off-by: Mengchi Zhang --- libopencl/opencl_runtime_api.cc | 4 +++- src/cuda-sim/cuda-sim.cc | 5 ++--- src/cuda-sim/cuda-sim.h | 4 +++- src/gpgpu-sim/gpu-sim.h | 4 ++-- src/gpgpusim_entrypoint.cc | 4 ++-- src/gpgpusim_entrypoint.h | 1 - 6 files changed, 12 insertions(+), 10 deletions(-) (limited to 'src/gpgpu-sim') diff --git a/libopencl/opencl_runtime_api.cc b/libopencl/opencl_runtime_api.cc index 03ec80c..d302ff8 100644 --- a/libopencl/opencl_runtime_api.cc +++ b/libopencl/opencl_runtime_api.cc @@ -877,6 +877,8 @@ clEnqueueNDRangeKernel(cl_command_queue command_queue, const cl_event * event_wait_list, cl_event * event) CL_API_SUFFIX__VERSION_1_0 { + gpgpu_context *ctx; + ctx = GPGPU_Context(); int _global_size[3]; int zeros[3] = { 0, 0, 0}; printf("\n\n\n"); @@ -956,7 +958,7 @@ clEnqueueNDRangeKernel(cl_command_queue command_queue, } kernel_info_t *grid = gpgpu_opencl_ptx_sim_init_grid(kernel->get_implementation(),params,GridDim,BlockDim,gpu); if ( g_ptx_sim_mode ) - gpgpu_opencl_ptx_sim_main_func( grid ); + ctx->func_sim->gpgpu_opencl_ptx_sim_main_func( grid ); else gpgpu_opencl_ptx_sim_main_perf( grid ); return CL_SUCCESS; diff --git a/src/cuda-sim/cuda-sim.cc b/src/cuda-sim/cuda-sim.cc index df0bbd7..12f38af 100644 --- a/src/cuda-sim/cuda-sim.cc +++ b/src/cuda-sim/cuda-sim.cc @@ -61,7 +61,6 @@ int g_debug_execution = 0; int g_debug_thread_uid = 0; addr_t g_debug_pc = 0xBEEF1518; // Output debug information to file options -int cp_count; int cp_cta_resume; unsigned g_ptx_sim_num_insn = 0; @@ -2146,7 +2145,7 @@ unsigned max_cta (const struct gpgpu_ptx_sim_info *kernel_info, unsigned threads This function simulates the CUDA code functionally, it takes a kernel_info_t parameter which holds the data for the CUDA kernel to be executed !*/ -void gpgpu_cuda_ptx_sim_main_func( kernel_info_t &kernel, bool openCL ) +void cuda_sim::gpgpu_cuda_ptx_sim_main_func( kernel_info_t &kernel, bool openCL ) { printf("GPGPU-Sim: Performing Functional Simulation, executing kernel %s...\n",kernel.name().c_str()); @@ -2312,7 +2311,7 @@ void functionalCoreSim::createWarp(unsigned warpId) void functionalCoreSim::execute(int inst_count, unsigned ctaid_cp) { - cp_count= m_gpu->checkpoint_insn_Y; + m_gpu->gpgpu_ctx->func_sim->cp_count= m_gpu->checkpoint_insn_Y; cp_cta_resume= m_gpu->checkpoint_CTA_t; initializeCTA(ctaid_cp); diff --git a/src/cuda-sim/cuda-sim.h b/src/cuda-sim/cuda-sim.h index 96d34f5..4c8d953 100644 --- a/src/cuda-sim/cuda-sim.h +++ b/src/cuda-sim/cuda-sim.h @@ -55,7 +55,6 @@ extern class kernel_info_t *gpgpu_opencl_ptx_sim_init_grid(class function_info * struct dim3 gridDim, struct dim3 blockDim, class gpgpu_t *gpu ); -extern void gpgpu_cuda_ptx_sim_main_func( kernel_info_t &kernel, bool openCL = false ); extern void print_splash(); extern void gpgpu_ptx_sim_register_const_variable(void*, const char *deviceName, size_t size ); extern void gpgpu_ptx_sim_register_global_variable(void *hostVar, const char *deviceName, size_t size ); @@ -137,8 +136,11 @@ class cuda_sim { public: //global variables char *opcode_latency_int; + int cp_count; //global functions void ptx_opcocde_latency_options (option_parser_t opp); + void gpgpu_cuda_ptx_sim_main_func( kernel_info_t &kernel, bool openCL = false ); + int gpgpu_opencl_ptx_sim_main_func( kernel_info_t *grid ); }; #endif diff --git a/src/gpgpu-sim/gpu-sim.h b/src/gpgpu-sim/gpu-sim.h index 7eeb7dd..b47ab16 100644 --- a/src/gpgpu-sim/gpu-sim.h +++ b/src/gpgpu-sim/gpu-sim.h @@ -494,6 +494,8 @@ public: */ simt_core_cluster * getSIMTCluster(); + // backward pointer + class gpgpu_context* gpgpu_ctx; private: // clocks @@ -511,8 +513,6 @@ private: void gpgpu_debug(); ///// data ///// -// backward pointer - 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 5018305..d9d1023 100644 --- a/src/gpgpusim_entrypoint.cc +++ b/src/gpgpusim_entrypoint.cc @@ -134,7 +134,7 @@ void *gpgpu_sim_thread_concurrent(void*) if( GPGPUsim_ctx_ptr()->g_the_gpu->is_functional_sim()) { kernel_info_t * kernel = GPGPUsim_ctx_ptr()->g_the_gpu->get_functional_kernel(); assert(kernel); - gpgpu_cuda_ptx_sim_main_func(*kernel); + GPGPUsim_ctx_ptr()->gpgpu_ctx->func_sim->gpgpu_cuda_ptx_sim_main_func(*kernel); GPGPUsim_ctx_ptr()->g_the_gpu->finish_functional_sim(kernel); } @@ -287,7 +287,7 @@ int gpgpu_opencl_ptx_sim_main_perf( kernel_info_t *grid ) /*! * This function call the CUDA PTX functional simulator */ -int gpgpu_opencl_ptx_sim_main_func( kernel_info_t *grid ) +int cuda_sim::gpgpu_opencl_ptx_sim_main_func( kernel_info_t *grid ) { //calling the CUDA PTX simulator, sending the kernel by reference and a flag set to true, //the flag used by the function to distinguish OpenCL calls from the CUDA simulation calls which diff --git a/src/gpgpusim_entrypoint.h b/src/gpgpusim_entrypoint.h index dfb82d0..887b3c8 100644 --- a/src/gpgpusim_entrypoint.h +++ b/src/gpgpusim_entrypoint.h @@ -83,6 +83,5 @@ struct GPGPUsim_ctx* GPGPUsim_ctx_ptr(); class stream_manager* g_stream_manager(); int gpgpu_opencl_ptx_sim_main_perf( kernel_info_t *grid ); -int gpgpu_opencl_ptx_sim_main_func( kernel_info_t *grid ); #endif -- 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') diff --git a/src/abstract_hardware_model.cc b/src/abstract_hardware_model.cc index 7755477..248e7a5 100644 --- a/src/abstract_hardware_model.cc +++ b/src/abstract_hardware_model.cc @@ -173,9 +173,10 @@ void gpgpu_functional_sim_config::ptx_set_tex_cache_linesize(unsigned linesize) m_texcache_linesize = linesize; } -gpgpu_t::gpgpu_t( const gpgpu_functional_sim_config &config ) +gpgpu_t::gpgpu_t( const gpgpu_functional_sim_config &config, gpgpu_context* ctx ) : m_function_model_config(config) { + gpgpu_ctx = ctx; m_global_mem = new memory_space_impl<8192>("global",64*1024); m_tex_mem = new memory_space_impl<8192>("tex",64*1024); diff --git a/src/abstract_hardware_model.h b/src/abstract_hardware_model.h index 68cb693..da29a11 100644 --- a/src/abstract_hardware_model.h +++ b/src/abstract_hardware_model.h @@ -31,6 +31,8 @@ // Forward declarations class gpgpu_sim; class kernel_info_t; +class gpgpu_context; + //Set a hard limit of 32 CTAs per shader [cuda only has 8] #define MAX_CTA_PER_SHADER 32 @@ -529,7 +531,9 @@ private: class gpgpu_t { public: - gpgpu_t( const gpgpu_functional_sim_config &config ); + gpgpu_t( const gpgpu_functional_sim_config &config, gpgpu_context* ctx ); + // backward pointer + class gpgpu_context* gpgpu_ctx; int checkpoint_option; int checkpoint_kernel; int checkpoint_CTA; diff --git a/src/cuda-sim/cuda-sim.cc b/src/cuda-sim/cuda-sim.cc index a143aa5..939358b 100644 --- a/src/cuda-sim/cuda-sim.cc +++ b/src/cuda-sim/cuda-sim.cc @@ -62,7 +62,6 @@ int g_debug_thread_uid = 0; addr_t g_debug_pc = 0xBEEF1518; // Output debug information to file options -unsigned g_ptx_sim_num_insn = 0; unsigned gpgpu_param_num_shaders = 0; char *opcode_latency_fp, *opcode_latency_dp,*opcode_latency_sfu,*opcode_latency_tensor; @@ -1629,7 +1628,7 @@ void ptx_thread_info::ptx_exec_inst( warp_inst_t &inst, unsigned lane_id) dim3 ctaid = get_ctaid(); dim3 tid = get_tid(); printf("%u [thd=%u][i=%u] : ctaid=(%u,%u,%u) tid=(%u,%u,%u) icount=%u [pc=%u] (%s:%u - %s) [0x%llx]\n", - g_ptx_sim_num_insn, + m_gpu->gpgpu_ctx->func_sim->g_ptx_sim_num_insn, get_uid(), pI->uid(), ctaid.x,ctaid.y,ctaid.z,tid.x,tid.y,tid.z, get_icount(), @@ -1687,7 +1686,7 @@ void ptx_thread_info::ptx_exec_inst( warp_inst_t &inst, unsigned lane_id) dump_regs(stdout); } update_pc(); - g_ptx_sim_num_insn++; + m_gpu->gpgpu_ctx->func_sim->g_ptx_sim_num_insn++; //not using it with functional simulation mode if(!(this->m_functionalSimulationMode)) @@ -1714,11 +1713,11 @@ void ptx_thread_info::ptx_exec_inst( warp_inst_t &inst, unsigned lane_id) if (space_type) StatAddSample( g_inst_classification_stat[g_ptx_kernel_count], ( int )space_type); StatAddSample( g_inst_op_classification_stat[g_ptx_kernel_count], (int) pI->get_opcode() ); } - if ( (g_ptx_sim_num_insn % 100000) == 0 ) { + if ( (m_gpu->gpgpu_ctx->func_sim->g_ptx_sim_num_insn % 100000) == 0 ) { dim3 ctaid = get_ctaid(); dim3 tid = get_tid(); DPRINTF(LIVENESS, "GPGPU-Sim PTX: %u instructions simulated : ctaid=(%u,%u,%u) tid=(%u,%u,%u)\n", - g_ptx_sim_num_insn, ctaid.x,ctaid.y,ctaid.z,tid.x,tid.y,tid.z ); + m_gpu->gpgpu_ctx->func_sim->g_ptx_sim_num_insn, ctaid.x,ctaid.y,ctaid.z,tid.x,tid.y,tid.z ); fflush(stdout); } diff --git a/src/cuda-sim/cuda-sim.h b/src/cuda-sim/cuda-sim.h index 76450dc..4566dc2 100644 --- a/src/cuda-sim/cuda-sim.h +++ b/src/cuda-sim/cuda-sim.h @@ -134,11 +134,15 @@ struct gpgpu_ptx_sim_info get_ptxinfo(); class cuda_sim { public: + cuda_sim() { + g_ptx_sim_num_insn = 0; + } //global variables char *opcode_latency_int; int cp_count; int cp_cta_resume; int g_ptxinfo_error_detected; + unsigned g_ptx_sim_num_insn; //global functions void ptx_opcocde_latency_options (option_parser_t opp); void gpgpu_cuda_ptx_sim_main_func( kernel_info_t &kernel, bool openCL = false ); diff --git a/src/gpgpu-sim/gpu-sim.cc b/src/gpgpu-sim/gpu-sim.cc index 39acdd9..4f9ccbf 100644 --- a/src/gpgpu-sim/gpu-sim.cc +++ b/src/gpgpu-sim/gpu-sim.cc @@ -699,7 +699,7 @@ void gpgpu_sim::stop_all_running_kernels(){ } gpgpu_sim::gpgpu_sim( const gpgpu_sim_config &config, gpgpu_context* ctx ) - : gpgpu_t(config), m_config(config) + : gpgpu_t(config, ctx), m_config(config) { gpgpu_ctx = ctx; m_shader_config = &m_config.m_shader_config; diff --git a/src/gpgpu-sim/gpu-sim.h b/src/gpgpu-sim/gpu-sim.h index b47ab16..119b934 100644 --- a/src/gpgpu-sim/gpu-sim.h +++ b/src/gpgpu-sim/gpu-sim.h @@ -62,6 +62,8 @@ #define SAMPLELOG 222 #define DUMPLOG 333 +class gpgpu_context; + extern tr1_hash_map address_random_interleaving; enum dram_ctrl_t { -- cgit v1.3 From 39910217c8137f5ff964d2f991a688020e2d5c8c Mon Sep 17 00:00:00 2001 From: Mengchi Zhang Date: Tue, 9 Jul 2019 09:53:55 -0400 Subject: Move opcode_latency_* Signed-off-by: Mengchi Zhang --- src/cuda-sim/cuda-sim.cc | 9 ++++----- src/cuda-sim/cuda-sim.h | 6 ++++-- src/gpgpu-sim/shader.cc | 8 ++++---- 3 files changed, 12 insertions(+), 11 deletions(-) (limited to 'src/gpgpu-sim') diff --git a/src/cuda-sim/cuda-sim.cc b/src/cuda-sim/cuda-sim.cc index 939358b..f8c21bc 100644 --- a/src/cuda-sim/cuda-sim.cc +++ b/src/cuda-sim/cuda-sim.cc @@ -64,7 +64,6 @@ addr_t g_debug_pc = 0xBEEF1518; unsigned gpgpu_param_num_shaders = 0; -char *opcode_latency_fp, *opcode_latency_dp,*opcode_latency_sfu,*opcode_latency_tensor; char *opcode_initiation_int, *opcode_initiation_fp, *opcode_initiation_dp,*opcode_initiation_sfu,*opcode_initiation_tensor; char *cdp_latency_str; unsigned cdp_latency[5]; @@ -668,15 +667,15 @@ void ptx_instruction::set_opcode_and_latency() sscanf(gpgpu_ctx->func_sim->opcode_latency_int, "%u,%u,%u,%u,%u", &int_latency[0],&int_latency[1],&int_latency[2], &int_latency[3],&int_latency[4]); - sscanf(opcode_latency_fp, "%u,%u,%u,%u,%u", + sscanf(gpgpu_ctx->func_sim->opcode_latency_fp, "%u,%u,%u,%u,%u", &fp_latency[0],&fp_latency[1],&fp_latency[2], &fp_latency[3],&fp_latency[4]); - sscanf(opcode_latency_dp, "%u,%u,%u,%u,%u", + sscanf(gpgpu_ctx->func_sim->opcode_latency_dp, "%u,%u,%u,%u,%u", &dp_latency[0],&dp_latency[1],&dp_latency[2], &dp_latency[3],&dp_latency[4]); - sscanf(opcode_latency_sfu, "%u", + sscanf(gpgpu_ctx->func_sim->opcode_latency_sfu, "%u", &sfu_latency); - sscanf(opcode_latency_tensor, "%u", + sscanf(gpgpu_ctx->func_sim->opcode_latency_tensor, "%u", &tensor_latency); sscanf(opcode_initiation_int, "%u,%u,%u,%u,%u", &int_init[0],&int_init[1],&int_init[2], diff --git a/src/cuda-sim/cuda-sim.h b/src/cuda-sim/cuda-sim.h index 4566dc2..977e3fc 100644 --- a/src/cuda-sim/cuda-sim.h +++ b/src/cuda-sim/cuda-sim.h @@ -47,8 +47,6 @@ extern int g_debug_thread_uid; extern void ** g_inst_classification_stat; extern void ** g_inst_op_classification_stat; extern int g_ptx_kernel_count; // used for classification stat collection purposes -extern char *opcode_latency_fp, *opcode_latency_dp,*opcode_latency_sfu,*opcode_latency_tensor; - extern class kernel_info_t *gpgpu_opencl_ptx_sim_init_grid(class function_info *entry, gpgpu_ptx_sim_arg_list_t args, @@ -139,6 +137,10 @@ class cuda_sim { } //global variables char *opcode_latency_int; + char *opcode_latency_fp; + char *opcode_latency_dp; + char *opcode_latency_sfu; + char *opcode_latency_tensor; int cp_count; int cp_cta_resume; int g_ptxinfo_error_detected; diff --git a/src/gpgpu-sim/shader.cc b/src/gpgpu-sim/shader.cc index 4d12068..69790fc 100644 --- a/src/gpgpu-sim/shader.cc +++ b/src/gpgpu-sim/shader.cc @@ -3022,15 +3022,15 @@ void shader_core_config::set_pipeline_latency() { sscanf(gpgpu_ctx->func_sim->opcode_latency_int, "%u,%u,%u,%u,%u", &int_latency[0],&int_latency[1],&int_latency[2], &int_latency[3],&int_latency[4]); - sscanf(opcode_latency_fp, "%u,%u,%u,%u,%u", + sscanf(gpgpu_ctx->func_sim->opcode_latency_fp, "%u,%u,%u,%u,%u", &fp_latency[0],&fp_latency[1],&fp_latency[2], &fp_latency[3],&fp_latency[4]); - sscanf(opcode_latency_dp, "%u,%u,%u,%u,%u", + sscanf(gpgpu_ctx->func_sim->opcode_latency_dp, "%u,%u,%u,%u,%u", &dp_latency[0],&dp_latency[1],&dp_latency[2], &dp_latency[3],&dp_latency[4]); - sscanf(opcode_latency_sfu, "%u", + sscanf(gpgpu_ctx->func_sim->opcode_latency_sfu, "%u", &sfu_latency); - sscanf(opcode_latency_tensor, "%u", + sscanf(gpgpu_ctx->func_sim->opcode_latency_tensor, "%u", &tensor_latency); //all div operation are executed on sfu -- cgit v1.3 From 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') 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') 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') 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') 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 87726d32ada00fcd93f2cf24ccae4ba593c4f9ec Mon Sep 17 00:00:00 2001 From: Mengchi Zhang Date: Wed, 10 Jul 2019 15:03:16 -0400 Subject: Move g_pc_to_finfo Signed-off-by: Mengchi Zhang --- src/abstract_hardware_model.cc | 3 ++- src/abstract_hardware_model.h | 13 ++++++++----- src/cuda-sim/cuda-sim.cc | 7 +++---- src/cuda-sim/cuda-sim.h | 5 +++-- src/cuda-sim/ptx_ir.cc | 12 +++++++----- src/cuda-sim/ptx_ir.h | 9 +++++++-- src/cuda-sim/ptx_parser.cc | 2 +- src/gpgpu-sim/shader.cc | 4 ++-- src/gpgpu-sim/shader.h | 2 +- 9 files changed, 34 insertions(+), 23 deletions(-) (limited to 'src/gpgpu-sim') diff --git a/src/abstract_hardware_model.cc b/src/abstract_hardware_model.cc index 248e7a5..ef09051 100644 --- a/src/abstract_hardware_model.cc +++ b/src/abstract_hardware_model.cc @@ -39,6 +39,7 @@ #include #include #include +#include "../libcuda/gpgpu_context.h" unsigned mem_access_t::sm_next_access_uid = 0; unsigned warp_inst_t::sm_next_uid = 0; @@ -945,7 +946,7 @@ void simt_stack::print (FILE *fout) const } else { fprintf(fout," " ); } - ptx_print_insn( stack_entry.m_pc, fout ); + m_gpu->gpgpu_ctx->func_sim->ptx_print_insn( stack_entry.m_pc, fout ); fprintf(fout,"\n"); } diff --git a/src/abstract_hardware_model.h b/src/abstract_hardware_model.h index da29a11..8ef8376 100644 --- a/src/abstract_hardware_model.h +++ b/src/abstract_hardware_model.h @@ -347,9 +347,11 @@ public: mutable bool volta_cache_config_set; }; -struct core_config { - core_config() - { +class core_config { + public: + core_config(gpgpu_context* ctx) + { + gpgpu_ctx = ctx; m_valid = false; num_shmem_bank=16; shmem_limited_broadcast = false; @@ -361,6 +363,8 @@ struct core_config { bool m_valid; unsigned warp_size; + // backward pointer + class gpgpu_context* gpgpu_ctx; // off-chip memory request architecture parameters int gpgpu_coalesce_arch; @@ -934,7 +938,7 @@ public: m_empty=true; m_config=NULL; } - warp_inst_t( const core_config *config ) + warp_inst_t( const core_config *config ) { m_uid=0; assert(config->warp_size<=MAX_WARP_SIZE); @@ -1105,7 +1109,6 @@ public: unsigned get_uid() const { return m_uid; } unsigned get_schd_id() const { return m_scheduler_id; } - protected: unsigned m_uid; diff --git a/src/cuda-sim/cuda-sim.cc b/src/cuda-sim/cuda-sim.cc index 3d1da62..fb9bc9e 100644 --- a/src/cuda-sim/cuda-sim.cc +++ b/src/cuda-sim/cuda-sim.cc @@ -220,7 +220,6 @@ void gpgpu_t::gpgpu_ptx_sim_unbindTexture(const struct textureReference* texref) } unsigned g_assemble_code_next_pc=0; -std::map g_pc_to_finfo; std::vector function_info::s_g_pc_to_insn; #define MAX_INST_SIZE 8 /*bytes*/ @@ -257,7 +256,7 @@ void function_info::ptx_assemble() const symbol *l = pI->get_label(); labels[l->name()] = n; } else { - g_pc_to_finfo[PC] = this; + gpgpu_ctx->func_sim->g_pc_to_finfo[PC] = this; m_instr_mem[n] = pI; s_g_pc_to_insn.push_back(pI); assert(pI == s_g_pc_to_insn[PC]); @@ -497,7 +496,7 @@ void gpgpu_t::gpu_memset( size_t dst_start_addr, int c, size_t count ) } } -void ptx_print_insn( address_type pc, FILE *fp ) +void cuda_sim::ptx_print_insn( address_type pc, FILE *fp ) { std::map::iterator f = g_pc_to_finfo.find(pc); if( f == g_pc_to_finfo.end() ) { @@ -509,7 +508,7 @@ void ptx_print_insn( address_type pc, FILE *fp ) finfo->print_insn(pc,fp); } -std::string ptx_get_insn_str( address_type pc ) +std::string cuda_sim::ptx_get_insn_str( address_type pc ) { std::map::iterator f = g_pc_to_finfo.find(pc); if( f == g_pc_to_finfo.end() ) { diff --git a/src/cuda-sim/cuda-sim.h b/src/cuda-sim/cuda-sim.h index 5c95100..c578524 100644 --- a/src/cuda-sim/cuda-sim.h +++ b/src/cuda-sim/cuda-sim.h @@ -61,8 +61,6 @@ unsigned ptx_sim_init_thread( kernel_info_t &kernel, bool functionalSimulationMode = false); const warp_inst_t *ptx_fetch_inst( address_type pc ); const struct gpgpu_ptx_sim_info* ptx_sim_kernel_info(const class function_info *kernel); -void ptx_print_insn( address_type pc, FILE *fp ); -std::string ptx_get_insn_str( address_type pc ); /*! @@ -162,6 +160,7 @@ class cuda_sim { void ** g_inst_op_classification_stat; std::set g_globals; std::set g_constants; + std::map g_pc_to_finfo; // backward pointer class gpgpu_context* gpgpu_ctx; //global functions @@ -181,6 +180,8 @@ class cuda_sim { struct rec_pts find_reconvergence_points( function_info *finfo ); address_type get_converge_point( address_type pc ); void gpgpu_ptx_sim_memcpy_symbol(const char *hostVar, const void *src, size_t count, size_t offset, int to, gpgpu_t *gpu ); + void ptx_print_insn( address_type pc, FILE *fp ); + std::string ptx_get_insn_str( address_type pc ); }; #endif diff --git a/src/cuda-sim/ptx_ir.cc b/src/cuda-sim/ptx_ir.cc index 1bd409e..c537091 100644 --- a/src/cuda-sim/ptx_ir.cc +++ b/src/cuda-sim/ptx_ir.cc @@ -83,8 +83,9 @@ symbol_table::symbol_table() assert(0); } -symbol_table::symbol_table( const char *scope_name, unsigned entry_point, symbol_table *parent ) +symbol_table::symbol_table( const char *scope_name, unsigned entry_point, symbol_table *parent, gpgpu_context* ctx ) { + gpgpu_ctx = ctx; m_scope_name = std::string(scope_name); m_reg_allocator=0; m_shared_next = 0; @@ -183,7 +184,7 @@ symbol_table* symbol_table::start_inst_group() { //previous added assert(m_inst_group_symtab.find(std::string(inst_group_name)) == m_inst_group_symtab.end()); - symbol_table *sym_table = new symbol_table(inst_group_name, 3/*inst group*/, this ); + symbol_table *sym_table = new symbol_table(inst_group_name, 3/*inst group*/, this, gpgpu_ctx ); sym_table->m_global_next = m_global_next; sym_table->m_shared_next = m_shared_next; @@ -221,7 +222,7 @@ bool symbol_table::add_function_decl( const char *name, int entry_point, functio *func_info = m_function_info_lookup[key]; prior_decl = true; } else { - *func_info = new function_info(entry_point); + *func_info = new function_info(entry_point, gpgpu_ctx); (*func_info)->set_name(name); (*func_info)->set_maxnt_id(0); m_function_info_lookup[key] = *func_info; @@ -232,7 +233,7 @@ bool symbol_table::add_function_decl( const char *name, int entry_point, functio *sym_table = m_function_symtab_lookup[key]; } else { assert( !prior_decl ); - *sym_table = new symbol_table( "", entry_point, this ); + *sym_table = new symbol_table( "", entry_point, this, gpgpu_ctx ); // Initial setup code to support a register represented as "_". // This register is used when an instruction operand is @@ -1373,8 +1374,9 @@ std::string ptx_instruction::to_string() const unsigned function_info::sm_next_uid = 1; -function_info::function_info(int entry_point ) +function_info::function_info(int entry_point, gpgpu_context* ctx ) { + gpgpu_ctx = ctx; m_uid = sm_next_uid++; m_entry_point = (entry_point==1)?true:false; m_extern = (entry_point==2)?true:false; diff --git a/src/cuda-sim/ptx_ir.h b/src/cuda-sim/ptx_ir.h index 1604551..babd54b 100644 --- a/src/cuda-sim/ptx_ir.h +++ b/src/cuda-sim/ptx_ir.h @@ -306,7 +306,7 @@ private: class symbol_table { public: symbol_table(); - symbol_table( const char *scope_name, unsigned entry_point, symbol_table *parent ); + symbol_table( const char *scope_name, unsigned entry_point, symbol_table *parent, gpgpu_context* ctx); void set_name( const char *name ); const ptx_version &get_ptx_version() const; unsigned get_sm_target() const; @@ -348,6 +348,9 @@ public: symbol_table* start_inst_group(); symbol_table* end_inst_group(); + // backward pointer + class gpgpu_context* gpgpu_ctx; + private: unsigned m_reg_allocator; unsigned m_shared_next; @@ -1233,7 +1236,7 @@ private: class function_info { public: - function_info(int entry_point ); + function_info(int entry_point, gpgpu_context* ctx ); const ptx_version &get_ptx_version() const { return m_symtab->get_ptx_version(); } unsigned get_sm_target() const { return m_symtab->get_sm_target(); } bool is_extern() const { return m_extern; } @@ -1403,6 +1406,8 @@ public: void set_maxnt_id(unsigned maxthreads) { maxnt_id = maxthreads;} unsigned get_maxnt_id() { return maxnt_id;} + // backward pointer + class gpgpu_context* gpgpu_ctx; private: unsigned maxnt_id; diff --git a/src/cuda-sim/ptx_parser.cc b/src/cuda-sim/ptx_parser.cc index 269ec4d..5a94679 100644 --- a/src/cuda-sim/ptx_parser.cc +++ b/src/cuda-sim/ptx_parser.cc @@ -112,7 +112,7 @@ symbol_table * gpgpu_context::init_parser( const char *ptx_filename ) { g_filename = strdup(ptx_filename); if (g_global_allfiles_symbol_table == NULL) { - g_global_allfiles_symbol_table = new symbol_table("global_allfiles", 0, NULL); + g_global_allfiles_symbol_table = new symbol_table("global_allfiles", 0, NULL, this); ptx_parser->g_global_symbol_table = ptx_parser->g_current_symbol_table = g_global_allfiles_symbol_table; } /*else { diff --git a/src/gpgpu-sim/shader.cc b/src/gpgpu-sim/shader.cc index 69790fc..6cd6d8f 100644 --- a/src/gpgpu-sim/shader.cc +++ b/src/gpgpu-sim/shader.cc @@ -1025,7 +1025,7 @@ void scheduler_unit::cycle() m_simt_stack[warp_id]->get_pdom_stack_top_info(&pc,&rpc); SCHED_DPRINTF( "Warp (warp_id %u, dynamic_warp_id %u) has valid instruction (%s)\n", (*iter)->get_warp_id(), (*iter)->get_dynamic_warp_id(), - ptx_get_insn_str( pc).c_str() ); + m_shader->m_config->gpgpu_ctx->func_sim->ptx_get_insn_str( pc).c_str() ); if( pI ) { assert(valid); if( pc != pI->pc ) { @@ -2690,7 +2690,7 @@ void warp_inst_t::print( FILE *fout ) const for (unsigned j=0; jwarp_size; j++) fprintf(fout, "%c", (active(j)?'1':'0') ); fprintf(fout, "]: "); - ptx_print_insn( pc, fout ); + m_config->gpgpu_ctx->func_sim->ptx_print_insn( pc, fout ); fprintf(fout, "\n"); } void shader_core_ctx::incexecstat(warp_inst_t *&inst) diff --git a/src/gpgpu-sim/shader.h b/src/gpgpu-sim/shader.h index e0cefac..2837f1b 100644 --- a/src/gpgpu-sim/shader.h +++ b/src/gpgpu-sim/shader.h @@ -1367,7 +1367,7 @@ const char* const pipeline_stage_name_decode[] = { class shader_core_config : public core_config { public: - shader_core_config(gpgpu_context* ctx){ + shader_core_config(gpgpu_context* ctx):core_config(ctx){ pipeline_widths_string = NULL; gpgpu_ctx = ctx; } -- cgit v1.3 From 57f8e9bd4e73757f4026e8b257fb625d465e0271 Mon Sep 17 00:00:00 2001 From: Mengchi Zhang Date: Wed, 10 Jul 2019 15:09:34 -0400 Subject: Move gpgpu_ptx_instruction_classification Signed-off-by: Mengchi Zhang --- src/cuda-sim/cuda-sim.cc | 3 +-- src/cuda-sim/cuda-sim.h | 1 + src/debug.h | 2 -- src/gpgpu-sim/gpu-sim.cc | 6 +++--- 4 files changed, 5 insertions(+), 7 deletions(-) (limited to 'src/gpgpu-sim') diff --git a/src/cuda-sim/cuda-sim.cc b/src/cuda-sim/cuda-sim.cc index fb9bc9e..b370400 100644 --- a/src/cuda-sim/cuda-sim.cc +++ b/src/cuda-sim/cuda-sim.cc @@ -53,7 +53,6 @@ typedef void * yyscan_t; #include "cuda_device_runtime.h" #include "../../libcuda/gpgpu_context.h" -int gpgpu_ptx_instruction_classification; int g_debug_execution = 0; int g_debug_thread_uid = 0; addr_t g_debug_pc = 0xBEEF1518; @@ -1681,7 +1680,7 @@ void ptx_thread_info::ptx_exec_inst( warp_inst_t &inst, unsigned lane_id) if(!(this->m_functionalSimulationMode)) ptx_file_line_stats_add_exec_count(pI); - if ( gpgpu_ptx_instruction_classification ) { + if ( m_gpu->gpgpu_ctx->func_sim->gpgpu_ptx_instruction_classification ) { m_gpu->gpgpu_ctx->func_sim->init_inst_classification_stat(); unsigned space_type=0; switch ( pI->get_space().get_type() ) { diff --git a/src/cuda-sim/cuda-sim.h b/src/cuda-sim/cuda-sim.h index c578524..89f67cf 100644 --- a/src/cuda-sim/cuda-sim.h +++ b/src/cuda-sim/cuda-sim.h @@ -161,6 +161,7 @@ class cuda_sim { std::set g_globals; std::set g_constants; std::map g_pc_to_finfo; + int gpgpu_ptx_instruction_classification; // backward pointer class gpgpu_context* gpgpu_ctx; //global functions diff --git a/src/debug.h b/src/debug.h index 1277494..1005bd5 100644 --- a/src/debug.h +++ b/src/debug.h @@ -83,8 +83,6 @@ private: unsigned m_value; }; -extern int gpgpu_ptx_instruction_classification ; - class ptx_thread_info; class ptx_instruction; bool thread_at_brkpt( ptx_thread_info *thd_info, const class brk_pt &b ); diff --git a/src/gpgpu-sim/gpu-sim.cc b/src/gpgpu-sim/gpu-sim.cc index 30411f1..93f041a 100644 --- a/src/gpgpu-sim/gpu-sim.cc +++ b/src/gpgpu-sim/gpu-sim.cc @@ -499,8 +499,8 @@ void gpgpu_sim_config::reg_options(option_parser_t opp) option_parser_register(opp, "-gpgpu_deadlock_detect", OPT_BOOL, &gpu_deadlock_detect, "Stop the simulation at deadlock (1=on (default), 0=off)", "1"); - option_parser_register(opp, "-gpgpu_ptx_instruction_classification", OPT_INT32, - &gpgpu_ptx_instruction_classification, + option_parser_register(opp, "-gpgpu_ptx_instruction_classification", OPT_INT32, + &(gpgpu_ctx->func_sim->gpgpu_ptx_instruction_classification), "if enabled will classify ptx instruction types per kernel (Max 255 kernels now)", "0"); option_parser_register(opp, "-gpgpu_ptx_sim_mode", OPT_INT32, &(gpgpu_ctx->func_sim->g_ptx_sim_mode), @@ -1200,7 +1200,7 @@ void gpgpu_sim::gpu_print_stat() spill_log_to_file (stdout, 1, gpu_sim_cycle); insn_warp_occ_print(stdout); } - if ( gpgpu_ptx_instruction_classification ) { + if ( gpgpu_ctx->func_sim->gpgpu_ptx_instruction_classification ) { StatDisp( gpgpu_ctx->func_sim->g_inst_classification_stat[gpgpu_ctx->func_sim->g_ptx_kernel_count]); StatDisp( gpgpu_ctx->func_sim->g_inst_op_classification_stat[gpgpu_ctx->func_sim->g_ptx_kernel_count]); } -- cgit v1.3 From e83a9c978ebd08eed616eb6fdbdef88160232076 Mon Sep 17 00:00:00 2001 From: Mengchi Zhang Date: Wed, 10 Jul 2019 16:09:54 -0400 Subject: Move cdp_latency Signed-off-by: Mengchi Zhang --- src/cuda-sim/cuda-sim.cc | 8 +++++--- src/cuda-sim/cuda-sim.h | 1 + src/gpgpu-sim/shader.cc | 7 +++---- 3 files changed, 9 insertions(+), 7 deletions(-) (limited to 'src/gpgpu-sim') diff --git a/src/cuda-sim/cuda-sim.cc b/src/cuda-sim/cuda-sim.cc index b370400..bd96f7e 100644 --- a/src/cuda-sim/cuda-sim.cc +++ b/src/cuda-sim/cuda-sim.cc @@ -58,7 +58,6 @@ int g_debug_thread_uid = 0; addr_t g_debug_pc = 0xBEEF1518; // Output debug information to file options -unsigned cdp_latency[5]; void cuda_sim::ptx_opcocde_latency_options (option_parser_t opp) { option_parser_register(opp, "-ptx_opcode_latency_int", OPT_CSTR, &opcode_latency_int, @@ -680,8 +679,11 @@ void ptx_instruction::set_opcode_and_latency() sscanf(gpgpu_ctx->func_sim->opcode_initiation_tensor, "%u", &tensor_init); sscanf(gpgpu_ctx->func_sim->cdp_latency_str, "%u,%u,%u,%u,%u", - &cdp_latency[0],&cdp_latency[1],&cdp_latency[2], - &cdp_latency[3],&cdp_latency[4]); + &gpgpu_ctx->func_sim->cdp_latency[0], + &gpgpu_ctx->func_sim->cdp_latency[1], + &gpgpu_ctx->func_sim->cdp_latency[2], + &gpgpu_ctx->func_sim->cdp_latency[3], + &gpgpu_ctx->func_sim->cdp_latency[4]); if(!m_operands.empty()){ std::vector::iterator it; diff --git a/src/cuda-sim/cuda-sim.h b/src/cuda-sim/cuda-sim.h index 89f67cf..0b6e84f 100644 --- a/src/cuda-sim/cuda-sim.h +++ b/src/cuda-sim/cuda-sim.h @@ -163,6 +163,7 @@ class cuda_sim { std::map g_pc_to_finfo; int gpgpu_ptx_instruction_classification; // backward pointer + unsigned cdp_latency[5]; class gpgpu_context* gpgpu_ctx; //global functions void ptx_opcocde_latency_options (option_parser_t opp); diff --git a/src/gpgpu-sim/shader.cc b/src/gpgpu-sim/shader.cc index 6cd6d8f..f380560 100644 --- a/src/gpgpu-sim/shader.cc +++ b/src/gpgpu-sim/shader.cc @@ -1084,12 +1084,11 @@ void scheduler_unit::cycle() if(pI->m_is_cdp && !warp(warp_id).m_cdp_dummy) { assert(warp(warp_id).m_cdp_latency == 0); - extern unsigned cdp_latency[5]; if(pI->m_is_cdp == 1) - warp(warp_id).m_cdp_latency = cdp_latency[pI->m_is_cdp - 1]; + warp(warp_id).m_cdp_latency = m_shader->m_config->gpgpu_ctx->func_sim->cdp_latency[pI->m_is_cdp - 1]; else //cudaLaunchDeviceV2 and cudaGetParameterBufferV2 - warp(warp_id).m_cdp_latency = cdp_latency[pI->m_is_cdp - 1] - + cdp_latency[pI->m_is_cdp] * active_mask.count(); + warp(warp_id).m_cdp_latency = m_shader->m_config->gpgpu_ctx->func_sim->cdp_latency[pI->m_is_cdp - 1] + + m_shader->m_config->gpgpu_ctx->func_sim->cdp_latency[pI->m_is_cdp] * active_mask.count(); warp(warp_id).m_cdp_dummy = true; break; } -- cgit v1.3