diff options
Diffstat (limited to 'libcuda')
| -rw-r--r-- | libcuda/cuda_api_object.h | 5 | ||||
| -rw-r--r-- | libcuda/cuda_runtime_api.cc | 14 | ||||
| -rw-r--r-- | libcuda/gpgpu_context.h | 13 |
3 files changed, 23 insertions, 9 deletions
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<cuobjdumpSection*> cuobjdumpSectionList; @@ -187,6 +188,8 @@ class cuda_runtime_api { std::map<void *, size_t> 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 eb645ce..59d2a60 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__); } @@ -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; @@ -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()); @@ -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 d819559..3c9f87c 100644 --- a/libcuda/gpgpu_context.h +++ b/libcuda/gpgpu_context.h @@ -4,15 +4,21 @@ #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" +#include "../src/cuda-sim/cuda_device_runtime.h" 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); + 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; @@ -22,6 +28,10 @@ class gpgpu_context { ptxinfo_data* ptxinfo; 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 ); @@ -32,6 +42,7 @@ class gpgpu_context { class symbol_table* init_parser(const char*); class gpgpu_sim *gpgpu_ptx_sim_init_perf(); struct _cuda_device_id *GPGPUSim_Init(); + void ptx_reg_options(option_parser_t opp); }; gpgpu_context* GPGPU_Context(); |