From 6e051c42b8c9c820af50a10025d182a96fdd12a6 Mon Sep 17 00:00:00 2001 From: Mengchi Zhang Date: Tue, 9 Jul 2019 10:06:31 -0400 Subject: Move m_ptx_save_converted_ptxplus Signed-off-by: Mengchi Zhang --- libcuda/cuda_runtime_api.cc | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) (limited to 'libcuda/cuda_runtime_api.cc') diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc index eb645ce..43c5e1f 100644 --- a/libcuda/cuda_runtime_api.cc +++ b/libcuda/cuda_runtime_api.cc @@ -2936,7 +2936,7 @@ void gpgpu_context::cuobjdumpParseBinary(unsigned int handle){ if(context->get_device()->get_gpgpu()->get_config().convert_to_ptxplus() ) { cuobjdumpELFSection* elfsection = api->findELFSection(ptx->getIdentifier()); assert (elfsection!= NULL); - char *ptxplus_str = gpgpu_ptx_sim_convert_ptx_and_sass_to_ptxplus( + char *ptxplus_str = ptxinfo->gpgpu_ptx_sim_convert_ptx_and_sass_to_ptxplus( ptx->getPTXfilename(), elfsection->getELFfilename(), elfsection->getSASSfilename()); -- 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 'libcuda/cuda_runtime_api.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 b3655bf28a7402db347f9d7f87049806b9315a05 Mon Sep 17 00:00:00 2001 From: Mengchi Zhang Date: Tue, 9 Jul 2019 15:03:18 -0400 Subject: Move g_global_name_lookup and g_constant_name_lookup Signed-off-by: Mengchi Zhang --- libcuda/cuda_runtime_api.cc | 4 ++-- src/cuda-sim/cuda-sim.cc | 14 ++++++-------- src/cuda-sim/cuda-sim.h | 6 ++++-- 3 files changed, 12 insertions(+), 12 deletions(-) (limited to 'libcuda/cuda_runtime_api.cc') diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc index 43c8bae..bbbaf23 100644 --- a/libcuda/cuda_runtime_api.cc +++ b/libcuda/cuda_runtime_api.cc @@ -767,9 +767,9 @@ void cudaRegisterVarInternal( ctx->cuobjdumpParseBinary((unsigned)(unsigned long long)fatCubinHandle); fflush(stdout); if ( constant && !global && !ext ) { - gpgpu_ptx_sim_register_const_variable(hostVar,deviceName,size); + ctx->func_sim->gpgpu_ptx_sim_register_const_variable(hostVar,deviceName,size); } else if ( !constant && !global && !ext ) { - gpgpu_ptx_sim_register_global_variable(hostVar,deviceName,size); + ctx->func_sim->gpgpu_ptx_sim_register_global_variable(hostVar,deviceName,size); } else cuda_not_implemented(__my_func__,__LINE__); } diff --git a/src/cuda-sim/cuda-sim.cc b/src/cuda-sim/cuda-sim.cc index 0ed125a..b3e2965 100644 --- a/src/cuda-sim/cuda-sim.cc +++ b/src/cuda-sim/cuda-sim.cc @@ -1948,18 +1948,16 @@ void print_splash() } } -std::map g_const_name_lookup; // indexed by hostVar -std::map g_global_name_lookup; // indexed by hostVar std::set g_globals; std::set g_constants; -void gpgpu_ptx_sim_register_const_variable(void *hostVar, const char *deviceName, size_t size ) +void cuda_sim::gpgpu_ptx_sim_register_const_variable(void *hostVar, const char *deviceName, size_t size ) { printf("GPGPU-Sim PTX registering constant %s (%zu bytes) to name mapping\n", deviceName, size ); g_const_name_lookup[hostVar] = deviceName; } -void gpgpu_ptx_sim_register_global_variable(void *hostVar, const char *deviceName, size_t size ) +void cuda_sim::gpgpu_ptx_sim_register_global_variable(void *hostVar, const char *deviceName, size_t size ) { printf("GPGPU-Sim PTX registering global %s hostVar to name mapping\n", deviceName ); g_global_name_lookup[hostVar] = deviceName; @@ -1972,14 +1970,14 @@ void gpgpu_ptx_sim_memcpy_symbol(const char *hostVar, const void *src, size_t co memory_space_t mem_region = undefined_space; std::string sym_name; - std::map::iterator c=g_const_name_lookup.find(hostVar); - if ( c!=g_const_name_lookup.end() ) { + std::map::iterator c=gpu->gpgpu_ctx->func_sim->g_const_name_lookup.find(hostVar); + if ( c!=gpu->gpgpu_ctx->func_sim->g_const_name_lookup.end() ) { found_sym = true; sym_name = c->second; mem_region = const_space; } - std::map::iterator g=g_global_name_lookup.find(hostVar); - if ( g!=g_global_name_lookup.end() ) { + std::map::iterator g=gpu->gpgpu_ctx->func_sim->g_global_name_lookup.find(hostVar); + if ( g!=gpu->gpgpu_ctx->func_sim->g_global_name_lookup.end() ) { if ( found_sym ) { printf("Execution error: PTX symbol \"%s\" w/ hostVar=0x%Lx is declared both const and global?\n", sym_name.c_str(), (unsigned long long)hostVar ); diff --git a/src/cuda-sim/cuda-sim.h b/src/cuda-sim/cuda-sim.h index e259f1f..25ebf7b 100644 --- a/src/cuda-sim/cuda-sim.h +++ b/src/cuda-sim/cuda-sim.h @@ -48,8 +48,6 @@ extern void ** g_inst_classification_stat; extern void ** g_inst_op_classification_stat; extern void print_splash(); -extern void gpgpu_ptx_sim_register_const_variable(void*, const char *deviceName, size_t size ); -extern void gpgpu_ptx_sim_register_global_variable(void *hostVar, const char *deviceName, size_t size ); extern void gpgpu_ptx_sim_memcpy_symbol(const char *hostVar, const void *src, size_t count, size_t offset, int to, gpgpu_t *gpu ); extern void read_sim_environment_variables(); @@ -147,6 +145,8 @@ class cuda_sim { unsigned g_ptx_sim_num_insn; char *cdp_latency_str; int g_ptx_kernel_count; // used for classification stat collection purposes + std::map g_global_name_lookup; // indexed by hostVar + std::map g_const_name_lookup; // indexed by hostVar //global functions void ptx_opcocde_latency_options (option_parser_t opp); void gpgpu_cuda_ptx_sim_main_func( kernel_info_t &kernel, bool openCL = false ); @@ -157,6 +157,8 @@ class cuda_sim { struct dim3 gridDim, struct dim3 blockDim, gpgpu_t *gpu ); + void gpgpu_ptx_sim_register_global_variable(void *hostVar, const char *deviceName, size_t size ); + void gpgpu_ptx_sim_register_const_variable(void*, const char *deviceName, size_t size ); }; #endif -- cgit v1.3 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 'libcuda/cuda_runtime_api.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