summaryrefslogtreecommitdiff
diff options
context:
space:
mode:
Diffstat
-rw-r--r--libcuda/cuda_api_object.h5
-rw-r--r--libcuda/cuda_runtime_api.cc2
-rw-r--r--libcuda/gpgpu_context.h2
-rw-r--r--libopencl/opencl_runtime_api.cc2
-rw-r--r--src/cuda-sim/cuda-sim.cc13
-rw-r--r--src/cuda-sim/cuda-sim.h14
-rw-r--r--src/gpgpu-sim/gpu-sim.cc4
7 files changed, 23 insertions, 19 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 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<unsigned> 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
generated by cgit v1.3 (git 2.53.0) at 2026-07-10 22:42:41 +0000