diff options
| -rw-r--r-- | libopencl/opencl_runtime_api.cc | 4 | ||||
| -rw-r--r-- | src/cuda-sim/cuda-sim.cc | 5 | ||||
| -rw-r--r-- | src/cuda-sim/cuda-sim.h | 4 | ||||
| -rw-r--r-- | src/gpgpu-sim/gpu-sim.h | 4 | ||||
| -rw-r--r-- | src/gpgpusim_entrypoint.cc | 4 | ||||
| -rw-r--r-- | src/gpgpusim_entrypoint.h | 1 |
6 files changed, 12 insertions, 10 deletions
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 |