summaryrefslogtreecommitdiff
diff options
context:
space:
mode:
Diffstat
-rw-r--r--libcuda/cuda_api_object.h5
-rw-r--r--libcuda/cuda_runtime_api.cc8
-rw-r--r--libcuda/gpgpu_context.h2
-rw-r--r--libopencl/opencl_runtime_api.cc6
-rw-r--r--src/abstract_hardware_model.cc3
-rw-r--r--src/abstract_hardware_model.h6
-rw-r--r--src/cuda-sim/cuda-sim.cc75
-rw-r--r--src/cuda-sim/cuda-sim.h42
-rw-r--r--src/cuda-sim/ptx_loader.cc5
-rw-r--r--src/cuda-sim/ptx_loader.h4
-rw-r--r--src/cuda-sim/ptx_parser.h1
-rw-r--r--src/cuda-sim/ptxinfo.l4
-rw-r--r--src/gpgpu-sim/gpu-sim.cc6
-rw-r--r--src/gpgpu-sim/gpu-sim.h6
-rw-r--r--src/gpgpu-sim/shader.cc8
-rw-r--r--src/gpgpusim_entrypoint.cc4
-rw-r--r--src/gpgpusim_entrypoint.h1
17 files changed, 100 insertions, 86 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..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__);
}
@@ -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 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 03ec80c..0a6eb3e 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");
@@ -954,9 +956,9 @@ 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 )
- 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/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 df0bbd7..b3e2965 100644
--- a/src/cuda-sim/cuda-sim.cc
+++ b/src/cuda-sim/cuda-sim.cc
@@ -56,20 +56,13 @@ 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;
// Output debug information to file options
-int cp_count;
-int cp_cta_resume;
-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;
-char *opcode_initiation_int, *opcode_initiation_fp, *opcode_initiation_dp,*opcode_initiation_sfu,*opcode_initiation_tensor;
-char *cdp_latency_str;
unsigned cdp_latency[5];
void cuda_sim::ptx_opcocde_latency_options (option_parser_t opp) {
@@ -671,30 +664,30 @@ 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",
+ sscanf(gpgpu_ctx->func_sim->opcode_initiation_int, "%u,%u,%u,%u,%u",
&int_init[0],&int_init[1],&int_init[2],
&int_init[3],&int_init[4]);
- sscanf(opcode_initiation_fp, "%u,%u,%u,%u,%u",
+ sscanf(gpgpu_ctx->func_sim->opcode_initiation_fp, "%u,%u,%u,%u,%u",
&fp_init[0],&fp_init[1],&fp_init[2],
&fp_init[3],&fp_init[4]);
- sscanf(opcode_initiation_dp, "%u,%u,%u,%u,%u",
+ sscanf(gpgpu_ctx->func_sim->opcode_initiation_dp, "%u,%u,%u,%u,%u",
&dp_init[0],&dp_init[1],&dp_init[2],
&dp_init[3],&dp_init[4]);
- sscanf(opcode_initiation_sfu, "%u",
+ sscanf(gpgpu_ctx->func_sim->opcode_initiation_sfu, "%u",
&sfu_init);
- sscanf(opcode_initiation_tensor, "%u",
+ sscanf(gpgpu_ctx->func_sim->opcode_initiation_tensor, "%u",
&tensor_init);
- sscanf(cdp_latency_str, "%u,%u,%u,%u,%u",
+ 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]);
@@ -1486,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() )
@@ -1631,7 +1624,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(),
@@ -1689,14 +1682,14 @@ 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))
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;
@@ -1712,15 +1705,15 @@ 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 ( (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);
}
@@ -1923,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,
@@ -1955,18 +1948,16 @@ void print_splash()
}
}
-std::map<const void*,std::string> g_const_name_lookup; // indexed by hostVar
-std::map<const void*,std::string> g_global_name_lookup; // indexed by hostVar
std::set<std::string> g_globals;
std::set<std::string> 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;
@@ -1979,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<const void*,std::string>::iterator c=g_const_name_lookup.find(hostVar);
- if ( c!=g_const_name_lookup.end() ) {
+ std::map<const void*,std::string>::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<const void*,std::string>::iterator g=g_global_name_lookup.find(hostVar);
- if ( g!=g_global_name_lookup.end() ) {
+ std::map<const void*,std::string>::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 );
@@ -2106,8 +2097,6 @@ void read_sim_environment_variables()
}
}
-ptx_cta_info *g_func_cta_info = NULL;
-
#define MAX(a,b) (((a)>(b))?(a):(b))
unsigned max_cta (const struct gpgpu_ptx_sim_info *kernel_info, unsigned threads_per_cta, unsigned int warp_size, unsigned int n_thread_per_shader, unsigned int gpgpu_shmem_size, unsigned int gpgpu_shader_registers, unsigned int max_cta_per_core)
@@ -2146,7 +2135,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());
@@ -2272,7 +2261,7 @@ void functionalCoreSim::initializeCTA(unsigned ctaid_cp)
assert(m_thread[i]!=NULL && !m_thread[i]->is_done());
char fname[2048];
snprintf(fname,2048,"checkpoint_files/thread_%d_0_reg.txt",i );
- if(cp_cta_resume==1)
+ if(m_gpu->gpgpu_ctx->func_sim->cp_cta_resume==1)
m_thread[i]->resume_reg_thread(fname,symtab);
ctaLiveThreads++;
}
@@ -2296,7 +2285,7 @@ void functionalCoreSim::createWarp(unsigned warpId)
char fname[2048];
snprintf(fname,2048,"checkpoint_files/warp_%d_0_simt.txt",warpId );
- if(cp_cta_resume==1)
+ if(m_gpu->gpgpu_ctx->func_sim->cp_cta_resume==1)
{
unsigned pc,rpc;
m_simt_stack[warpId]->resume(fname);
@@ -2312,8 +2301,8 @@ void functionalCoreSim::createWarp(unsigned warpId)
void functionalCoreSim::execute(int inst_count, unsigned ctaid_cp)
{
- cp_count= m_gpu->checkpoint_insn_Y;
- cp_cta_resume= m_gpu->checkpoint_CTA_t;
+ m_gpu->gpgpu_ctx->func_sim->cp_count= m_gpu->checkpoint_insn_Y;
+ m_gpu->gpgpu_ctx->func_sim->cp_cta_resume= m_gpu->checkpoint_CTA_t;
initializeCTA(ctaid_cp);
int count=0;
@@ -2406,8 +2395,6 @@ unsigned translate_pc_to_ptxlineno(unsigned pc)
extern std::map<unsigned,const char*> get_duplicate();
-int g_ptxinfo_error_detected;
-
static char *g_ptxinfo_kname = NULL;
static struct gpgpu_ptx_sim_info g_ptxinfo;
static std::map<unsigned,const char*> g_duplicate;
diff --git a/src/cuda-sim/cuda-sim.h b/src/cuda-sim/cuda-sim.h
index 96d34f5..25ebf7b 100644
--- a/src/cuda-sim/cuda-sim.h
+++ b/src/cuda-sim/cuda-sim.h
@@ -46,19 +46,8 @@ 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 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,
- 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 );
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();
@@ -135,10 +124,41 @@ struct gpgpu_ptx_sim_info get_ptxinfo();
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;
+ char *opcode_latency_fp;
+ char *opcode_latency_dp;
+ char *opcode_latency_sfu;
+ char *opcode_latency_tensor;
+ char *opcode_initiation_int;
+ char *opcode_initiation_fp;
+ char *opcode_initiation_dp;
+ char *opcode_initiation_sfu;
+ char *opcode_initiation_tensor;
+ int cp_count;
+ int cp_cta_resume;
+ 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
+ std::map<const void*,std::string> g_global_name_lookup; // indexed by hostVar
+ std::map<const void*,std::string> 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 );
+ 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 );
+ 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
diff --git a/src/cuda-sim/ptx_loader.cc b/src/cuda-sim/ptx_loader.cc
index e0d9a11..2bf464c 100644
--- a/src/cuda-sim/ptx_loader.cc
+++ b/src/cuda-sim/ptx_loader.cc
@@ -58,7 +58,6 @@ extern int ptxinfo_lex_destroy(yyscan_t scanner);
static bool g_save_embedded_ptx;
static int g_occupancy_sm_number;
-bool m_ptx_save_converted_ptxplus;
bool ptxinfo_data::keep_intermediate_files() {return g_keep_intermediate_files;}
@@ -71,7 +70,7 @@ void gpgpu_context::ptx_reg_options(option_parser_t opp)
"keep intermediate files created by GPGPU-Sim when interfacing with external programs",
"0");
option_parser_register(opp, "-gpgpu_ptx_save_converted_ptxplus", OPT_BOOL,
- &m_ptx_save_converted_ptxplus,
+ &(ptxinfo->m_ptx_save_converted_ptxplus),
"Saved converted ptxplus to a file",
"0");
option_parser_register(opp, "-gpgpu_occupancy_sm_number", OPT_INT32, &g_occupancy_sm_number,
@@ -106,7 +105,7 @@ void gpgpu_context::print_ptx_file( const char *p, unsigned source_num, const ch
fflush(stdout);
}
-char* gpgpu_ptx_sim_convert_ptx_and_sass_to_ptxplus(const std::string ptxfilename, const std::string elffilename, const std::string sassfilename)
+char* ptxinfo_data::gpgpu_ptx_sim_convert_ptx_and_sass_to_ptxplus(const std::string ptxfilename, const std::string elffilename, const std::string sassfilename)
{
printf("GPGPU-Sim PTX: converting EMBEDDED .ptx file to ptxplus \n");
diff --git a/src/cuda-sim/ptx_loader.h b/src/cuda-sim/ptx_loader.h
index aa3e2f3..decfdc8 100644
--- a/src/cuda-sim/ptx_loader.h
+++ b/src/cuda-sim/ptx_loader.h
@@ -43,14 +43,14 @@ class ptxinfo_data{
const char *g_ptxinfo_filename;
class gpgpu_context* gpgpu_ctx;
bool g_keep_intermediate_files;
+ bool m_ptx_save_converted_ptxplus;
void ptxinfo_addinfo();
bool keep_intermediate_files();
+ char* gpgpu_ptx_sim_convert_ptx_and_sass_to_ptxplus(const std::string ptx_str, const std::string sass_str, const std::string elf_str);
};
extern bool g_override_embedded_ptx;
extern int no_of_ptx; //counter to track number of ptx files to be extracted in an application.
-char* gpgpu_ptx_sim_convert_ptx_and_sass_to_ptxplus(const std::string ptx_str, const std::string sass_str, const std::string elf_str);
-
#endif
diff --git a/src/cuda-sim/ptx_parser.h b/src/cuda-sim/ptx_parser.h
index e09aab1..ef422e5 100644
--- a/src/cuda-sim/ptx_parser.h
+++ b/src/cuda-sim/ptx_parser.h
@@ -30,7 +30,6 @@
#include "../abstract_hardware_model.h"
#include "ptx_ir.h"
-extern int g_error_detected;
#ifdef __cplusplus
const class ptx_instruction *ptx_instruction_lookup( const char *filename, unsigned linenumber );
diff --git a/src/cuda-sim/ptxinfo.l b/src/cuda-sim/ptxinfo.l
index 3a152b0..51371e3 100644
--- a/src/cuda-sim/ptxinfo.l
+++ b/src/cuda-sim/ptxinfo.l
@@ -88,13 +88,11 @@ OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
%%
-extern int g_ptxinfo_error_detected;
-
int ptxinfo_error(yyscan_t yyscanner, ptxinfo_data* ptxinfo, const char* msg)
{
struct yyguts_t * yyg = (struct yyguts_t*)yyscanner;
int i;
- g_ptxinfo_error_detected = 1;
+ ptxinfo->gpgpu_ctx->func_sim->g_ptxinfo_error_detected = 1;
fflush(stdout);
printf("GPGPU-Sim: ERROR while parsing output of ptxas (used to capture resource usage information)\n");
if( msg != NULL )
diff --git a/src/gpgpu-sim/gpu-sim.cc b/src/gpgpu-sim/gpu-sim.cc
index 39acdd9..9f47067 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;
@@ -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
diff --git a/src/gpgpu-sim/gpu-sim.h b/src/gpgpu-sim/gpu-sim.h
index 7eeb7dd..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<new_addr_type,unsigned> address_random_interleaving;
enum dram_ctrl_t {
@@ -494,6 +496,8 @@ public:
*/
simt_core_cluster * getSIMTCluster();
+ // backward pointer
+ class gpgpu_context* gpgpu_ctx;
private:
// clocks
@@ -511,8 +515,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/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
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
generated by cgit v1.3 (git 2.53.0) at 2026-07-11 07:25:58 +0000